Intel® FPGA SDK for OpenCL™ Pro Edition: Programming Guide

ID 683846
Date 12/19/2022
Public
Document Table of Contents

11.2. Memory Attributes for Configuring Kernel Memory Systems

The Intel® FPGA SDK for OpenCL™ provides kernel memory attributes that you can apply to constant, local, and private variables or arrays in your design to customize the on-chip memory architecture of the local and private memory systems.
Attention: Only apply these local memory kernel attributes to constant, local or private variables.
Table 21.  OpenCL Attributes for Configuring Local Memory
Attribute Description
register Specifies that the variable or array must be carried through the pipeline in registers. Registers can be implemented either exclusively in FFs or in a combination of FFs and RAM-based FIFOs.
memory("impl_type") Specifies that the variable or array must be implemented in a memory system. Including the memory kernel attribute is equivalent to declaring the variable or array with the __local qualifier.

You can pass an optional string argument to specify the memory implementation type. Specify impl_type as either BLOCK_RAM or MLAB to implement the memory using memory blocks (such as M20K) or memory logic array blocks (MLABs), respectively.

numbanks(N) Specifies that the memory system implementing the variable or array must have N banks, where N is a power-of-2 integer value greater than zero.
bankwidth(N) Specifies that the memory system implementing the variable or array must have banks that are N bytes wide, where N is a power-of-2 integer value greater than zero.
singlepump Specifies that the memory system implementing the variable or array must be clocked at the same rate as the component accessing it.
doublepump Specifies that the memory system implementing the variable or array must be clocked at twice the rate as the component accessing it.
merge("label", "direction") Forces two or more variables or arrays to be implemented in the same memory system.

label is an arbitrary string. Assign the same label to all variables that you want to merge.

Specify direction as either width or depth to identify whether the memories should be merged width-wise or depth-wise, respectively.

bank_bits(b 0 , b 1 , ..., b n ) Forces the memory system to split into 2n banks, with {b 0 , b 1 , ..., b n } forming the bank-select bits.
Important: b 0 , b 1 , ..., b n must be consecutive, positive integers.
Note: If you specify the numbanks(n) attribute without the bank_bits attribute, the compiler automatically infers the bank-select bits based on the memory access pattern.
private_copies(N)

Specifies that the variable or array declared or accessed inside a pipelined loop has a maximum of N private copies to allow N simultaneous iterations of the loop at any given time, where N is an unsigned integer value.

Apply this attribute when the scope of a variable (through its declaration or access pattern) is limited to a loop. If the loop also has a #pragma max_concurrency M , the number of private copies created is min(M,N).

max_replicates(N) Specifies that the memory implementing the variable or array has no more than N replicates, where N is an integer value greater than 0, to enable simultaneous reads from the datapath.
simple_dual_port Specifies that the memory implementing the variable or array should have no port that services both reads and writes.
force_pow2_depth(N) Specifies that the memory implementing the variable or array has a power-of-2 depth. This option is enabled if N is 1 and disabled if N is 0. The default value is 1.
Table 22.  Code Examples for Memory Attributes
Example Use Case Syntax
Implements a variable in a register
int __attribute__((register)) a[12];
Implements a memory system with eight banks, each with a width of 8 bytes
int __attribute__((memory,
                   numbanks(8),
                   bankwidth(8)) b[16];
Implements a double-pumped memory system with one 128-byte wide bank, and a maximum of two replicates.
int __attribute__((memory,
                   numbanks(1),
                   bankwidth(128),
                   doublepump,
                   max_replicates(2)))  c[32];

You can also apply memory attributes to data members of a struct. Specify attributes for struct data members in the struct declaration. If you apply attributes to an object instantiation of a struct, then those attributes override the attributes specified in the declaration for struct data members. For example, consider the following code:

struct State { 
     int array[100] __attribute__((__memory__)); 
     int reg[4] __attribute__((__register__)); 
}; 
__kernel void sum(...) { 
     struct State S1; 
     struct State S2 __attribute__((__memory__)); 
     // some uses 
}

The offline compiler splits S1 into two variables as S1.array[100] (implemented in memory) and S1.reg[4] (implemented in registers). However, the compiler ignores attributes applied at struct declaration for object S2 and does not split it as the S2 has the attribute memory applied to it.