FP8 fine-grained scaling on Hopper vs Blackwell: per-tensor vs per-block scales and accuracy implications

FP8 (E4M3/E5M2) was introduced on Hopper with per-tensor scaling: a single FP32 scale factor for the entire tensor, applied before the FP8 GEMM and after in the epilogue. This coarse scaling is problematic for tensors with high dynamic range or outliers, as a single scale cannot adequately represent both large and small values. Blackwell extends FP8 support with block-scaled formats: MXFP8 (E4M13 + UE8M0 scale per 32-value block) and NVFP8 (E4M3 + E4M3 scale per 16-value block + per-tensor FP32 scale). The tcgen05.mma instruction supports these via .kind::mxf8f6f4 (block-scaled FP8/FP6 with UE8M0 scale). The finer-grained scaling (per-16 or per-32 values instead of per-tensor) reduces quantization error by allowing different scale factors for different regions of the tensor. The accuracy implication: block-scaled FP8 can achieve close to FP16/BF16 accuracy, while per-tensor FP8 may suffer noticeable degradation on tensors with outliers. The CUTLASS block-scaled tensor core record notes that DeepGEMM requires packed UE8M0 scales (four per int) on SM100, while SM90 wants FP32 scales - this changes the memory layout the engine must produce. The existing block-scaled-tensor-core-fp4-fp8 record covers the PTX ISA shape tables; this record focuses on the Hopper-to-Blackwell evolution of FP8 scaling granularity and its accuracy implications.

The evolution from per-tensor FP8 (Hopper) to block-scaled FP8 (Blackwell) is analogous to the evolution from per-tensor INT8 to per-channel INT8 in earlier generations. The key accuracy implication: block-scaled FP8 can maintain close to FP16 accuracy on most tensors, while per-tensor FP8 may suffer 1-5% degradation depending on the tensor distribution. The FlashAttention-3 record notes that FP8 doubles tensor-core throughput but leaves SFU exp unchanged, making the softmax bottleneck worse - this is true on both Hopper and Blackwell. The block-scaled FP8 formats are consumed by tcgen05.mma via .kind::mxf8f6f4. DeepGEMM's requirement for packed UE8M0 scales on SM100 (vs FP32 on SM90) is a concrete example of how the scale format changes the engine's required data layout. 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.