Skip to content

Commit 273034a

Browse files
[SYCL] [FPGA] Update the experimental latency control API to use property list (#5993)
As planned before, this patch deprecates the template argument approach in the experimental latency control API, and use property list instead. I'm still updating the experimental API rather than the formal API because the frontend support (clang and SPIR-V) for latency control is not ready yet, so I'd like to wait for them before fully moving on to deprecate the experimental API. --- This patch contains three parts: - Update the extension documents for pipes and LSUs, change the experimental latency control to use property list in the API. - sycl/doc/extensions/supported/sycl_ext_intel_dataflow_pipes.asciidoc - sycl/doc/extensions/supported/sycl_ext_intel_fpga_lsu.md - Create properties for latency control. - sycl/include/sycl/ext/oneapi/latency_control/properties.hpp - sycl/include/sycl/ext/oneapi/properties/property.hpp - sycl/test/extensions/properties/properties_latency_control.cpp - Update headers for the experimental latency control, change the API to use property list. - sycl/include/sycl/ext/intel/experimental/fpga_utils.hpp - sycl/include/sycl/ext/intel/experimental/fpga_lsu.hpp - sycl/include/sycl/ext/intel/experimental/pipes.hpp Test: intel/llvm-test-suite#982
1 parent b2d4d67 commit 273034a

File tree

8 files changed

+503
-381
lines changed

8 files changed

+503
-381
lines changed

sycl/doc/extensions/supported/sycl_ext_intel_dataflow_pipes.asciidoc

Lines changed: 87 additions & 63 deletions
Original file line numberDiff line numberDiff line change
@@ -110,9 +110,9 @@ A pipe type is a specialization of the pipe class:
110110

111111
[source,c++,Pipe type def,linenums]
112112
----
113-
template <typename name,
114-
typename dataT,
115-
size_t min_capacity = 0>
113+
template <typename Name,
114+
typename DataT,
115+
size_t MinCapacity = 0>
116116
class pipe;
117117
----
118118

@@ -131,7 +131,7 @@ using pipe<class bar, float, 5>;
131131

132132
The interface of a pipe is through static member functions, and instances of a pipe class cannot be instantiated. Allowing instances of pipe objects, when their type defines connectivity, would introduce an error prone secondary mechanism of reference.
133133

134-
The first template parameter, `name`, can be any type, and is typically expected to be a user defined class in a user namespace. The type only needs to be forward declared, and not defined.
134+
The first template parameter, `Name`, can be any type, and is typically expected to be a user defined class in a user namespace. The type only needs to be forward declared, and not defined.
135135

136136
Above this basic mechanism of {cpp} type being used to identify a pipe, additional layers can be built on top to contain the type in an object that can be passed by value. Because such mechanisms (e.g. `boost::hana::type`) can layer on top of the fundamental type-based mechanism described here, those mechanisms are not included in the extension specification.
137137

@@ -166,37 +166,37 @@ myQueue.submit([&](handler& cgh) {
166166

167167
== Read/write member functions, and pipe template parameters
168168

169-
The pipe class exposes static member functions for writing a data word to a pipe, and for reading a data word from a pipe. A data word in this context is the data type that the pipe contains (`dataT` pipe template argument).
169+
The pipe class exposes static member functions for writing a data word to a pipe, and for reading a data word from a pipe. A data word in this context is the data type that the pipe contains (`DataT` pipe template argument).
170170

171171
Blocking and non-blocking forms of the read and write members are defined, with the form chosen based on overload resolution.
172172

173173
[source,c++,Read write members,linenums]
174174
----
175-
template <typename name,
176-
typename dataT,
177-
size_t min_capacity = 0>
175+
template <typename Name,
176+
typename DataT,
177+
size_t MinCapacity = 0>
178178
class pipe {
179179
// Blocking
180-
static dataT read();
181-
static void write( const dataT &data );
180+
static DataT read();
181+
static void write( const DataT &Data );
182182
183183
// Non-blocking
184-
static dataT read( bool &success_code );
185-
static void write( const dataT &data, bool &success_code );
184+
static DataT read( bool &Success );
185+
static void write( const DataT &Data, bool &Success );
186186
187187
// Static members
188-
using value_type = dataT;
189-
size_t min_capacity = min_capacity;
188+
using value_type = DataT;
189+
size_t min_capacity = MinCapacity;
190190
}
191191
----
192192

193193
The read and write member functions may be invoked within device code, or within a SYCL host program. Some interfaces may not be available on all devices/implementations, but the pipe definition itself does not gate availability. Instead, implementations should error if an unsupported pipe is used. See section <<device_queries>> for information on querying the availability of specific pipe features relative to a device.
194194

195195
The template parameters of the device type are defined as:
196196

197-
* `name`: Type that is the basis of pipe identification. Typically a user-defined class, in a user namespace. Forward declaration of the type is sufficient, and the type does not need to be defined.
198-
* `dataT`: The type of data word/packet contained within a pipe. This is the data type that is read during a successful `pipe::read` operation, or written during a successful `pipe::write` operation. The type must be standard layout and trivially copyable. This template parameter can be queried by using the `value_type` type alias.
199-
* `min_capacity`: User defined minimum number of words in units of `dataT` that the pipe must be able to store without any being read out. A minimum capacity is required in some algorithms to avoid deadlock, or for performance tuning. An implementation can include more capacity than this parameter, but not less. This template parameter can be queried by using the `min_capacity` static member.
197+
* `Name`: Type that is the basis of pipe identification. Typically a user-defined class, in a user namespace. Forward declaration of the type is sufficient, and the type does not need to be defined.
198+
* `DataT`: The type of data word/packet contained within a pipe. This is the data type that is read during a successful `pipe::read` operation, or written during a successful `pipe::write` operation. The type must be standard layout and trivially copyable. This template parameter can be queried by using the `value_type` type alias.
199+
* `MinCapacity`: User defined minimum number of words in units of `DataT` that the pipe must be able to store without any being read out. A minimum capacity is required in some algorithms to avoid deadlock, or for performance tuning. An implementation can include more capacity than this parameter, but not less. This template parameter can be queried by using the `min_capacity` static member.
200200

201201
== Pipe types and {cpp} scope
202202

@@ -260,12 +260,12 @@ Pipes expose two additional static member functions that are available within ho
260260

261261
[source,c++,Read write members,linenums]
262262
----
263-
template <typename name,
264-
typename dataT,
265-
size_t min_capacity = 0>
263+
template <typename Name,
264+
typename DataT,
265+
size_t MinCapacity = 0>
266266
class pipe {
267267
template <pipe_property::writeable host_writeable>
268-
static dataT* map(size_t requested_size, size_t &mapped_size);
268+
static DataT* map(size_t requested_size, size_t &mapped_size);
269269
270270
static size_t unmap(T* mapped_ptr, size_t size_to_unmap);
271271
}
@@ -284,11 +284,11 @@ The APIs are defined as:
284284
|Function |Description
285285
|`template <pipe_property::writeable host_writeable> +
286286
dataT* map(size_t requested_size, size_t &mapped_size);`
287-
|Returns a _dataT *_ in the host address space. The host can write data to this pointer for reading by a device pipe endpoint, if it was created with template parameter `host_writeable = true`. Alternatively, the host can read data from this pointer if it was created with template parameter `host_writeable = false`.
287+
|Returns a DataT *_ in the host address space. The host can write data to this pointer for reading by a device pipe endpoint, if it was created with template parameter `host_writeable = true`. Alternatively, the host can read data from this pointer if it was created with template parameter `host_writeable = false`.
288288

289289
The value returned in the mapped_size argument specifies the maximum number of bytes that the host can access. The value specified by _mapped_size_ is less than or equal to the value of the _requested_size_ argument that the caller specifies. _mapped_size_ does not impact the _min_capacity_ property of the pipe.
290290

291-
After writing to or reading from the returned _dataT *_, the host must execute one or more `unmap` calls on the same pipe, to signal to the runtime that data is ready for transfer to the device (on a write), and that the runtime can reclaim the memory for reuse (on a read or write). If `map` is called on a pipe before `unmap` has been used to unmap all memory mapped by a previous `map` operation, the buffer returned by the second `map` call will not overlap with that returned by the first call.
291+
After writing to or reading from the returned DataT *_, the host must execute one or more `unmap` calls on the same pipe, to signal to the runtime that data is ready for transfer to the device (on a write), and that the runtime can reclaim the memory for reuse (on a read or write). If `map` is called on a pipe before `unmap` has been used to unmap all memory mapped by a previous `map` operation, the buffer returned by the second `map` call will not overlap with that returned by the first call.
292292

293293
|`static size_t unmap(T* mapped_ptr, size_t size_to_unmap);`
294294
|Signals to the runtime that the host is no longer using _size_to_unmap_ bytes of the host allocation that was returned previously from a call to `map`. In the case of a writeable host pipe, calling `unmap` allows the unmapped data to become available to the kernel. If the _size_to_unmap_ value is smaller than the _mapped_size_ value specified to `map`, then multiple `unmap` function calls are necessary to unmap the full capacity of the host allocation. It is legal to perform multiple `unmap` function calls to unmap successive bytes in the buffer returned by `map`, up to the _mapped_size_ value defined in the `map` operation.
@@ -382,22 +382,22 @@ The pipe class described above exposes both read and write static member functio
382382

383383
[source,c++,iopipes,linenums]
384384
----
385-
template <typename name,
386-
typename dataT,
387-
size_t min_capacity = 0>
385+
template <typename Name,
386+
typename DataT,
387+
size_t MinCapacity = 0>
388388
class kernel_readable_io_pipe {
389389
public:
390-
static dataT read(); // Blocking
391-
static dataT read( bool &success_code ); // Non-blocking
390+
static DataT read(); // Blocking
391+
static DataT read( bool &Success ); // Non-blocking
392392
};
393393
394-
template <typename name,
395-
typename dataT,
396-
size_t min_capacity = 0>
394+
template <typename Name,
395+
typename DataT,
396+
size_t MinCapacity = 0>
397397
class kernel_writeable_io_pipe {
398398
public:
399-
static void write( dataT data ); // Blocking
400-
static void write( dataT data, bool &success_code ); // Non-blocking
399+
static void write( DataT Data ); // Blocking
400+
static void write( DataT Data, bool &Success ); // Non-blocking
401401
}
402402
----
403403

@@ -642,55 +642,76 @@ Automated mechanisms are possible to provide uniquification across calls, and co
642642

643643
*NOTE*: The APIs described in this section are experimental. Future versions of this extension may change these APIs in ways that are incompatible with the versions described here.
644644

645-
In the experimental API version, read/write methods take template arguments, which can contain the latency control properties `latency_anchor_id` and/or `latency_constraint`.
645+
The Intel FPGA experimental `pipe` class is implemented in `sycl/ext/intel/experimental/pipes.hpp` which is included in `sycl/ext/intel/fpga_extensions.hpp`.
646+
647+
In the experimental API version, read/write methods take in a property list as function argument, which can contain the latency control properties `latency_anchor_id` and/or `latency_constraint`.
646648

647649
* `sycl::ext::intel::experimental::latency_anchor_id<N>`, where `N` is an integer: An ID to associate with the current read/write function call, which can then be referenced by other `latency_constraint` properties elsewhere in the program to define relative latency constaints. ID must be unique within the application, and a diagnostic is required if that condition is not met.
648650
* `sycl::ext::intel::experimental::latency_constraint<A, B, C>`: A tuple of three values which cause the current read/write function call to act as an endpoint of a latency constraint relative to a specified `latency_anchor_id` defined by a different instruction.
649651
** `A` is an integer: The ID of the target anchor defined on a different instruction through a `latency_anchor_id` property.
650-
** `B` is an enum value: The type of control from the set {`type::exact`, `type::max`, `type::min`}.
652+
** `B` is an enum value: The type of control from the set {`latency_control_type::exact`, `latency_control_type::max`, `latency_control_type::min`}.
651653
** `C` is an integer: The relative clock cycle difference between the target anchor and the current function call, that the constraint should infer subject to the type of the control (exact, max, min).
652654

653-
The template arguments above don't have to be specified if user doesn't want to apply latency controls. The template arguments can be passed in arbitrary order.
654-
655-
=== Implementation
655+
=== Synopsis
656656

657657
[source,c++]
658658
----
659659
// Added in version 2 of this extension.
660660
namespace sycl::ext::intel::experimental {
661-
enum class type {
661+
enum class latency_control_type {
662662
none, // default
663663
exact,
664664
max,
665665
min
666666
};
667667
668-
template <int32_t _N> struct latency_anchor_id {
669-
static constexpr int32_t value = _N;
670-
static constexpr int32_t default_value = -1;
668+
struct latency_anchor_id_key {
669+
template <int Anchor>
670+
using value_t =
671+
oneapi::experimental::property_value<latency_anchor_id_key,
672+
std::integral_constant<int, Anchor>>;
671673
};
672674
673-
template <int32_t _N1, type _N2, int32_t _N3> struct latency_constraint {
674-
static constexpr std::tuple<int32_t, type, int32_t> value = {_N1, _N2, _N3};
675-
static constexpr std::tuple<int32_t, type, int32_t> default_value = {
676-
0, type::none, 0};
675+
struct latency_constraint_key {
676+
template <int Target, latency_control_type Type, int Cycle>
677+
using value_t = oneapi::experimental::property_value<
678+
latency_constraint_key, std::integral_constant<int, Target>,
679+
std::integral_constant<latency_control_type, Type>,
680+
std::integral_constant<int, Cycle>>;
677681
};
678682
679-
template <typename name,
680-
typename dataT,
681-
size_t min_capacity = 0>
683+
template <int Anchor>
684+
inline constexpr latency_anchor_id_key::value_t<Anchor> latency_anchor_id;
685+
686+
template <int Target, latency_control_type Type, int Cycle>
687+
inline constexpr latency_constraint_key::value_t<Target, Type, Cycle>
688+
latency_constraint;
689+
690+
template <typename Name,
691+
typename DataT,
692+
size_t MinCapacity = 0>
682693
class pipe {
683694
// Blocking
684-
template <class... _Params>
685-
static dataT read();
686-
template <class... _Params>
687-
static void write( const dataT &data );
695+
static DataT read();
696+
697+
template <typename PropertiesT>
698+
static DataT read( PropertiesT Properties );
699+
700+
static void write( const DataT &Data);
701+
702+
template <typename PropertiesT>
703+
static void write( const DataT &Data, PropertiesT Properties );
688704
689705
// Non-blocking
690-
template <class... _Params>
691-
static dataT read( bool &success_code );
692-
template <class... _Params>
693-
static void write( const dataT &data, bool &success_code );
706+
static DataT read( bool &Success );
707+
708+
template <typename PropertiesT>
709+
static DataT read( bool &Success, PropertiesT Properties );
710+
711+
static void write( const DataT &Data, bool &Success );
712+
713+
template <typename PropertiesT>
714+
static void write( const DataT &Data, bool &Success, PropertiesT Properties );
694715
}
695716
} // namespace sycl::ext::intel::experimental
696717
----
@@ -709,17 +730,20 @@ using Pipe3 = ext::intel::experimental::pipe<class PipeClass2, int, 8>;
709730
myQueue.submit([&](handler &cgh) {
710731
cgh.single_task<class foo>([=] {
711732
// The following Pipe1::read is anchor 0
712-
int value = Pipe1::read<ext::intel::experimental::latency_anchor_id<0>>();
733+
int value = Pipe1::read(
734+
ext::oneapi::experimental::properties(latency_anchor_id<0>));
713735
714736
// The following Pipe2::write is anchor 1
715737
// The following Pipe2::write occurs exactly 2 cycles after anchor 0
716-
Pipe2::write<ext::intel::experimental::latency_anchor_id<1>,
717-
ext::intel::experimental::latency_constraint<
718-
0, ext::intel::experimental::type::exact, 2>>(value);
738+
Pipe2::write(value,
739+
ext::oneapi::experimental::properties(
740+
latency_anchor_id<1>,
741+
latency_constraint<0, latency_control_type::exact, 2>));
719742
720743
// The following Pipe3::write occurs at least 2 cycles after anchor 1
721-
Pipe3::write<ext::intel::experimental::latency_constraint<
722-
1, ext::intel::experimental::type::min, 2>>(value);
744+
Pipe3::write(value,
745+
ext::oneapi::experimental::properties(
746+
latency_constraint<1, latency_control_type::min, 2>));
723747
});
724748
});
725749
----

0 commit comments

Comments
 (0)