(Advanced) CUDA Programming Course

Welcome to the CUDA Programming Course. This course covers high-performance kernel development for modern NVIDIA GPUs, from core concepts to the newer features in Ampere, Hopper, and Blackwell architectures.

The exercises are downloaded and run against your own GPU: how the exercises work. No GPU of your own? The gpu-submit toolkit runs your CUDA on a datacentre GPU from your laptop.

Course Outline

Part 1 — Introduction

• 1.1 GPUs vs CPUs
• 1.2 The CUDA runtime
• 1.3 Moving data
• 1.4 The memory hierarchy
• † Unified addressing
• 1.5 The toolchain
• 1.6 Recap
• 1.7 Exercise — Copying a rectangle

Part 2 — Basic Kernels

• 2.1 A first kernel
• 2.2 Threads, Blocks, and Grids
• † Thread block clusters
• 2.3 Automatic scaling
• 2.4 Launching a kernel
• 2.5 What nvcc actually does
• 2.6 Debugging kernel errors
• 2.7 Exercise — SAXPY

Part 3 — Divergence and Coalescing

• 3.1 Single instruction, multiple threads
• 3.2 Divergence
• Case — Mandelbrot
• 3.3 Coalescing
• Case — RGB to greyscale
• 3.4 Profiling
• 3.5 Exercise — Embedding bags
• † SIMD and SIMT

Part 4 — Pipelining and Occupancy

• 4.1 Pipelining and latency hiding
• 4.2 Occupancy
• 4.3 Register reuse
• 4.4 Exercise — Matrix multiplication

Part 5 — Shared Memory

• 5.1 What we can already do
• 5.2 Making the cooperation local
• 5.3 Shared memory
• 5.4 Exercise — Parallel Reduction
• † Dynamic shared memory
• 5.5 Banks and bank conflicts
• 5.6 Matrix transpose
• 5.7 Exercise — Matrix Transpose
• 5.8 Recap

Part 6 — Thread Coarsening and Vectorized Memory Access

• 6.1 Per-thread overhead
• † Instruction breakdown
• 6.2 Bytes in flight
• 6.3 Thread coarsening
• 6.4 __launch_bounds__, __restrict__ and #pragma unroll
• † The effect of restrict at the assembly level
• 6.5 Persistent Kernels
• 6.6 Vectorized load and store operations
• 6.7a Exercise — AXPY with bfloat16
• 6.7b Exercise — AXPY with misaligned data

Part 7 — Warp Shuffles, Reductions, and Cooperative Groups

• 7.1 Reduction: from Shared Memory to Warp Shuffles
• † Floating-point reductions
• 7.2 Masks, __activemask, and Why It Is Not Enough
• 7.3 Cooperative Groups
• 7.4 Exercise — Row-wise Softmax in BF16

Part 8 — Asynchronous Data Movement: LDGSTS

• 8.1 The Problem with Synchronous Loads
• 8.2 Asynchronous Copy to Shared Memory
• 8.3 Software Pipelining
• 8.4 Putting It Together — Pipelined SAXPY
• † L1 Bypass
• 8.5 Exercise — Max pooling

Glossary.