[BugFix][CUDA] Advance tcgen05 ld/st segment pointers in b32 columns (#2952)
The multi-segment recursion in tcgen05_{ld,st}_core advanced the
register pointer in dst_t/src_t elements while segment lengths count
b32 columns, so a 16-bit copy whose column count is not a power of two
(e.g. 48 columns = one x32 + one x16 segment) accessed the tail
segments' registers at half the intended offset -- a silent
miscompilation. Advance the pointer in b32 units instead.
Add asymmetric whole-store/whole-load roundtrips at (128, 96) bf16: a
symmetric roundtrip cancels the unit error, so exactly one side is
split into power-of-two (single-segment) slices, catching each
direction independently.
Co-authored-by: Claude Fable 5 <noreply@anthropic.com> Y
Yongqi Zhuo committed
c2d847f4970343d0b773d14ec44a6c0f6407731c
Parent: 4d15031
Committed by GitHub <noreply@github.com>
on 8/13/2026, 5:35:42 AM