Take CUDA C++ from syntax to hardware intuition
__global__ and threadIdx take half an hour. The real bar is looking at a kernel and knowing where it is stuck. This course goes write-it-correct, then write-it-fast, then read the compiler output, then ship it. Every optimization names the bottleneck it removes.
You can finish the whole course without a GPU: the four labs simulate the execution model in the browser.
__global__ void add(const float* a, const float* b, float* c, int n) { int i = blockIdx.x * blockDim.x + threadIdx.x; if (i < n) c[i] = a[i] + b[i];} add<<<(n + 255) / 256, 256>>>(d_a, d_b, d_c, n);Learning path
Take them in order. Each stage builds on the last.Getting started
Write, compile, and run your first kernel3 lessons · 60 minYour first kernel: the CPU tells the GPU to work
How host and device split the work, what __global__ means, and why a kernel launch is asynchronous
Thread hierarchy: index arithmetic for grid / block / thread
Map 1D and 2D data onto threads, and why every kernel needs a bounds check
Device memory and error checks: CUDA that does not hide failures
The right way to call cudaMalloc / cudaMemcpy, unified memory, and the macro you should copy once
Execution model and memory
Warps, the memory hierarchy, coalescing, shared memory: where performance comes from4 lessons · 92 minWarp and SM: the GPU's real unit of execution
The SIMT model, warp divergence, latency hiding, and why the number 32 shows up everywhere
Memory hierarchy: the table that sets your performance ceiling
Latency and bandwidth of registers, shared memory, L1/L2, and HBM, plus arithmetic intensity as a diagnostic
Memory coalescing: the same kernel, a 5x gap
How a warp issues memory transactions, what stride access is, and the AoS vs SoA choice
Shared memory and bank conflicts: using on-chip storage correctly
A full derivation of tiled matmul, and that mysterious [TILE][TILE + 1]
Optimization in practice
Seven-step reduction, occupancy tradeoffs, pipeline overlap3 lessons · 76 minSeven-step reduction: 30x from one kernel
From the naive version all the way to warp shuffle, with a clear reason for each speedup
Occupancy: why 'higher is better' is a misunderstanding
The three occupancy limiters, how to compute them, and when you should deliberately lower occupancy
Streams and async: overlap copies with compute
Stream semantics, event timing, a three-stage pipeline, and CUDA Graph
Going lower
Warp primitives, PTX and SASS: see what the compiler emitted2 lessons · 46 minProduction and AI
Profiling methodology, PyTorch ops, Tensor Cores and Triton3 lessons · 66 minProfiling: replace guesses with data
How Nsight Compute and Nsight Systems split the work, the metrics that matter, and how to read a roofline
Plug into PyTorch: turn your kernel into an operator
From load_inline for a quick check, to a proper setuptools extension and autograd
Next steps: from writing kernels to knowing the GPU
Tensor Core, CUTLASS, Triton, FlashAttention: reading order and how to choose
Interactive labs
Turn abstractions into something you can dragThread index visualizer
How grid / block / thread map onto data
Drag gridDim and blockDim, then click any thread to see which element it owns. 1D and 2D grids, with out-of-range threads shown live.
Memory coalescing simulator
How many sector transactions one warp produces
Adjust stride, start offset, and access width. Watch the 32 addresses, the sector count, and the bandwidth efficiency.
Bank conflict simulator
How a warp hits the 32 shared-memory banks
Pick an access pattern and padding, then see how many threads land on each bank, the worst conflict degree, and how many cycles it takes.
Occupancy calculator
Is it registers, shared memory, or thread count that is capping you
Pick an architecture, set block size, registers per thread, and shared memory per block. Get theoretical occupancy and the limiting resource.
Every important kernel in the lessons has a full runnable copy under cuda/, with timing, result checks, and multi-version performance comparisons. On a machine with a GPU, make builds all of them. Without a GPU you can still finish the course.
cd cudamake # compile every sample for the detected architecturemake run # run them in order./bin/06_reduction # seven reduction versions, timed