Intel® FPGA SDK for OpenCL™ Standard Edition: Best Practices Guide

ID 683176
Date 9/24/2018
Public
Document Table of Contents

2.3.2. Changing the Memory Access Pattern Example

The following is an example code of a simple OpenCL kernel:
kernel void big_lmem_4r_4w_nosplit (global int* restrict in, 
                                    global int* restrict out) {  
    local int lmem[4][1024]; 
 
    int gi = get_global_id(0);  
    int gs = get_global_size(0);  
    int li = get_local_id(0);  
    int ls = get_local_size(0);  
    int res = in[gi];              
    
    #pragma unroll   
    for (int i = 0; i < 4; i++) {          
         lmem[i][(li*i) % ls] = res;    
         res >>= 1;  }    

    // Global memory barrier
    barrier(CLK_GLOBAL_MEM_FENCE); 

    res = 0;  
    #pragma unroll   
    for (int i = 0; i < 4; i++) {    
        res ^= lmem[i][((ls-li)*i) % ls];  }     
    out[gi] = res;
}

The system viewer report of this example highlights the stallable loads and stores.

Figure 16. System View of the Example
Figure 17. Area Report of the Example
Figure 18. Kernel Memory Viewer of the Example

Observe that only two memory banks are created, with high arbitration on the first bank between load and store operations. Now, switch the banking indices to the second dimension, as shown in the following example code, :

kernel void big_lmem_4r_4w_nosplit (global int* restrict in, 
                                    global int* restrict out) {  
  local int lmem[1024][4];  

  int gi = get_global_id(0);  
  int gs = get_global_size(0);  
  int li = get_local_id(0);  
  int ls = get_local_size(0);  
  int res = in[gi];              
    
  #pragma unroll   
  for (int i = 0; i < 4; i++) {          
    lmem[(li*i) % ls][i] = res;    
    res >>= 1; 
  }    
    
  // Global memory barrier
  barrier(CLK_GLOBAL_MEM_FENCE); 
    
  res = 0;  
  #pragma unroll   
  for (int i = 0; i < 4; i++) {    
    res ^= lmem[((ls-li)*i) % ls][i];  
  }     
  out[gi] = res;
}

In the kernel memory viewer, you can observe that now four memory banks are created, with separate load store units. All load store instructions are stall-free.

Figure 19. Kernel Memory Viewer of the Example After Changing the Banking Indices
Figure 20. Area Report of the Example After Changing the Banking Indices