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

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.

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

You can use the annotated_arg class to annotate kernel arguments with either lambda kernels or functor kernels.

Functor Example
#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;
     }
};
Lambda Example
#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;
      });
}