MKPipe: A Compiler Framework for Optimizing Multi-Kernel Workloads in OpenCL for FPGA

We first study the multi-kernel workloads from the existing OpenCL for FPGA benchmark suites, including Rodinia(Zohouri et al., 2016), Spector(Gautier et al., 2016) and OpenDwarf(Verma et al., 2016). These multi-kernel workloads in our study share a common implementation that all the different kernels in the same workload are synthesized into a single bitstream, thereby co-residing on the same FPGA chip. The main reason is to eliminate the FPGA re-programming overhead. If there is data dependence among the kernels, the kernel invocations are sent to the same command queue, which imposes global synchronization among the invocations. The advantage of this approach is that it ensures correctness easily. The disadvantage is that the sequential execution of multiple kernels may lead to poor resource utilization since only the hardware corresponding to one kernel is active at a time.

To quantify the FPGA resource utilization, we propose a metric, effective resource utilization (ERU), for each kernel. As shown in Eq. 3, it is defined as the maximum usage among different types of resources of the FPGA chip, including both the static resources: adaptive loop-up tables (ALUTs), dedicated logic registers (FFs), DSP blocks, RAMs blocks, and the dynamic resource: DRAM bandwidth. The static resource usage is computed as the percentage of the resource consumed by the kernel. The DRAM bandwidth usage is the ratio of the utilized bandwidth over the peak one when the kernel is active. The reason for the maximum is to capture the effect of the critical resource. This way, low effective resource utilization indicates that there is room available in the hardware for performance improvement.

Here we use a case study on the benchmark CFD to illustrate the low ERU problem. CFD contains three kernels with the kernel data flow graph shown in Figure 1. Each kernel needs data from the previous one.

Figure 2a shows the ERU over time of the CFD benchmark. As the kernels are executed sequentially, we can visualize ERU as a stepwise function based on the order of kernel invocation and the execution time of each kernel. Although these kernels have been optimized individually, the overall ERU is low. In comparison, when we enable concurrent execution between K2 and K3, the execution time is reduced as shown in Figure 2b. Furthermore, CKE using kernel fusion or channels may free up some hardware resource. For example, as shown in Figure 2b, the RAM usage of K2&3 is less than the aggregated usage of K2 and K3. The reasons will be discussed in Section 5.4. Such freed up resource enables more opportunities for single-kernel optimizations.

The GPU multi-kernel pipeline programming frameworks (Steinberger et al., 2014; Zheng et al., 2017) have been developed recently. These frameworks leverage multiple Streaming Multiprocessors (SMs) on a GPU and may schedule different pipeline stages on different SMs. This is possible as each SM is a programmable processing unit and can execute different kernels. In contrast, an FPGA has fixed functionality once it is synthesized. Therefore, the multi-kernel programming frameworks for GPU are not directly applicable for FPGAs. Here we dissect the five GPU pipeline execution models used in VersaPipe (Zheng et al., 2017) and analyze their similarities and differences compared to pipeline execution models for FPGA.

The first GPU pipeline execution model is ”Run to completion (RTC)”. This execution model combines all stages of a pipeline into a single kernel, which is similar to the kernel fusion method discussed in Section 5.4.1. The limitation of this model is that it does not support global synchronization between stages.

The second GPU execution model is ”Kernel by kernel (KBK)”. In this model, multiple kernels are used and the kernels are executed one after the other. This model is the same as the baseline in the benchmark suites we studied. The limitation is that there is no concurrent kernel execution as discussed in Section 4.1.

The third GPU execution model is ”Megakernel”. Megakernel organizes pipeline computations into a huge kernel and each stage is scheduled with a software scheduler. The persistent thread technique (Aila and Laine, 2009; Gupta et al., 2012) is used to implement Megakernel. These persistent threads fetch data from a shared queue and run the corresponding pipeline stage upon the data. After each stage, the produced data are sent back to the same queue for subsequent processing. We tried to implement the Megakernel design on FPGA. However, the OpenCL for FPGA compiler is unable to handle this type of kernel and cannot construct hardware based on the OpenCL code. The main problem is the switch statement which chooses among different pipeline stages. The compiler regards these kernels as ”FPGA-unfriendly”. Furthermore, there are additional drawbacks for this model. First, the scheduler requires extra hardware resources. Second, the data communication between each stage is based on a shared queue. Although we can implement this shared queue using local memory on FPGA, this queue becomes a bottle neck due to its high number of read and write requests.

The fourth and fifth GPU pipeline execution models in VersaPipe are ”Coarse pipeline” and ”Hybrid pipeline”. In ”Coarse pipeline”, each pipeline stage is bounded to one SM. In ”Hybrid pipeline”, each pipeline stage is assigned to multiple thread blocks on a few SMs. As discussed before, a synthesized FPGA is not able to perform different functions based on the SM id. Therefore these two execution models are not feasible to FPGA.

Our compiler framework is shown in Figure 3. The input to our compiler includes naive kernel code, host code, and profiling data. The naive kernel code means that there is no device-specific or platform-specific optimization. In our implementation, the naive kernel is the same as the one from the benchmark suite with all the optimization #pragma and attributes stripped. The profiling data include the execution time and throughput of each naive kernel, where the throughput of one kernel is computed as the ratio of the output data size over the execution time. The compiler generates the kernel data flow graph from the host code and determines how the kernels can be executed concurrently while satisfying data dependency. Then the compiler analyzes the producer-consumer relationship between workitems/loop iterations in different kernels (workitems for NDRange kernels or loop iterations for single-workitem kernels) and uses different ways to enable CKE. Next, kernel balancing is performed to either balance the throughput in a multi-kernel pipeline or adjust the resource allocation among the kernels which require global synchronization. After kernel balancing, the compiler explores the option of bitstream splitting. Finally, the compiler produces the optimized kernel code and the host code.

The compiler derives the kernel data flow graph from the host code. The kernels are invoked in the host code using clEnqueueTask or clEnqueueNDRangeKernel functions. Their inputs and outputs arguments are explicitly set in clSetKernelArg functions. Among the kernels, the compiler excludes the kernels that can not be executed concurrently using the condition that they have dependency carried over through CPUs or CPU memory.

For kernels with data dependency, the compiler analyzes the kernel code to identify the producer-consumer relationship among their workitems/loop iterations for NDRange/single-workitem kernels, respectively. As the data dependency is carried over through the variables with the same name, for each global-memory variable, the compiler searches all the kernels to see which one(s) uses it as live-in/live-out. As the array indices in OpenCL workloads are typically affine functions of workitem ids or loop iteration indices, the compiler performs polyhedra analysis (Aho et al., 2006) to determine the exact dependency between the workitems/iterations of the producer kernel and those of the consumer kernel. Based on this producer-consumer relationship, the dependency between two kernels at the workitem/iteration level is classified into the following categories: few-to-few, few-to-many, many-to-many and many-to-few.

For example, in the code of the two single-workitem naive kernels of the CFD benchmark as shown in Figure 4, the compiler finds that the global variable ’fluxes_energy’ is produced in the kernel ’compute_flux’ and consumed in the kernel ’time_step’. When the iteration index variable i == j, these two iterations in the two kernels access the same global memory address. Therefore, the compiler identifies this producer-consumer relationship as one-to-one (or few-to-few), since one loop iteration in the producer kernel produces the data for one iteration in the consumer kernel.

Another example is shown in Figure 10, where the two NDRange kernels in the LUD benchmark access a common array ’m’. In each workitem of the kernel ’lud_perimeter’, a number (BSIZE) of elements in array ’m’ are updated as shown in line 6 of Figure 10 and their array indices are a linear function of its workitem id and its workgroup id. In the kernel ’lud_internal’, each workitem reads two elements from the same array for computation and their indices are also linear functions of their workitem ids and workgroup ids. Through polyhedral analysis, the compiler determines the dependency relationship between the producer workitem/workgroup ids in the kernel ’lud_perimeter’ and the consumer workitem/workgroup ids in the kernel ’lud_internal’ and classifies the producer-consumer relationship as one-to-many (or few-to-many).

Besides the dependency relationship, the compiler also produces a constant queue structure, id_queue, which is used to determine the desired execution order for the workitems in the consumer kernel. This id_queue is used in the workitem/workgroup id remapping step (Section 5.4.4). As the workitems in the producer kernel are dispatched in the sequential order based on their workitem ids, the compiler mimics this order to process the workitems of the producer kernel. For each producer workitem, the compiler checks its dependent workitems in the consumer kernel. If a dependent workitem has its dependency completely resolved, its workitem id will be pushed into the id_queue. If there are multiple dependent workitems in the consumer kernel and they are ready at the same time, all their workitem ids will be pushed in the id_queue. The compiler also builds a similar queue at the workgroup granularity, i.e., the queue contains the consumer workgroup ids.

Based on the producer-consumer relationship between the kernels, we propose a systematic decision tree approach to enable multi-kernel pipelines through CKE. Our approach is shown in Figure 5. First, the compiler checks if there is a dominant kernel in the workload. We define a kernel as dominant if its execution time is over 95% of total execution time. The reason for such a check is that as long as this dominant kernel has high resource utilization, the overall utilization is high and CKE would have very limited impact. Then, the compiler checks if there is a need for global synchronization among the kernels as a result of the producer-consumer relationship. For many-to-many or many-to-few producer-consumer relationship, the consumer workitems/loop iterations have to wait for almost all the producer workitems/iterations to finish. Therefore, the gains from CKE typically is not high enough to offset the potential overhead of CKE. As a result, global synchronizations are justified in such cases.

For multi-kernels not requiring global synchronization, the compiler explores different ways to enable multi-kernel pipelines through CKE. If the kernels exhibit few-to-many producer-consumer relationship, we propose to use global memory for data communication and this approach is referred to as CKE through global memory. If they exhibit few-to-few producer-consumer relationship, the compiler chooses between kernel fusion and CKE through channels. It estimates the overall execution time. When the execution time is high, the compiler chooses kernel fusion and CKE with channel otherwise for the reason discussed in Section 5.4.2.

With multiple kernels sharing the FPGA device, we need to coordinate the optimizations upon them. Our compiler considers two different scenarios for kernel balancing. The first is for kernels with CKE, i.e., the kernels form a pipeline. The second is that there are global synchronizations separating kernels and making them run sequentially.

For kernels with CKE, each kernel becomes a stage in a multi-kernel pipeline. Therefore, the goal is to balance the throughput among the stages. However, if kernels are separated by global synchronizations, throughput imbalance is not an issue. The goal then is to balance resource allocation such that more resources can be allocated for the kernels, which can lower the most execution time from such additional resources. A workload may also contain both cases at the same time. For example, the CFD benchmark has three kernels. Our compiler determines that it is beneficial to enable CKE between K2 and K3 while K1 should be ended with a global synchronization. In such a case, the compiler considers the K2&3 pipeline as a single kernel and allocates resources to K1 and K2&3 accordingly using the algorithm discussed in Section 5.5.2. Then, the allocated resources for K2&3 are further distributed between K2 and K3 using the algorithm discussed in Section 5.5.1 for throughput balancing.

In our compiler framework, we consider three parameters for single-kernel optimizations: compute unit replication (CU) factor, SIMD factor, and loop unroll (Unroll) factor. As these factors have similar performance impacts, i.e., increasing any factor by N times can potentially increase the throughput by N times, our compiler first determines a unified performance factor, denoted as N_uni, for each kernel, and then realize this factor by adjusting the unroll, SIMD, and CU factors as discussed in Section 5.5.3.

The bitstream splitting optimization explores the option of placing kernels in multiple bitstream files. This way, more resources are available for each kernel such that more aggressive single-kernel optimizations can be performed. However, using multiple bitstreams has to pay the penalty of device reprogramming and data transfer between the device and the host. Therefore, we limit the maximal number of bitstreams as 2. As a result, if there are more than two kernels in a workload, our compiler decouples them into two virtual kernels. Such decoupling is essentially the same as bi-partitioning the kernel data flow graph.

Our compiler employs the following criteria for bi-partitioning the graph: (a) loops cannot be partitioned unless each iteration of the loop has very high execution latency compared to reprogramming overhead; (b) a multi-kernel pipeline can not be broken by partitioning; and (c) the difference between the accumulated critical resource utilization over time in either partition needs to be minimized. A loop in the kernel data flow graph means that the kernels will be invoked multiple times. If we break kernels in a loop into different bitstreams, we have to pay the reprogramming overhead for each iteration. Therefore, unless the execution time of each iteration is high, the loop should not be partitioned. The last criterion aims to isolate the long-running kernels, which are resource constrained due to co-residence with other kernels. Such kernels are more likely to benefit from more resources. Using the notation of Equation 2, the condition can be expressed as find a partition to minimize $|\textit{T${}_{1}$}\times\textit{ERU${}_{1}$}-\textit{T${}_{2}$}\times\textit{ERU${}_{2}$}|$. As the number of the kernels in multi-kernel pipelines are small, our compiler exhaustively goes through all the possible partitions to find one that meets the criteria.

With the two virtual kernels, our compiler uses Equation 2 to determine whether to put them into separate bitstreams or to let them co-reside in the same one. In the equation, K₁ and K₂ are the two virtual kernels. The ERU of them are ERU₁ and ERU₂ and their execution times are T₁ and T₂. The reprogramming and data transfer overhead are T_r and T_d, respectively. We consider kernel co-residence in a single bitstream beneficial if:

The LHS of Equation 2 is the execution time if both kernels reside on the same device. The RHS is an estimate of the execution time if they are separated into two different bitstreams. When one kernel monopolizes the device, its execution time can be reduced with more aggressive optimizations. Such reduced execution time is estimated with a factor of the kernel’s ERU, i.e, the utilization of its critical resource. For example, if one kernel uses 80% of the DSP blocks, when the entire chip, i.e., 100% DSP blocks are available to it, the potential performance improvement would be 100%/80%. The corresponding execution time is $80\%\times$T, i.e., $ERU\times$T. If LHS is less than RHS, co-residence is preferred. Otherwise, the compiler produces two source code files, one for each virtual kernel, which will be used by the OpenCL compiler to synthesize into separate bitstreams.

We tested the reprogramming overhead T_r using kernels with different complexities. We found that the reprogramming overhead is around 1400ms for different kernels and it is independent upon the complexity or resource requirement of the kernel.

After kernel optimizations, the host code is adjusted accordingly. For kernel fusion, unnecessary kernel invocations and allocations for the global memory data that are used for cross kernel communication would be removed. For CKE with channel and CKE with global memory, kernel arguments are adjusted. The compiler also allocates global memory for the global ’flags’ array and the ’id_order’ array. All the clFinish functions between concurrent executing kernels are removed since they are synchronization points.

The main cost of MKPipe is the OpenCL to FPGA compilation overhead. We will compile the kernels during two steps: the profiling step and auto-tuning step. Since the naive kernel has no optimization pragma and attributes, the compilation time in profiling step is usually much smaller than compiling the baseline kernels as they enable these pragma/attributes. During the auto-tuning step, each kernel needs to be compiled for $2p+1$ times, where $p$ is the user-defined parameter discussed in Section 5.5.1. Nonetheless, all these compilations in the tuning step can be performed in parallel.

Figure 14 reports the normalized performance of the multi-kernel workloads. We use the following notations: ’KBK’ represents the kernels from the benchmark suites which use the KBK model. KBK is our baseline. For the benchmarks derived from the GPU benchmark suite, we applied the SIMD and CU attributes to optimize the kernel as our baseline. ’Fusion’ represents kernels executed with a hybrid model of ’KBK’ and ’RTC’. The reason is that the kernel fusion model or RTC does not support global synchronization between stages. As a result, the kernels with global synchronization in between are executed with the KBK model. ’Channel’ represents kernels executed with the hybrid model of ’KBK’ and ’CKE with channel’. ’Global Memory’ represents kernels executed with the hybrid model of ’KBK’ and ’CKE with global memory’. ’Kernel Balancing’ shows the speedup of the kernels with the best pipeline model and kernel balancing optimization. ’Bitstream splitting’ is for the kernels optimized with kernel balancing and bitstream splitting.

Among all the benchmarks, BP uses single workitem kernels. CFD has kernel implementations in both single-workitem (labeled ’CFD _SI’) and NDRange mode (labeled ’CFD_NDR’), and Hist has the original implementation using an NDRange producer kernel and a single workitem consumer kernel (labeled ’Hist_MIX’). We found that this NDRange kernel results in low frequency and rewrote it as a single workitem kernel. This Hist version is labeled ’Hist_SI’. All the remaining benchmarks use NDRange kernels.

From Figure 14, we can see that the CKE optimization and the kernel balancing optimization contribute the most performance improvement. Overall, the multi-kernel workloads optimized by MKPipe achieve up to 3.6x (1.4x on average) speedup over the baseline. Among the multi-kernel workloads, BFS has a dominant kernel, which takes 95.8% of the overall execution time. MKPipe identifies this dominant kernel and performs kernel balancing optimization. Our optimized kernel achieves a speedup of 1.1x as our compiler balances the optimizations on the kernels more judiciously. The Histogram benchmark has one producer kernel and one consumer kernel, their dependency relationship is identified as one-to-one. For the single workitem implementation (Hist_SI), MKPipe generates both fused design and CKE with channel design. As the fused design forms a longer loop body, the OpenCL compiler optimizes the code more effectively and the synthesized design achieves a speedup of 1.7x over the baseline. For HIST_MIX, due to different numbers of workitems in the producer and consumer kernels, MKPipe chooses to enable CKE using channels. The benchmark Tdm benefits the most from the kernel balancing optimization as it efficiently searches a large design space of the optimization parameters. The main benefit of LUD comes from CKE with global memory and workgroup mapping as discussed in Section 5. Color benefits from kernel fusion. Dijkstra benefits from CKE with channel due to the low execution time of its kernels. For the remaining benchmarks, CFD and BP, we analyze them in Section 7.3.

Besides the performance impact, we present the resource consumption and the frequency of different designs for each benchmark in Table 2. From the table, we can see that for most benchmarks our optimized design utilizes resources more aggressively and one side effect is the slightly lower frequency due to longer critical paths.

In this experiment, we compare the FPGA performance with NVIDIA RTX 2080 GPUs. For benchmarks BFS and Hist, the CUDA kernels from Parboil (Stratton et al., 2012) benchmark suite are used as they are optimized for GPUs. Similarly, for benchmarks CFD, LUD and BP, the CUDA kernels from Rodinia benchmark suite are used. For benchmark Tdm, the OpenCL kernels is used as the OpenDwarf benchmark suite does not have the CUDA version. The benchmarks from Pannotia benchmark suite are not included since they require AMD drivers and SDK support. The results are shown in Figure 15. Given the bandwidth difference between our FPGA board (25.6 GB/s for Stratix V) and GPU (448 GB/s on RTX2080), the performance of OpenCL kernels for FPGA is not competitive.

To make a more fair comparison, we include a performance projection for the state-of-art Stratix 10 MX FPGA. Compared to Stratix V GX, Stratix 10 MX (Corporation, 2019b) has 6x DSP capability, 2.6x memory blocks 20x memory bandwidth(512GB/s). Taking the advantage of 14nm manufacturing node and HyperFlex technology (Hutton, 2015), Stratix 10 family FPGA boards are expected to reach an operating frequency $f_{max}$ of up to 1 GHz. However, when $f_{max}$ is limited by the critical path, HyperFlex will have a limited impact. Therefore, we only assume a 150MHz increase in $f_{max}$ compared to Stratix V, which is in accordance with prior study (Zohouri et al., 2018) (Chung et al., 2018). As can be seen in Table 2, most of the optimized benchmarks are bandwidth limited. Based on the existing performance estimation model (Wang et al., 2016), the speedup of benchmarks on Stratix 10 MX can be predicted as:

$freq_{proj}$ and $freq$ are the frequencies of Stratix 10 MX and Stratix V. $\#Banks_{proj}$ and $\#Banks$ are the number of memory banks which are 32 and 2 for Stratix 10 MX and Stratix V, respectively. #Mem _trans _width is the maximum transaction width and it is 64Byptes for both devices. Based on these data, the average speedup (geometric mean) of all six benchmarks is 26.8x. As can be found in figure 15, the average speedup of kernels on Stratix 10 MX FPGA is comparable with the average speedup of the kernels on the state-of-the-art GPU. Such results are also consistent with existing works (Zohouri et al., 2016; Zohouri et al., 2018) that FPGAs deliver inferior performance but superior energy efficiency to the same generation GPUs.