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 in the following ways:
- Using the streaming_interface kernel property
- (Deprecated)Using the streaming_interface macro
The macro is deprecated and will be removed in a future release. Convert any uses of this macro to use the streaming_interface kernel property instead.
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.
(Deprecated) The streaming_interface Macro
Deprecated: The streaming_interface macro is deprecated and might be removed in a future release. Use The streaming_interface Kernel Property instead.
To have the compiler implement the IP component invocation interface with a "ready-valid" handshake:
- Implement the IP kernel as a functor.
- Include the following header file:
sycl/ext/intel/prototype/interfaces.hpp
- Add one of the following options to the compiler command (icpx -fsycl):
- Linux:-I/$INTELFPGAOCLSDKROOT/include
- Windows:/I %INTELFPGAOCLSDKROOT%\include
- Add the streaming_interface macro to the functor operator()().
The following code shows an example of implementing a streaming interface:
#include <sycl/sycl.hpp>
#include <sycl/ext/intel/prototype/interfaces.hpp>
using namespace sycl;
struct MyIP {
int *input_a, *input_b, *input_c;
int n;
streaming_interface void operator()() const {
for (int i = 0; i < n; i++) {
input_c[i] = input_a[i] + input_b[i];
}
}
};
The resulting IP core kernel is invoked with a ready/valid handshake. Compiling the example code generates the start signal, the done signal, the ready_in signal, and ready_out signals as conduits. The compilation of the example code also generates conduits for the base addresses of the three pointers as well the value of N.
The streaming handshaking is similar to the Avalon® Streaming (ST) protocol. The IP kernel consumes the arguments on the clock cycle that the start and ready_out signals are asserted. The IP component kernel invocation is finished on the clock cycle that the done and ready_in signals are asserted.