Developer Guide

Intel oneAPI FPGA Handbook

ID 785441
Date 2/07/2024
Public

A newer version of this document is available. Customers should click here to go to the newest version.

Document Table of Contents

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.

The annotated_arg template class is provided in the fpga_extension.hpp header file and uses the sycl::ext::intel::experimental namespace. To use the annotated_arg template class, use the following declarations in your code:

#include <sycl/ext/intel/fpga_extension.hpp>
using namespace sycl::ext::intel::experimental;
using namespace sycl::ext::oneapi::experimental;

The annotated_arg template class has the following syntax:

annotated_arg<data type,decltype(properties { (conduit | register_map),
                                                [stable,]
                                                [buffer_location<id>,]
                                                [dwidth<data_bus_width,]
                                                [awidth<address_bus_wdith>,]
                                                [latency<value>,]
                                                [maxburst<value>,]
                                                [alignment<value>,]
                                                [read_write_mode<mode>] } ) >

When passing pointers to a SYCL* kernel with the annotated_arg template class, the pointer needs memory allocated using USM memory allocation APIs such as malloc_shared or alligned_alloc_shared.

RESTRICTION:
  • For aggregate data types (like struct), the annotated_arg template class does not support the dot (.) and the arrow (->) member-access operators. You must use an explicit or implicit conversion to the underlying date type before accessing the members of the aggregate data type.
  • The annotated_arg template class does not support using variable-precision data types (like ac_int) operations (like addition, multiplication...) directly. If you want to use the annotated_arg template class with a variable-precision data type, you must cast the annotated_arg object to the underlying data type object and perform operations on that object directly. For more details, refer to Recommended Method for Annotating struct Data Types.

Recommended Method for Annotating struct Data Types

When you annotate struct data types, create a local variable that stores the kernel argument of struct type so that member functions and data members of the struct types can be accessed directly.

The following example code demonstrates how to annotate a kernel argument of ac_int data type:

using ac_int_type = ac_int<17, true>;

struct MyIP {
    annotated_arg<ac_int_type, decltype(properties{conduit})> a;
    annotated_arg<ac_int_type, decltype(properties{conduit})> b;
    annotated_arg<int*, decltype(properties{conduit})> c;
    void operator()() const {
        // Make local variables of the struct type before using them
        ac_int_type n1 = a;
        ac_int_type n2 = b;
        *c = n1 + n2;
    }
};

The following example code demonstrates how to annotate a kernel argument of struct data type:

struct DataT{
    float f1;
    float f2;
};

struct MyIP {
    annotated_arg<DataT, decltype(properties{conduit})> inp;
    annotated_arg<float*, decltype(properties{conduit})> sum;
    void operator()() const {
        // Make local variables of the struct type from the kernel argument
        // before using the kernel argument
        DataT data = inp;
        // This allows access to struct members with the 'dot' operator
        *c = data.f1 + data.f2;
    }
};

The annotated_arg Template Class Properties Summary
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

sycl::ext::intel::experimental::conduit Declaration

Syntax
sycl::ext::intel::experimental:conduit
Description
Directs the compiler to create a dedicated input conduit on the kernel for the kernel arguments. That input conduit is associated with the kernel start and busy signals.

This parameter is mutually exclusive with the sycl::ext::intel::experimental::register_map declaration.

sycl::ext::intel::experimental::register_map Declaration

Syntax
sycl::ext::intel::experimental::register_map
Description
Directs the compiler to create a register map to store the input.

This declaration is mutually exclusive with the sycl::ext::intel::experimental:conduit declaration.

sycl::ext::intel::experimental::stable Declaration

Syntax
sycl::ext::intel::experimental::stable
Description

While the SYCL* software model makes kernel arguments read-only, the RTL IP component generated by the compiler can be plugged into external systems where kernel arguments can change while the kernel executes.

This property specifies that the input to the kernel does not change between pipelined invocations of the kernel. The input can still change after all active kernel invocations have finished.

If the input changes while the pipelined kernel invocations are executing, the behavior is undefined.

sycl::ext::intel::experimental::buffer_location Property

Syntax

template <unsigned id>
sycl::ext::intel::experimental::buffer_location<id>

Description
Specifies a global memory identifier for the pointer interface. You must specify this argument to use any of the other Avalon® MM Host interface properties. Pointers with this argument specified are often referred to as annotated pointers.

Each unique ID results in a separate Avalon® MM Host interface on your RTP IP core. All hosts with the same address space are arbitrated within the kernel to a single interface. As such, these hosts must share the same template parameters that describe the interface.

If this property is not specified, the pointer argument is considered to be an unannotated pointer.

If you mix annotated pointers (that is, pointers with specified buffer locations) and unannotated pointers, ensure that you review Mixing Annotated and Unannotated Pointers in Your Kernel to understand the implications of mixing pointer types.

If all pointers in your kernel are unannotated, the compiler infers an Avalon® MM Host interface with a buffer location ID of 0.

IMPORTANT:
The caller is responsible for ensuring the correct buffer location is specified; otherwise, functional failures might occur. If you specify a buffer_location property on your kernel argument, specify the same buffer location on the USM allocation API call that allocates memory on the test bench (that is, the host code).

RESTRICTION:
If you use annotated pointers, you cannot mix sycl::accessor and annotated_arg kernel arguments.

sycl::ext::intel::experimental::dwidth Property

Syntax

template <unsigned width>
sycl::ext::intel::experimental::dwidth<width>

Valid Values
8, 16, 32, 64, 128, 256, 512, or 1024
Default Value
64
Description
The width of the memory-mapped data bus in bits.

To specify the dwidth property, you must also specify the buffer_location property.

sycl::ext::intel::experimental::awidth Property

Syntax

template <unsigned width>
sycl::ext::intel::experimental::awidth<width>

Valid Values
Integer value in the range 11-41
Default Value
41
Description
The width of the address bus. This value affects only the width of the Avalon® MM host interface.

RESTRICTION:
The lowest-numbered buffer location must have a minimum address width of 11-bits (awidth<11>). There is no minimum width for higher-numbered buffer locations.

To specify the awidth property, you must also specify the buffer_location property.

sycl::ext::intel::experimental::latency Property

Syntax

template <unsigned value>
sycl::ext::intel::experimental::latency<value>

Valid Values
Non-negative integer value
Default Value
1
Description
The guaranteed latency from when a read command exits the kernel to when the external memory returns valid read data.

If this latency is variable (such as when accessing DRAM), set it to 0.

To specify the latency property, you must also specify the buffer_location property.

sycl::ext::intel::experimental::maxburst Property

Syntax

template <unsigned value>
sycl::ext::intel::experimental::maxburst<value>

Valid Values
Integer value in the range 1 – 1024
Default Value
1
Description
The maximum number of data transfers that can associate with a read or write transaction. This value controls the width of the burstcount signal.

For fixed latency interfaces, this value must be set to 1.

For more details, review information about burst signals and the burstcount signal role in "Avalon Memory-Mapped Interface Signal Roles" in Avalon Interface Specifications.

To specify the maxburst property, you must also specify the buffer_location property and set the latency property to 0 (latency<0>).

sycl::ext::intel::experimental::alignment Property

Syntax

template <unsigned value>
sycl::ext::intel::experimental::alignment<value>

Valid Values
Integer value greater than the alignment of the data type
Default Value
1
Description
The alignment of the base pointer address in bytes.

The compiler uses this information to determine how many simultaneous loads and stores this pointer can permit.

For example, if you have a bus with 4 32-bit integers on it, you should use sycl::ext::intel::experimental::dwidth<128> (bits) and sycl::ext::intel::experimental::align<16> (bytes). This means that up to 16 contiguous bytes (or 4 32-bit integers) can be loaded or stored as a coalesced memory word per clock cycle.

IMPORTANT:
The caller is responsible for aligning the data to the set value for the align argument; otherwise, functional failures might occur. If you specify this property, ensure that you specify the same alignment value the following locations:
  • The sycl::ext::intel::experimental::alignment<value> property
  • The alignment argument of the aligned_malloc_shared template function in your host code. For example,
    aligned_malloc_shared<int>(value, 1, q );

sycl::ext::intel::experimental::read_write_mode Property

Syntax

template <unsigned mode>
sycl::ext::intel::experimental::read_write_mode<value>

Valid Values
read_write, read, or write
Default Value
readwrite
Description
The port direction of the memory interface associated with the input pointer. Only the relevant Avalon host signals are generated.

The available modes are as follows:

  • read_write: The interface can used for read-write operations. You can also specify the sycl::ext::intel::experimental::read_write_mode_readwrite property to enable this mode.
  • read: The interface is read-only. It be used only for read operations. You can also specify the sycl::ext::intel::experimental::read_write_mode_read property to enable this mode.
  • write: The interface is write-only. It can be used only for write operations. You can also specify the sycl::ext::intel::experimental::read_write_mode_write property to enable this mode.

To specify the read_write_mode property, you must also specify the buffer_location property.