Tensor-core instruction generations: wmma/mma (Ampere) -> wgmma (Hopper) -> tcgen05 (Blackwell)

This is the layer that determines what a kernel can possibly do. Three generations exist and they change how a kernel must be STRUCTURED, not just how fast it runs. (1) Ampere and earlier: warp-level wmma (implicit load/store of fragments, the compiler handles element distribution) and mma (explicit fragment layout, the programmer handles it). Both are warp-synchronous - the whole warp participates, the accumulators live in registers, and the instruction blocks until it is issued in the warp's own issue stream. (2) Hopper: wgmma, issued by a warpgroup of four contiguous warps (the first warp's warp-rank must be a multiple of 4). It is ASYNCHRONOUS and runs in a separate async proxy, so it needs wgmma.fence, fence.proxy.async, commit_group and wait_group to be correct; M is always 64 (shapes are .m64nNk16 for f16/bf16/tf32 and .m64nNk32 for fp8 and int8), the accumulator can be in registers or shared memory, and sparse variants exist. The payoff is concrete: per the FA-3 author blog, without wgmma the older mma.sync reaches only about two thirds of Hopper tensor-core peak. (3) Blackwell: tcgen05, where the MMA is launched by a SINGLE THREAD rather than a warp or warpgroup, accumulates in dedicated on-chip Tensor Memory instead of registers, and can be issued across TWO CTAs in a cluster. TMEM is a 2D array - on sm_100a/sm_100f, 512 columns x 128 lanes per CTA with 32-bit cells - dynamically allocated by a single warp in units of 32 columns with power-of-2 counts, and everything allocated must be deallocated before the kernel exits. The largest single-CTA BF16/FP16 UMMA tile is 128x256x16, about twice the largest Hopper WGMMA atom, and 2-CTA mode reaches 256x256x16.

Why this is a real decision input and not trivia: a Hopper-only kernel library tells you a consumer Blackwell part gets a worse path, and CUTLASS's README states the trap explicitly - sm100a-compiled kernels are NOT compatible with RTX 50 series (SM120) because SM120 has a different compute capability from datacenter SM100. Hopper's extended static shared memory is 228 KB per CTA (vs the 48 KB default elsewhere), which is what makes the large Hopper tiles possible; Blackwell entries go further with 328 KB shared-memory allocations. The accumulator-location change is the deeper generational difference: registers force a serial schedule once pressure bites, TMEM does not, which is why FA-4's backward pass can keep MMAs in flight while CUDA cores do elementwise work where FA-3 could not. HARDWARE-PARAMETER ROWS ADDED 2026-10-05 (schema authority pass, agent w5-schema). `affected_by_hardware` names PRODUCTS (accelerator record ids); these rows name the hardware QUANTITIES that move this class's cost, which had nowhere to live until flop.hardware_parameters was added in this commit. Every effect below is lifted verbatim from this record's own existing prose and no figure here is new - `magnitude` reproduces numbers the record already quoted. `basis` is 'unverified' where the effect was asserted with nothing concrete behind it, rather than being dressed up. The rows were pulled out of free-text notes; they are not a change to what this record claims.