feat(cuda-device): derive the m16n8k16 accumulator fragment layout

[honeyspoon]

wmma::mma_m16n8k16_f32_f16 takes and returns [f32; 4], and those four values are not four adjacent elements of the output tile — each lane holds a scattered quarter of a 16x8 tile. The intrinsic's own docs give the mapping as index arithmetic the caller must reproduce, and its safety section warns that a different array order computes a different matrix operation.

That failure is silent: it produces a wrong product rather than an error. This derives the mapping once so call sites do not each re-derive it.

The layout is a mixed-radix numeral

Substituting the PTX row/col formulas into the row-major offset and regrouping:

offset = (group + 8*(j>>1))*8 + 2*tig + (j&1)
       = (j&1)*1 + tig*2 + group*8 + (j>>1)*64

digit    range   place   source
j & 1      2       1     fragment index, low bit
tig        4       2     lane % 4
group      8       8     lane / 4
j >> 1     2      64     fragment index, high bit

Place values are the running product of the lower ranges (1, 2, 8, 64), span 128. So the numeral is uniquely decodable, (lane, j) -> offset is injective, and since domain and codomain both have 128 points it is a bijection — the warp's fragments tile the output exactly, no overlap and no gap.

What makes this awkward to write by hand is that the fragment index is split across two non-adjacent digit positions, at places 1 and 64, interleaved with the lane digits. That is why j 0 and 1 differ by one column while j 0 and 2 differ by eight rows.

API

All const fn, no unsafe, no new dependencies.

Function	Purpose
acc_coords(lane, j)	position in the 16x8 tile
acc_offset(lane, j)	row-major offset, in numeral form
acc_matrix_offset(..., tile_row, tile_col, row_stride)	strided destination a GEMM epilogue writes; None on overflow
acc_decode(offset)	the inverse

acc_decode is exported deliberately: a caller with an existing hand-rolled layout can round-trip it in a test and catch a disagreement cheaply, rather than discovering it as wrong numerics.

Verified against hardware, not only against its own tests

A bijection onto the wrong positions would pass every unit test here. So the check that matters ran real mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32 on sm_86 (RTX A2000, CUDA 13.3), gathering A and B from the PTX ISA layouts, scattering the result with acc_coords, and comparing against a host matmul. Inputs are asymmetric and non-repeating so a transposed or half-swapped layout could not coincidentally match:

LAYOUT CORRECT  (0/128 mismatched, worst abs diff 0.00000)

That also validates the A and B fragment layouts indirectly — a wrong gather would have produced a wrong product.

Tests

Bijection via an ownership map (fails on both overlap and gaps), agreement between the numeral and row/col forms, decode as a true inverse over all 128 points, the two-position split of the fragment index, disjointness preserved under strided placement, and overflow rejection. 6 tests, fmt --check clean, zero warnings.

Scope

The C/D accumulator only. A and B are different numerals and would each need the same treatment. Only sm_86 was available to test on; the formulas come from the PTX ISA rather than from hardware probing, and the layout is part of the instruction's ISA contract, so I would not expect divergence — but it is untested elsewhere.