GPU Programming with CUDA

How GPU code really runs: grids of blocks of threads, warps of 32 and waves over SMs; why most kernels are memory-bound; coalescing, occupancy, shared memory and tiling; fusion and launch overhead; and how to read Nsight Compute and Nsight Systems.

An interactive AI Infrastructure lesson: 24 steps, about 35 minutes, on a live simulation in your browser.

The ML team wants to write their own CUDA kernels, and the platform team owns the box they will run on: ws-1, one H100. To review their work you need to know how GPU code runs. Start with the smallest kernel there is: c = a + b over 64 million floats.

A kernel is a function the CPU launches on the GPU. The launch names a grid of thread blocks: here 250,000 blocks of 256 threads, one thread per element. Every thread runs the same code and works out which element is its own from blockIdx.x, blockDim.x and threadIdx.x. The if (i < n) guard stops the threads of the last block that fall off the end.

What you will learn

1. Grids, blocks and warps

   • A first kernel: add two vectors: A kernel launch is a grid of blocks of threads, all running the same function. Each thread finds its own element from its block and thread index.
   • Warps of 32, waves of blocks: A warp is 32 threads executing one instruction together. An SM hides memory latency by switching between many resident warps, so it needs enough of them.
   • Drill: the global thread index

2. Memory-bound by nature

   • 768 MB moved for 64 MFLOP: Arithmetic intensity (FLOPs per byte of DRAM traffic) against the ridge point (peak FLOPS / peak bandwidth) tells you whether a kernel is limited by memory or by compute before you profile it.
   • Twice the maths per element: In a memory-bound kernel, time is bytes moved divided by bandwidth. To make it faster, move fewer bytes; doing less maths changes nothing.

3. Coalescing

   • Break it: complex numbers, real parts: A warp's loads are served in 32-byte sectors. When neighbouring lanes read neighbouring addresses the bytes are all used; any stride wastes the rest of each sector.
   • One field out of 32: The waste of a strided access is capped by the sector: for 4-byte values, stride 2 costs 2×, stride 4 costs 4×, stride 8 or more costs 8×.

4. Occupancy

   • Break it: 128 registers per thread: Registers, shared memory and block size each cap how many warps an SM holds. Too few warps means too few loads in flight, and the kernel waits on latency with the bandwidth idle.
   • 768 threads per block: Occupancy is a means, not the goal. A kernel needs enough work in flight to hide latency; beyond that, higher occupancy does not make it faster.
   • Sixteen blocks for 132 SMs: A grid with fewer blocks than the GPU has slots leaves SMs idle, however good each block is. Size grids from the SM count, not from a constant.

5. Shared memory and tiling

   • Break it: a transpose with bank conflicts: Shared memory serves a warp at full speed only when its lanes hit different banks. Column access to a power-of-two-wide tile serialises the warp.
   • One extra column: Padding a shared-memory tile by one column turns a 32-way conflict into none: the classic fix, worth knowing by sight in a code review.
   • Matrix multiply: naive, then tiled: Tiling moves data reuse from DRAM into on-chip memory. Each level you fix exposes the next roof: L2, then shared memory, then the arithmetic units.
   • Register tiles: 77% of FP32: Fast kernels keep data in the fastest level that can hold it: registers, then shared memory, then L2, then HBM. Reuse at each level multiplies the FLOPs per byte.
   • Tensor cores: cuBLAS in bf16: Tensor cores raise the compute roof about 15× for matrix multiplies in low precision. Real GEMMs reach 70–75% of that peak through libraries, not hand-written loops.
   • The tail effect: Work is scheduled in waves of blocks; a partial last wave leaves most SMs idle for a whole tile's time. Efficiency depends on the shape, not only the size.

6. Fusion and launch overhead

   • Four kernels or one: For memory-bound ops, every kernel boundary is a round trip to HBM. Fusing a chain of N elementwise ops cuts the traffic, and the time, about N-fold.
   • When the launches are the cost: A kernel launch costs a few microseconds of CPU time. When kernels are shorter than that, the GPU waits on the CPU: fuse them or replay them as a CUDA Graph.

7. Errors, profilers, libraries

   • Break it: launches that never run: A kernel launch returns immediately and reports nothing. Check cudaGetLastError() after it and the result of the next synchronise, or errors surface far from their cause.
   • Reading Nsight Compute: Profile top-down: nsys for the timeline and the expensive kernels, then ncu Speed Of Light for the bound, then the section it points to.
   • Drill: profile one kernel with ncu
   • Libraries first, custom kernels when: Use the library until a profile shows a hot op it does badly; write a custom kernel to fuse memory traffic away, not to beat cuBLAS at GEMM.

8. Recap & playground

   • Cheat sheet
   • Playground