Developer Guide

Intel® oneAPI DPC++/C++ Compiler Handbook for FPGAs

ID 785441
Date 10/24/2024
Public
Document Table of Contents

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:

IMPORTANT:
If the buffer location ID specified with the annotated_ptr does not match the actual buffer location of the pointer at run time the behavior is undefined. Ensure that the pointers that are being annotated inside the kernel are allocated with the matching buffer location via USM memory allocation APIs such as malloc_shared or aligned_alloc_shared.

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
These annotated_arg arguments direct the compiler to build an interconnect that allows accessing the data in the following ways:
  • 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>).
For more details on the annotated_arg class, refer to Memory-Mapped Host Interfaces Using Unified Shared Memory Pointers and the annotated_arg Class.

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];
       }
    }
};