Volta m8n8k4: e_zero for C in f16 depends on the output type (-19 for f16 outputs, not -20); V100 measurements #22

Thanks for picking up #19 in v2.1.0. Here are measurements of the other Volta m8n8k4 parameter, e_zero, on four Tesla V100-SXM2-16GB (driver 580.178.04, VBIOS 88.00.8D.00.03, CUDA 12.8, sm_70).

Summary. On Volta, e_zero for a zero C in f16 depends on the output type, as you noted for Hopper in #16. sim.py currently selects it by c_type (-20 if self.c_type == torch.float16 else -131). What the V100 matches:

So selecting by d_type would give −131 for f32 outputs (as now for C in f32) and −19 for f16 outputs.

Setup.

• Our harness is independent of mmasim-kernels. It uses mma.sync.aligned.m8n8k4.row.col with the PTX fragment layout, and each quad-pair computes its own 8×8×4 tile.
• 16,384 random tiles per instruction (1,048,576 results), with raw bit patterns and scaled operands, plus directed cases.
• The four GPUs give byte-identical outputs on everything below.

Random tiles: 0 mismatches between the GPU and MMA-Sim on all three instructions. e_zero only matters for very small products, which random bits almost never produce.

Directed cases (C = +0, all other terms zero):

• In the first case the GPU keeps every product bit; with e_zero = −20 and F = 23, MMA-Sim drops one.
• In the second case the GPU drops the 2⁻⁴³ product, which leaves an exact tie that rounds to even (+0). MMA-Sim keeps it and rounds up. With D in f32 (f32.f16.f16.f16), the GPU keeps the same 2⁻⁴³ (0x33000020).

Probes (C = +0 in f16):

1. D in f32: one product per tile with significand 2047 × 1025 (ends in binary 1111) and exponent −14 to −28. The GPU returns every product exactly, so every e_zero ≤ −25 is consistent. This is the same probe and result as Turing's m8n8k4 on an RTX 2060.
2. D in f16: a sum placed exactly on an f16 rounding tie, plus a small product 2⁻ᵏ below the tie bit (k = 1..25). It is run at two levels (1 + 2⁻¹¹ and 2⁻²⁵) and with normal×normal, normal×subnormal and subnormal×subnormal operands.
   • What the GPU does: at the normal level everything matches F = 23. At the subnormal level, only products of 2⁻⁴³ and below are lost (k = 18..23), whatever the operand kinds.
   • The fit: scanning e_zero from −14 to −45, only −19 reproduces all 86 tiles, and with it only F = 23 does.

With e_zero = −131 for f32.f16.f16.f16 and −19 for f16.f16.f16.f16, our integer T-FDPA model gives 0 mismatches against the V100 on all random tiles, directed cases and probes. We have not probed the f16-output case on Turing.