Visible to Intel only — GUID: GUID-E7605FCA-7734-4FF9-BCC3-5EAE9B496DFE
Visible to Intel only — GUID: GUID-E7605FCA-7734-4FF9-BCC3-5EAE9B496DFE
Ready/Valid Handshaking Kernel Invocation Interface
While the default kernel invocation interface is a memory-mapped (MM) agent interface, you can also have the Intel® oneAPI DPC++/C++ Compiler implement the kernel invocation interface with a ready/valid handshake, similar to what an Avalon® streaming interface uses. A kernels with a ready/valid handshake kernel invocation interface is often referred to as a streaming kernel.
Streaming kernels can also be pipelined. For details, refer to Pipelined Kernels.
You can indicate that a kernel uses a "ready-valid" handshake interface as its invocation interface by using the streaming_interface kernel property.
The streaming_interface Kernel Property
Use the streaming_interface kernel property to request the compiler to implement the kernel invocation interface with a "ready-valid" handshake.
The streaming_interface kernel property supports the following template parameters:
- accept_downstream_stall
If the streaming_interface property with accept_downstream_stall (streaming_interface<accept_downstream_stall>) is specified on a kernel, the compiler generates a "ready-valid" kernel interface at both input and output such that the "ready-valid" handshaking happens both at kernel invocation and kernel completion.
A streaming_interface_accept_downstream_stall property is also provided for convenience.
- remove_downstream_stall
If the streaming_interface property with remove_downstream_stall (streaming_interface<remove_downstream_stall>) is specified on a kernel, the compiler generates a "ready-valid" kernel interface only at the input to the kernel. The "ready-valid" handshaking happens only at kernel invocation. There is no interface to gate the done signals emitted by the kernel. That is, there is no "back-pressure" on the kernel.
A streaming_interface_remove_downstream_stall property is also provided for convenience.
If no template parameter is provided, the accept_downstream_stall value is inferred.
To use the streaming_interface kernel property:
- Include the following header file in your code:
sycl/ext/intel/fpga_extensions.hpp
- Label your kernel with the streaming_interface property as follows:
- Functor Model:
- Add a member function named get to the functor. Have the get function take an argument of type properties_tag and a return type auto.
- Create a properties object in the new function with the streaming_interface property and return it.
Functor Model streaming_interface Kernel Property Code Example#include <sycl/sycl.hpp> #include <sycl/ext/intel/fpga_extensions.hpp> using namespace sycl; using namespace sycl::ext::intel::experimental; using namespace sycl::ext::oneapi::experimental; struct MyFunctorIP { int *input_a, *input_b, *input_c; int n; void operator()() const { for (int i = 0; i < n; i++) { input_c[i] = input_a[i] + input_b[i]; } } auto get(properties_tag) { return properties{streaming_interface<>}; } }; ... q.single_task(MyFunctorIP{functor_input_A, functor_input_B, functor_input_C, kN}).wait();
- Lambda Model:
- Pass a properties object that contains the sreaming_interface property to your q.single_task call.
Lambda Model streaming_interface Kernel Property Code Example#include <sycl/sycl.hpp> #include <sycl/ext/intel/fpga_kernel_properties.hpp> using namespace sycl; using namespace sycl::ext::intel::experimental; using namespace sycl::ext::oneapi::experimental; class MyLambdaIP; // Create a properties object containing the kernel invocation // interface property properties kernel_properties{streaming_interface<>}; ... q.single_task<MyLambdaIP>(kernel_properties, [=] { for (int i = 0; i < n; i++) { lambda_input_C[i] = lambda_input_A[i] + lambda_input_B[i]; } }).wait();
- Functor Model:
Limitations of Streaming Kernels
The following actions are not supported when using a streaming kernel:
Using streaming kernels as SYCL NDRange kernels.
Profiling of streaming kernels.