100 DAYS OF GPU Course preview
Day 78Module 8
draft

Which Blackwell do you have: wgmma, tcgen05 and what your card cannot do

Compile this lesson's Hopper GEMM for the RTX 5090 target and ptxas reports:

ptxas /app/example.compute_120a.ptx, line 146; error   : Instruction 'wgmma.mma_async with floating point types' not supported on .target 'sm_120a'

Captured 2026-09-01 with NVCC 12.8.1; the repository contains the full transcript. The RTX 5090 is a Blackwell GPU, but consumer and datacenter Blackwell support different tensor core instructions.

This lesson maps those capabilities from primary sources. It also explains a warp-specialized wgmma GEMM and the compiler errors for unsupported targets.

Three tensor core instruction sets

The earlier tensor core lessons issued work by warp: WMMA's portable API on day 72, then mma.sync with explicit fragments on day 73. Hopper changed the issuing unit. wgmma.mma_async is issued by a warpgroup, four consecutive warps, 128 threads, and it is asynchronous: the issue returns, the tensor core works, and a separate wait retires the result.

Datacenter Blackwell changed it again: tcgen05 instructions compute out of a dedicated tensor memory instead of registers, and can use two SMs for one matmul. Consumer Blackwell uses neither instruction set: its tensor cores use day 73's mma.sync, with new input types.

Here is the map, every cell from a primary source (compute-capabilities appendix Tables 29 to 31, https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/compute-capabilities.html ; wgmma and tcgen05 target notes in the PTX ISA; CUTLASS's Blackwell functionality page. All checked 2026-08-29, the errors re-captured 2026-09-01):

Feature H100 (9.0, sm_90a) B200 (10.0, sm_100a) RTX 50 (12.0)
Thread block clusters yes yes yes
Distributed shared memory yes yes yes
TMA, single CTA yes yes yes
TMA multicast across a cluster yes yes no hardware; CUTLASS pins GeForce clusters to 1x1x1
wgmma yes no no
tcgen05, tensor memory no yes no
2-SM MMA (cta_group::2) no yes no
Shared memory per SM 228 KB (227 per block) 228 KB (227 per block) 100 KB (99 per block)
Resident warps per SM 64 64 48

Two rows need explicit sources. wgmma "Requires sm_90a" (https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#asynchronous-warpgroup-level-matrix-instructions-wgmma-mma , checked 2026-08-29): Hopper only, and no Blackwell of any kind. And the tcgen05 target notes list sm_100a, sm_101a and sm_103a with their family variants, never sm_120 (https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-mma-instructions-mma , checked 2026-08-29).

Hopper's wgmma instructions do not carry forward to datacenter Blackwell, and tcgen05 does not exist on consumer Blackwell. Each GPU in the table supports one of the three instruction sets.

Diagram: three cards, one feature ladder. Three columns, H100 sm_90a, B200 sm_100a, RTX 5090 CC 12.0, each a stack of feature bands; a band is filled when the hardware has it and hatched when the instruction only exists elsewhere. Shared bottom bands on all three: clusters, DSMEM, single-CTA TMA. Caption "the 9.0-and-up floor: all three, including consumer." H100-only band: wgmma. Caption "sm_90a only; carried to no successor." B200-only bands: tcgen05 plus tensor memory, 2-SM MMA. Caption "sm_100a family only; never sm_120." RTX 5090 column capped by two numbers where the others keep going: 100 KB shared per SM against 228, 48 resident warps against 64. Alt text: "Three GPUs share clusters, distributed shared memory and TMA. Only the H100 has wgmma, only the B200 has tcgen05 and tensor memory, and the RTX 5090 tops out at 100 kilobytes of shared memory per SM against their 228."

Compute capability is not a feature ranking

Compute capability 12.0 does not include every feature from 9.0. Compute capability is a family label, not a feature ranking. An architecture-specific target like sm_90a enables instructions that exist on one architecture and removes forward portability: the binary carries no PTX a newer card could JIT, and the instructions were never promised to exist again.

The a suffix marks an architecture-specific target.

Day 69 built fat binaries on the assumption that PTX flows forward; a targets are the documented exception.

The other limit is capacity. A Hopper GEMM pipeline is written against 227 KB of shared memory per block and 64 resident warps per SM. The consumer card offers 99 KB and 48.

Even the Hopper techniques that an RTX 50 can express, such as deep multi-stage double buffering with TMA feeding it, do not fit at Hopper sizes. That is why CUTLASS ships different kernels per family rather than one kernel with flags ("On Geforce series graphics card, there is no multicast feature therefore the cluster shape is fixed to 1x1x1", https://github.com/NVIDIA/cutlass/blob/main/media/docs/cpp/blackwell_functionality.md#cluster-size , checked 2026-08-29).

The GEMM you annotate

Full program in code/day78-blackwell/wgmma_gemm.cu. It is a small Hopper GEMM: 256 cubed, FP16 inputs, and FP32 accumulation. The exercise asks you to annotate it, and the listing plus the PTX ISA contain every quiz answer.

The block uses two warpgroups with different jobs. Threads 0 to 127 are the consumer: they issue every wgmma.mma_async and end the kernel holding the 64x64 output tile in their registers, 32 floats per thread. Threads 128 to 255 are the producer: they copy the next K-tile of A and B into shared memory while the tensor cores compute the current one.

That split is warp specialization, the pattern day 75's producer-consumer barriers rehearsed.

wgmma reads its operands from shared memory through 64-bit matrix descriptors (an encoded address plus two byte offsets and a swizzle mode), built by tileDesc() in the listing.

Three synchronization operations have separate roles. wgmma.fence orders register access before the issues, wgmma.commit_group and wgmma.wait_group retire the asynchronous work, and plain __syncthreads() is what lets the two warpgroups swap buffers.

    if (wg == 1) {
        loadStage(0, 0);  // prologue: stage 0 filled before anyone computes
    }
    __syncthreads();

    for (int kt = 0; kt < kTiles; ++kt) {
        const int s = kt & 1;
        if (wg == 1) {
            // Producer: fill the other stage while the consumer eats
            // this one. Nothing here waits on the tensor cores.
            if (kt + 1 < kTiles) {
                loadStage(s ^ 1, kt + 1);
            }
        } else {
            // Consumer: 4 wgmma issues walk the 64-deep tile in k16
            // steps. wgmma.fence first (register accesses, mandatory),
            // then the issues, then one commit_group, then wait_group 0,
            // which blocks until every committed wgmma has finished
            // reading shared memory and writing acc.
            asm volatile("wgmma.fence.sync.aligned;");
            for (int ks = 0; ks < kTileK / kWgmmaK; ++ks) {
                const uint64_t off = uint64_t(ks) * (2 * 128 >> 4);
                wgmmaM64n64k16(acc, tileDesc(smA[s]) + off,
                               tileDesc(smB[s]) + off);
            }
            asm volatile("wgmma.commit_group.sync.aligned;");
            asm volatile("wgmma.wait_group.sync.aligned 0;");
        }
        // Both warpgroups: the consumer is done reading stage s and the
        // producer is done writing stage s^1, so the next iteration may
        // swap them. Move the wait_group after this barrier and the
        // producer can overwrite a tile the tensor cores are still
        // reading: a race with no error message.
        __syncthreads();
    }

And the host-side gate that makes the hardware story explicit rather than a driver error:

    // wgmma exists on sm_90a and nowhere else, and an sm_90a binary
    // carries no PTX a newer card could JIT. Every non-Hopper card,
    // including all of Blackwell, fails here by design; the page's table
    // says what each card lacks.
    cudaDeviceProp prop;
    CUDA_CHECK(cudaGetDeviceProperties(&prop, 0));
    if (prop.major != 9 || prop.minor != 0) {
        std::fprintf(stderr,
                     "%s is CC %d.%d. This program needs CC 9.0 (Hopper): "
                     "wgmma is sm_90a-only and no Blackwell card has it.\n",
                     prop.name, prop.major, prop.minor);
        return EXIT_FAILURE;
    }

What this listing leaves out: a production Hopper GEMM would use TMA instead of ordinary loads, mbarriers instead of __syncthreads(), more stages, and swizzled shared memory layouts. Every one of those raises the annotation difficulty without changing the answers to today's questions, which are about who issues, who waits, and what each wait retires.

Results

Not yet run. The program needs an H100 to execute sm_90a code. The verified compile matrix was captured 2026-09-01 through the Compiler Explorer API into evidence/compile-2026-09-01.txt, using the compiler whose nvcc matches the course toolkit:

Toolkit Target Outcome
NVCC 12.6.2 (V12.6.85) sm_90a compiles, exit 0
NVCC 12.6.2 (V12.6.85) sm_90 8 ptxas errors, exit 255
NVCC 12.6.2 (V12.6.85) sm_120a nvcc fatal, target unknown, exit 1
NVCC 12.8.1 sm_120a ptxas errors, exit 255

What a Hopper session adds

An H100 run checks three claims:

  1. The harness passes. The two matrix-descriptor constants and the accumulator fragment mapping were written from the PTX ISA's tables and Figure 149, not from a run, and they are the two most common ways a hand-written wgmma kernel is wrong. A failure prints the first bad index; whole wrong 8x8 blocks point at the descriptors, scrambling inside blocks points at the epilogue.
  2. The runtime rejection string. Launching this sm_90a binary on a non-Hopper card should fail at load with a no-kernel-image error; the exact wording belongs in Pitfalls and awaits that capture. This page quotes only errors it has actually collected.
  3. The overlap, on a timeline. An nsys trace would show the producer warpgroup's loads running under the consumer's tensor-core work, the same evidence day 51 demanded for streams. No throughput number is promised: a 4x4 grid of CTAs cannot fill 132 SMs, and this program optimizes for readability, not occupancy.

Run it yourself

This is a reading lesson. The exercise and quiz use the listing, and the table lets you check a GPU's documented capabilities. To reproduce the compile matrix, compile for each listed target, including -arch=sm_90a.

An H100 is required for the runtime and nsys checks; an RTX 50 cannot execute wgmma. The setup guide lists remote GPU options, and the planning snapshot is in FACT-SHEET.md.

Exercise

Annotate the pipeline listing above. For each of its six synchronization points (the prologue __syncthreads(), wgmma.fence, the four-issue loop, commit_group, wait_group 0, the loop-end __syncthreads()), write one line: which warpgroup executes it, and what has been guaranteed the moment it completes. Then answer the same question for fence.proxy.async.shared::cta inside loadStage.

Time: 25 to 40 minutes. Submit: the seven annotation lines, then the quiz.

Check: the quiz in content/quizzes/day78.toml, marked in the browser, five questions with an explanation on every option. Every answer is derivable from the listing plus the PTX ISA sections linked under Go deeper; none needs a run. A wrong answer costs nothing and opens all four explanations, which is where the teaching is.

Hint 1

Sort the seven lines into two piles first: instructions that order threads against each other, and instructions that order work that has already been issued and left the threads behind. The two piles never substitute for each other, and one pile is executed by both warpgroups.

Hint 2

For each wait, ask what physically becomes safe after it: reading acc? Overwriting a shared memory stage? Reading a tile another warpgroup wrote through a different proxy?

The listing's loop-end comment tells you what goes wrong when wait_group and the barrier swap; the same reasoning answers the other five.

Solution

The prologue barrier publishes stage 0 to the consumer before the first issue. wgmma.fence (consumer only) orders the warpgroup's prior register accesses against the coming asynchronous issues; the PTX ISA makes it mandatory, "Otherwise, the behavior is undefined." The four issues put work into flight and guarantee nothing yet.

commit_group (consumer) closes the batch; wait_group 0 (consumer) is the only line after which acc holds the sums and the stage's shared memory is no longer being read. The loop-end __syncthreads() (both warpgroups) is buffer handover: consumer done reading stage s, producer done writing s^1.

The proxy fence (producer) makes generic-proxy stores visible to the async proxy wgmma reads through; the barrier then carries that guarantee across warpgroups. The obvious wrong answer, "the barrier covers all of it", fails because __syncthreads() orders threads, and an in-flight wgmma.mma_async is no longer attached to any thread; only its own wait retires it.

On Hopper, every synchronization instrument answers one question, who may touch this memory next, and the compiler checks none of them.

Pitfalls

You built with -arch=sm_90 and got eight errors from a compiler you did not invoke. nvcc's front end accepts the source; ptxas, the second compiler day 46 introduced, rejects each wgmma site: ptxas /app/example.ptx, line 148; error : Instruction 'wgmma.mma_async with floating point types' not supported on .target 'sm_90', ending in ptxas fatal : Ptx assembly aborted due to errors (captured 2026-09-01, NVCC 12.6.2). The fix is -arch=sm_90a, and the cost of the fix is a binary that runs on Hopper alone.

You targeted your RTX 50 and the toolkit refused the target itself. Under CUDA 12.6: nvcc fatal : Value 'sm_120a' is not defined for option 'gpu-architecture', three spaces before the colon. That is a toolkit-age error, not a feature verdict; the verdict arrives under 12.8.1, which knows the target and still rejects every wgmma instruction on it (both captured 2026-09-01).

You expected the sm_90a binary to run on a newer card. a targets embed no forward-portable PTX, so there is nothing for the driver to JIT. Day 58 built with plain sm_90 for exactly this reason. This program's host gate reports the capability mismatch before any launch; the driver's own load-failure string awaits the hardware session.

You read "5th generation tensor cores" on a consumer spec sheet and went looking for tcgen05. The PTX target notes for every tcgen05 instruction list the sm_100a/sm_101a/sm_103a families and never sm_120. No tensor memory, no 2-SM MMA, on any GeForce card. The consumer path onto its tensor cores is day 73's mma.sync, and day 80's capstone uses it.

Your cluster code compiles with .multicast::cluster and runs slower than without it. The qualifier compiles on any sm_90+ target, but the PTX ISA warns it is "optimized for" specific targets and "may have substantially reduced performance on other targets"; on GeForce there is no multicast hardware and CUTLASS pins the cluster shape to 1x1x1 (sources in the table above). Compiling is not having.

You ported a 4-stage Hopper pipeline and launch fails or occupancy collapses on a consumer card. The budget rows do it: 99 KB of shared memory per block against 227, and 48 resident warps per SM against 64. A pipeline sized for Hopper's shared memory does not fit, and one sized to fit may not have the warps to hide latency. Resize the stages; do not port the constants.

Go deeper

Next

Day 79 covers cuTile, where the compiler picks the instruction for each target. Day 80 returns to a hand-written WMMA GEMM that must beat the measured baseline day 44 set, 67.7 percent of cuBLAS FP32 on the course T4 at 2048, using tensor cores supported by every card in the table.