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.
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.
