Abstract:

This document provides a systematic introduction to General-Purpose GPU (GPGPU) programming using NVIDIA's CUDA platform. It begins by refreshing the architectural differences between CPUs and GPUs, highlighting the GPU's suitability for data-parallel processing despite the PCIe data transfer bottleneck. The core of the CUDA programming model is explained, including the concepts of the host (CPU) and device (GPU), and the use of kernels (__global__) for device-side execution.

The guide details CUDA's fundamental abstractions: the thread hierarchy (grids, blocks, threads, warps) and the memory hierarchy (global, shared, constant, and local memory). It covers practical aspects of memory management, such as dynamic allocation (cudaMalloc), data transfer (cudaMemcpy), and the use of specialized memory types like __shared__ and __constant__ for performance optimization. Advanced topics are then explored, including the performance implications of branching (warp divergence), efficient inter-thread communication using warp shuffle operations, and maximizing GPU occupancy by overlapping computation and data transfers with CUDA Streams. Finally, the document outlines strategies for utilizing multiple GPUs, from manual device selection and data partitioning to simplified memory management using Unified Memory.

An Introduction to GPGPU Programming with CUDA

• Chapter 21: GPU vs. CPU Architecture: GPUs excel at data-parallel tasks due to their architecture, which features many cores sharing a single control unit. In contrast, CPUs are more versatile, with each core having its own control unit, making them better for serial or complex decision-making tasks. A key bottleneck in GPGPU programming is the data transfer speed between the CPU and GPU via PCIe.
• Chapter 21: CUDA Programming Basics: CUDA is NVIDIA's platform for GPGPU programming. It defines the CPU as the host and the GPU as the device.
  • Functions executed on the device are called kernels and are defined with the __global__ specifier.
  • Functions called from and executed on the device use the __device__ specifier.
  • Code is compiled with the NVIDIA CUDA compiler (nvcc).
• Chapter 21 & 22: Thread Hierarchy: CUDA organizes execution into a three-level hierarchy to manage thousands of threads.
  • Threads: The fundamental unit of execution.
  • Thread Blocks: Groups of threads (up to 1024) that execute on a single Streaming Multiprocessor (SM) and can share memory and be synchronized.
  • Grids: Composed of multiple thread blocks. A single kernel launch creates one grid.
  • Threads are executed in groups of 32 called warps. For optimal performance, all threads in a warp should execute the same instruction to avoid "warp divergence."
• Chapter 22: Memory Management: Because the host and device have separate memory spaces, data must be explicitly managed.
  • cudaMalloc(): Allocates memory on the device's global memory.
  • cudaMemcpy(): Copies data between the host and device. Minimizing these transfers is crucial for performance.
  • cudaFree(): Deallocates device memory.
• Chapter 22: Memory Hierarchy: CUDA provides different types of memory with varying scope, speed, and size.
  • Shared Memory (__shared__): A small, fast memory space shared by all threads within a block. It is useful for coordinating and sharing data to reduce reliance on slower global memory.
  • Constant Memory (__constant__): Read-only memory optimized for when all threads in a warp read from the same location.
• Chapter 22: Error Handling: All CUDA runtime API calls return a cudaError_t value, which should be checked against cudaSuccess to ensure operations complete without issues. cudaDeviceSynchronize() is used to block the host until a kernel has finished, allowing for proper error checking.
• Chapter 23: Performance and Advanced Features:
  • Branching: Conditional statements (if-else) can cause threads within a warp to take different execution paths (warp divergence), serializing their execution and significantly reducing performance.
  • Warp Shuffling: Allows threads within the same warp to directly exchange data without using shared memory, offering a highly efficient communication method in a single instruction.
  • CUDA Streams: Enable asynchronous execution by placing operations (memory copies, kernel launches) into a queue. This allows for overlapping data transfers with computation, hiding latency and maximizing GPU utilization.
• Chapter 23: Using Multiple GPUs: Problems can be divided across multiple devices to increase computational power.
  • Unified Memory (cudaMallocManaged()): Simplifies multi-GPU programming by creating a single memory space accessible by both the CPU and all GPUs. The CUDA runtime automatically manages data migration between devices, hiding the complexity of manual memory transfers.