Visible to Intel only — GUID: GUID-5C02C6E1-D81F-471E-ABED-F75E9B012105
Visible to Intel only — GUID: GUID-5C02C6E1-D81F-471E-ABED-F75E9B012105
Memory-Mapped Host Interfaces Using Unified Shared Memory Pointers and the annotated_arg Class
A kernel can interface with an external memory over an Avalon® Memory-mapped (MM) Host interface. You can specify the Avalon® MM host interface with the annotated_arg class.
Describe a customized Avalon® MM host interface in your code by adding an annotated_arg object in your kernel functor definition or by capturing an annotated_arg object with a lambda kernel function.
Each annotated_arg argument of a kernel can be configured with a conduit or register map input type.
Each pointer type annotated_arg argument can be customized to different port configuration attaching to a specific buffer location. An Avalon® MM Host interface is created for each unique buffer location. Host interfaces that share the same buffer location are arbitrated on the same interface.
Learn how to configure memory-mapped host interfaces in your RTL IP kernel by reviewing the Memory-Mapped Host Interfaces sample from the oneAPI Samples for FPGA GitHub repository.
The annotated_arg class has the following arguments. For full descriptions and details, refer to The annotated_arg Template Class.
Template Object or Parameter | Description | Valid Values |
---|---|---|
Interface Type Properties | ||
sycl::ext::intel::experimental::conduit | Create a dedicated input conduit on the kernel. | N/A |
sycl::ext::intel::experimental::register_map | Creates an input register map for the input instead of a dedicated input port. | N/A |
sycl::ext::intel::experimental::stable | Specifies that the input to the kernel does not change between pipelined invocations of the kernel. The input can change after all active kernel invocations have finished. |
N/A |
Pointer Kernel Argument Properties | ||
sycl::ext::intel::experimental::buffer_location | The global memory identifier for the pointer interface. To use any of the other Avalon® MM Host interface properties, you must also specify the buffer_location property. |
N/A |
sycl::ext::intel::experimental::dwidth | The width of the memory-mapped data bus in bits | 8, 16, 32, 64, 128, 256, 512, or 1024 Default: 64 |
sycl::ext::intel::experimental::awidth | The width of the memory-mapped address bus in bits. | Integer 11 – 41 Default: 41 |
sycl::ext::intel::experimental::latency | The guaranteed latency from when a read command exits the kernel until the external memory returns valid read data. | Non-negative integer value Default: 1 |
sycl::ext::intel::experimental::maxburst | The maximum number of data transfers that can associate with a read or write transaction. | Integer 1 – 1024 |
sycl::ext::intel::experimental::alignment | The alignment of the base pointer address in bytes. | See description |
sycl::ext::intel::experimental::read_write_mode | The port direction of the interface. | read_write, read, or write Default: read_write |
You can use the annotated_arg class to annotate kernel arguments with either lambda kernels or functor kernels.
#include <sycl/sycl.hpp> #include <sycl/ext/intel/fpga_extension.hpp> using namespace sycl; using namespace ext::intel::experimental; using namespace ext::oneapi::experimental; struct kernel { annotated_arg<int*, decltype(properties{ conduit, buffer_location<1>, dwidth<32>, awidth<16>, latency<0>, read_write_mode_readwrite })> mm_a; void operator()() const { *mm_a = 1; } };
#include <sycl/sycl.hpp> #include <sycl/ext/intel/fpga_extension.hpp> using namespace sycl; using namespace ext::intel::experimental; using namespace ext::oneapi::experimental; class kernel; void launch_kernel(queue& q, int* a) { auto mm_a = annotated_arg(a, conduit, buffer_location<1>, dwidth<32>, awidth<16>, latency<0>, read_write_mode_readwrite); q.single_task<class kernel>([=] { *mm_a = 1; }); }
- Buffer Locations in Memory-Mapped Host Interfaces
- The annotated_arg Template Class
Use the annotated_arg template class to annotate your kernel arguments with properties that direct the compiler to customize the kernel argument interface.