Glossary
Short definitions of every term this guide uses, in alphabetical order, with a pointer to the chapter where it matters most.
Arithmetic intensity
FLOPs performed per byte moved from memory. For a BF16 projection it is about m, the number of tokens. Below about 206 FLOP/byte an H200 kernel is limited by memory, above it by math. Chapter 1.7
Atomic operation
A read-modify-write in global memory that no other thread can interrupt, for example atomicAdd in CUDA or tl.atomic_add in Triton. It returns the old value. Used for split-K counters and for the router's atomic epilogue. Chapters 6.7, 9.4
Bank conflict
Shared memory is divided into 32 banks. When several lanes of a warp access different addresses in the same bank, the accesses are serialized. Padding an array row by one element (ecorr[16][NB·8 + 1]) avoids it. Chapter 8.1
BF16 (bfloat16)
A 16-bit float with 1 sign bit, 8 exponent bits and 7 mantissa bits: FP32's range with less precision. All weights except the router's, and all activations, are BF16. Chapter 7.1
Coalescing
When the 32 lanes of a warp access consecutive addresses, the hardware serves them as a few wide transactions. All our weight loads are 32 × 16 contiguous bytes. Chapter 2.4
CTA (thread block)
A group of up to 1,024 threads that runs on one SM, shares its shared memory and can synchronize with __syncthreads(). Our GEMVs use one CTA per 16 weight rows (times S K-slices). Chapter 2.2
cuBLAS
NVIDIA's linear-algebra library, which SGLang uses for these projections. Its Hopper kernels appear in profiles as nvjet_sm90_…. Chapter 3
CUDA graph
A recording of a sequence of GPU launches that can be replayed with almost no CPU overhead. The judge times replays of our kernel_fn, so it must not sync with the CPU or branch on data values. Chapters 1.4, 2.9
Decode and verify loads
Decode: m = 1, 16 or 32 requests generating one token each. Verify: 8 speculative draft tokens per request, so m = 8, 128 or 256. Chapter 1.2
Exception (12-bit format)
A weight whose top exponent bits fall outside the 8-value window (0.014 %: tiny weights and zeros). It is stored clamped, and the exact FP32 difference goes into a per-row list that an extra warp applies. Chapters 7, 8.2
Floor (speed of light)
The judge's lower bound for a shape: max(bytes ÷ 4.8 TB/s, FLOPs ÷ 989 TFLOP/s). Our compressed kernels can beat it, because it assumes 16-bit weights. Chapter 1.7
Fragment
The part of an mma operand tile that one lane holds in its registers. For m16n8k16, a lane holds 8 values of A, 4 of B and 4 accumulators. Chapter 2.5
GEMV / GEMM
General matrix-vector product, and general matrix-matrix product. With m ≤ 32 tokens, our "GEMV" kernels are really skinny GEMMs that treat x as a few vectors. Chapter 3
Grid
All the CTAs of one kernel launch, numbered by blockIdx. Ours is (N/16, S) for the GEMVs, and (N/BN, SPLIT, ⌈M/BM⌉) for Triton. Chapters 2.2, 4.5
HBM
High-bandwidth memory, the GPU's 141 GB main memory, at 4.8 TB/s. Every weight byte comes from here once per call, because the judge clears L2. Chapter 2.1
Lane
One of the 32 threads of a warp (lane = threadIdx.x & 31). In mma layouts, lane = 4·g + t. Chapter 2.5
__launch_bounds__
A hint to the compiler about CTA size and the minimum number of CTAs per SM. It limits registers per thread. Our PIPE 2 and 3 modes set it to 2 and 1. Chapter 6.2
L1 / L2 cache
L1 is small and private to each SM. L2 is 50 MB, shared by the whole GPU, and flushed by the judge before every timed call. __ldcg and .cg loads bypass L1 to read other SMs' fresh writes from L2. Chapter 2.3
Little's law
Bytes in flight = bandwidth × latency. Keeping 4.8 TB/s flowing with about 0.7 µs latency needs about 3.4 MB in flight, roughly 25 KB per SM. Chapter 2.4
load_inline
A PyTorch utility that compiles C++/CUDA source code from a Python string into an importable extension module, at run time. Chapter 4.3
lop3
A PTX instruction that computes any bitwise function of three 32-bit inputs, chosen by an 8-bit truth table (0xF8 = a | (b & c)). Chapter 7.4
mma.sync.m16n8k16
A warp-wide tensor-core instruction: D (16×8, FP32) = A (16×16, BF16) · B (16×8, BF16) + C. Our GEMVs put 16 weight rows in A and up to 8 tokens in B. Chapter 2.5
NB, NW, U, UG, PIPE, S
The GEMV knobs:
- NB: 8-token column blocks (1, 2 or 4);
- NW: streaming warps per CTA;
- U / UG: 32-wide K blocks / 64-wide K groups loaded per loop trip;
- PIPE: the loop style (0–3);
- S: K-splits across CTAs.
Occupancy
How many warps or CTAs an SM holds at once. It is limited by registers, shared memory and launch bounds. More resident warps hide more latency.
prmt
A PTX "permute bytes" instruction: builds a 32-bit word from any 4 of the 8 bytes of two inputs. We use it as an 8-entry lookup table and to interleave bytes. Chapter 7.4
PTX / SASS
PTX is NVIDIA's virtual instruction set, which asm volatile emits directly. SASS is the real machine code the driver compiles it to (see it with cuobjdump -sass). Chapter 2.8
Register
The fastest storage, private to a thread: up to 255 32-bit registers per thread, 65,536 per SM. Our GEMVs keep weights, x values and sums in registers. Chapter 2.3
Roofline
A chart of best possible throughput against arithmetic intensity: a sloped memory roof, a flat math roof, and the "ridge" where they meet. Chapter 1.7
Shared memory
Fast on-chip memory shared by the threads of one CTA (up to 228 KB per SM). Our GEMVs use it only for the final reduction across warps. Chapter 6.6
SM (streaming multiprocessor)
One of the H200's 132 processing units. Each runs one or more CTAs. A grid with fewer CTAs than SMs leaves hardware idle. Chapter 2.1
Software pipelining
Issuing the loads for the next loop iteration before computing the current one, so memory latency overlaps with work. These are our PIPE 1/2 loops, and Triton's num_stages. Chapter 6.5
Split-K
Dividing the K (reduction) dimension among several CTAs and adding their partial sums. It gives more parallelism to shapes with few rows, at the cost of combining the partials. Chapters 6.7, 9.4
Swap A and B
Putting the weight rows on the MMA's M side and the tokens on its N side (out ᵀ = W · xᵀ). With tiny m this wastes less of the tile than the usual orientation. Chapter 2.5
Template instantiation
Compiling a C++ template for one specific set of parameters. Every (NW, U, NB, DUAL, PIPE) combination our tables use must be listed (C3/Z3 lines), or the run fails with "unsupported … config". Chapter 6.9
Tensor core
Matrix-multiply hardware inside each SM, driven by mma.sync or wgmma: 989 TFLOP/s of BF16 on the H200. Chapter 2.5
TMA / TensorDescriptor
Hopper's Tensor Memory Accelerator copies whole 2-D tiles between global and shared memory with one instruction. A TensorDescriptor describes the tensor and tile shape for it. Chapter 9
Triton
A Python-embedded language for GPU kernels, where you program one tile ("program") with block operations. Used for our split-K GEMM. Chapter 9.2
Warp
32 threads that execute one instruction at a time together. The unit of scheduling, and of mma.sync. Chapter 2.2
wgmma
Hopper's warpgroup MMA: 4 warps together, operands in shared memory, much higher throughput for large tiles. cuBLAS and Triton use it; our shipped GEMVs don't. Chapter 2.5