RT-CUDA: A Software Tool for CUDA Code Restructuring

Abstract

Recent development in graphic processing units (GPUs) has opened a new challenge in harnessing their computing power as a new general purpose computing paradigm. However, porting applications to CUDA remains a challenge to average programmers, which have to package code in separate functions, explicitly manage data transfers between the host and device memories, and manually optimize GPU memory utilization. In this paper, we propose a restructuring tool (RT-CUDA) that takes a C-like program and some user directives as compiler hints to produce an optimized CUDA code. The tool strategy is based on efficient management of the memory system to minimize data motion by managing the transfer between host and device, maximizing bandwidth for device memory accesses, and enhancing data locality and re-use of cached data using shared-memory and registers. Enhanced resource utilization is implemented by re-writing code as parametric kernels and use of efficient auto-tuning. The tool enables calling numerical libraries (CuBLAS, CuSPARSE, etc.) to help implement applications in science simulation like iterative linear algebra solvers. For the above applications, the tool implement an inter-block global synchronization which allow the execution overall among a few iterations which is helpful to balance load and to avoid polling. Evaluation of RT-CUDA has been performed using a variety of basic linear algebra operators (Madd, MM, MV, VV, etc.) as well as the programming of iterative solvers for systems of linear equations like Jacobi and Conjugate Gradient algorithms. Significant speedup has been achieved over other compilers like PGI OpenACC and GPGPU compilers for the above applications. Evaluation shows that generated kernels efficiently call math libraries and enable implementing complete iterative solvers. The tool help scientists developing parallel simulators like reservoir simulators, molecular dynamics, etc. without exposing to complexity of GPU and CUDA programming. We have partnership with a group of researchers at the Saudi Aramco, a national company in Saudi Arabia. RT-CUDA is currently explored as a potential development tool for applications involving linear algebra solvers by the above group. In addition, RT-CUDA is being used by Senior and Graduate students at King Fahd University of Petroleum and Minerals in their projects as part of RT-CUDA continuous enhancement.

Appendix: CUDA Kernel Optimizations

1.1 Manual Optimizations

Vectorization improves the bandwidth utilization by using one of the vector data types such as float2, float4, float8 in CUDA as structs with some special data alignment [32, 44]. Global memory transactions in GPU are aligned to 128 bytes even if the actual data load is less than 128 bytes. So, if a specific thread within a warp performs 8 loads of float data type with sequence of load instructions then the GPU perform 8 different memory transactions. On the other hand, if a float8 type is used and perform single float8 load operation then it can be done by only one global memory transaction. So, to improve global memory bandwidth utilization, the programmer should use vector data types in the case when consecutive data loads are not aligned to 128 bytes. Texture Fetching utilized the texture memory that is a read-only portion of memory in device memory (DRAM) that has been cached (off-chip cache) on access [32, 44]. It has been accessible by all threads and host. It is optimized for 2D spatial locality, so threads of the same warp that read texture addresses that are close together achieve best performance. Texture references that are bound to CUDA arrays can be written to via surface-write operations by binding a surface to the same underlying CUDA array storage. Reading from a texture while writing to its underlying global memory array in the same kernel launch should be avoided because the texture caches are read-only and are not invalidated when the associated global memory is modified. So, texture fetches from addresses that have been written via global stores in the same kernel call returned undefined data. The data in texture can consist of 1, 2, or 4 elements of any of the following types: (1) Signed or unsigned 8, 16, or 32-bit integers, (2) 16-bit floating point values, and (3) 32-bit floating point values. Arrays declared in texture memory can be used in kernels by invoking texture intrinsic provided in CUDA such as tex1D(), tex2D(), and tex3D() for 1D, 2D, and 3D CUDA arrays respectively. Before invoking a kernel that uses texture memory, the texture must be bound to a CUDA array or device memory by calling cudaBindTexture(), cudaBindTexture2D(), or cudaBindTexturetoArray(). Coalesced Global Memory Access refers to combining multiple memory accesses into a single transaction [13]. Global memory is the slowest memory on the GPU. Simultaneous global memory accesses by each thread of a half-warp (16 threads) during the execution of a single read and write instruction are coalesced into a single access. This is achieved based on the following conditions: (1) the size of the memory element accessed by each thread is either 4, 8, or 16 bytes, (2) the elements to be accessed form a contiguous block of memory, (3) the Nth element is accessed by the Nth thread in the half-warp, does not affect if any thread in between not accessing the global memory that is divergent warp, and (4) the address of the first element is aligned to 16 times the element’s size. nVidia Kepler’s Shuffle Instructions perform data exchange between threads within a warp [30]. It is more faster than the use of shared memory. This feature allows the threads of a warp to exchange data with each other directly without going through shared (or global) memory. So, this can present an attractive way for applications to rapidly interchange data among threads. There are four variants of shuffle instructions in CUDA that are __shfl(), __shfl_up(), __shfl_down(), and __shfl_xor(). Shuffle instructions can be used to free up shared memory to be used for other data or to increase warp occupancy and to perform warp-synchronous optimizations (removing __syncthreads()). All the __shfl() intrinsics take an optional width parameter which permits sub-division of the warp into segments. For example, to exchange data between 4 groups of 8 lanes in a SIMD manner. If width is less than 32 then each subsection of the warp behaves as a separate entity with a starting logical lane ID of 0. A thread may only exchange data with others in its own subsection. Width must have a value which is a power of 2 so that the warp can be subdivided equally; results are undefined if width is not a power of 2, or is a number greater than warpSize.

1.2 User Driven Optimizations

Loop Collapsing is a technique to transform some nested loops into a single-nested loop to reduce loop overhead and improve runtime performance [18] specifically for irregular applications such as sparse matrix vector multiplication (spMV). Such applications pose challenges in achieving high performance on GPU programs because stream architectures are optimized for regular program patterns. It improves the performance of the application in three ways: (1) the amount of parallel work (the number of iterations, to be executed by GPU threads) is increased (2) inter-thread locality is increased (3) control flow divergence is eliminated, such that adjacent threads can be executed concurrently in an SIMD manner [24]. Thread and Thread-Block Merging enhance the data sharing among thread blocks to reduce the number of global memory accesses [47, 48]. Thread-Block Merge determines the workload for each thread block while Thread Merge decides the workload for each thread. If data sharing among neighbouring blocks is due to a global to shared memory (G2S) access, Thread-Block Merge should be preferred to better utilization of the shared memory. When data sharing is from a global to register (G2R) access, Thread Merge from neighbouring blocks should be preferred due to the reuse of registers. If there are many G2R accesses, which lead to data sharing among different thread blocks, the register file is not large enough to hold all of the reused data. In this case, Thread-Block Merge should be used and shared memory variables should be introduced to hold the shared data. In addition, if a block does not have enough threads, Thread-Block Merge instead of Thread Merge should also be used to increase the number of threads in a block even if there is no data sharing. Thread Merge achieves the effects of loop unrolling. It combines several threads’ workload into one thread (combining N neighbouring blocks along column direction into one). By doing this, they can share not only shared memory but also the registers in the register file. Furthermore, some control flow statements and address computation can be reused, thereby further reducing the overall instruction count. The limitation is that an increased workload typically requires a higher number of registers, which may reduce the number of active threads that can fit in the hardware. Parallel Loop Swap is used to improve the performance of regular data accesses in nested loops [3, 24]. It uses to transform non-continuous memory accesses within the loop nest to a continuous memory access which is a candidate for the coalesced global memory access optimization. Strip Mining splits a loop into two nested loops [16, 41–43]. The outer loop has stride equal to the strip size and the inner loop has strides of the original loop within a strip. This technique is also used in loop tiling. In loop tiling or loop blocking, loops are also interchanged after performing strip mining to improve the locality of memory references that is why loop tiling is also called strip-mine-and-interchange. Bank Conflict Free Shared Memory Access improves performance by reordering the data into shared memory such that the memory addresses requested by the consecutive threads in a half-warp should be mapped to different memory banks of shared memory [10]. Shared memory banks are organized such that successive 32-bit words are assigned to successive banks and the bandwidth is 32 bits per bank per clock cycle. In GPUs, the warp size is 32 threads and the number of banks is 16. So, a shared memory request for a warp is split into one request for the first half of the warp and one request for the second half of the warp. However, no bank conflict occurs if only one memory location per bank is accessed by a half warp of threads. Using Read-Only Data Cache, introduced in nVidia Kepler in addition to L1 cache, can benefit the performance of bandwidth-limited kernels [30, 39]. This is the same cache used by the texture pipeline via a standard pointer without the need to bind a texture beforehand and without the sizing limitations of standard textures. This feature is automatically enabled and managed by the compiler, access to any variable or data structure that is known to be constant through programmer use of the C99-standard “const __restrict__” keyword tagged by the compiler to be loaded through constant data cache.

1.3 Compiler Optimizations

Common Sub-Expression Elimination is a compiler optimization technique that searches for instances of identical expressions, evaluates to the same value, and replace them with a single variable holding the computed value [9, 49]. It enhances the application performance by reducing the number of floating point operations. In CUDA, common sub-expression elimination can be used to avoid redundant calculations for the initial address of an array. Loop Invariant Code Motion (also called hoisting or scalar promotion) is a compiler optimization that has been performed automatically [17, 36]. Loop invariant code is a set of statements or expressions within the body of a loop that can be moved outside of the body without affecting the semantics of the program. It makes loops faster by reducing the amount of code that executes in each iteration of the loop. The CUDA C compiler automatically applies this optimization technique to the PTX code. Loop Unrolling is a compiler optimization technique that is applied for the known trip counts at the compile time either by using the constants or templating the kernel [28]. NVIDIA compiler also provides a directive ‘#pragma unroll’ to explicitly activate the loop unrolling on a particular loop.