11.2. Memory Attributes for Configuring Kernel Memory Systems

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

Download PDF

ID 683846

Date 12/13/2021

Version

Public

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

Visible to Intel only — GUID: ewa1456330445366

Ixiasoft

View Details

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 ₀, b ₁, ..., b _n)`	Forces the memory system to split into 2ⁿ banks, with {`b` ₀, `b` ₁, ..., `b` _n} forming the bank-select bits. Important: `b` ₀, `b` ₁, ..., `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.

Select Your Language

Using Intel.com Search

Quick Links

Recent Searches

Advanced Search

Only search in

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

11.2. Memory Attributes for Configuring Kernel Memory Systems