ID 683176
Date 9/24/2018
Public

## 8.3. Memory Access Considerations

recommends kernel programming strategies that can improve memory access efficiency and reduce area usage of your OpenCL™ kernel.
1. Minimize the number of access points to external memory.

If possible, structure your kernel such that it reads its input from one location, processes the data internally, and then writes the output to another location.

2. Instead of relying on local or global memory accesses, structure your kernel as a single work-item with shift register inference whenever possible.
3. Instead of creating a kernel that writes data to external memory and a kernel that reads data from external memory, implement the channels extension between the kernels for direct data transfer.
4. If your OpenCL application includes many separate constant data accesses, declare the corresponding pointers using __constant instead of __global const. Declaration using __global const creates a private cache for each load or store operation. On the other hand, declaration using __constant creates a single constant cache on the chip only.
CAUTION:
If your kernel targets a Cyclone® V device (for example, Cyclone V SoC), declaring __constant pointer kernel arguments might degrade FPGA performance.

5. If your kernel passes a small number of constant arguments, pass them as values instead of pointers to global memory.

For example, instead of passing __constant int * coef and then dereferencing coef with index 0 to 10, pass coef as a value (int16 coef). If coef was the only __constant pointer argument, passing it as a value eliminates the constant cache and the corresponding load and store operations completely.

6. Conditionally shifting large shift registers inside pipelined loops leads to the creation of inefficient hardware. For example, the following kernel consumes more resources when the if (K > 5) condition is present:
#define SHIFT_REG_LEN 1024
__kernel void bad_shift_reg (__global int * restrict src,
__global int * restrict dst,
int K)
{
float shift_reg[SHIFT_REG_LEN];
int sum = 0;

for (unsigned i = 0; i < K; i++)
{
sum += shift_reg[0];
shift_reg[SHIFT_REG_LEN-1] = src[i];

// This condition will cause sever area bloat.
if (K > 5)
{
#pragma unroll
for (int m = 0; m < SHIFT_REG_LEN-1 ; m++)
{
shift_reg[m] = shift_reg[m + 1];
}
}
dst[i] = sum;
}
}
Attention: Conditionally accessing a shift register does not degrade hardware efficiency. If it is necessary to implement conditional shifting of a large shift register in your kernel, consider modifying your code so that it uses local memory.