Enabling OpenMP Task Parallelism on Multi-FPGAs

The OpenMP specification [17] defines a programming model based on code annotations that use minimal modifications of the program to expose potential parallelism. OpenMP has been successfully used to parallelize programs on CPUs and GPUs residing on a single computer node.

Tasks were introduced to OpenMP in version 3.0 [18], to expose a higher degree of parallelism through the non-blocking execution of code fragments annotated with task directives. The annotated code is outlined as tasks that are dispatched to the OpenMP runtime by the control thread. Dependencies between tasks are specified by the depend clause, which the programmer uses to define the set of input variables that the task depends on and the output variables that the task writes to. Whenever the dependencies of a given task are satisfied, the OpenMP runtime dispatches it to a ready queue which feeds a pool of worker threads. The OpenMP runtime manages the task graph, handles data management, creates and synchronizes threads, among other activities.

Listing 1 shows an example of a simple program that uses OpenMP tasks to perform a sequence of increments in the vector V, similarly as discussed in Figure 1. First, the program creates a pool of worker threads. This is done using the #pragma omp parallel directive (line 3). It is with this set of threads that an OpenMP program runs in parallel. Next, the #pragma omp single directive (line 4) is used to select one of the worker threads. This selected thread is called control thread and is responsible for creating the tasks. The #pragma omp task directive (line 6) is used to effectively build N worker tasks, each one encapsulating a call to function foo(V,i) (line 8) that performs some specific operation on the vector V depending on the task (loop) index. Worker threads are created by the OpenMP runtime to run on the cores of the system. In order to assure the execution order of the tasks the depend clause (line 7) is used to specify the input and output dependencies of each task. In the specific case of the example of Listing 1 we have to introduce a deps vector which assures that the foo(V,i) tasks are executed in a pipeline order.

Now assume the user wants to accelerate the above described tasks in a GPU. For this purpose, consider the code in Listing 2 which is the version of Listing 1 using a GPU as an acceleration device. In Listing 2 the target directive (line 6) has a similar role as the task directive from Listing 1, but has for goal to offload the computation to an acceleration device instead of a CPU. The device clause (line 6) receives as an argument an integer which specifies the device that will perform the computation. In the case of Listing 2 this number corresponds to a GPU. Similar to the task directive, the target directive also accepts the depend clause (lines 7). The map clause (line 8) specifies the direction of the transfer operation that the host must execute with the given data. The most common transfer operations are to, from, and tofrom. For example, in line 8 of Listing 2 vector V is annotated with tofrom given that it is first sent from the host memory to the GPU memory, and then received back after processed by function foo. Notice that clause nowait also appears in line 8 of Listing 1. This clause is necessary because, by default, the target directive is blocking. In other words, the nowait clause allows the control thread to create all tasks without waiting for the previous ones to complete. Finally in line 9 of Listing 2 the worker thread foo function is called.

The current LLVM OpenMP implementation of task offloading is not enough to handle devices like Multi-FPGA. This paper proposes an OpenMP plugin implementation for Multi-FPGA devices. Moreover, it extends the LLVM OpenMP runtime to substitute software tasks for target tasks running on hardware IPs of a Multi-FPGA architecture. This is done by implementing the declare variant pragma which is already defined in the OpenMP standard. Also, the depend and the map clauses are used to dynamically create communication paths between the IPs of the various FPGAs in the architecture.

The Xilinx VC709 Connectivity Kit provides a hardware environment for developing and evaluating designs targeting the Virtex-7 FPGA. The VC709 board provides features common to many embedded processing systems, including dual DDR3 memories, an 8-lane PCI Express interface [19], and small SFP connectors for fiber optical (or Ethernet) transceivers. In addition to the board’s physical components, the kit also comes with a Target Reference Design (TRD) featuring a PCI Express IP, a DMA IP, a Network Module, and a Virtual FIFO memory controller interfacing to DDR3 memory. Figure 2 shows the schematic of the main board components and their corresponding components.

This kit was selected to explore the ideas discussed in this paper due to its reduced cost and the fact that the TRD has components ready for inter-FPGAs communication. Each one of the TRD components is explained in more detail below.

To explain how the proposed system works, please consider the code fragment shown in Listing 3. The declare variant directive (line 1), which is part of the OpenMP standard, declares a specialized hardware (FPGA IP-core) variation (hw_laplace2d, in line 4) of a C function (do_laplace2d, in line 2), and specifies the context in which that variation should be called. For example, line 3 of Listing 3 states that the variant hw_laplace2d should be selected when the proper vc709 device flag is provided to the compiler at compile time. This compiler flag is matched by the match (device=arch(vc709)) clause and when the call to function do_laplace2d(&V, x, y) (line 16) is to be executed a call to the IP-core hw_laplace2d(int*, int, int) is performed instead.

As shown in line 11 of Listing 3, the main function creates a pipeline of N tasks. Each task receives a vector V containing h*w (height and width) grid elements that are used to calculate a Laplace 2D stencil. As the tasks are created within a target region and the vc709 device flag was provided to the compiler, the hw_laplace2d variant is selected to run each task in line 16. The compiler then uses this name to specify and offload the hardware IP that will run the task. As a result, at each loop iteration, a hardware IP task is created inside the FPGA. Details of the implementation of the Laplace 2D IP and other stencil IPs used in this work are provided in Section IV.

The map clause (line 12) specifies that the data is mapped back and forth to the host at each iteration. However, the implemented mapping algorithm concludes that vector V is sent to the IP from the host memory and its output forwarded to the next IP in the following iteration. The interconnections between these IPs are defined according to the depend clauses (line 13). In the particular example of Listing 3, given that the dependencies between the tasks follow the loop iteration order, a simple pipelined task graph is generated.

A careful comparison of Listing 3 with Listing 2 reveals that in terms of syntax, the proposed approach does not require the user to change anything concerning the OpenMP standard, besides specifying the vc709 flag to the compiler. This gives the programmer a powerful verification flow. He/she can write the software version of do_laplace2d for algorithm verification purpose, and then switch to the hardware (FPGA) version hw_laplace2d by just using the vc709 compiler flag.

To achieve this level of alignment with the OpenMP standard, two extensions were required to the OpenMP runtime implementation: (a) a modification in the task graph construction mechanism and; (b) the design of the VC709 plugin in the libomptarget library. All these changes have been done very carefully to ensure full compatibility with the current implementation of the OpenMP runtime.

As shown in Figure 3, the plugin receives the task graph generated by the runtime and maps these tasks to the available IPs in the cluster. The cluster configuration is passed through a conf.json file, which contains: (a) the location of the bitstream files, (b) the number of FPGAs, (c) the IPs available in each FPGA, and (d) the addresses of IPs and FPGAs. As in our experiments, the FPGAs are connected in a ring topology, a round-robin algorithm is used to map tasks to IPs. Each task is mapped in a circular order to the free IP that is closest to the host computer.

Besides the extensions to OpenMP, an entire hardware infrastructure was designed to support OpenMP programming of Multi-FPGA architectures. This infrastructure leverages on the Target Reference Design (TRD) presented in Section II-B, but could also be ported to other modern Alveo FPGAs. To facilitate its understanding, the design is described below using two perspectives: Single-FPGA execution and Multi-FPGA execution.

Consider, for example, the hardware infrastructure in Figure 1, but this time with a single FPGA board and 4 task IPs (IP0-IP3). These IPs are separately designed using a standard FPGA toolchain (e.g. Vivado), which is not addressed in this paper. To connect the IPs into the infrastructure, the designer just needs to ensure that they use the AXI-Stream interface. The infrastructure can be easily changed to accept other interfaces, although for the purpose of this paper AXI-Stream suffices.

According to the proposed programming model, FPGA IPs can execute tasks that have dependencies between each other. In order to implement such dependencies an AXI4-Stream Interconnet (AXIS-IC) module [22] was inserted to the infrastructure (Figure 1). This enables the IPs to communicate directly to each other, based on the OpenMP dependencies programmed between them, thus avoiding unnecessary communication through the host memory.

The AXI4-Stream Interconnect is an IP core that enables the connection of heterogeneous master/slave AMBA® AXI4-Stream protocol compliant endpoint IP. The AXI4-Stream Interconnect routes data from one or more AXI4-Stream master channels to one or more AXI4-Stream slave channels. The VC709 plugin uses the CONF register (Figure 1) bank to program the source and destination ports of each IP according to their specified task dependencies.

To enable that, a MAC Frame Handler module (MFH) was designed and inserted into the hardware infrastructure, as shown in Figure 1. This module is required because the Network Subsystem that routes packages through the optical fibers that connects the boards receives data in the form of MAC Frames, which contain four fields: (a) destination, (b) source, (c) type/length and (d) payload. Therefore, to use the optical fibers, a module that can assemble and disassemble MAC frames is required.

The MFH module is responsible for inserting and removing the source and destination MAC addresses and type/lengh fields whenever the IPs need to send/receive data through the Network Subsystem. MAC addresses are extracted from the dependencies in the task graph while the type/lengh fields are extracted from the map clause. The VC709 plugin uses this information to set up the CONF registers, which in turn configure the MFH module.

With all of these components in place, the proposed OpenMP runtime can to distribute tasks IPs across a cluster of FPGAs and map the dependency graph so that FPGA IPs communicate directly.

Figure 5 shows an overview of a typical stencil IP implementation using cell and iteration parallelism. Figure 5(a) shows the grid to be computed, and Figure 5(b) the components that implement the stencil, namely: (a) a shift-register that stores the grid data in processing order; and (b) the processing element (PE), which does the actual stencil computation. The cells in Figure 5(a) are computed by the architecture in Figure 5(a) from left-to-right and top-to-bottom one after the other. At each clock cycle, data in the shift-registers are shifted to the left in Figure 5(a), and a new cell value is pushed into the input of the first shift-register (i.e. $cell^{t}_{3,2}$). The computation starts after all neighboring data of a cell are available in the shift-register array. In the example of Figure 5, $cell^{t+1}_{1,1}$ is computed while input data is stored into $cell^{t}_{3,2}$. In the next clock cycle, the data at $cell^{t}_{0,0}$ at the output of the shift-register is discarded (shifted out), and the data of $cell^{t}_{4,2}$ is pushed into the input of the shift-register. Notice that the data at $cell^{t}_{0,0}$ is no longer required for any computation at this stage.

Each stencil IP has a shift-register and eight processing elements and is thus capable of processing up to eight elements at a time until the end of an iteration. Each IP works with a 256-bit AXI4-Stream interface, as each cell in the matrix is a 32-bit float.

The A-SWT switch in the architecture of Figure 1 can be configured so that the IPs can be reused, thus expanding the system’s capacity to deal with larger grids and iteration counts. By doing so, the stencil pipeline can be scaled in both space and time. Such scaling is required to leverage the processing power of the multiple FPGAs, and to enable the computation of large-size problems that could not be done by a single FPGA due to the lack of resources. Unfortunately, as discussed in the Section V, the size and number of IPs in an FPGA is constrained by the ability of the synthesis tool and designer to make efficient use of the FPGA resources, and this sometimes can become a bottleneck.

The FPGA’s scalability experiments were executed with the settings shown in Table II, and varying the number of FPGAs from 1 to 6. The Grid Size column shows the dimensions of the initial grid for each kernel. The more computation a kernel does, the more difficult it was for Vivado to synthesize the design respecting the time constraints. For this reason, the dimensions of the grid at each kernel were adjusted to avoid negative slacks. The Iterations column was set at 240 so that it was possible to execute with all 6 FPGAs. The # IPs column specifies the number of IPs at each FPGA. The number of IPs varies for the same reason as the dimensions of the initial grid. The larger the kernel computation, the smaller the number of IPs Vivado could synthesize within the time constraints. On the other hand, as discussed in Section V-C there is still plenty of hardware to be used before the FPGA runs out of resources, which reinforces the long term potential of the model proposed herein if using a more efficient FPGA design flow like (e.g. U250 and Vitis).

The graph of Figure 6 shows the speedup concerning the execution on a single FPGA, achieved by the various stencil kernels as the number of FPGAs varies on the x-axis. The speedup grows almost linearly with the number of FPGAs for all five kernels. This result shows that it is possible to scale applications using Multi-FPGA architectures by using programming models like the one proposed in this paper to facilitate the design of such systems. The graph of Figure 7 shows, on the y-axis, the number of floating-point operations (GFLOPs) for each kernel as the number of FPGA varies on the x-axis. The Laplace-2D kernel (yellow line) executes more GFLOPs than the other kernels. This is because, although the computation of this kernel is the simplest one, during synthesis, it was possible to insert more IPs (four) per FPGA, which allowed more iteration parallelism as discussed in Section IV. Just below the Laplace-2D is the Laplace-3D (green line), with only 2 IPs per FPGA still managed to sustain a linear performance growth. For the remaining kernels, as they all have only one IP per FPGA, the number of GFLOPs is related to the number of operations executed and the grid’s dimensions. Notice that Diffusion-3D (red line) and Diffusion-2D (blue line) perform less computation than the Jacobi 9-pts (orange line). However, they achieve better GFLOP numbers due to their higher grid dimension, enabling them to take advantage of increased iteration parallelism.

A second experiment was performed to evaluate the IPs’ scalability concerning the number of iterations. The Laplace-2D kernel was used as an example, although similar results have been achieved for the other kernels. The graph in Figure 8 shows, on the y-axis, the number of GFLOPs produced by the system, as the number of iterations varies on the x-axis. The yellow, blue, red, and green lines represent executions with 1, 2, 3, and 4 IPs, respectively. As shown, the execution with a single IP (yellow line) remains practically constant.

On the other hand, the execution with 4 IPs shows an increase in performance until reaching a plateau. The executions with 2 and 3 IPs also show a gradual performance increase. This experiment reveals that by increasing the number of IPs, it is possible to improve the system’s scalability in terms of iterations.

The graph of Figure 9 shows on the y-axis the number of GFLOPs for the Laplace-2D kernel as the number of IPs increase (x-axis). Each line in the graph is a different number of iterations. The graph reinforces the insight revealed in Figure 8: as more IPs are added to the system, the more significantly the increase in the number of iterations improve performance. This can be confirmed by looking at the distances between the lines in Figure 9, which grow larger as the number of IPs increase. This experiment also supports the case for Multi-FPGA architectures.

Regarding resource utilization, the graph in Figure 10 shows the percentage of occupancy of the FPGA main components of the proposed infra-structure (not considering the IPs). Remarkably, the DMA/PCIe component occupies 30.2% of the available LUTs. This large utilization is because the DMA/PCIe was designed to support a board with four communication channels, although the proposed approach just requires one. Components MFH, SWITCH, VFIFO, and Network occupy, respectively, 1.7%, 11.5%, 13.2% , and 6.1% of the available LUTs. BRAMs are used by the DMA/PCIe (5.5%), VFIFO (18.3%), and NET (2.4%). The most significant usage of BRAMs comes from VFIFO, which uses it to multiplex and demultiplex the four channels of the virtual FIFO. DSP is the least used component (1%).

Table III shows the quantity and percentage of the FPGA components used by each IP from the free region (gray area) of Figure 10. The percentage of the available LUTs effectively used by the stencil IPs varies from 7.5% to 28.3%, depending on the complexity of the kernel. As for BRAM, the utilization ranges from 0.7% to 6.0%. This is directly linked to the size of the shift-registers, and is impacted by the size of the grid to be calculated. The number of DSP components used by the IPs varies from 0.4% to 4.0%, and is related to the number of multiplications performed by each kernel. The small utilization of the FPGA resources by the kernels has been previously discussed in the beginning of Section V. Additional work will be done to address these shortcomings.