This document assumes that you are familiar with OpenCL concepts and application programming interfaces (APIs), as described in the OpenCL Specification version 1.0 by the Khronos Group™. It also assumes that you have experience in creating OpenCL applications.

To achieve the highest performance of your OpenCL™ application for FPGAs, familiarize yourself with details of the underlying hardware. In addition, understand the compiler optimizations that convert and map your OpenCL application to FPGAs.

For more information on the OpenCL Specification version 1.0, refer to the OpenCL Reference Pages on the Khronos Group website. For detailed information on the OpenCL APIs and programming language, refer to the OpenCL Specification version 1.0.

An FPGA consists of several small computational units. Custom datapaths can be built directly into the fabric by programming the compute units and connecting them as shown in the following figure. Data flow is programmed directly into the architecture.

Figure 1. FPGA Architecture

With FPGAs, low-level operations like bit masking, shifting, and addition are all configurable. Also, you can assemble these operations in any order. To implement computation pipelines, FPGAs integrate combinations of lookup tables (LUTs), registers, on-chip memories, and arithmetic hardware (for example, digital signal processor (DSP) blocks) through a network of reconfigurable connections. As a result, FPGAs achieve a high level of programmability. LUTs are responsible for implementing various logic functions. For example, reprogramming a LUT can change an operation from a bit-wise AND logic function to a bit-wise XOR logic function.

The key benefit of using FPGAs for algorithm acceleration is that they support wide, heterogeneous and unique pipeline implementations. This characteristic is in contrast to many different types of processing units such as symmetric multiprocessors, DSPs, and graphics processing units (GPUs). In these types of devices, parallelism is achieved by replicating the same generic computation hardware multiple times. In FPGAs, however, you can achieve parallelism by duplicating only the logic that your algorithm exercises.

A processor implements an instruction set that limits the amount of work it can perform each clock cycle. For example, most processors do not have a dedicated instruction that can execute the following C code:

E = (((A + B) ^ C) & D) >> 2;

Without a dedicated instruction for this C code example, a CPU, DSP, or GPU must execute multiple instructions to perform the operation. In contrast, you may think of an FPGA as a hardware platform that can implement any instruction set that your software algorithm requires. You can configure an FPGA to perform a sequence of operations that implements the code example above in a single clock cycle. An FPGA implementation connects specialized addition hardware with a LUT that performs the bit-wise XOR and AND operations. The device then leverages its programmable connections to perform a right shift by two bits without consuming any hardware resources. The result of this operation then becomes a part of subsequent operations to form complex pipelines.

When a kernel describes a single work item, the Intel^® FPGA SDK for OpenCL™ host can execute the kernel as a single work-item, which is equivalent to launching a kernel with an NDRange size of (1, 1, 1). The compiler tries to accelerate the single work item for best performance.

The OpenCL Specification version 1.0 describes this mode of operation as task parallel programming. A task refers to a kernel executed with one work-group that contains one work-item.

Generally, the host launches multiple work-items in parallel. However, this data parallel programming model is not suitable for situations where fine-grained data must be shared among parallel work-items. In these cases, you can maximize throughput by expressing your kernel as a single work-item. Unlike NDRange kernels, single work-item kernels follow a natural sequential model similar to C programming. Particularly, you do not have to partition the data across work-items.

To ensure high-throughput single work-item-based kernel execution on the FPGA, the Intel^® FPGA SDK for OpenCL™ Offline Compiler must process multiple pipeline stages in parallel at any given time. This parallelism is realized by pipelining the iterations of loops.

Consider the following simple example code that shows accumulating with a single-work item:

1 kernel void accum_swg (global int* a, 
                         global int* c, 
                         int size, 
                         int k_size) {
2     int sum[1024];
3     for (int k = 0; k < k_size; ++k) {
4        for (int i = 0; i < size; ++i) {
5            int j = k * size + i;
6            sum[k] += __prefetching_load(&a[j]);
7        }
8     }
9     for (int k = 0; k < k_size; ++k) {
10       c[k] = sum[k];
11    }
12 }

Figure 6. Loop Analysis Report

Figure 7. System View of Single-Work Item Kernel

The following figure illustrates how each iteration of i enters into the block:

When you observe the outer loop, having an II value of 1 also means that each iteration of the thread can enter at every clock cycle. In the example, k_size of 20 and size of 4 is considered. This is true for the first eight clock cycles as outer loop iterations 0 to 7 can enter without any downstream stalling it. Once thread 0 enters into the inner loop, it takes four iterations to finish. Threads 1 to 8 cannot enter into the inner loop and they are stalled for four cycles by thread 0. Thread 1 enters into the inner loop after thread 0's iterations are completed. As a result, thread 9 enters into the outer loop on clock cycle 13. Threads 9 to 20 enters into the loop at every four clock cycles, which is the value of size. Through this example, you can observe that the dynamic initiation interval of the outer loop is greater than the statically predicted initiation interval of 1 and it is a function of the trip count of the inner loop.

Figure 9. Single Work-Item Execution

Consider the same accumulate example written in NDRange:

kernel void accum_ndr (global int* a, 
                       global int* c, 
                       int size) {
   int k = get_global_id(0);

   int sum[1024];
   for (int i = 0; i < size; ++i) {
     int j = k * size + i;
     sum[k] += a[j];
   }
   c[k] = sum[k];
}

Figure 10. Loop Analysis Report

Figure 11. System View of the NDRange Kernel

The report summary gives you a quick overview of the results of compiling your design including a summary of each kernel in your design and a summary of the estimated resources that each kernel in your design uses.

The report summary is divided into the following sections: Info, Quartus Fit Summary, Kernel Summary , Estimated Resource Usage, and Compile Warnings.

Use the f_max II report to help you with the following tasks:

• Identify f_max bottleneck in your design
• Estimate the execution latency
• Determine where to use loop pragmas

Access the f_max II report, by selecting Throughput Analysis→f_max II Report on the high-level design reports page (report.html).

Following is an OpenCL design example that includes a loop:

kernel void lowered_fmax (global int *dst, int N) {
    int res = N;
    #pragma unroll  9
    for (int i = 0; i < N; i++) {
	    res += 1;
	    res ^= i;
    }
    dst[0] = res;
}

Figure 13. F_max II Report for the OpenCL Design Example

You can observe that block B1 is the f_max limiting block of the design. The scheduler can schedule this block only at 90.5 MHz due to the loop feedback path.

You can interact with the Graph Viewer in the following ways:

• Use the mouse wheel to zoom in and out within the Graph Viewer.

• Hover over any node to view information on that node in the tool-tip and in the Details pane.

For example, a logic block that has a pipelined loop with a high initiation interval (II) value might be highlighted in red because the high II value might affect design throughput.

You can hide certain types of connections in the system view of the Graph Viewer by unchecking that type of connection. By default, both Control and Memory are checked in the Graph Viewer. Control refers to connections between blocks and loops. Memory refers to connections to and from global or local memories. If your design includes connections to and from read or write channels, you also have a Channels option in the Graph Viewer.

The block view of the Graph Viewer provides a more granular graph view of the kernel. This view shows the following:

• Fine grained details within kernels (including instructions and dependencies of the instructions) of the generated datapath of computations. The Intel^® FPGA SDK for OpenCL™ Offline Compiler encapsulates maximum instructions in clusters for better QoR. The Graph Viewer shows clusters, instructions outside clusters and their connections.
• Linking from the instruction back to source line by clicking the instruction node.
• Various information about the instructions, such as data width, node’s schedule information in start cycle and latency are provided, if applicable.

The following image is an example of the block view of the Graph Viewer:

Figure 15. Granular Graph View of a Block

If your design has loops, the Intel^® FPGA SDK for OpenCL™ Offline Compiler encapsulates the loop control logic into loop orchestration nodes and the initial condition of the loops into loop input node and their connection to the datapath.

Inside a block, there are often stallable channel RD/WR or memory LD/ST nodes connecting to computation nodes or clusters. You can click on different nodes and view the Details pane (or hover over the nodes) to see detailed information about the instruction. For example, you can click on the LD/ST nodes to see attributes such as instruction type, width, LSU style, stall-free, global memory, scheduled start cycle, and estimated latency. For stallable nodes, the latency value provided is an estimate. Perform a simulation or hardware run for more accurate latency values.

Figure 16. Instruction Node with Details in Pop-up Box

A cluster has a FIFO in its exit node to store any pipelined data in-flight. You can click on the cluster exit node to find the exit FIFO width and depth attribute. The cluster exit FIFO size is also available in the cluster view of the Graph Viewer. Refer to the following image for cluster exit FIFO details in a block view example.

Figure 17. Cluster Node With FIFO Details

This view shows the following:

• Fine grained details within clusters (including instructions and dependencies of the instructions) of the generated datapath of computations.
• Linking from the instruction back to source line by clicking the instruction node.
• Various information about the instructions, such as data width, node’s schedule information in start cycle and latency are provided, if applicable.

A cluster starts with an entry node and ends with an exit node. The cluster exit node has a FIFO of depth greater than or equal to the latency of the cluster to store any data in-flight. You can find the size of the cluster exit FIFO by clicking on the exit node. The cluster exit FIFO size information is also available in the block view of the Graph Viewer.

Figure 18. Granular Graph View of a Cluster

Figure 19. Cluster Exit Node Example

Besides computation nodes, when your design contains loops, you can see loop orchestration nodes and variable nodes along with their Feedback nodes. The compiler generates the loop orchestration logic for loops in your design. This logic is represented by loop orchestration nodes in the cluster view of the Graph Viewer. A variable node corresponds to a variable that has loop-carried dependency in your design. A variable node goes through various computation logic and finally feeds to a Feedback node that connects back to the variable node. This back edge means that the variable is passed to the next iteration after the new value is evaluated. Scan for loop-carried variables that have a long latency to the Feedback nodes as they can be the II bottlenecks. You can cross-check by referring to the Loop Analysis report for more information about the II bottleneck. The Feedback node has a FIFO to store any data in-flight for the loop and is sized to d*II where d is the dependency distance and II is the initiation interval. You can find the size of the cluster exit FIFO by clicking on the feedback node and looking at the Details pane or the pop-up box. Refer to Figure 19 for cluster exit FIFO details in the cluster view example.

Data movement is a bottleneck in many algorithms. The Kernel Memory Viewer in the High Level Design Report (report.html) shows you how the Intel^® FPGA SDK for OpenCL™ Offline Compiler interprets the data connections and synthesizes memory for your kernel. Use the Kernel Memory Viewer to help you identify data movement bottlenecks in your kernel design.

Some patterns in memory accesses can cause undesired arbitration in the load-store units (LSUs), which can affect the throughput performance of your kernel. Use the Kernel Memory Viewer to identify unwanted arbitration in the LSUs.

Access the Kernel Memory Viewer by clicking System Viewers > Kernel Memory Viewer.

The following image illustrates the layout of the Kernel Memory Viewer:

The Kernel Memory Viewer has the following panes:

Kernel Memory List

Lists all memories present in your design. When you select a memory name, you can view its graphical representation in the Kernel Memory Viewer pane.

Kernel Memory Viewer

Shows a graphical representation of the memory system or memory bank selected in the Kernel Memory List pane.

Code View

Shows the source code file for which the reports are generated.

Details

Shows the details of the memory system or memory bank selected in the Kernel Memory List pane.

Kernel Memory List

The Kernel Memory List pane displays a hierarchy of kernels with memories synthesized (RAMs, ROMs, and registers) and optimized away in that kernel.

Figure 21. Features and Details of the Kernel Memory List Pane

The following table describes each numbered feature highlighted in the above image:

Table 5.  Kernel Memory List Pane Icons and Labels

	Icon or Label	Name	Description
1		Kernel name	The list of memories in your kernel can be expanded or collapsed. Memories that do not belong to any kernel are shown under (Other).
2		RAM	A RAM is a memory that has at least one write to it. The name of the RAM memory is same as its name in your design. When you select a memory name, you can view a logical representation of the RAM in the Kernel Memory Viewer pane. By default, only the first bank of the memory system is displayed. To select banks that you want the Kernel Memory Viewer pane to display: Expand the memory name. Clear the memory name check box to collapse all memory banks in the view. Select the memory name check box to show all memory banks in the view.
3		ROM	A ROM is a memory that is read-only. The name of the ROM memory is same as its name in your design. When you select a memory name, you can view a logical representation of the ROM in the Kernel Memory Viewer pane. By default, only the first bank of the memory system is displayed. To select banks that you want the Kernel Memory Viewer pane to display: Expand the memory name. Clear the memory name check box to collapse all memory banks in the view. Select the memory name check box to show all memory banks in the view.
4	Bank #num	Bank	A memory bank is always associated with a RAM or a ROM. Each bank is named as Bank #num, where #num is the ID of the memory bank starting from 0. Click on the bank name to display the bank view in the Kernel Memory Viewer pane, which displays a graphical representation of the bank with all of its replicates and private copies. This view can help you focus on specific memory banks of a complex memory design. Clear the memory bank name check box to collapse the bank in the logical representation of the memory. Select the memory bank name check box to display the bank in the logical representation of the memory.
5		Register	A register is a kernel variable that is carried through the pipeline in registers (rather than being stored in a RAM or ROM). The name of the register is same as its name in your design. You can implement a register variable either exclusively in FFs or in a combination of FFs and RAM-based FIFOs.
6	text label	Optimized Away	A kernel variable can be optimized away because it is unused in your design, or compiler optimizations have transformed all uses of the variable such that it is unnecessary. The name of the optimized away variable is same as its name in your design.
7		Filter	Use the Kernel Memory List filter to selectively view the list of RAMs, ROMs, registers, and optimized away variables in your design. When you clear the check box associated with an item in the filter, you hide all occurrences of that kind of item in the Kernel Memory List. Filter your Kernel Memory List to help you focus on a specific type of memory in your design.

Access the Schedule Viewer by clicking System Viewers > Schedule Viewer.

In the Schedule Viewer graph, the vertical axis depicts the order of execution and the horizontal axis depicts the clock cycles. The Schedule List (left-hand pane) displays a list of kernels, blocks, and clusters.

OpenCL design example that performs matrix square AxA:

 // performs matrix square A*A
 // A is a square LEN*LEN matrix
 // A = [ r0       : [ c[0], c[1], ... c[LEN-1] ],
 //       r1       : [ c[0], ...                ],
 //       ...                                   ],
 //       r[LEN-1] : [                          ] ]
 
 // LEN = 100
 
 kernel void matrix_square (global float* restrict A, global float* restrict out)
 {
 	for( unsigned oi = 0 ; oi < LEN*LEN ; oi++ ) 
 	{
 		float sum = 0.f;
 		int row = oi / LEN;
 		int col = oi % LEN;
 
 		#pragma unroll
 		for ( int stride = 0 ; stride < LEN ; stride++ )
 		{
 			unsigned i = (row * LEN) + stride;
 			unsigned j = (stride * LEN) + col;
 			sum += A[i] * A[j];
 		}
 		
 		out[oi] = sum;
 	}
 }

The area analysis of the kernel matrix_square indicates that the estimated usages of flipflops (FF) and RAMs are high.

Figure 24. Area Report of the Unoptimized Kernel matrix_square

Further examination of block 2 in the Graph viewer shows that Block2 also has a high latency value.

Figure 25. Graph Viewer Results for the Unoptimized Kernel matrix_square

The cause for these performance bottlenecks is the system is loading data from global memory from inside a loop. Therefore, the first optimization step you can take is to preload the data into local memory, as shown in the following modified code:

 // 1. preload the data into local memory
 
 kernel void matrix_square_v1 (global float* restrict A, global float* restrict out)
 {
 	local float cache_a[LEN*LEN];
 	for( unsigned k = 0 ; k < LEN*LEN ; k++ )
 	{
 		cache_a[k] = A[k];
 	}
 
 	for( unsigned oi = 0 ; oi < LEN*LEN ; oi++ ) 
 	{
 		float sum = 0.f;
 		int row = oi / LEN;
 		int col = oi % LEN;
 		
     #pragma unroll
     for( unsigned stride = 0 ; stride < LEN ; stride++ )
 		{
 			unsigned i = (row * LEN) + stride;
 			unsigned j = (stride * LEN) + col;
 			sum += cache_a[i] * cache_a[j];
 		}
 		
 		out[oi] = sum;
 	}
 }

Figure 26. Area Report Results for the Modified Kernel matrix_square_v1

Figure 27. Cluster View of the Modified Kernel matrix_square_v1

If you remove the modulus computation and replace it with a column counter, as shown in the modified kernel matrix_square_v2, you can reduce the amount of adaptive look-up table (ALUT) and FF use.

 // 1. preload the data into local memory
 // 2. remove the modulus computation
 
 kernel void matrix_square_v2 (global float* restrict A, global float* restrict out)
 {
 	local float cache_a[LEN*LEN];
 	for( unsigned k = 0 ; k < LEN*LEN ; k++ )
 	{
 		cache_a[k] = A[k];
 	}
 
 	unsigned row = 0;
 	unsigned col = 0;
 	
 	for( unsigned oi = 0 ; oi < LEN*LEN ; oi++ )
 	{
 		float sum = 0.f;
 
 		// keep a column counter to know when to increment row
 		if( col == LEN - 1 )
 		{
 			col = 0;
 			row += 1;
 		}
 		else
 		{
 			col += 1;
 		}
 		
 		#pragma unroll
 		for( unsigned stride = 0 ; stride < LEN ; stride++ )
 		{
 			unsigned i = (row * LEN) + stride;
 			unsigned j = (stride * LEN) + col;
 			sum += cache_a[i] * cache_a[j];
 		}
 		
 		out[oi] = sum;
 	}
 }

Figure 28. Area Report of Kernel matrix_square_v2

Figure 29. Cluster View of Kernel matrix_square_v2

Further examination of the area report of matrix_square_v2 reveals that the computations for indexes i and j (that is, unsigned i = (row * LEN) + stride and unsigned j = (stride * LEN) + col, respectively) have very different ALUT usage estimations.

A way to optimize DSP and RAM block usages for index calculation is to remove the multiplication computation and simply keep track of the addition, as shown in the modified kernel matrix_square_v3 below.

// 1. preload the data into local memory
 // 2. remove the modulus computation
 // 3. remove DSP and RAM blocks for index calculation helps reduce the latency
 
 kernel void matrix_square_v3 (global float* restrict A, global float* restrict out)
 {
 
   local float cache_a[LEN*LEN];
 	for( unsigned k = 0 ; k < LEN*LEN ; k++ )
 	{
 		cache_a[k] = A[k];
 	}
 
 	unsigned row_i = 0;
 	unsigned row_j = 0;
 
 	for( unsigned oi = 0 ; oi < LEN*LEN ; oi++ )
 	{
 		unsigned i, j;
 
 		// keep a column base counter to know when to increment the row base
 		if( row_j == LEN - 1 )
 		{
 			row_i += LEN;
 			row_j = 0;
 		}
 		else
 		{
 			row_j += 1;
 		}
 		
 		// initialize i and j
 		i = row_i;
 		j = row_j;
 		
 		float sum = 0.f;
 		#pragma unroll
 		for( unsigned stride = 0 ; stride < LEN ; stride++ )
 		{
 			i += 1;       // 0, 1, 2, 3, 0,...
 			j += LEN;     // 0, 3, 6, 9, 1,...
 			sum += cache_a[i] * cache_a[j];
 		}
 		
 		out[oi] = sum;
 	}
 }

By removing the multiplication step, you can reduce DSP usage as shown in the area report below. In addition, the modification helps reduce latency.

Figure 30. Area Report of Kernels matrix_square_v3

Figure 31. Graph Viewer of the Kernel matrix_square_v3

To resolve the loop-carried dependency, unroll the sum-product for complete parallelism and create registers to avoid multiple copies of cache_a, as shown in the code below in the modified kernel matrix_square_v4.

// 1. preload the data into local memory
 // 2. remove the modulus computation
 // 3. remove DSP and RAM blocks for index calculation helps reduce the latency
 // 4. unroll the sum-product for full parallelism, create registers to avoid many copies of cache_a
 
 kernel void matrix_square_v4 (global float* restrict A, global float* restrict out)
 {
 
   local float cache_a[LEN*LEN];
 	for( unsigned k = 0 ; k < LEN*LEN ; k++ )
 	{
 		cache_a[k] = A[k];
 	}
 
 	unsigned row_i = 0;
 	unsigned row_j = 0;
 
 	for( unsigned oi = 0 ; oi < LEN*LEN ; oi++ )
 	{
 		unsigned i, j;
 		// keep a column base counter to know when to increment the row base
 		if( row_j == LEN - 1 )
 		{
 			row_i += LEN;
 			row_j = 0;
 		}
 		else
 		{
 			row_j += 1;
 		}
 		
 		// initialize i and j
 		i = row_i;
 		j = row_j;
 		
 		float r_buf[LEN];
 		float c_buf[LEN];
 		for( int stride = 0 ; stride < LEN ; stride++ )
 		{
 			i += 1;       // 0, 1, 2, 3, 0,...
 			j += LEN;     // 0, 3, 6, 9, 1,...
 			r_buf[stride] = cache_a[i];
 			c_buf[stride] = cache_a[j];
 		}
 
 		// uses harder floating point when -fp-relaxed is used
 		float sum = 0.f;
 		#pragma unroll
 		for(unsigned idx = 0; idx &lt; LEN; idx++)
 		{
 			sum += r_buf[idx] * c_buf[idx];
 		}
 
 		out[oi] = sum;
 	}
 }

As shown in the cluster view results below, by breaking up the computation steps, you can achieve higher throughput at the expense of increased area usage. The modification also reduces the latency by 50%.

Figure 32. Cluster View of the Modified Kernel matrix_square_v4

The following cluster view provides an alternative with -fp-relaxed to show dot product instead of chain:

Figure 33. Cluster View of the Modified Kernel matrix_square_v4

The following table provides the throughput comparison for all five versions of the kernel:

The JSON files containing the HLD FPGA Report report data are available in the <your_kernel_filename>/reports/lib/json directory. The directory provides the following .json files:

You can read the following .json files without a special parser:

• area.json
• area_src.json
• loops.json
• quartus.json
• summary.json

For example, if you want to identify all of the values and bottlenecks for the initiation interval (II) of a loop, you can find the information in the children section in the loops.json file, as shown below:

“name”:”<block name|Kernel: kernel name>  # Find the loops which does not begin with “Kernel:”

“data”:[<Yes|No>, <#|n/a>, <II|n/a>]      # The data field corresponds to “Pipelined”, “II”, “Bottleneck”

Each kernel in your OpenCL system is represented by a set of blocks. Inside each block is a set of non-branching instructions that implement your algorithm and the offline compiler's loop orchestration logic. The block shows the execution flow of your kernel. When there are loops, there is a back edge to the block or its previous block(s), depending on the loop structure, for example, nested loops. Loops usually impose II bottlenecks and are a main focus for optimization.

A block has three main parts — an input or loop input node, a set of instructions and a branch node. The input and branch nodes may not be present depending on if there is branching in or out of the block. The input or loop input node determines the initial value for variables depending on where the branch into this block originated. The rest of the block contains stallable and non-stallable instructions, and clusters. A well-optimized design should contain a minimal number of stallable instructions, such as stallable I/O or memory accesses.

The non-stallable instructions within in block are grouped into clusters to reduce handshaking overheads with stallable instructions. A cluster has an entry and an exit node. There is only a stall-free cluster. You can find the exit FIFO information in the cluster’s exit node. Finally, the branch node informs the next block to go to, under which condition.

In the HLD report, you can find different views of your kernel, under the System Viewers drop-down menu. For more information about System Viewers, refer to Using System Viewers. The Graph Viewer shows the branching, stallable operations and the clusters, and their connections. The cluster view shows the content inside clusters.

The Intel^® FPGA SDK for OpenCL™ Offline Compiler compiles a kernel that does not use any built-in work-item functions, such as get_global_id() and get_local_id(), as a single work-item kernel. Otherwise, the offline compiler compiles the kernel as an NDRange kernel. For more information on built-in work-item functions, refer to section 6.11.1: Work-Item Functions of the OpenCL Specification version 1.0.

For single work-item kernels, the offline compiler attempts to pipeline every loop in the kernel to allow multiple loop iterations to execute concurrently. Kernel performance might degrade if the compiler cannot pipeline some of the loops effectively, or if it cannot pipeline the loops at all.

The offline compiler cannot pipeline loops in NDRange kernels. However, these loops can accept multiple work-items simultaneously. A kernel might have multiple loops, each with nested loops. If you tabulate the total number of iterations of nested loops for each outer loop, kernel throughput is usually reduced by the largest total iterations value that you have tabulated. To execute an NDRange kernel efficiently, there usually needs to be a large number of threads.

Unlike a GPU, an FPGA can build any custom LSU that is best suited for the memory access pattern that the compiler infers for your application. As a result, your ability to write OpenCL code that selects the ideal LSU types for your application might help improve the performance of your design significantly.

When reviewing the HTML area report of your design, the values in the Global interconnect entry at the system level represents the size of the global memory interconnect.

Figure 35. HTML Area Report Showing the Size of the Global Memory Interconnect in an OpenCL Design

In the HTML report, the system view of the Graph Viewer depicts global memory interconnects as loads (LD), stores (ST), and connections (gray lines).

Figure 36. Graph Viewer (System View) Result of Global Memory Interconnects in an OpenCL Design

The Intel^® FPGA SDK for OpenCL Offline Compiler selects the appropriate type of LSU for your OpenCL system based on the memory access pattern of your design. Example LSU types include contiguous access (or consecutive access) and burst-interleaved access. Figure 91 and Figure 90 illustrate the difference in access patterns between contiguous and burst-interleaved memory accesses, respectively.

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;
}

In the Graph Viewer report, the system view of this example highlights the stallable loads and stores.

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.

The orig kernel has nested loops to a depth of four. The nested loops created extra blocks (Block 2, 3, 4, 6, 7 and 8) that consume area due to the variables being carried, as shown in the following reports:

Due to loop coalescing, you can see the reduced latency in the lc_test. The Block 5 of orig kernel and Block 12 of lc_test kernel are the inner most loops.

When declaring a channel in your kernel code, precede the declaration with the keyword channel.

For example:

channel long16 myCh __attribute__((depth(16)));

In the HTML report, the area report maps the channel area to the declaration line in the source code. Channels and channel arrays are reported with their width and depth.

The depth attribute is treated as the minimum depth specification. The offline compiler may increase the depth for the following reasons:

• Instruction scheduling requirements. This may happen due to the following reasons:
  • To balance reconverging paths through multiple kernels.

    When multiple paths exist from one kernel to another via channels and other kernels, the depths on these channels may be increased to balance latencies among all these paths. This is a throughput optimization, as unbalanced paths are likely to lead to pipeline stalls.

  • Due to the compiler optimization that can relax a loop-carried dependency at the expense of deepening the channel with a non-blocking write.

    Consider the following code that reads next_val from the global memory only if the non-blocking write to my_channel succeeded in the previous iteration.

    bool write_valid = true;
    int next_val = 0;
    while (not_done) {
       if (write_valid) {
          next_val = *global_mem_ptr;
          global_mem_ptr++;
       }
       write_valid = write_channel_nb_intel(my_channel, next_val);
       not_done = some_fn(next_val);
    }

    With a naive implementation, this loop has a very high II because the high-latency global memory read must complete before the write into the channel can begin and the next value of write_valid is computed. To remove the global read from the II-critical path, the compiler can instead check if my_channel has space to accept whatever value is read from the global memory before doing the actual global read. The check for channel fullness takes one clock cycle and hence, the next loop iteration can start as soon as the channel fullness check is complete, giving II=1 for this loop. To make the resulting hardware functionally correct, the channel must be deepened by the latency of the global read or, to be precise, schedule distance between the channel fullness check and the actual write, which may be slightly greater than the global read. If you do not want to have your channel deepened in this situation, identify and remove the loop-carried dependency involving a valid value from the write_channel_nb_intel() call.

• The nature of the underlying FIFO implementation. This happens if the chosen underlying implementation cannot support the exact depth required and must be increased to the next supported size.

• Burst-Coalesced Load-Store Units
• Prefetching Load-Store Units
• Pipelined Load-Store Units
• Constant-Pipelined Load-Store Units
• Atomic-Pipelined Load-Store Units

kernel void burst_coalesced (global int * restrict in, 
                             global int * restrict out) {
  int i = get_global_id(0);
  int value = in[i/2];      // Burst-coalesced LSU
  out[i] = value;
}

kernel void prefetching (global int * restrict in, 
                         global int * restrict out, 
                         int N) {
  int res = 1;
  for (int i = 0; i < N; i++) {
    int v = in[i];             // Prefetching LSU
    res ^= v;
  }
  out[0] = res;
}

__attribute((reqd_work_group_size(1024,1,1)))
kernel void local_pipelined (global int* restrict in, 
                             global int* restrict out) {
  local int lmem[1024];
  int gi = get_global_id(0);
  int li = get_local_id(0);

  int res = in[gi];
  for (int i = 0; i < 4; i++) {
    lmem[li - i] = res;                 // pipelined LSU
    res >>= 1;
  }

  barrier(CLK_GLOBAL_MEM_FENCE);

  res = 0;
  for (int i = 0; i < 4; i++) {
    res ^= lmem[li - i];                // pipelined LSU
  }

  out[gi] = res;
}

kernel void global_infrequent (global int * restrict in, 
                               global int * restrict out, 
                               int N) {
  int a = 0;
  if (get_global_id(0) == 0)
      a = in[0];                   // Pipelined LSU
  for (int i = 0; i < N; i++) {
    out[i] = in[i] + a;
  }
}

 kernel void constant_pipelined (constant int *src, 
                                 global int *dst) {
  int i = get_global_id(0);
  dst[i] = src[i];              // Constant pipelined LSU
}

kernel void atomic_pipelined (global int* restrict out) {
  atomic_add(&out[0], 1);  // Atomic LSU
}

1. Pipelined LSU (load/store): It is area efficient but it can be slower than other LSUs. You should use this LSU if you are constricted on area or if your access patterns are not necessarily consecutive.
2. Prefetching LSU (only for loads): It is also area efficient but it is perfect for fully consecutive access patterns. There is a throughput penalty for using it for non-consecutive access patterns, so, use it only if you know that the addresses accessed are strictly consecutive.
3. Burst-coalesced LSU (load/store): It is expensive in area but can process consecutive access patterns very efficiently. There is an area penalty for checking whether the access patterns are consecutive or not. The LSU dynamically attempts to combine several kernel requests into one big burst spanning multiple memory words, if possible.
4. Burst-coalesced cached LSU (only for loads): It is the most expensive in area because it contains an extra cache that is local to the LSU. It can help the throughput in cases where you intend to read the same location in memory multiple times, especially across multiple ND-range threads.

In general, you should optimize a kernel that targets a single compute unit first. After you optimize this compute unit, increase the performance by scaling the hardware to fill the remainder of the FPGA. The hardware footprint of the kernel correlates with the time it takes for hardware compilation. Therefore, the more optimizations you can perform with a smaller footprint (that is, a single compute unit), the more hardware compilations you can perform in a given amount of time.

In addition to data processing and memory access optimizations, consider implementing the following design practices, if applicable, when you create your kernels.

Sometimes, FPGA-to-global memory bandwidth constrains the data transfer efficiency between kernels. The theoretical maximum FPGA-to-global memory bandwidth varies depending on the number of global memory banks available in the targeted Custom Platform and board. To determine the theoretical maximum bandwidth for your board, refer to your board vendor's documentation.

In practice, a kernel does not achieve 100% utilization of the maximum global memory bandwidth available. The level of utilization depends on the access pattern of the algorithm.

If global memory bandwidth is a performance constraint for your OpenCL kernel, first try to break down the algorithm into multiple smaller kernels. Secondly, as shown in the figure below, eliminate some of the global memory accesses by implementing the SDK's channels or OpenCL pipes for data transfer between kernels.

For more information on the usage of channels, refer to the Implementing Intel^® FPGA SDK for OpenCL™ Channels Extension section of the Intel^® FPGA SDK for OpenCL™ Programming Guide.

For more information on the usage of pipes, refer to the Implementing OpenCL Pipes section of the Intel^® FPGA SDK for OpenCL™ Programming Guide.

Consider the following code examples:

__kernel void
consumer (__global uint*restrict dst) {
  for (int i = 0; i < 5; i++) {
    dst[2*i] = read_channel_intel(c0);
    dst[2*i+2] = read_channel_intel(c1);
  }
}

__kernel void
consumer (__global uint*restrict dst,
  read_only pipe uint 
    __attribute__((blocking)) c0,
  read_only pipe uint
    __attribute__((blocking)) c1)
{
  for (int i = 0; i < 5; i++) {
    read_pipe (c0, &dst[2*i]);
    read_pipe (c1, &dst[2*i+2]);
  }
}

The code example on the left makes two read channel calls. The code example on the right makes two read pipe calls. In most cases, the kernel executes these channel or pipe calls in parallel; however, channel and pipe call executions might occur out of sequence. Out-of-sequence execution means that the read operation from c1 can occur and complete before the read operation from c0.

During compilation, the offline compiler computes scheduling mismatches between interacting channels or pipes. These mismatches might cause imbalances between read and write operations. The offline compiler performs buffer inference optimization automatically to correct the imbalance.

Consider the following examples:

__kernel void producer (
  __global const uint * restrict src,
  const uint iterations)
{
  for(int i = 0; i < iteration; i++)
  {
    write_channel_intel(c0,src[2*i]);
    write_channel_intel(c1,src[2*i+1]);
  }
}

__kernel void consumer (
  __global uint * restrict dst,
  const uint iterations)
{
  for(int i = 0; i < iterations; i++)
  {
    dst[2*i] = read_channel_intel(c0);
    dst[2*i+1] = read_channel_intel(c1);
  }
}

__kernel void producer (
  __global const uint * restrict src,
  const uint iterations,
  write_only pipe uint
    __attribute__((blocking)) c0,
  write_only pipe uint
    __attribute__((blocking)) c1)
{
  for(int i = 0; i < iteration; i++)
  {
    write_pipe(c0,&src[2*i]);
    write_pipe(c1,&src[2*i+1]);
  }
}

__kernel void consumer (
  __global uint * restrict dst,
  const uint iterations,
  read_only pipe uint
    __attribute__((blocking)) c0,
  read_only pipe uint
    __attribute__((blocking)) c1)
{
  for(int i = 0; i < iterations; i++)
  {
    read_pipe(c0,&dst[2*i]);
    read_pipe(c1,&dst[2*i+1]);
  }
}

The offline compiler performs buffer inference optimization if channels or pipes between kernels cannot form a cycle. A cycle between kernels is a path that originates from a kernel, through a write channel or a write pipe call, and returns to the original kernel. For the example, assume that the write channel or write pipe calls in the kernel producer are scheduled 10 cycles apart and the read channel or read pipe calls are scheduled 15 cycles apart. There exists a temporary mismatch in the read and write operations to c1 because five extra write operations might occur before a read operation to c1 occurs. To correct this imbalance, the offline compiler assigns a buffer size of five cycles to c1 to avoid stalls. The extra buffer capacity decouples the c1 write operations in the producer kernel and the c1 read operations in the consumer kernel.

• Use single-threaded kernels over multi-threaded kernels.
• Consider how the design model can be represented with a feed forward datapath, for example, back-to-back loops or discrete processing steps. Determine whether you should split the design into multiple kernels connected by channels.
• Aggregate data on channels only when the entire data is used at the same point of kernel.
• Attempt to keep the number of channels per kernel reasonable.
• Do not use non-blocking channels or pipes if you are using a looping structure waiting for the data. Non-blocking channels consume more resources than the blocking channels.

Consider the OpenCL code for a parallel application in which each work-item is responsible for computing the accumulation of four elements in an array:

__kernel void example ( __global const int * restrict x,
                        __global int * restrict sum ) {
   int accum = 0;

   for (size_t i = 0; i < 4; i++) {
      accum += x[i + get_global_id(0) * 4];
   }

   sum[get_global_id(0)] = accum;
}

Notice the following three main operations that occur in this kernel:

• Load operations from input x
• Accumulation
• Store operations to output sum

The offline compiler arranges these operations in a pipeline according to the data flow semantics of the OpenCL kernel code. For example, the offline compiler implements loops by forwarding the results from the end of the pipeline to the top of the pipeline, depending on the loop exit condition.

The OpenCL kernel performs one loop iteration of each work-item per clock cycle. With sufficient hardware resources, you can increase kernel performance by unrolling the loop, which decreases the number of iterations that the kernel executes. To unroll a loop, add a #pragma unroll directive to the main loop, as shown in the code example below. Keep in mind loop unrolling significantly changes the structure of the compute unit that the offline compiler creates.

__kernel void example ( __global const int * restrict x,
                        __global int * restrict sum ) {
  int accum = 0;

  #pragma unroll
  for (size_t i = 0; i < 4; i++) {
    accum += x[i + get_global_id(0) * 4];
  }

  sum[get_global_id(0)] = accum;
}

In this example, the #pragma unroll directive causes the offline compiler to unroll the four iterations of the loop completely. To accomplish the unrolling, the offline compiler expands the pipeline by tripling the number of addition operations and loading four times more data. With the removal of the loop, the compute unit assumes a feed-forward structure. As a result, the compute unit can store the sum elements every clock cycle after the completion of the initial load operations and additions. The offline compiler further optimizes this kernel by coalescing the four load operations so that the compute unit can load all the necessary input data to calculate a result in one load operation.

Unrolling the loop and coalescing the load operations from global memory allow the hardware implementation of the kernel to perform more operations per clock cycle. In general, the methods you use to improve the performance of your OpenCL kernels should achieve the following results:

• Increase the number of parallel operations
• Increase the memory bandwidth of the implementation
• Increase the number of operations per clock cycle that the kernels can perform in hardware

The offline compiler might not be able to unroll a loop completely under the following circumstances:

• You specify complete unrolling of a data-dependent loop with a very large number of iterations. Consequently, the hardware implementation of your kernel might not fit into the FPGA.
• You specify complete unrolling and the loop bounds are not constants.
• The loop consists of complex control flows (for example, a loop containing complex array indexes or exit conditions that are unknown at compilation time).

For the last two cases listed above, the offline compiler issues the following warning:

Full unrolling of the loop is requested but the loop bounds cannot be determined. The loop is not unrolled.

To enable loop unrolling in these situations, specify the #pragma unroll <N> directive, where <N> is the unroll factor. The unroll factor limits the number of iterations that the offline compiler unrolls. For example, to prevent a loop in your kernel from unrolling, add the directive #pragma unroll 1 to that loop.

Refer to Good Design Practices for Single Work-Item Kernel for tips on constructing well-structured loops.

The OpenCL^® standard does not support fixed-point representation; you must implement fixed-point representations using integer data types. Hardware developers commonly achieve hardware savings by using fixed-point data representations and only retain a data resolution required for performing calculations. You must use an 8, 16, 32, or 64-bit scalar data type because the OpenCL standard supports only these data resolutions. However, you can incorporate the appropriate masking operations in your source code so that the hardware compilation tools can perform optimizations to conserve hardware resources.

For example, if an algorithm uses a fixed-point representation of 17-bit data, you must use a 32-bit data type to store the value. If you then direct the Intel^® FPGA SDK for OpenCL™ Offline Compiler to add two 17-bit fixed-point values together, the offline compiler must create extra hardware to handle the addition of the excess upper 15 bits. To avoid having this additional hardware, you can use static bit masks to direct the hardware compilation tools to disregard the unnecessary bits during hardware compilation. The code below implements this masking operation:

__kernel fixed_point_add (__global const unsigned int * restrict a,
                          __global const unsigned int * restrict b,
                          __global unsigned int * restrict result)
{
	   size_t gid = get_global_id(0);

   	unsigned int temp;
   	temp = 0x3_FFFF & ((0x1_FFFF & a[gid]) + ((0x1_FFFF & b[gid]));

   	result[gid] = temp & 0x3_FFFF;
}

In this code example, the upper 15 bits of inputs a and b are masked away and added together. Because the result of adding two 17-bit values cannot exceed an 18-bit resolution, the offline compiler applies an additional mask to mask away the upper 14 bits of the result. The final hardware implementation is a 17-bit addition as opposed to a full 32-bit addition. The logic savings in this example are relatively minor compared to the sheer number of hardware resources available in the FPGA. However, these small savings, if applied often, can accumulate into a larger hardware saving across the entire FPGA.

To set up aligned memory allocations, add the following source code to your host program:

• For Windows:

  #define AOCL_ALIGNMENT 64
  #include <malloc.h>
  void *ptr = _aligned_malloc (size, AOCL_ALIGNMENT);

  To free up an aligned memory block, include the function call _aligned_free(ptr);

• For Linux:

  #define AOCL_ALIGNMENT 64
  #include <stdlib.h>
  void *ptr = NULL;
  posix_memalign (&ptr, AOCL_ALIGNMENT, size);

  To free up an aligned memory block, include the function call free(ptr);

typedef struct {
	char r,g,b,alpha;
} Pixel_s;

typedef union {
     Pixel_s p;
     int not_used;
} Pixel;

typedef struct {
       char r,g,b,alpha;
} __attribute__((aligned(4))) Pixel;

The offline compiler conforms with the ISO C standard that requires the alignment of a struct to satisfy all of the following criteria:

• The alignment must be an integer multiple of the lowest common multiple between the alignments of all struct members.
• The alignment must be a power of two.

You may set the struct alignment by including the aligned(N) attribute in your kernel code. Without an aligned attribute, the offline compiler determines the alignment of each struct in an array of struct based on the size of the struct. Consider the following example:

__kernel void test (struct mystruct* A,
                    struct mystruct* B)
{
    A[get_global_id(0)] = B[get_global_id(0)];
}

If the size of mystruct is 101 bytes, each load or store access is 1-byte aligned. If the size of mystruct is 128 bytes, each load or store access is 128-byte aligned, which generates the most efficient hardware.

When the struct fields are not aligned within the struct, the offline compiler inserts padding to align them. Inserting padding between struct fields affects hardware efficiency in the following manner:

• Increases the size of the struct
• Might affect the alignment

To prevent the offline compiler from inserting padding, include the packed attribute in your kernel code. The aforementioned ISO C standard applies when determining the alignment of a packed or unpacked struct. Consider the following example:

struct mystruct1
{
    char a;
    int b;
};

The size of mystruct1 is 8 bytes. Therefore, the struct is 8-byte aligned, resulting in efficient accesses in the kernel. Now consider another example:

struct mystruct2
{
    char a;
    int b;
    int c;
};

The size of mystruct2 is 12 bytes and the struct is 4-byte aligned. Because the struct fields are padded and the struct is unaligned, accesses in the kernel are inefficient.

Following is an example of a struct that includes the packed attribute:

struct __attribute__((packed)) mystruct3
{
    char a;
    int b;
    int c;
};

The size of mystruct3 is 16 bytes. Because mystruct3 is aligned and there is no padding between struct fields, accesses in this kernel are more efficient than accesses in mystruct3.

To include both the aligned(N) and packed attributes in a struct, consider the following example:

struct __attribute__((packed)) __attribute__((aligned(16))) mystruct5
{
    char a;
    int b;
    int c;
};

The size of mystruct5 is 9 bytes. Because of the aligned(16) attribute, the struct is stored at 16-byte aligned addresses in an array. Because mystruct5 is 16-byte aligned and has no padding, accesses in this kernel is efficient.

For more information on struct alignment and the aligned(N) and packed attributes, refer to the following documents:

• Section 6.11.1 of the OpenCL Specification version 1.2
• Disabling Insertion of Data Structure Padding section of the Intel^® FPGA SDK for OpenCL™ Programming Guide
• Specifying the Alignment of a Struct section of the Intel^® FPGA SDK for OpenCL™ Programming Guide

The following code example illustrates a scenario where you should update a vector element:

__kernel void update (__global const float4 * restrict in,
                      __global const float4 * restrict out)
{
	  size_t gid = get_global_id(0);

	  out[gid].x = process(in[gid].x);
	  out[gid].y = process(in[gid].y);
	  out[gid].z = process(in[gid].z);
	  out[gid].w = 0; //Update w even if that variable is not required.
}

The restrict keyword informs the offline compiler that the pointer does not alias other pointers. For example, if your kernel has two pointers to global memory, A and B, that never overlap each other, declare the kernel in the following manner:

__kernel void myKernel (__global int * restrict A,
                        __global int * restrict B)

The following functions are expensive:

• Integer division and modulo (remainder) operators
• Most floating-point operators except addition, multiplication, absolute value, and comparison
  Note: For more information on optimizing floating-point operations, refer to the Optimize Floating-Point Operations section.
• Atomic functions

In contrast, inexpensive functions have minimal effects on kernel performance, and their implementation consumes minimal hardware.

The following functions are inexpensive:

• Binary logic operations such as AND, NAND, OR, NOR, XOR, and XNOR
• Logical operations with one constant argument
• Shift by constant
• Integer multiplication and division by a constant that is a power of two

If an expensive function produces a new piece of data for every work-item in a work-group, it is beneficial to code it in a kernel. On the contrary, the code example below shows a case of an expensive floating-point operation (division) executed by every work-item in the NDRange:

__kernel void myKernel (__global const float * restrict a,
                        __global float * restrict b,
                        const float c, const float d)
{
   size_t gid = get_global_id(0);
   
   //inefficient since each work-item must calculate c divided by d
   b[gid] = a[gid] * (c / d); 
}

The result of this calculation is always the same. To avoid this redundant and hardware resource-intensive operation, perform the calculation in the host application and then pass the result to the kernel as an argument for all work-items in the NDRange to use. The modified code is shown below:

__kernel void myKernel (__global const float * restrict a,
                        __global float * restrict b,
                        const float c_divided_by_d)
{
   size_t gid = get_global_id(0);

   /*host calculates c divided by d once and passes it into  
   kernel to avoid redundant expensive calculations*/   
   b[gid] = a[gid] * c_divided_by_d;  						
}

The Intel^® FPGA SDK for OpenCL™ Offline Compiler consolidates operations that are not work-item-dependent across the entire NDRange into a single operation. It then shares the result across all work-items. In the first code example, the offline compiler creates a single divider block shared by all work-items because division of c by d remains constant across all work-items. This optimization helps minimize the amount of redundant hardware. However, the implementation of an integer division requires a significant amount of hardware resources. Therefore, it is beneficial to off-load the division operation to the host processor and then pass the result as an argument to the kernel to conserve hardware resources.

For example, the following code fragment illustrates branching that involves work-item ID such as get_global_id or get_local_id:

for (size_t i = 0; i < get_global_id(0); i++)
{
	  // statements
}

Consider the following OpenCL kernel program:

__kernel void add (__global int * a,
                   __global int * b, 
                   __global int * c)
{
    int gid = get_global_id(0);
    c[gid] = a[gid]+b[gid];
}

As shown in the figure below, the Profiler instruments and connects performance counters in a daisy chain throughout the pipeline generated for the kernel program. The host then reads the data collected by these counters. For example, in PCI Express^® (PCIe^®)-based systems, the host reads the data via the PCIe control register access (CRA) or control and status register (CSR) port.

Work-item execution stalls might occur at various stages of an Intel^® FPGA SDK for OpenCL™ pipeline. Applications with large amounts of memory accesses or load and store operations might stall frequently to enable the completion of memory transfers. The Profiler helps identify the load and store operations or channel accesses that cause the majority of stalls within a kernel pipeline.

For usage information on the Intel^® FPGA Dynamic Profiler for OpenCL™ , refer to the Profiling Your OpenCL Kernel section of the Intel^® FPGA SDK for OpenCL™ Programming Guide.

• Include the -profile Intel^® FPGA SDK for OpenCL™ Offline Compiler command option in your aoc command during development to insert performance counters into your kernel.
• Periodically check kernel f_max and performance without using the Profiler.
• Run the host application from a local folder to reduce profiler overhead. Avoid running your host from a remote or NAS folder.
• Ensure that kernel runtime is longer than 20 ms; otherwise, the Profiler's overhead takes over.
• Understand how all the load and store operations and channels are connected in the data flow.

For example, if a kernel has a read channel call to an Ethernet I/O and the Profiler identifies a stall, it implies that the write channel is not writing data to the Ethernet I/O at the same rate as the read rate of the kernel.

For kernel-to-kernel channels, stalls occur if there is an imbalance between the read and write sides of the channel, or if the read and write kernels are not running concurrently.

For example, if the kernel that reads is not launched concurrently with the kernel that writes, or if the read operations occur much slower than the write operations, the Profiler identifies a stall for the write_channel_intel call in the write kernel.

In the Source Code tab of the Intel^® FPGA Dynamic Profiler for OpenCL™ GUI, the tool tip on the Occupancy% column might specify an Activity percentage. Activity differs from occupancy in that activity relates to predication, as explained below.

Each LSU has a predicate signal besides the ivalid signal. The ivalid signal indicates that upstream logic is providing valid data to the LSU. The predicate signal indicates that the LSU should act on the data that it receives. A work-item or loop iteration can occupy a memory instruction even if it is predicated. If the branch statements do not contain loops, the offline compiler converts the branches to minimize control flow, which leads to more efficient hardware. As part of the conversion, memory and channel instructions must be predicated and the output results much be selected through multiplexer logic.

Consider the following code example:

int addr = compute_address();
int x = 0;
if (some_rare_condition)
    x = src[addr];

The offline compiler will modify the code as follows:

int addr = compute_address();
int x = 0;
x = src[addr] if some_rare_condition;

In this case, src[] receives a valid address every clock cycle. Assuming src[] itself does not generate stalls into the pipeline, the ivalid signal for src[] will be high most of the time. In actuality, src[] only performs loading if the predicate signal some_rare_condition is true. Therefore, for this load operation, occupancy will be high but activity will be low.

Because activity percentages available in the tool tips do not account for predicated accesses, you can identify predicated instructions based on low activity percentages. Despite having low activity percentages, these instructions might have high occupancies.

In the Source Code tab of the Intel^® FPGA Dynamic Profiler for OpenCL™ GUI, the tool tip on the Attributes column might specify a Cache Hit rate. For some global load units, the Intel^® FPGA SDK for OpenCL™ Offline Compiler might instantiate a private cache. In this case, the offline compiler creates an additional hardware counter to measure the effectiveness of this cache. Note that details of this private cache is available in the HTML area report.

Memory instructions stall often whenever bandwidth usage is inefficient or if a large amount of data transfer is necessary during the execution of your application. Inefficient memory accesses lead to suboptimal bandwidth utilization. In such cases, analyze your kernel memory accesses for possible improvements.

Channel instructions stall whenever there is a strong imbalance between read and write accesses to the channel. Imbalances might be caused by channel reads or writes operating at different rates.

For example, if you find that the stall percentage of a write channel call is high, check to see if the occupancy and activity of the read channel call are low. If they are, the performing speed of the kernel controlling the read channel call is too slow for the kernel controlling the write channel call, leading to a performance bottleneck.

If a memory or channel access is causing high percentage pipeline stalls, the line in the source code that instructs the memory or channel is highlighted in red. A stall percentage of 20% or higher results in a high stall identification. The higher the stall percentage, the darker the red highlight will be. To easily traverse through high stall percentage values, right and left arrows can be found at the bottom right corner of the Source Code tab.

Figure 83. High Stall Percentage Identification

Consider the following code example:

__kernel void proc (__global int * a, ...) {
  for (int i = 0; i < N; i++) {
    for (int j = 0; j < 1000; j++) {
      write_channel_intel (c0, data0);
    }
    for (int k = 0; k < 3; k++) {
      write_channel_intel (c1, data1);
    }
  }
}    

Assuming all the loops are pipelined, the first inner loop with a trip count of 1000 is the critical loop. The second inner loop with a trip count of three will be executed infrequently. As a result, you can expect that the occupancy and activity percentages for channel c0 are high and for channel c1 are low.

Also, occupancy percentage might be low if you define a small work-group size, the kernel might not receive sufficient work-items. This is problematic because the pipeline is empty generally for the duration of kernel execution, which leads to poor performance.

Review your memory accesses to see if you can rewrite them such that accesses to memory sites address consecutive memory regions.

Usually, the sum of the stall and occupancy percentages approximately equals 100%. If a load and store operation or channel has a high stall percentage, it means that the load and store operation or channel has the ability to execute every cycle but is generating stalls.

Solutions for stalling global load and store operations:

• Use local memory to cache data.
• Reduce the number of times you read the data.
• Improve global memory accesses.
  • Change the access pattern for more global-memory-friendly addressing (for example, change from stride accessing to sequential accessing).
  • Compile your kernel with the -no-interleaving=default Intel^® FPGA SDK for OpenCL™ Offline Compiler command option, and separate the read and write buffers into different DDR banks.
  • Have fewer but wider global memory accesses.
• Acquire an accelerator board that has more bandwidth (for example, a board with three DDRs instead of 2 DDRs).

Solution for stalling local load and store operations:

• Review the HTML area report to verify the local memory configuration and modify the configuration to make it stall-free.

Solutions for stalling channels:

• Fix stalls on the other side of the channel. For example, if channel read stalls, it means that the writer to the channel is not writing data into the channel fast enough and needs to be adjusted.
• If there are channel loops in your design, specify the channel depth.

Figure 84. Example OpenCL Kernel and Profiler Analysis

In this example, dst[] is executed once every 20 iterations of the FACTOR2 loop and once every four iterations of the FACTOR1 loop. Therefore, FACTOR2 loop is the source of the bottleneck.

Solutions for resolving loop bottlenecks:

• Unroll the FACTOR1 and FACTOR2 loops evenly. Simply unrolling FACTOR1 loop further will not resolve the bottleneck
• Vectorize your kernel to allow multiple work-items to execute during each loop iteration

Figure 85. Example OpenCL Kernel and Profiler Analysis

In this example, the accelerator board can provide a bandwidth of 25600 megabytes per second (MB/s). However, the vector_add kernel is requesting (2 reads + 1 write) x 4 bytes x 294 MHz = 12 bytes/cycle x 294 MHz = 3528 MB/s, which is 14% of the available bandwidth. To increase the bandwidth, increase the number of tasks performed in each clock cycle.

Solutions for low bandwidth:

• Automatically or manually vectorize the kernel to make wider requests
• Unroll the innermost loop to make more requests per clock cycle
• Delegate some of the tasks to another kernel

Kernels that are marked with the autorun attribute are referred to as autorun kernels. Hence, an autorun kernel starts executing automatically before other kernels are launched explicitly by the host, and restarts automatically on completion. For more information about the autorun attribute, refer to Omit Communication Hardware between the Host and the Kernel topic.

Since autorun kernels never complete, you must call the host library call clGetProfileDataDeviceIntelFPGA to capture the autorun profiler data. You can instruct the host application to make this call at any point during execution.

aocl report <filename>.aocx profile.mon <filename>.source

• The Profiler can only extract one set of profile data from a kernel while it is running.

  If the Profiler collects the profile data after kernel execution completes, you can call the host API to generate the profile.mon file multiple times.

  For more information on how to collect profile data during kernel execution, refer to the Collecting Profile Data During Kernel Execution section of the Intel^® FPGA SDK for OpenCL™ Programming Guide.

• Profile data is not persistent across OpenCL programs or multiple devices.

  You can request profile data from a single OpenCL program and on a single device only. If your host swaps a new kernel program in and out of the FPGA, the Profiler will not save the profile data.

• Instrumenting the Verilog code with performance counters increases hardware resource utilization (that is, FPGA area usage) and typically decreases performance.

  For information on instrumenting the Verilog code with performance counters, refer to the Instrumenting the Kernel Pipeline with Performance Counters section of the Intel^® FPGA SDK for OpenCL™ Programming Guide.

The following flowchart outlines the approach you can take to iterate on your design and optimize your single work-item kernel. For usage information on the Intel^® FPGA SDK for OpenCL™ Emulator and the Profiler, refer to the Emulating and Debugging Your OpenCL Kernel and Profiling Your OpenCL Kernel sections of the Intel^® FPGA SDK for OpenCL™ Programming Guide, respectively. For information on the Intel^® FPGA Dynamic Profiler for OpenCL™ GUI and profiling information, refer to the Profile Your Kernel to Identify Performance Bottlenecks section.

Intel^® recommends the following optimization options to address single work-item kernel loop-carried dependencies, in order of applicability: removal, relaxation, simplification, and transfer to local memory.

The offline compiler assumes a default work-group size for your kernel depending on certain constraints imposed during compilation time and runtime .

The offline compiler imposes the following constraints at compilation time:

• If you specify a value for the reqd_work_group_size attribute, the work-group size must match this value.
• If you specify a value for the max_work_group_size attribute, the work-group size must not exceed this value.
• If you do not specify values for reqd_work_group_size and max_work_group_size, and the kernel contains a barrier, the offline compiler defaults to a maximum work-group size of 256 work-items.
• If you do not specify values for both attributes and the kernel does not contain any barrier, the offline compiler does not impose any constraint on the work-group size at compilation time.

The OpenCL™ standard imposes the following constraints at runtime:

• The work-group size in each dimension must divide evenly into the requested NDRange size in each dimension.
• The work-group size must not exceed the device constraints specified by the CL_DEVICE_MAX_WORK_GROUP_SIZE and CL_DEVICE_MAX_WORK_ITEM_SIZES queries to the clGetDeviceInfo API call.

If you do not specify values for both the reqd_work_group_size and max_work_group_size attributes, the runtime determines a default work-group size as follows:

• If the kernel contains a barrier or refers to the local work-item ID, or if you use the clGetKernelWorkGroupInfo and clGetDeviceInfo API calls in your host code to query the work-group size, the runtime defaults the work-group size to one work-item.
• If the kernel does not contain a barrier or refer to the local work-item ID, or if your host code does not query the work-group size, the default work-group size is the global NDRange size.

When queuing an NDRange kernel (that is, not a single work-item kernel), specify an explicit work-group size under the following conditions:

• If your kernel uses memory barriers, local memory, or local work-item IDs.
• If your host program queries the work-group size.

If your kernel uses memory barriers, perform one of the following tasks to minimize hardware resources:

• Specify a value for the reqd_work_group_size attribute.
• Assign to the max_work_group_size attribute the smallest work-group size that accommodates all your runtime work-group size requests.

Specifying a smaller work-group size than the default at runtime might lead to excessive hardware consumption. Therefore, if you require a work-group size other than the default, specify the max_work_group_size attribute to set a maximum work-group size. If the work-group size remains constant through all kernel invocations, specify a required work-group size by including the reqd_work_group_size attribute. The reqd_work_group_size attribute instructs the offline compiler to allocate exactly the correct amount of hardware to manage the number of work-items per work-group you specify. This allocation results in hardware resource savings and improved efficiency in the implementation of kernel compute units. By specifying the reqd_work_group_size attribute, you also prevent the offline compiler from implementing additional hardware to support work-groups of unknown sizes.

For example, the code fragment below assigns a fixed work-group size of 64 work-items to a kernel:

__attribute__((reqd_work_group_size(64,1,1)))
__kernel void sum (__global const float * restrict a,
                   __global const float * restrict b,
                   __global float * restrict answer)
{
  size_t gid = get_global_id(0);

  answer[gid] = a[gid] + b[gid];
}

Include the num_simd_work_items attribute in your kernel code to direct the offline compiler to perform more additions per work-item without modifying the body of the kernel. The following code fragment applies a vectorization factor of four to the original kernel code:

__attribute__((num_simd_work_items(4)))
__attribute__((reqd_work_group_size(64,1,1)))
__kernel void sum (__global const float * restrict a,
                   __global const float * restrict b,
                   __global float * restrict answer)
{
   size_t gid = get_global_id(0);

   answer[gid] = a[gid] + b[gid];
}

To use the num_simd_work_items attribute, you must also specify a required work-group size of the kernel using the reqd_work_group_size attribute. The work-group size you specify for reqd_work_group_size must be divisible by the value you assign to num_simd_work_items. In the code example above, the kernel has a fixed work-group size of 64 work-items. Within each work-group, the work-items are distributed evenly among the four SIMD vector lanes. After the offline compiler implements the four SIMD vector lanes, each work-item now performs four times more work.

The offline compiler vectorizes the code and might coalesce memory accesses. You do not need to change any kernel code or host code because the offline compiler applies these optimizations automatically.

You can vectorize your kernel code manually, but you must adjust the NDRange in your host application to reflect the amount of vectorization you implement. The following example shows the changes in the code when you duplicate operations in the kernel manually:

__kernel void sum (__global const float * restrict a,
                   __global const float * restrict b,
                   __global float * restrict answer)
{
   size_t gid = get_global_id(0);

   answer[gid * 4 + 0] = a[gid * 4 + 0] + b[gid * 4 + 0];
   answer[gid * 4 + 1] = a[gid * 4 + 1] + b[gid * 4 + 1];
   answer[gid * 4 + 2] = a[gid * 4 + 2] + b[gid * 4 + 2];
   answer[gid * 4 + 3] = a[gid * 4 + 3] + b[gid * 4 + 3];
}

In this form, the kernel loads four elements from arrays a and b, calculates the sums, and stores the results into the array answer. Because the FPGA pipeline loads and stores data to neighboring locations in memory, you can manually direct the offline compiler to coalesce each group of four load and store operations.

To increase overall kernel throughput, the hardware scheduler in the FPGA dispatches work-groups to additional available compute units. A compute unit is available for work-group assignments as long as it has not reached its full capacity.

Assume each work-group takes the same amount of time to complete its execution. If the offline compiler implements two compute units, each compute unit executes half of the work-groups. Because the hardware scheduler dispatches the work-groups, you do not need to manage this process in your own code.

The offline compiler does not automatically determine the optimal number of compute units for a kernel. To increase the number of compute units for your kernel implementation, you must specify the number of compute units that the offline compiler should create using the num_compute_units attribute, as shown in the code sample below.

__attribute__((num_compute_units(2)))
__kernel void sum (__global const float * restrict a,
                   __global const float * restrict b,
                   __global float * restrict answer)
{
    size_t gid = get_global_id(0);

    answer[gid] = a[gid] + b[gid];
}

Increasing the number of compute units achieves higher throughput. However, as shown in the figure below, you do so at the expense of increasing global memory bandwidth among the compute units. You also increase hardware resource utilization.

Both the num_compute_units and num_simd_work_items attributes increase throughput by increasing the amount of hardware that the Intel^® FPGA SDK for OpenCL™ Offline Compiler uses to implement your kernel. The num_compute_units attribute modifies the number of compute units to which work-groups can be scheduled, which also modifies the number of times a kernel accesses global memory. In contrast, the num_simd_work_items attribute modifies the amount of work a compute unit can perform in parallel on a single work-group. The num_simd_work_items attribute duplicates only the datapath of the compute unit by sharing the control logic across each SIMD vector lane.

Generally, using the num_simd_work_items attribute leads to more efficient hardware than using the num_compute_units attribute to achieve the same goal. The num_simd_work_items attribute also allows the offline compiler to coalesce your memory accesses.

Multiple compute units competing for global memory might lead to undesired memory access patterns. You can alter the undesired memory access pattern by introducing the num_simd_work_items attribute instead of the num_compute_units attribute. In addition, the num_simd_work_items attribute potentially offers the same computational throughput as the equivalent kernel compute unit duplication that the num_compute_units attribute offers.

You cannot implement the num_simd_work_items attribute in your kernel under the following circumstances:

• The value you specify for num_simd_work_items is not 2, 4, 8 or 16.
• The value of reqd_work_group_size is not divisible by num_simd_work_items.

  For example, the following declaration is incorrect because 50 is not divisible by 4:

  __attribute__((num_simd_work_items(4)))
  __attribute__((reqd_work_group_size(50,0,0)))

• Kernels with complex control flows. You cannot vectorize lines of code within a kernel in which different work-items follow different control paths (for example, the control paths depend on get_global_ID or get_local_ID).

During kernel compilation, the offline compiler issues messages informing you whether the implementation of vectorization optimizations is successful. Kernel vectorization is successful if the reported vectorization factor matches the value you specify for the num_simd_work_items attribute.

Consider a case where a kernel with a num_simd_work_items attribute set to 16 does not fit in the FPGA. The kernel might fit if you modify it by duplicating a narrower SIMD kernel compute unit. Determining the optimal balance between the number of compute units and the SIMD width might require some experimentation. For example, duplicating a four lane-wide SIMD kernel compute unit three times might achieve better throughput than duplicating an eight lane-wide SIMD kernel compute unit twice.

The following example code shows how you can combine the num_compute_units and num_simd_work_items attributes in your OpenCL™ code:

__attribute__((num_simd_work_items(4)))
__attribute__((num_compute_units(3)))
__attribute__((reqd_work_group_size(8,8,1)))
__kernel void matrixMult(__global float * restrict C,
                         __global float * restrict A,
. . .

The figure below illustrates the data flow of the kernel described above. The num_compute_units implements three replicated compute units. The num_simd_work_items implements four SIMD vector lanes.

An interconnect topology connects shared global, constant, and local memory systems to their underlying memory. Interconnect includes access arbitration to memory ports.

Memory accesses compete for shared memory resources (that is, global, local, and constant memories). If your OpenCL kernel performs a large number of memory accesses, the Intel^® FPGA SDK for OpenCL™ Offline Compiler must generate complex arbitration logic to handle the memory access requests. The complex arbitration logic might cause a drop in the maximum operating frequency (f_max), which degrades kernel performance.

The following sections discuss memory access optimizations in detail. In summary, minimizing global memory accesses is beneficial for the following reasons:

• Typically, increases in OpenCL kernel performance lead to increases in global memory bandwidth requirements.
• The maximum global memory bandwidth is much smaller than the maximum local memory bandwidth.
• The maximum computational bandwidth of the FPGA is much larger than the global memory bandwidth.
  Attention: Use local, private or constant memory whenever possible to increase the memory bandwidth of the kernel.

In most circumstances, the default burst-interleaved configuration leads to the best load balancing between the memory banks. However, in some cases, you might want to partition the banks manually as two non-interleaved (and contiguous) memory regions to achieve better load balancing.

The figure below illustrates the differences in memory mapping patterns between burst-interleaved and non-interleaved memory partitions.

Consider the following code example:

__kernel void sum ( __global const float * restrict a,
                    __global const float * restrict b,
                    __global float * restrict c )
{
	size_t gid = get_global_id(0);

	c[gid] = a[gid] + b[gid];
}

The load operation from array a uses an index that is a direct function of the work-item global ID. By basing the array index on the work-item global ID, the offline compiler can direct contiguous load operations. These load operations retrieve the data sequentially from the input array, and sends the read data to the pipeline as required. Contiguous store operations then store elements of the result that exits the computation pipeline in sequential locations within global memory.

The following figure illustrates an example of the contiguous memory access optimization:

Contiguous load and store operations improve memory access efficiency because they lead to increased access speeds and reduced hardware resource needs. The data travels in and out of the computational portion of the pipeline concurrently, allowing overlaps between computation and memory accesses. If possible, use work-item IDs that index consecutive memory locations for load and store operations that access global memory. Sequential accesses to global memory increase memory efficiency because they provide an ideal access pattern.

If your FPGA board offers heterogeneous global memory types, keep in mind that they handle different memory accesses with varying efficiencies.

For example:

• Use DDR SDRAM for long sequential accesses.
• Use QDR SDRAM for random accesses.
• Use on-chip RAM for random low latency accesses.

For more information on how to allocate buffers in global memory and how to modify your host application to leverage heterogeneous buffers, refer to the Specifying Buffer Location in Global Memory and Allocating OpenCL Buffer for Manual Partitioning of Global Memory sections of the Intel^® FPGA SDK for OpenCL™ Programming Guide.

• To improve kernel f_max on board support packages with four or more banks of global memory, use the -global-ring option in your aoc command. For more information about the -global-ring option, refer to Forcing Ring Interconnect for Global Memory (-global-ring).
• In conjunction with the -global-ring option, you can improve the write throughput to global memory by using the -duplicate-ring option in your aoc command. For more information about the -duplicate-ring option, refer to Duplicating the Store Ring to Improve the Write Throughput to Global Memory (-duplicate-ring).
• If global memory throughput is still insufficient, to maximize throughput for both reads and writes, Intel recommends using the -no-interleaving option in conjunction with -global-ring and -duplicate-ring options. For more information about the -no-interleaving option, refer to Disabling Burst-Interleaving of Global Memory (-no-interleaving=<global_memory_type>).

To minimize global memory accesses, you must first preload data from a group of computations from global memory to constant, local, or private memory. You perform the kernel computations on the preloaded data, and then write the results back to global memory.

By default, the constant cache size is 16 kB. You can specify the constant cache size by including the -const-cache-bytes=<N> option in your aoc command, where <N> is the constant cache size in bytes.

Unlike global memory accesses that have extra hardware for tolerating long memory latencies, the constant cache suffers large performance penalties for cache misses. If the __constant arguments in your OpenCL™ kernel code cannot fit in the cache, you might achieve better performance with __global const arguments instead. If the host application writes to constant memory that is already loaded into the constant cache, the cached data is discarded (that is, invalidated) from the constant cache.

For more information on the -const-cache-bytes=<N> option, refer to the Configuring Constant Memory Cache Size (-const-cache-bytes=<N>) section of the Intel^® FPGA SDK for OpenCL™ Programming Guide.

For more information on the implementation of private memory using registers, refer to the Inferring a Register section of the Intel^® FPGA SDK for OpenCL™ Programming Guide.

The following code example depicts an 8 x 4 local memory system that is implemented in a single bank. As a result, no two elements in the system can be accessed in parallel.

local int lmem[8][4];

#pragma unroll
for(int i = 0; i<4; i+=2) {
    lmem[i][x] = …; 
}

To improve performance, you can add numbanks(N) and bankwidth(M) in your code to define the number of memory banks and the bank widths in bytes. The following code implements eight memory banks, each 16-bytes wide. This memory bank configuration enables parallel memory accesses down the 8 x 4 array.

local int __attribute__((numbanks(8), 
                        bankwidth(16)))
                        lmem[8][4]; 
#pragma unroll
for (int i = 0; i < 4; i+=2) {
    lmem[i][x & 0x3] = …; 
}

By specifying different values for the numbanks(N) and bankwidth(M) kernel attributes, you can change the parallel access pattern. The following code implements four memory banks, each 4-bytes wide. This memory bank configuration enables parallel memory accesses across the 8 x 4 array.

local int __attribute__((numbanks(4), 
                        bankwidth(4)))
                        lmem[8][4]; 

#pragma unroll
for (int i = 0; i < 4; i+=2) {
    lmem[x][i] = …; 
}

The following code examples illustrate how the bank geometry changes based on the values you assign to numbanks and bankwidth.

local int
__attribute__((numbanks(2),
               bankwidth(16)))
               lmem[2][4]; 

local int 
__attribute__((numbanks(2), 
               bankwidth(8))) 
               lmem[2][4]; 

local int 
__attribute__((numbanks(2), 
               bankwidth(4))) 
               lmem[2][4]; 

local int 
__attribute__((numbanks(4), 
               bankwidth(8))) 
               lmem[2][4]; 

local int
__attribute__((numbanks(4),
               bankwidth(4)))
               lmem[2][4]; 

Intel^® 's M20K memory blocks have two physical ports. The number of logical ports that are available in each M20K block depends on the degree of pumping. Pumping is a measure of the clock frequency of the M20K blocks relative to the rest of the design.

Consider the following code example where the singlepump attribute is applied to a local memory system, lmem, which has three read accesses and one write access. The singlepump attribute indicates that the M20K blocks runs at the same frequency as the rest of the design.

__kernel void three_copies(int raddr, int waddr) {
    int __attribute__((memory,
                       numbanks(1),
                       singlepump,
                       max_replicates(3)))
                       lmem[16];

    lmem[waddr] = lmem[raddr] + lmem[raddr + 1] + lmem[raddr + 2];
    // do something with lmem
}

If you include the doublepump kernel attribute in your local variable declaration, you specify that the M20K memory blocks runs at double the frequency as the rest of the design.

__kernel void three_copies(int raddr, int waddr) {
    int __attribute__((memory,
                       numbanks(1),
                       doublepump))
                       lmem[16];

    lmem[waddr] = lmem[raddr] + lmem[raddr + 1] + lmem[raddr + 2];
    // do something with lmem
}

Each double-pumped M20K block has four logical ports available. As such, there only needs to be one M20K block to implement three read ports and one write port in lmem.

The figure below shows a common case where kernel performance might benefit from static memory coalescing:

Consider the following vectorized kernel:

__attribute__((num_simd_work_items(4)))
__attribute__((reqd_work_group_size(64,1,1)))
__kernel void sum (__global const float * restrict a,
                   __global const float * restrict b,
                   __global float * restrict answer)
{
   size_t gid = get_global_id(0);

   answer[gid] = a[gid] + b[gid];
}

The OpenCL™ kernel performs four load operations that access consecutive locations in memory. Instead of performing four memory accesses to competing locations, the offline compiler coalesces the four loads into a single wider vector load. This optimization reduces the number of accesses to a memory system and potentially leads to better memory access patterns.

Although the offline compiler performs static memory coalescing automatically when it vectorizes the kernel, you should use wide vector loads and stores in your OpenCL code whenever possible to ensure efficient memory accesses. To implement static memory coalescing manually, you must write your code in such a way that a sequential access pattern can be identified at compilation time. The original kernel code shown in the figure above can benefit from static memory coalescing because all the indexes into buffers a and b increment with offsets that are known at compilation time. In contrast, the following code does not allow static memory coalescing to occur:

__kernel void test (__global float * restrict a,
		          __global float * restrict b,
                    __global float * restrict answer;
		          __global int * restrict offsets)
{
 size_t gid = get_global_id(0);

 answer[gid*4 + 0] = a[gid*4 + 0 + offsets[gid]] + b[gid*4 + 0];
 answer[gid*4 + 1] = a[gid*4 + 1 + offsets[gid]] + b[gid*4 + 1];
 answer[gid*4 + 2] = a[gid*4 + 2 + offsets[gid]] + b[gid*4 + 2];
 answer[gid*4 + 3] = a[gid*4 + 3 + offsets[gid]] + b[gid*4 + 3];
}

The value offsets[gid] is unknown at compilation time. As a result, the offline compiler cannot statically coalesce the read accesses to buffer a.

Optimizing kernel performance generally requires additional FPGA resources. In contrast, area optimization often results in performance decreases. During kernel optimization, Intel^® recommends that you run multiple versions of the kernel on the FPGA board to determine the kernel programming strategy that generates the best size versus performance trade-off.

1. To review the estimated resource usage summary on-screen, compile your kernel by including the -report flag in your aoc command. To review kernel-specific area usage information, refer to the <your_kernel_filename>/reports/report.html file.
2. If possible, perform floating-point computations by compiling your OpenCL kernel with the -fpc or -fp-relaxed option of the aoc command.

For more usage information on the -report, -fp-relaxed and -fpc options, refer to the Displaying Estimated Resource Usage Summary (-report), Relaxing Order of Floating-Point Operations (-fp-relaxed), and Reducing Floating-Point Operations (-fpc) sections of the Intel^® FPGA SDK for OpenCL™ Programming Guide.

For more information on floating-point operations, refer to Optimize Floating-Point Operations.

For example, if your kernel requires one external memory bank, target a board variant that only supports a single external memory bank. Targeting a board with multiple external memory banks increases the area usage of your kernel unnecessarily.

If your Custom Platform does not provide a board variant that meets your needs, consider creating a board variant. Consult the Intel^® FPGA SDK for OpenCL™ Custom Platform Toolkit User Guide for more information.

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 Intel^® FPGA SDK for OpenCL™ 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. However, 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.

1. Introduce floating-point arithmetic operations only when necessary.
2. The Intel^® FPGA SDK for OpenCL™ Offline Compiler defaults floating-point constants to double data type. Add an f designation to the constant to make it a single precision floating-point operation.

   For example, the arithmetic operation sin(1.0) represents a double precision floating-point sine function. The arithmetic operation sin(1.0f) represents a single precision floating-point sine function.

3. If you do not require full precision result for a complex function, compute simpler arithmetic operations to approximate the result. Consider the following example scenarios:
   1. Instead of computing the function pow(x,n) where n is a small value, approximate the result by performing repeated squaring operations because they require much less hardware resources and area.
   2. Ensure you are aware of the original and approximated area usages because in some cases, computing a result via approximation might result in excess area usage. For example, the sqrt function is not resource-intensive. Other than a rough approximation, replacing the sqrt function with arithmetic operations that the host has to compute at runtime might result in larger area usage.
   3. If you work with a small set of input values, consider using a LUT instead.

4. If your kernel performs a complex arithmetic operation with a constant that the offline compiler computes at compilation time (for example, log(PI/2.0)), perform the arithmetic operation on the host instead and pass the result as an argument to the kernel at runtime.

1. Select the most appropriate data type for your application.

   For example, do not define your variable as float if the data type short is sufficient.

2. Ensure that both sides of an arithmetic expression belong to the same data type.

   Consider an example where one side of an arithmetic expression is a floating-point value and the other side is an integer. The mismatched data types cause the Intel^® FPGA SDK for OpenCL™ Offline Compiler to create implicit conversion operators, which can become expensive if they are present in large numbers.

3. Take advantage of padding if it exists in your data structures.

   For example, if you only need float3 data type, which has the same size as float4, you may change the data type to float4 to make use of the extra dimension to carry an unrelated value.

Intel^® assumes that you are familiar with the information presented in the Intel^® FPGA SDK for OpenCL™ Pro Edition Best Practices Guide and know how to apply the general optimization strategies described in the document.

The following topics outline strategies you can implement to reduce channel overhead.

The following example shows a producer kernel and a consumer kernel communicating via channels:

kernel producer(unsigned N) {
   int result;
   for (unsigned int i = 0; i < N; i++) {
      write_channel_intel(Produce(i));
   }
} 
 
kernel consumer(unsigned N) {
   for (unsigned int i = 0; i < N; i++) {
      Consume(i, read_channel_intel());
   }
}

The optimized code below merges the two kernels in the example above into a single kernel, which uses the computation results directly without channel accesses:

kernel fused(unsigned N) {
   for (unsigned int i = 0; i < N; i++) {
      Consume(i, Produce(i));
   }
}

Unoptimized multi-kernel systolic array pseudocode:

// data distribution network over an array of channels

channel int c[ROWS][COLS]; 
channel int d[ROWS][COLS]; 

attribute((num_compute_units(ROWS,COLS)) 
kernel void PE() {
   // get data values from my neighbors 
   while(1){
      x = read_channel_intel(c[ROWS-1][COLS]); 
      y = read_channel_inel(d[ROWS][COLS-1]);
 
 
      // some code that uses x and y 
      ... 
      // send the same data values to the next neighbors 
      write_channel_intel(c[ROWS][COLS], x); 
      write_channel_intel(d[ROWS][COLS], y); 
   }
}

Optimized single-kernel pseudocode:

kernel void allPEs() {
   while(1){
      int c[ROWS], d[COLS]; 
   
      #pragma unroll 
      for (int i = 0; i < ROWS; i++) 
         #pragma unroll 
         for (int j = 0; j < COLS; j++) {
            PE(c[i], d[j]); 
         } 
      }
   } 
}

Optimized pseudocode with the __fpga_reg() function:

kernel void allPEs() {
   int c[ROWS], d[COLS]; 
   
   while(1){
      #pragma unroll 
      for (int i = 0; i < ROWS; i++) 
         #pragma unroll 
         for (int j = 0; j < COLS; j++) {
            // compute and store outputs 
            PE(c[i], d[j]); 
            c[i] = __fpga_reg(c[i]); 
            d[j] = __fpga_reg(d[j]); 
         } 
      }
   } 
}

After the offline compiler unrolls the loop, there will be one more register before every PE on both c and d, allowing the Intel^® Quartus^® Prime Pro Edition software to place the PEs apart. You may add more than one register by inserting multiple __fpga_reg() calls in your code. For example, the call __fpga_reg(__fpga_reg(x)) adds two registers on the data path. However, having excessive __fpga_reg() calls in your kernel will increase the design area, and the congestion might result in Fmax degradation.

The example code below has a blocking channel read:

while (cond) {
   val = read_channel_intel (my_ch);
   <do_compute (val)>
}

To switch to a non-blocking channel read that is functionally equivalent to the blocking channel read, modify the code in the following manner:

bool have_data = true;
while (cond) {
   val = read_channel_nb_intel (my_ch, &have_data);
   if (have_data) <do_compute (val)>
} 

For this example, the downside of changing from a blocking channel read to a non-blocking channel read is that the loop control logic becomes more complex. If you transform multiple channel accesses this way, the loop control logic might limit your performance.

To achieve high performance, Intel^® recommends achieving a loop initiation interval (II) of 1. An II value of 1 indicates that a loop is able to start a new iteration of a loop data path every clock cycle. Doing so helps your design consume the available FPGA resources efficiently.

Intel^® has established a new loop control scheme specifically for the Intel^® Stratix^® 10 architecture.

kernel loop_overhead(unsigned N) {
   for (unsigned int i = 0; i < N; i++) {
      for (unsigned int j = 0; j < N; j++) {
         //do work
         //total iterations: i * (j + s)
      }
   }
}

kernel void good_loop(global int * restrict A, 
                      global int * restrict result,
					  unsigned N) {
  
   unsigned int sum = 0;
  
   for (unsigned int i = 0; i < N; i++) {
      sum += A[i];
   }
   *result = sum; 
}

channel unsigned int c0; 

kernel void producer() {
 
   for (unsigned int i = 0; i < 10; i++) {
      write_channel_intel(c0, i); 
   }   
}

kernel void consumer (global int * restrict A, 
                      global int * restrict result, 
                      unsigned N) {
   unsigned int sum = 0;
   for (unsigned int i = 0; 
        i < N && read_channel_intel(c0) != 5; i++) {
      sum += A[i];
   }
   *result = sum;
}

The offline compiler cannot pipeline computations used for loop-carried dependencies. Loops that contain many complex computations limit the amount of retiming optimizations that the compiler can perform because the compiler cannot make any functional changes to the loop path. Even if II=1, the HTML report identifies the Fmax bottleneck. Use this information in conjunction with the information presented in the Fmax analysis report pane to assess the most critical paths in your design.

If a loop-carried dependency contains logic that the offline compiler cannot compute in one clock cycle, one mitigation approach is to lengthen the dependency distance. The dependency distance is the number of loop iterations that occur from when the compiler reads the value to when the next value becomes available. The Loop analysis report within the High Level Design Report identifies the most complex loop dependency.

int i = 0;
int N = 256;
while (!done) {
   i++;
   if (i == N) i = 0;
   <use i for some computation…>
}

Multiple concurrent stores often require more than two ports to implement. Multiple concurrent stores also cause the compiler to create stallable memories where memory accesses must be arbitrated on every clock cycle. To determine if a memory is arbitrated, examine the load and store units in the system or memory viewer of the High Level Design Report. In the viewer, load and store units that are highlighted red are stallable memories. The report will present this information when you hover over each highlighted unit.

The unoptimized code below creates a cached burst-coalesced LSU that consumes more resources and disables other optimizations:

kernel void cached (global int * restrict in,
                    global int * restrict out) {
   int i = get_global_id(0);
   int idx = out[i];
   int cached_value = in[idx]; // Burst-coalesced cached LSU
   out[i] = cached_value;
}

To prevent the caching, mark the memory pointer as volatile, as shown in the optimized code below:

kernel void not_cached (global volatile int * restrict in,
                        global int * restrict out) {
   int i = get_global_id(0);
   int idx = out[i];
   int not_cached_value = in[idx];
   out[i] = not_cached_value;
}

For more information on optimizing load-store units, refer to the Load-Store Units section.

Intel^® also recommends that you leverage on-chip storage to achieve optimal results. Note that compared to non- Intel^® Stratix^® 10 devices, Intel^® Stratix^® 10 has a larger M20K to ALM ratio, allowing you to create larger local memory systems.

The unoptimized code below has a function that receives a pointer from the array in global memory. In this case, the offline compiler modifies the array and then stores it back to memory. Then, in a subsequent iteration of the outer loop, the array is reused.

void cached_array(int N, int M, int BUF_SIZE, 
                  global float2 * global_buf)
{
   for (uint i = 0; i< N; i++) {
      float2 data_array[BUF_SIZE];
      for (uint j= 0; j < M; j++) {
	          
         data_array[i] = global_buf [j]; //Load value
		 
         ... //do work to modify data_array 
		 
         global_buf[j] = data_array[i]; 
      } 
   } 
}

To prevent unnecessary global memory accesses, define a private array in the kernel to declare the array on chip. The function will access the private array instead of the one declared on chip. As a result, the array will receive on-chip local memory storage. Accessing this on-chip local memory does not require accesses to global memory.

Optimized code:

void local_array(int N, int M, int BUF_SIZE, 
                 global float2 * global_buf)
{
   float2 local_buf[BUF_SIZE];
   populate_buf(local_buf, global_buf); 

   for (uint i = 0; i< N; i++) {
      float2	data_array[BUF_SIZE];			
      for (uint j= 0; j < M; j++) {

         data_array[i] = local_buf[j];//Load value 

		 ... //do work to modify data_array

         local_buf[j] = data_array[i]; 

      }
   } 
}

The following constructs or features prevent Intel^® Stratix^® 10 data path optimization:

• NDRange designs with loops
• Stallable LSUs with the exception of burst-coalesced LSUs

  Burst-coalesced LSUs are the default type of LSUs that the offline compiler instantiates. Example of a burst coalesced LSU instantiation:

  kernel void burst_coalesced (global int * restrict in,
                               global int * restrict out){
     int i = get_global_id(0);
     int value = in[i/2];  //Burst-coalesced LSU
     out[i] = value;
  }

  You can view the LSU type of various instructions in the High Level Design Report by hovering over the load or store operation in the Graph Viewer. Refer to the Load-Store Units section for more information on the types of LSUs and how you can influence the compiler on which type of LSUs to instantiate.

• Channels with multiple call sites
• Stallable RTL library calls

  Refer to the Create RTL Modules section for more information.

• Reconvergent control flow in the optimized control flow graph, with the exception of loops that use the new control optimization

  The following pseudocode example of a simple reconvergent control flow shows that the flow of the code goes in one of two paths. The offline compiler implements different control logic for each path. It also implements logic to reconverge the control flow after the two paths are completed.

  while (some_some condition){
     if (some_other_condition){
        for(...){ }
     } else{
        for(...){ }
     }
  }

• Loops that do not use the new loop control scheme

  Refer to the Loop Control Optimization section for more information on what loops are affected by this restriction.

• Basic block structures with the exception of the following:
  • Basic block with only one predecessor, as shown in Figure 98
  • Basic block with exactly two predecessors, where one predecessor is the back-edge of a loop, as shown in Figure 99
  Note: The majority of optimized designs belong to one of the two supported basic block structures. You may review images of these basic blocks in the Graph Viewer of the High Level Design Report.

  The following code example generates the two types of supported basic block structures:

  __attribute__((max_global_work_dim(0)))
  void kernel basic_block(global unsigned int *myvar,
                              unsigned int insize)
  {
     for(int i=0; i < insize; i++){
        myvar[i] += insize;
     }
  }

Figure 98. Basic Block Structure with One PredecessorThe basic block in question refers to gzip.B2, which is outlined in red. Its predecessor is the previous basic block, gzip.B1.

Figure 99. Basic Block Structure with Two PredecessorsThe basic block in question is gzip.B1, which is outlined in red. This block refers to the body of the loop; its predecessors are itself and the previous basic block, gzip.B0. The purple line in gzip.B1 denotes the back end of the loop.

There are stall-free and stallable RTL modules. A stall-free RTL module is a fixed-latency module for which the offline compiler can optimize away stall logic. Refer to the Stall-Free RTL section in the Intel^® FPGA SDK for OpenCL™ Pro Edition Programming Guide for more information.

A stallable RTL module has variable latency and relies on backpressured input and output interfaces to function correctly. Implementing stallable interfaces in Intel^® Stratix^® 10 designs consumes a lot of FPGA resources because of the handshake logic, which limits retiming. Using stallable interfaces in Intel^® Stratix^® 10 designs also disables the data path control optimization scheme.

Intel^® strongly recommends that you use stall-free RTL modules because the offline compiler can incorporate them into your Intel^® Stratix^® 10 designs more effectively.

In traditional FPGA RTL for non- Intel^® Stratix^® 10 devices, it is common practice to reset every register indiscriminately for easy implementation without negative effects on performance. However, to improve the performance of your Intel^® Stratix^® 10 design, you must minimize the number of resets.

Avoid unnecessary resets in your Intel^® Stratix^® 10 design for the following reasons:

• The resulting high-fanout signal prevents the Intel^® Quartus^® Prime Pro Edition software's retimer to find a satisfactory solution.

  For more in-depth explanation, refer to the Avoid Broadcast Signals section in the Intel^® Stratix^® 10 High-Performance Design Handbook. Note that the term "broadcast signals" refers to high-fanout signals.

• In some situations, simply having a reset on a register, regardless of whether it is a high-fanout reset signal, is enough to degrade the performance of your Intel^® Stratix^® 10 design.

  For more in-depth explanation, refer to the Synchronous Resets and Limitations section in the Intel^® Stratix^® 10 High-Performance Design Handbook.

Intel^® recommends the following design guidelines for resets:

• To improve Intel^® Stratix^® 10 design performance, do not reset registers that do not hold internal states to reduce reset fanout.
• The reset signal is guaranteed to remain asserted for at least 50 clock cycles. Leverage this guaranteed assertion by "flushing" chains of registers that have internal states, which further reduces the reset fanout.
• You have the option to pipeline the reset signal inside the RTL module for fanout management, to a depth of no more than 15 pipeline registers for stall-free RTL modules, or no more than 25 pipeline registers for stallable RTL modules. If the RTL module is sufficiently large, pipelining the reset signal might improve design performance.

The figure below illustrates how a single-threaded host application might process parallel, independent data paths between kernel executions:

The OpenCL runtime is thread safe and supports multithreaded applications. Thus, you can perform work tasks on the host in parallel threads while still allowing those threads to access the OpenCL APIs in a thread-safe way.

However, thread safety is enforced at the OpenCL API boundary by synchronizing threads immediately after any OpenCL host API is called, which means, only one thread is ever active in the runtime while the other threads wait. As such, if one thread calls clFlinish on a queue full of work, another thread calling clSetKernelArg may have to wait for all of that work to complete before executing.

If possible, process data on CPU on multiple-threads, and use single dedicated thread to interact with the OpenCL API.

The figure below illustrates how a multi-threaded host application processes parallel, independent data paths between kernel executions:

As illustrated in the following figure, when the invocation queue is used, system and OpenCL runtime environment overheads (from responding to the finish and sending in the next set of invocation arguments) are overlapped with the kernel executions. This allows kernels to execute continuously, maximizing the system level throughput.

Kernel invocations are queued in hardware when another enqueued kernel with same kernel function name and same program is already running on the device, and the following are true:

• OpenCL kernel was not compiled with hardware kernel invocation buffer disabled (-no-hardware-kernel-invocation-buffer).
• OpenCL kernel was not compiled with performance counters (-profile)
• Enqueued OpenCL kernel does not have printf.
• All event objects queued earlier in the command queue have execution status equal to CL_COMPLETE.

  If the status is CL_SUBMITTED or CL_RUNNING, then that status is for the enqueue kernel with the same kernel function name in the same program.

• All event objects in the event wait list have execution status equal to CL_COMPLETE.

  If the status is CL_SUBMITTED or CL_RUNNING, then that status is for the enqueue kernel on the same device with the same kernel function name in the same program.

• If the OpenCL kernel uses heterogeneous memory, kernel currently running on the device and the one getting enqueued did not set the same memory object on different memory types.

Consider the following two example host code snippets where kernel invocation can be queued on hardware kernel invocation queue:

Example 1

int main()
{	…
  clEnqueueNDRangeKernel(queue, kernel, …, NULL);
  clEnqueueNDRangeKernel(queue, kernel, …, NULL);
  …
}

As soon as the first enqueue kernel is running, the second enqueue kernel can be queued on hardware.

Example 2

int main()
{	…
  clEnqueueNDRangeKernel(queue0, kernel0, …, NULL);
  clEnqueueNDRangeKernel(queue1, kernel1, …, NULL);
  clEnqueueNDRangeKernel(queue0, kernel0, …, NULL);
  clEnqueueNDRangeKernel(queue1, kernel1, …, NULL);
  …
}

As soon as the first enqueue of kernel0 is running, the second enqueue of kernel0 can be queued on the hardware irrespective of status of kernel1. Similarly, as soon as the first enqueue of kernel1 is running, the second enqueue of kernel1 can be queued on hardware irrespective of the status of kernel0.

Now, consider the following two examples where kernel invocation cannot be queued on hardware:

Example 1

int main()
{	…
  clEnqueueNDRangeKernel(queue, kernel0, …, NULL);
  clEnqueueNDRangeKernel(queue, kernel1, …, NULL);
  clEnqueueNDRangeKernel(queue, kernel0, …, NULL);
  …
}

Since the queue is in-order, enqueue kernel1 prevents the second enqueue of kernel0 from being queued on the hardware invocation queue.

Example 2

int main()
{	…
  clEnqueueNDRangeKernel(queue0, kernel0, …, NULL);
  clEnqueueNDRangeKernel(queue1, kernel1, …, &event);
  clFlush(queue1);
  clEnqueueNDRangeKernel(queue0, kernel0, …, 1, &event, NULL);
  …
}

Since the second enqueue of kernel0 is waiting on enqueue of kernel1 to complete, it will only get queued on the hardware kernel invocation queue if kernel1 finishes execution before first enqueue of kernel0 finishes.

To utilize hardware kernel invocation queue while double buffering, write your host code as shown in the following code snippet:

int main()
{	…
  cl_event dependencies[2];
  for (int i=0; i<MAX_ITERATIONS; i++) {
    if (i < 2) {
      clEnqueueWriteBuffer(writeQ,  inputBufferD[i%2],  CL_FALSE,  …,  inputBufferH[i],  0,  NULL,  &writeEvent[i]);
      clFlush(writeQ);
      clSetKernelArg(kernel,  0,  sizeof(cl_mem *),  &inputBufferD[i%2]);
      clSetKernelArg(kernel,  1,  sizeof(cl_mem *),  &outputBufferD[i%2]);
      clEnqueueNDRangeKernel(kernelQ,  kernel,  …,  1,  &writeEvent[i],  &kernelEvent[i]);
      clFlush(kernelQ);
    } else {
      clEnqueueWriteBuffer(writeQ,  inputBufferD[i%2],  CL_FALSE,  …,  inputBufferH[i],  1,  &kernelEvent[i-2],  &writeEvent[i]); 
      clFlush(writeQ);
      dependencies[0] = writeEvent[i];
      dependencies[1] = readEvent[i-2];
      clSetKernelArg(kernel,  0,  sizeof(cl_mem *),  &inputBufferD[i%2]);
      clSetKernelArg(kernel,  1,  sizeof(cl_mem *),  &outputBufferD[i%2]);
      clEnqueueNDRangeKernel(kernelQ,  kernel,  …,  2,  dependencies,  &kernelEvent[i]);
      clFlush(kernelQ);
    }
    clEnqueueReadBuffer(readQ,  output_device[i%2],  CL_FALSE,  …,  outputBufferH[i],  1,  &kernelEvent[i],  &readEvent[i]);
    clFlush(readQ);
  }
  …
}

The following diagram will help you visualize the event dependency:

The following figure illustrates the order the commands would be executed on the device assuming kernel execution is longer than reads and writes, and the device supports concurrent reads and writes: