Turing runs m8n8k4 with F = 23 (not 24); kind::f8f6f4 reads E2M1's low padding bits as mantissa (RTX 2060 / RTX 5070 Ti measurements) #19

Thanks for MMA-Sim. We have been checking it against our own GPUs with an independent harness, and found two places where hardware and MMA-Sim (commit 54294d4) disagree. Neither instruction has a hardware kernel in mmasim-kernels, which is probably why they slipped through. Everything else we measured matched bit for bit. Details and minimal cases below.

Setup. Our harness is a raw-bit mma.sync kernel written independently of mmasim-kernels: one independent tile per warp (per quad-pair for m8n8k4), with fragment layouts from the PTX ISA. Inputs:

• half of the tiles are uniform random words;
• half have operands within a few binades of 1, so that every fractional bit of the fused sum shows;
• the bits a format does not use (the low 13 bits of TF32, the padding around FP4/FP6) are random on purpose;
• a set of directed cases.

GPUs and software:

• RTX 2060 (sm_75), driver 610.62;
• RTX 5070 Ti (GB203, sm_120a), driver 616.64;
• A100-PCIE-40GB (sm_80), driver 580.178.04;
• CUDA 12.9.

1. Turing executes Volta's m8n8k4 with F = 23, not 24

MMA("Turing", "m8n8k4.*") applies Turing's parameters (F = 24; e_zero −21 / −132). On an RTX 2060 the native HMMA.884 (it is native: cuobjdump -sass shows HMMA.884.F32.F32.STEP0..3) behaves like Volta's sum width.

Statistics:

• Against MMA-Sim: m8n8k4.f32.f16.f16.f32 differs in 49,758 of 262,784 results, and m8n8k4.f32.f16.f16.f16 in 47,429 of 262,720. All differences are ±1 ulp and spread evenly over quad-pairs, rows and columns.
• With F = 23: a T-FDPA with L = 4, F = 23 and RZ gives 0 mismatches on all three m8n8k4 variants. No other F in 22..26, and no L = 2 chaining, does.
• Turing's own instructions agree: m16n8k8 (f32 and f16 accumulators) matches MMA-Sim exactly, with 0 mismatches in about 1.05 M results.

Minimal case, the discrepancy example from your Table 6:

• inputs: a = (−2^13, −0.5, −0.25, −0.125) f16, b = (2^10, 1, 1, 1) f16, c = 2^23 f32, m8n8k4.f32.f16.f16.f32;
• MMA-Sim (Turing): 0xbf000000 (−0.5);
• RTX 2060: 0x0, Volta's answer.

Zero accumulator with C in f16. We probed single products a·b with C = +0 (f16) and product exponents from −14 to −28. The significand 2047·1025 ends in binary 1111, so any bits lost in alignment would show. No bits are lost at any exponent. So with F = 23, every e_zero ≤ −25 is consistent: a zero f16 C does not affect the alignment. MMA-Sim's −21 truncates tiny products. Example, m8n8k4.f32.f16.f16.f16, c = +0 (f16):

Ampere, for comparison. On sm_80, ptxas lowers m8n8k4 to FFMA/FADD/FMUL with no HMMA at all. An A100 then computes (((p0+p1)+p2)+p3)+c in f32 with round-to-nearest-even: 0 mismatches on 233,452 finite f32-output results. So m8n8k4 is a tensor-core instruction only on Volta and Turing. That matches MMA-Sim not listing it for Ampere, but it may be worth a note. The eight Ampere mma.sync instructions MMA-Sim models match on the A100 with 0 mismatches.

2. kind::f8f6f4 reads the low padding bits of an E2M1 byte as mantissa bits

The PTX ISA says only that ".e2m1 resides in central 4 bits of the 8-bit container with padding in the upper 2 bits and lower 2 bits of the container". It does not say whether the padding is ignored. MMA-Sim's decode_uint8_to_fp6_fp4 drops bits 1:0. On an RTX 5070 Ti, the byte is read as E2M3 in bits 5:0: bits 1:0 act as two more mantissa bits, and bits 7:6 are ignored. With zero padding both readings agree, which is presumably the common case.

One-hot probe with m16n8k32.f32.e2m1.e4m3.f32 (mma.sync.aligned.m16n8k32.row.col.kind::f8f6f4): one non-zero byte in A, B = 1.0 (e4m3 0x38), C = 0. 512 tiles per row of the table, every (row, k) position:

With E2M1 operands decoded as E2M3 from bits 5:0 and everything else as in your T-FDPA (F = 25, L = 32), we get 0 mismatches on all 18 kind::f8f6f4 instructions involving e2m1, f32 and f16 accumulators. That is 14.2 M results across two runs with different seeds, random padding included. On the 10 of these instructions that MMA-Sim models, its current decoding mismatches 58–96 % of the results of the first run.

3. A data point for the e3m2 / e2m3 TODO

For kind::f8f6f4 with e3m2 and e2m3 (in bits 5:0, bits 7:6 random), we used T-FDPA with F = 25 and L = 32, as for the other RTX Blackwell 8-bit formats. It gives 0 mismatches on all 24 combinations without e2m1 (e3m2/e2m3 with each other and with e4m3/e5m2, f32 and f16 accumulators), 18.9 M results over two runs. So adding them to rtx_blackwell_mma_shape_and_type with the existing model looks safe.

Happy to share the raw inputs and outputs (case files, GPU words and SHA-256) or to run specific probes on these GPUs if useful.