WGMMA Part 2

Why lifecycle control matters

wgmma.mma_async returns immediately. That is the whole reason Hopper can overlap tensor math with TMA loads, pointer work, and staging for the next tile. But once math is asynchronous, the program also needs a way to bundle launches, cap how much is in flight, and know when accumulators are actually safe to consume.

Commit, wait, and fence define the WGMMA lifecycle

Grouping gives WGMMA pipeline-friendly granularity. The current batch of async MMAs stays “open” until you close it with wgmma.commit_group.sync.aligned. Later, wgmma.wait_group lets the warpgroup stall only when too many committed batches are still in flight.

// One output-tile worth of async math
wgmma.mma_async ...
wgmma.mma_async ...
wgmma.commit_group.sync.aligned;

// Prepare the next tile while some math still runs
// TMA, pointer math, shared-memory staging, etc.
wgmma.wait_group.sync.aligned 1;

// Before reading accumulators or exporting them
wgmma.wait_group.sync.aligned 0;

The lesson is careful about that last distinction. wgmma.fence is about the WGMMA pipeline itself, especially accumulators and register-resident A fragments. If shared-memory contents were produced through the async proxy path, you still need the right cross-proxy publication rule before consuming them.

stmatrix is how opaque accumulator fragments become a usable tile

When WGMMA finishes, the result is not sitting in registers as a clean row-major matrix. It is fragmented across 128 threads in the exact layout tensor cores use internally. stmatrix is the cooperative store instruction that writes those fragments back to shared memory in a normal matrix layout.

// From the repo's matmul_12.cuh epilogue flow
// convert/pack fragments
// stmatrix.sync.aligned.m8n8.x4.trans.shared::cta.b16 [...]
// bar.sync across the warpgroup
// issue TMA store for the completed tile

FP32 accumulators are the important caveat

stmatrix does not support .f32 data. It supports packed .b16 and .b8 forms. So if WGMMA accumulated into FP32 registers, you cannot feed those registers directly into stmatrix. You first downcast and pack the values, which is why FP32 output paths often include an explicit conversion loop before the shared-memory store.

FP8 is a transport and layout problem as much as a precision choice

Hopper's FP8 path matters because it changes both math throughput and the operand contract. The lesson frames FP8 as a compressed training format that still accumulates in FP32, but whose range, scaling, packing, and layout decisions can make or break correctness and performance.

Scaling is the escape hatch that keeps 8-bit floating point useful. Tensor-wise, vector-wise, and block-wise schemes all exist because the raw FP8 range is too cramped to trust without a quantization policy. The notes also call out .satfinite conversion forms so out-of-range values clamp instead of exploding into unusable results.

// Representative PTX conversion style
cvt.rn.satfinite.e4m3x2.f16x2 ...
cvt.rn.satfinite.e5m2x2.f32 ...

Hopper FP8 WGMMA has stricter operand rules

The course notes call out the fixed m64nNk32 FP8 shape and a crucial Hopper limitation: the obvious ldmatrix ... .b8 path is not what SM90a exposes here. That is why practical Hopper FP8 kernels tend to use either shared-memory descriptors for both operands or an RS path where A is loaded into registers with ordinary shared loads already packed in the layout WGMMA expects.

The same notes also stress the strict K-major rule. Unlike the FP16/BF16 forms, FP8 WGMMA does not give you the same transpose escape hatch. If the staged operand tiles are not already K-major, you pay for extra repacking and synchronization before the tensor cores can consume them.

Sparse WGMMA adds a second contract: metadata correctness

Sparse WGMMA is still matrix multiply-accumulate, but the hardware assumes a very specific structured sparsity pattern. The supported model is sparse A times dense B. A carries the compressed data plus the metadata that tells tensor cores where the surviving values belong.

The physical sparsity rule is 2:4. You do not just drop two values and hope for the best. You must store the survivors plus enough metadata for hardware to reconstruct their original positions inside each quartet. Wrong metadata does not degrade gracefully. It simply makes the tensor-core read the wrong sparse pattern.

Practical guidance

1. Think in issue windows. Batch async MMAs into logical groups, then choose a wait depth that preserves overlap instead of collapsing back to full serialization.
2. Drain only when you consume. Use wait_group 0 before reading accumulators, exporting them with stmatrix, or otherwise treating the results as complete.
3. Separate register ordering from memory visibility. wgmma.fence is not a substitute for the async proxy publication rules used by TMA-backed pipelines.
4. Treat FP8 as a layout contract. Scaling, packing, K-major staging, and conversion policy matter as much as the nominal data type.
5. Never hand-wave sparse metadata. Sparse tensor cores only work when the payload and selector logic match exactly.

Glossary

Continue the course

Lesson 7 completes the WGMMA instruction model. The next lesson zooms out from one primitive to the whole kernel: arithmetic intensity, warp specialization, circular buffers, persistent scheduling, and the pipeline shapes that keep Hopper busy for full GEMMs.