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

Buffer Locations in Memory-Mapped Host Interfaces

Each unique identifier specified by the buffer_location<id> property of a pointer kernel argument represents a unique Avalon® host interface for the kerne. Ensure that you avoid using the same buffer location id for interfaces with different properties.

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).

Similarly, the caller is responsible for aligning the data to the set value for the align argument; otherwise, functional failures might occur. If you specify the alignment<value> 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 );

Memory-Mapped Interface Unified-Shared-Memory Virtual Address Space

The compiler encodes certain information regarding the virtual address space in the top bits of a 64-bit pointer address as follows:

Pointer-Address Bit-Range Descriptions

Bit Range

Description

40:0

Used for addressing within the memory system

63:41

Stores the virtual address space information that is derived from the buffer location

In some cases, the compiler cannot determine which buffer location a pointer corresponds to and it creates logic in the generated RTL that inspects the top bits of the pointer at runtime to detect the buffer location and route the memory transaction to the correct external memory interface.

You do not need to encode the buffer location information yourself in most cases. Exceptions are outlined in Buffer Location Virtual Address Spaces.

The compiler automatically generates logic to embed this information from the buffer location specified on the pointer kernel argument in the source file.

Mixing Annotated and Unannotated Pointers in Your Kernel

The compiler infers Avalon host interfaces as follows:

  • If annotated kernel arguments are present, one interface is inferred for each unique buffer location.
  • If no annotated pointer kernel arguments are present, a single Avalon Host interface is inferred for all pointer arguments.

If your design has a mix of annotated kernel arguments and unannotated pointer arguments, the unannotated pointers can access any of the inferred Avalon host interfaces depending on the addresses passed to the kernel via the pointer arguments.

When your designs mix annotated and unannotated pointers, there are situations where the compiler is unable to determine if an error condition exists. In these situations, functional errors or hangs can appear when you try to simulate your design or use the generated RTL.

If you have a design that mixes annotated and unannotated pointers and your design fails in simulation, check for the following conditions:

Buffer Locations and Unannotated Pointers

Ensure that your unannotated pointer accesses are to interfaces that support the access. For example, an unannotated pointer attempting a read access to a write-only interface is an error that the compiler cannot always catch.

Consider the following design with two kernels:

// Kernel1: Write-only MM Host Interface
// Kernel1 defines a write-only memory-mapped host interface
struct Kernel1 {
  annotated_arg<int*, decltype(properties{buffer_location<1>, read_write_mode_write})> arg_a;
  annotated_arg<int*, decltype(properties{buffer_location<2>, read_write_mode_read})> arg_b;

  ...
  void operator()() const {
    ...
  }
};

//Kernel2: Unannotated Pointer
// Kernel2 uses an unannotated pointer and does a read access
struct Kernel2 {
  int* arg_a;
  ...
  void operator()() const {
    ... = *arg_a;
  }
};

The kernel Kernel1 specifies two memory-mapped host interfaces. The compiler infers the two interfaces, and the unannotated pointer argument in kernel Kernel2 can access either of them, depending on the memory addresses that are passed to it.

Because the buffer_location<1> memory has a write-only interface, passing an address for the buffer_location<1> memory to the unannotated pointer in kernel Kernel2 is an illegal operation.

If your program has such mixed usage, ensure that correct addresses are passed via kernel arguments.

Buffer Location Virtual Address Spaces

If the kernel has at least one kernel argument where the buffer location is specified (annotated argument) and at least one argument where the buffer location is not specified (unannotated argument), then you must embed the buffer location information in the top bits of any unannotated kernel arguments.

Consider the following code example:

// This struct defines the IP that is generated
struct MyIPComponent{
  // struct members are kernel arguments 
  int* a; // no buffer location specified
  annotated_arg<int*, decltype(properties{buffer_location<1>})> b; // buffer location 1
  annotated_arg<int*, decltype(properties{buffer_location<2>})> c; // buffer location 2
  annotated_arg<int*, decltype(properties{buffer_location<3>})> d; // buffer location 3

// operator()() defines the device/IP code 
void operator()() const {
    *a *= 2;
    *b *= 2;
    *c *= 2;
    *d *= 2;
   }
};

In this example, the compiler does not know what external memory pointer a will point to, so the compiler creates logic to check the top bits of the pointer to determine, at run time, which buffer location to access. Therefore, you must set those top bits for argument a (but not for the other kernel arguments).

In such cases, if the unannotated pointer argument has a conduit interface then the port is 64 bits wide. And, if the interface is register-map based, all 64 bits are passed to the kernel.

Simulation Exception:
When you simulate your kernel you do not need to write any host code (even in this use case) to embed information in the pointer bits. The buffer locations are all taken care of by the runtime stack that allocates the pointers in the host code.

When the compiler can deduce which buffer location that a pointer argument points to (for example, when there is only one mm_host interface), the compiler embeds the buffer location automatically. Buffer locations need to be specified manually only when your kernels have a mix of annotated and unannotated kernel arguments.

The compiler infers one global memory (with inferred buffer location 0) when there are no annotated pointer kernel arguments in your entire design (that is, across all kernels). In the following code example, because there are only unannotated pointer arguments, the compiler infers only one global memory and so it can embed the correct information in the top bits of the pointer kernel arguments.

// This struct defines the IP that will be generated
struct MyIPComponent{
    // struct members are kernel arguments
    int* a; // no buffer location specified
    int* b; // no buffer location specified
    // no other annotated kernel argument is present

    // operator()() defines the device/IP code
    void operator()() const {
        *a = ...
        *b = ...
    }
};

Determining Virtual Address Space Information

If you need to embed the virtual address space information in the top bits of the pointer kernel arguments because the compiler cannot do so, get the information to embed from the HTML reports:

  1. If you have not already done so, compile your kernel to obtain the HTML reports
  2. Open the HTML reports.
  3. Go to Views > System Viewer
  4. In the left pane, expand System and then expand Global memory

    Under Global memory, you see entries for all external memories for your kernel.

  5. Click on a memory to display that memory in the System Viewer pane.
  6. In the System Viewer pane, find the box that represents the memory and click the node inside the box.

    This node represents the “interface” of that global memory.

  7. In the Details pane, find the Start Address of the memory.

    The top bits of your unannotated pointer argument must match the top bits of this start address if you want the pointer to access this buffer location.

The following screen capture shows an example of determining the top bit of the start address needed to address a buffer location 1.