CUDA Programming for NVIDIA H100s – Comprehensive Course

Solid C++ and CUDA foundations required

You must understand C++ syntax, CUDA thread hierarchies, and shuffle operations before attempting Hopper's asynchronous programming model.

Matrix tiling knowledge is mandatory

Understanding how matrices are tiled and multiplied forms the essential foundation for WGMMA operations and memory layout optimizations.

H100 introduces asynchronous execution paradigm

The architecture represents a significant shift from synchronous models, requiring new mental models for warp-group level computation and memory management.

Tensor Memory Accelerator handles bulk transfers

TMA uses tensor map descriptors to execute asynchronous global-to-shared memory copies without consuming SM execution resources.

cp.async.bulk instruction powers data movement

This PTX instruction enables structured and unstructured bulk copies with features like L2 cache hinting and multicast capabilities unique to Hopper.

M-barriers synchronize compute and memory

Hardware barrier primitives manage the massive speed gap between fast tensor cores and relatively slow memory accesses.

WGMMA operates at warpgroup granularity

Warpgroup Matrix Multiply Accumulate instructions replace warp-level operations, requiring careful register management and descriptor-based matrix fetching.

FP8 precision requires specific memory layouts

Hopper's FP8 implementation demands K-major layouts and specific packing strategies when accumulating in FP32.

Sparse operations support 2:4 structured sparsity

WGMMA can operate on structurally sparse tensors to effectively double computational throughput for compatible matrices.

Persistent scheduling maximizes tensor core utilization

Modern kernel design uses warp specialization, circular buffering, and persistent warps to eliminate idle cycles and maintain near-100% occupancy.

Production code analysis includes CUTLASS and fast.cu

The course examines SM90 pipeline implementations and a from-scratch kernel achieving 107% of cuBLAS performance.

Multi-GPU scaling covers NCCL and parallelism strategies

Training trillion-parameter models requires understanding NCCL primitives, NVLink topologies, and data/model/pipeline parallelism techniques.