Visible to Intel only — GUID: GUID-2E32ED61-CE7E-4715-AB89-5554DB1BD93A
Visible to Intel only — GUID: GUID-2E32ED61-CE7E-4715-AB89-5554DB1BD93A
The annotated_ptr Template Class (Beta)
Use the annotated_ptr template class to annotate variable pointers in your kernel with a specific buffer location. This annotation directs the compiler to constrain memory accesses and build more efficient FPGA hardware.
The annotated_ptr template class is provided in the sycl.hpp header file. To use the annotated_ptr template class, use the following declarations in your code:
#include <sycl/sycl.hpp> using namespace sycl::ext::intel::experimental; using namespace sycl::ext::oneapi::experimental;
The annotated_ptr template class has the following syntax:
annotated_ptr<data type,decltype(properties { buffer_location<id> } ) >
Where the data type indicates the data type of the underlying pointer and id in buffer_location<id> corresponds to a memory location defined as follows, depending on your FPGA development flow:
-
- SYCL* HLS Flow
- The memory location must be a buffer location that you defined in the kernel interface with the annotated_arg template class. For details about defining buffer location in a kernel interface, refer to sycl::ext::intel::experimental::buffer_location in The annotated_arg Template Class.
-
- FPGA Acceleration Flow
- The memory location must be one of the global memories defined in board_spec.xml file of your FPGA platform. To identify the index of the buffer_location<index> property, refer to Partitioning Buffers Across Different Memory Types (Heterogeneous Memory).
Construction of an annotated_ptr Instance
Before using an instance of annotated_ptr, it must be explicitly constructed in one of the following ways:
datatype *p = …; annotated_ptr<datatype,decltype(properties { buffer_location<id> } )> ap {p};
datatype *p = …; annotated_ptr ap { p, properties { buffer_location<id> } };
Example of Using the annotated_ptr Template Class
The following example demonstrates how to constrain external memory accesses via pointers inside a kernel. The kernel in the example takes two annotated_arg arguments:
- pX with buffer location 1
- Y with buffer location 2
- Data that the pX pointer points to is accessed only through memory-mapped host interface 1 (buffer_location<1>).
- Data that the Y pointer points to is accessed only through memory-mapped host interface 2 (buffer_location<2>).
When dereferencing the pX pointer to get another pointer temp, the buffer location for the temp pointer is ambiguous. Use the annotated_ptr class to specify the buffer location for the temp pointer, so that the compiler needs to build an interconnect only between the load unit and memory-mapped host interface 1 when reading the data that the temp pointer points to.
struct MyIP { annotated_arg<int **, decltype(properties{buffer_location<1>})> pX; annotated_arg<int *, decltype(properties{buffer_location<2>})> Y; void operator()() const { for (int i = 0; i < N; i++) { int *temp = pX[i]; // The buffer location of temp is unknown // Use annotated_ptr to constrain accesses via temp to buffer location 1 annotated_ptr<int, decltype(properties{buffer_location<1>})> X{temp}; Y[i] = X[i]; } } };