Beating cuBLAS on an H100

Beating cuBLAS on an H100

NVIDIA’s Hopper H100 microarchitecture introduced two revolutionary hardware subsystems:

1. Tensor Memory Accelerator (TMA)

In pre-Hopper architectures, threads in a warp explicitly load data from global memory into registers, then write from registers into shared memory. TMA executes multidimensional tensor transfers (e.g. 2D/3D tiles) directly between global HBM and shared memory asynchronously in hardware, completely bypassing the register file and consuming zero SM ALU cycles.

2. Warp Specialization

Instead of every thread in a block performing both memory loading and arithmetic: - Producer Warps (1-2 warps): Issue asynchronous TMA copy instructions to fetch upcoming tiles into double-buffered shared memory. - Consumer Warps (Warp Groups): Execute continuous WGMMA matrix multiply instructions without ever pausing for memory latency. - Hardware barrier arrive/wait primitives (cuda::barrier) synchronize producers and consumers with cycle-level precision.