Learn CUDA: The Best GPU Programming Books, in Order
This curriculum is designed for expert-level programmers who already have strong C/C++ and systems programming backgrounds, diving straight into GPU architecture, CUDA kernel design, and high-performance computing. The three stages move from CUDA's core programming model and memory hierarchy, through advanced optimization and parallel algorithms, to production-grade HPC and numerical computing — each stage building directly on the mental models and vocabulary of the last.
CUDA Core Model & Memory Mastery
IntermediateEstablish a rigorous mental model of the CUDA execution model, thread hierarchy, memory spaces, and kernel optimization fundamentals — the vocabulary every advanced topic depends on.
▸ Study plan for this stage
Pace: 8–10 weeks, ~40–50 pages/day (alternating between both books; ~2 weeks per major topic cluster)
- Thread hierarchy: grids, blocks, threads, and how they map to GPU hardware (warps, SMs)
- CUDA memory spaces: global, shared, constant, local, and registers—access patterns and latency implications
- Kernel execution model: how threads are scheduled, synchronized, and how divergence impacts performance
- Memory coalescing and bandwidth optimization—the difference between naive and optimized memory access
- Synchronization primitives: __syncthreads(), atomic operations, and race condition avoidance
- Occupancy and resource constraints: registers, shared memory, and how they limit thread block size
- Compute capability and hardware differences—why code written for one GPU may not scale to another
- Performance profiling and analysis: using NVIDIA tools to measure bottlenecks and validate optimizations
- Explain the relationship between threads, blocks, and grids. How does a 2D block layout map to a 1D warp structure on the GPU?
- Why does memory coalescing matter, and what patterns of global memory access are efficient vs. inefficient? Give a concrete example.
- What is occupancy, and why is it not always the best metric for kernel performance? How do register pressure and shared memory constraints affect it?
- Describe the differences between global, shared, constant, and local memory in terms of scope, latency, size, and typical use cases.
- What happens when threads in a warp diverge (take different code paths)? How does this affect performance and how can you minimize it?
- How would you use __syncthreads() safely? What are the pitfalls of synchronization in conditional code?
- Implement a simple vector addition kernel and measure bandwidth; then optimize memory access patterns and re-measure to see the difference (from 'CUDA by Example' Chapter 4–5 exercises)
- Write a matrix multiplication kernel using global memory, then rewrite it with shared memory tiling; profile both and explain the performance gap
- Create a kernel with intentional warp divergence (e.g., if-else based on thread ID), measure performance, then refactor to eliminate divergence
- Implement a parallel reduction kernel (sum, max, etc.) using shared memory and __syncthreads(); test correctness and measure occupancy vs. performance
- Write a kernel that reads from constant memory vs. global memory in a broadcast pattern; measure latency and bandwidth differences
- Analyze a provided inefficient kernel using NVIDIA Nsight or nvprof; identify bottlenecks (memory, compute, occupancy) and propose optimizations
Next up: Mastering the execution model and memory hierarchy here provides the foundation to tackle advanced optimization techniques (instruction-level parallelism, tensor operations, graph-based scheduling) and domain-specific libraries in the next stage.
A concise, hands-on entry into CUDA's programming model that quickly covers threads, blocks, shared memory, and streams — ideal for experts who want to get productive fast without padding.
The definitive academic-and-practitioner textbook on CUDA; covers the full memory hierarchy, warp execution, tiling, and performance reasoning in rigorous depth. Read this second to solidify and formalize what Sanders introduced.
Parallel Algorithms & Advanced Optimization
ExpertDesign and analyze high-performance parallel algorithms — reductions, scans, sorting, stencils — and master advanced CUDA features like dynamic parallelism, streams, and occupancy tuning.
▸ Study plan for this stage
Pace: 8–10 weeks, ~40–50 pages/day with code walkthroughs and implementation labs
- Parallel reduction algorithms: tree-based, warp-level, and atomic-free implementations for correctness and performance
- Scan (prefix sum) algorithms: inclusive/exclusive scans, work-efficient implementations, and their role in sorting pipelines
- Sorting networks and GPU sorting: bitonic sort, merge sort, and radix sort with memory access patterns and bank conflict avoidance
- Stencil computations: halo exchange, boundary conditions, and tiling strategies for iterative solvers and image processing
- Dynamic parallelism: nested kernel launches, recursion, and when to use it versus flattening kernels
- Stream management: concurrent kernel execution, asynchronous memory transfers, and dependency graphs for multi-GPU scaling
- Occupancy tuning: register pressure, shared memory allocation, thread block sizing, and profiling with NVIDIA tools
- Advanced memory optimization: coalescing patterns for non-uniform access, texture caching, and unified memory trade-offs
- How do tree-based reductions minimize thread divergence and synchronization overhead compared to naive implementations?
- What is the difference between inclusive and exclusive scans, and why is work-efficiency critical for large-scale parallel scans?
- How does bitonic sort leverage GPU memory hierarchy and what are its complexity and communication costs?
- Explain the halo exchange pattern in stencil computations: how do you minimize redundant communication and synchronization?
- When is dynamic parallelism beneficial versus harmful, and how do you flatten recursive algorithms for better performance?
- How do CUDA streams enable overlapping computation and communication, and what synchronization primitives ensure correctness?
- What metrics (registers per thread, shared memory per block, warp occupancy) determine kernel occupancy, and how do you trade off occupancy for performance?
- How do you profile and optimize memory access patterns to achieve near-peak bandwidth on modern GPUs?
- Implement a tree-based parallel reduction for sum, max, and min; profile with nsys and compare against a naive reduction.
- Code an inclusive and exclusive scan using the Blelloch algorithm; measure work efficiency and memory bandwidth utilization.
- Implement bitonic sort for small arrays and radix sort for larger datasets; benchmark sorting time and memory throughput.
- Build a 2D stencil kernel (Jacobi iteration or Laplacian) with halo exchange; optimize for bank conflicts and measure speedup.
- Write a recursive algorithm (e.g., quicksort or tree traversal) using dynamic parallelism; flatten it and compare performance and code complexity.
- Create a multi-kernel pipeline with CUDA streams: overlap a compute kernel with asynchronous H2D and D2H transfers; measure latency hiding.
- Profile a kernel with varying block sizes and register usage; construct an occupancy vs. performance trade-off curve using Nsight Compute.
- Optimize a non-coalesced memory access pattern (e.g., transpose or gather) using shared memory tiling; measure bandwidth improvement.
Next up: This stage equips you with the algorithmic foundations and optimization discipline to tackle domain-specific applications (machine learning, scientific simulation, graphics) where you'll apply these patterns at scale and integrate them into larger systems.
Bridges the gap between basic CUDA and real optimization work, covering profiling, memory bandwidth analysis, and algorithmic patterns that expert programmers need to write production kernels.
Focuses on the full application lifecycle — profiling with Nsight, multi-GPU strategies, and integrating CUDA into larger software systems — building directly on the algorithmic foundations of Cook.
High-Performance & Scientific Computing on GPUs
ExpertApply GPU programming expertise to large-scale HPC workloads: numerical linear algebra, FFTs, scientific simulations, and multi-GPU/cluster-scale parallelism.
▸ Study plan for this stage
Pace: 12–14 weeks, ~40–50 pages/day with 2–3 days/week for hands-on implementation
- Optimizing memory hierarchies (global, shared, texture, constant memory) for bandwidth-bound and compute-bound kernels in large-scale workloads
- Numerical linear algebra primitives (matrix multiplication, LU decomposition, SVD) and their GPU-accelerated implementations using CUDA
- Fast Fourier Transform (FFT) algorithms, in-place vs. out-of-place computation, and batched FFT operations for scientific computing
- Warp-level and block-level synchronization patterns, atomic operations, and reduction algorithms for correctness at scale
- Multi-GPU programming: peer-to-peer transfers, unified virtual addressing (UVA), and domain decomposition for distributed HPC workloads
- Performance profiling and optimization workflows: identifying bottlenecks, roofline analysis, and achieving peak bandwidth/compute utilization
- Numerical stability and precision considerations in GPU-accelerated scientific simulations (floating-point error accumulation, convergence criteria)
- Cluster-scale parallelism: combining CUDA with MPI, overlapping computation and communication, and scaling to thousands of GPUs
- How do you optimize a memory-bound kernel like matrix transpose or reduction to achieve near-peak bandwidth on modern GPUs, and what role do shared memory and coalescing play?
- Describe the algorithmic differences between row-major and column-major matrix multiplication on GPUs, and explain how tiling strategies reduce global memory traffic
- What are the key numerical challenges when implementing iterative solvers (e.g., conjugate gradient) on GPUs, and how do you ensure convergence and stability across multiple GPU nodes?
- Explain the in-place FFT algorithm and how bit-reversal permutation, butterfly operations, and memory access patterns affect GPU performance for large-scale transforms
- How does unified virtual addressing (UVA) simplify multi-GPU programming, and what are the trade-offs between UVA and explicit peer-to-peer transfers?
- Design a domain decomposition strategy for a 3D finite-difference stencil computation across 4 GPUs, including halo exchange and communication overlap
- Implement a tiled matrix multiplication kernel (e.g., 32×32 or 64×64 tiles) and profile it against cuBLAS; measure bandwidth utilization and identify memory bottlenecks using Nsight Compute
- Optimize a reduction kernel (sum, max, or dot product) using warp-level primitives (__shfl_down, __shfl_xor) and compare performance against a naive global-memory version
- Implement a batched LU decomposition using Gaussian elimination with partial pivoting; validate numerical accuracy against LAPACK and measure speedup over CPU
- Write an in-place radix-2 FFT kernel for 1D signals (power-of-2 length); verify correctness against FFTW and benchmark on various sizes (2^10 to 2^20)
- Implement a 2D or 3D finite-difference stencil (e.g., Laplacian, heat equation) with shared-memory optimization; profile memory access patterns and compute utilization
- Develop a multi-GPU Jacobi or conjugate-gradient solver using NCCL or MPI+CUDA for peer-to-peer communication; measure weak and strong scaling across 2–4 GPUs
Next up: This stage equips you with production-grade GPU optimization techniques and distributed-memory parallelism, preparing you to tackle specialized domains (e.g., deep learning frameworks, graph processing, or real-time physics simulation) or to architect custom HPC solutions for domain-specific problems.
A curated collection of expert-authored chapters covering real HPC applications — molecular dynamics, image processing, financial computing — showing how CUDA is applied at scale across domains.
The canonical reference for numerical algorithms (FFTs, linear solvers, ODEs); reading this alongside GPU Gems lets you map classical HPC algorithms onto GPU architectures with full mathematical grounding.
The most comprehensive low-level reference for CUDA internals — hardware architecture details, PTX, memory subsystem, and concurrency — essential for squeezing peak performance out of NVIDIA GPUs at the expert level.
Discussion
Keep reading
Paths that share books, cover the same subject, or open a related topic.
