Lecture 15: CUDA
This lecture mainly introduces the programming model of CUDA and discusses how to develop programs for CUDA based on this programming model. Specifically, the lecture covers the following topics:
- CUDA bandwidth
- CUDA overview
- CUDA development process
Most slides are courtesy Massimiliano Fatica (NVIDIA).
CUDA Bandwidth on the GeForce 8800GTX
CUDA provides a huge amount of computation and memory bandwidth to support massive parallel computation. As an example, the following table summaries the bandwidth on GeForce 8800 GTX
| Operating Frequency | Operation/Data per clock cycle | Bandwidth | |
|---|---|---|---|
| ALU | 1.35GHz | 16SM*(8SP*2MADD+3 SFU) | ~388 GFLOPs |
| Register | 1.35GHz | 16SM*8SP*4Word | 2.8 TBytes/s |
| Shared Memory | 575MHz | 16SM*16Banks*1Word | 588 GBytes/s |
| Device Memory | 1.8GHz | 384 Bits | 86.4 GBytes/s |
| Host Memory | --- | --- | 1.5 GB/s or 3GB/s |
Apparently, the host memory bandwidth is the bottleneck of the overall performance. In other words, the overall performance could be significantly limited by the IO interface between GPU and CPU. To highlight this idea, let’s look at the single-precision general matrix multiplication(SGEMM) shown in the following figure.We have the follwing observations:
- As the dimension of the matrix increases, the percentage of IO activity versus the computational activity decreases. Therefore, the performance limiting effect of IO becomes less prominent, leading to the increase of computation bandwidth.
- The performance of GPU only algorithm consistently delivers the highest bandwidth, followed by GPU+IO Pinned and GPU+IO.
Therefore, to fully exploit the bandwidth CUDA has to offer, we need to minimize the IO demands of the program.
CUDA Overview
Compute Unified Device Architecture is a programming system for utilizing the G80 processor for compute. It supports general purpose programming model that allows users to:
- Run batches of threads on GPU
- Turn GPU into dedicated super-thread, massively data parallel coprocessor
In this section, we will briefly introduce CUDA programming system and CUDA API, then go through CUDA programming model in detail.
CUDA Programming System
CUDA programming system is essentailly a software stack consists of a hardware driver, an application programming interface and its runtime, and two higher-level mathematical libraries of common usage, CUFFT and CUBLAS. Application can call functions from CUDA libraries or talk to the CUDA Runtime directly. The CUDA runtime eases device code management by providing implicit initialization, context management, and module management.CUDA driver is responsible for loading the computation programs to GPU, and has the following features:
- Standalone Driver -Optimized for computation
- Interface designed for compute -graphics free API
- Data sharing with OpenGL buffer objects
- Guaranteed maximum download & readback speeds
- Explicit GPU memory management
CUDA API: an extension to the ANSI C
The CUDA API comprises an extension to the C programming language for a minimum learning curve. It consists of:
- A minimal set of extensions to the C language that allow the programmer to target portions of the source code for execution on the device
- A runtime library split into:
- A host component that runs on the host and provides functions to control and access one or more compute devices from the host;
- A device component that runs on the device and provides device-specific functions;
- A common component that provides built-in vector types and a subset of the C standard library that are supported in both host and device code.
The following table gives a snapshot of the extensions CUDA makes to ANSI C.
| Categories | Extensions | Code examples |
|---|---|---|
| Declaration Specs | global, device, shared, local, host | __host__ float hostFunc(); __global__ void KernelFunc(); __device__ float DeviceFunc(); __shared__ float region[M]; |
| Keywords | threadIdx, blockIdx | region[threadIdx] = image[i]; |
| Intrinsic | _syncthreads | __syncthreads() |
| Runtime API | memory, symbol,execution management | cudaMalloc((void**)&Md.elements, size); cudaFree(Md.elements); cudaMemcpy( d_A, h_A, N * sizeof(float), cudaMemcpyHostToDevice)); |
| Function launch | --- | //500blocks, 128 threads per block KernelFunc«< 500, 128 »>(…); |
CUDA Programming Model
The basic idea of GPU programming is that GPU is treated as a computing coprocessor to the CPU or host. The data parallel portions of an application are executed on GPU as kernels, which run in parallel with many threads. However, GPU thread is different with CPU thread in the sense that:
- GPU thread is extremely lightweight, which means it has very little overhead in creating a thread
- GPU needs thousands of threads to fully exploit its efficiency, while Multi-core CPU only needs a few
The following sub-sections would briefly talk about how threads are organized in GPU, what is the memory hierarchy of GPU, and how to leverage the features of GPU to high performance parallel computing.
Thread Batching
The threads in CUDA are organized with grids and blocks.
- A grid is a collection of thread blocks. All threads in the grid share data memory space. One kernel is mapped to only one grid.
- A thread block is a batch of threads that can cooperate with each other by either execution synchronization or data sharing through shared memory.
- Two threads from different thread blocks can not cooperate with each other.
As shown in the following figure, the thread blocks in a grid are organized as 1D or 2D arrays. One can refer to a particular thread block by using the block index, or in other words, block ID. Similarly, the threads in a thread block is organized as 1D, 2D or 3D arrays. Each thread can be indexed with its thread ID. This thread organization practically simplifies memory addressing when processing multi-dimensional data.
CUDA provides _syncthreads() API for thread synchronization within thread block. When writing to the common resource, multiple threads in a warp would be serialized. However the thread behavior is undefined between warps. In addition, there is no explicit way to synchronize the thread blocks. You can use kernal boundary as the implicit barrier for all blocks of threads.
Memory Space
The following figure shows the thread’s accessibility to different levels of memory hierarchy. As a review of CUDA memory hierarchy, we highlight a few things here:
- Register and shared memory have high bandwidth
- Local memory is logically partitioned among different thread blocks, yet physically sharing offchip memory, therefore it has high access latency.
- Since local memory is private to each thread, it can not be clobbered by other threads. However, it is possible for multiple threads writing to the same location in gobal memory. To prevent race condition, you have to create locks on top of the shared resource.
- Both constant and texture memory are read only from GPU’s perspective. Constant memory has the capability of broadcasting data to different multi-processors.
- The contents in constant memory and texture memory are loaded by the host. The host is also able to read and write the global memory.
It should be noted that different thread blocks could be mapped to the same streaming multiprocessor. That means shared memory is divided among different thread blocks. Therefore, the larger the number of blocks being mapped to the same streaming multi-processor, the smaller the size of the shared memory each thread block could have.
The following table summaries the memory access latency of different memory hierarchy
| Memory Hierarchy | Type | Speed |
|---|---|---|
| Register | Dedicated HW | Single cycle |
| Shared memory | Dedicated HW | Two cycle |
| Local memory | DRAM, no cache | slow |
| Global memory | DRAM, no cache | slow |
| Constant memory | DRAM, cached | Depending on cache locality. 1… 10s.. 100s of cycles. |
| Texture memory | DRAM, cached | Depending on cache locality. 1… 10s.. 100s of cycles. |
| Instruction memory | DRAM, cached | — |
Since local memory and global memory reside in device memory, which is DRAM, their access speed is much slower than shared memory. In order to take advantage of the fast shared memory, a common way is to partition the data set into subsets that can fit into shared memory. Each data set can be operated with one thread block by:
- Loading the subset from global memory to shared memory
- Performing the computation on the subset from shared memory; each thread can efficiently multi-pass over any data element
- Copying results from shared memory to global memory
Example: Square matrix multiplication
This example demonstrates how to use data partitioning to save memory bandwidth. Assume we are going to perform matrix multiplication of two square matrixes M and N with size of WIDTH*WIDTH, and store the results to matrix P. There are two ways to do this.
- One straightforward way is that one thread handles one element in P without matrix blocking. That means each thread has to fetch one row of M and one column of N from global memory. In total, M and N would be loaded with WIDTH times from global memory.
- Another way is to partition the matrix into several data blocks, with each thread block operating on each data block, as shown in the following figure. The data blocks are small enough so that one block from M and one block from N can both fit into the shared memory. Once the blocks are in shared memory, the multiplication of these two blocks can be operated on top of shared memory, and each element in these blocks could be reused for BLOCK_SIZE times. Therefore, M and N only need to be loaded with WIDTH/BLOCK_SIZE time from the global memory
CUDA Development Process
Parameterize the application
Parameterization is important during the application development process since it improves the portability of the application to different GPUs. Note that GPU programs are often optimized to specific GPU configurations. However, GPUs may vary in many ways:
- Number of multiprocessor
- Shared memory size
- Register file size
- Threads per block
- Memory bandwidth
With parameterization, the programs can be easily adapted to different GPU architectures by changing the values of hardware specific parameters.
Compilation
Any source file containing CUDA language extensions must be compiled with nvcc.Nvcc is a compiler driver that invokes all necessary tools and compilers. It can generate:
- Either C code (CPU code), which must be compiled with another tool, like g++.
- Or PTX (Parallel Thread eXecution ) object code directly.
Also, any executable with CUDA code requires two dynamic libraries:
- The CUDA runtime library (cudart)
- The CUDA core library (cuda).
The following figure shows the compilation flow. The steps before compiling PTX to GPU code are called virtual steps, since these steps do not need information about the configuration of the target GPU; the step of compiling PTX to target GPU code is called physical step, since this step does target specific optimization and resource allocation.
Another thing need to mention about NVCC is that NVCC is composed of EDG and Open64. The former separates the GPU and CPU code, while the latter generates the GPU PTX assembly. More importantly, Open64 provides:
- A complete C/C++ compiler framework, which allows us not to change the infrastructure framework as the micro-architecture advances over time.
- A good collection of high level architecture independent optimizations. Compiler infrastructure that interacts well with other related standardized tools.
Debugging with Emulation
Device emulation mode allows the executable to run completely on the host using CUDA runtime without any device and CUDA drivers. When running with CUDA driver, one can:
- Use host native debug support (breakpoints, inspection, etc.)
- Access any device-specific data from host code and vice-versa
- Call any host function from device code (e.g. printf) and vice-versa
- Detect deadlock situations caused by improper usage of __syncthreads
However, there are some pitfalls when using device emulation mode:
- Emulated device threads execute sequentially, so simultaneous accesses of the same memory location by multiple threads potentially produce different results
- Dereferencing device pointers on the host or host pointers on the device can produce correct results in device emulation mode, but will generate an error in device execution mode
- Results of floating-point computations will slightly differ because of:
- Different compiler outputs
- Different instruction sets
- Use of extended precision for intermediate results
Summary
This lecture gives an introdution of CUDA programming system as well as its program development process.
