Summary
The ck_tile 2:4-structured-sparse SWMMAC path (MmaOpFamily::SPARSE) produces incorrect GEMM results on gfx1201 (RDNA4) for real (arbitrary-position) 2:4 data. Three independent bugs in compress_a_impl / the iu4 packed path; canonical test patterns (always-exactly-2 nonzeros, zeros in fixed slots) mask all three, so the existing tests pass while real magnitude-pruned weights fail.
All three were found by device-side dumps + hand verification, and a full sparse+dense regression passes after the fixes (int8/fp8/fp16 sparse, all MX converts, dense iu4 K16/K32 — all max_abs_err=0).
Environment
- GPU: Radeon AI PRO R9700 (gfx1201, RDNA4), ROCm 7.2.4
- CK:
develop @ b6759456103a6e78137ecbae673c69865884f7e3 (2026-07-18)
- Files:
include/ck_tile/core/arch/mma/sparse/{sparse_transforms.hpp, sparse_mma_pipeline.hpp}
Bug 1 — compress_a_impl <2-nonzero fallback plants phantom values (all sparse dtypes)
When a 4-element group has fewer than 2 nonzeros (legal under 2:4 = "at most 2 per 4"), the fallback seeds an output survivor slot from a fixed input position (slot 3) while the index default points at slot 2 — injecting a phantom duplicate value into the compressed operand. Only visible with real data (exact2frac < 1.0); canonical exactly-2 patterns never hit it.
Fix: initialize unassigned survivor slots to true zero {0, 0} rather than copying a fixed input position.
Bug 2 — iu4/pk_int4 compressed over PHYSICAL bytes instead of LOGICAL nibbles
compress_a_impl groups 4 physical pk_int4_t bytes and does a whole-byte nonzero test, but 2:4 for int4 is over logical nibbles. Real per-nibble data yields >2 "nonzero" bytes per group → out-of-bounds write into the 2-slot survivor array (UB; observed as register-aliasing corruption). Compounded by vector_traits<ext_vector_t<pk_int4_t,N>>::scalar_type erasing packedness to int8_t, and the pipeline's index buffer being sized by physical bytes rather than logical nibbles (2× too small). RDNA4 ISA §7.12.3 ("8 index values per lane" for K=32 iu4) confirms the index must track logical nibbles (InternalAVecSize * APackedSize).
Fix: logical-nibble compression path + corrected idx sizing + a LogicalADataType param threaded through SparseCompressTransform::exec → execExtVec → compress_a_impl.
Bug 3 — iu4 index encoding does not match the swmmac_iu4 hardware read order
With Bugs 1–2 fixed and the compressed operand byte-correct, the end-to-end swmmac_i32_16x16x{32,64}_iu4_w32 result is still wrong (~280–530 abs err). Empirically, the hardware reads the two per-group index fields swapped relative to the survivor they govern, with an XOR-1 transform:
pos(HIGH-nibble survivor) = idx1_field XOR 1
pos(LOW-nibble survivor) = idx0_field XOR 1
Found via a position-pair sweep; the formula predicts all test cases and explains an earlier "always position 3" single-survivor anomaly (default field 2, XOR 1 = 3). This appears to be undocumented gfx1201 swmmac_iu4 behavior; CK's compress currently emits the index in the un-swapped/un-XORed convention.
Fix: swap the scan-slot→field mapping and XOR the written index value with 1 in the iu4 packed-nibble branch.
Repro / patch
Minimal standalone probes (warp-tile correctness vs CPU reference, single-nonzero and position-pair sweeps, K32/K64) reproduce each bug. A local patch against b6759456 fixing all three is available; happy to open a PR if useful — let me know the preferred form.
Bugs 1 and 2 are unambiguous library defects (phantom values / OOB write). Bug 3 is a hardware-convention mismatch worth documenting regardless, since it currently blocks iu4 2:4 bring-up on RDNA4.
Summary
The
ck_tile2:4-structured-sparse SWMMAC path (MmaOpFamily::SPARSE) produces incorrect GEMM results on gfx1201 (RDNA4) for real (arbitrary-position) 2:4 data. Three independent bugs incompress_a_impl/ the iu4 packed path; canonical test patterns (always-exactly-2 nonzeros, zeros in fixed slots) mask all three, so the existing tests pass while real magnitude-pruned weights fail.All three were found by device-side dumps + hand verification, and a full sparse+dense regression passes after the fixes (int8/fp8/fp16 sparse, all MX converts, dense iu4 K16/K32 — all
max_abs_err=0).Environment
develop@b6759456103a6e78137ecbae673c69865884f7e3(2026-07-18)include/ck_tile/core/arch/mma/sparse/{sparse_transforms.hpp, sparse_mma_pipeline.hpp}Bug 1 —
compress_a_impl<2-nonzero fallback plants phantom values (all sparse dtypes)When a 4-element group has fewer than 2 nonzeros (legal under 2:4 = "at most 2 per 4"), the fallback seeds an output survivor slot from a fixed input position (slot 3) while the index default points at slot 2 — injecting a phantom duplicate value into the compressed operand. Only visible with real data (
exact2frac < 1.0); canonical exactly-2 patterns never hit it.Fix: initialize unassigned survivor slots to true zero
{0, 0}rather than copying a fixed input position.Bug 2 — iu4/
pk_int4compressed over PHYSICAL bytes instead of LOGICAL nibblescompress_a_implgroups 4 physicalpk_int4_tbytes and does a whole-byte nonzero test, but 2:4 for int4 is over logical nibbles. Real per-nibble data yields >2 "nonzero" bytes per group → out-of-bounds write into the 2-slot survivor array (UB; observed as register-aliasing corruption). Compounded byvector_traits<ext_vector_t<pk_int4_t,N>>::scalar_typeerasing packedness toint8_t, and the pipeline's index buffer being sized by physical bytes rather than logical nibbles (2× too small). RDNA4 ISA §7.12.3 ("8 index values per lane" for K=32 iu4) confirms the index must track logical nibbles (InternalAVecSize * APackedSize).Fix: logical-nibble compression path + corrected idx sizing + a
LogicalADataTypeparam threaded throughSparseCompressTransform::exec → execExtVec → compress_a_impl.Bug 3 — iu4 index encoding does not match the
swmmac_iu4hardware read orderWith Bugs 1–2 fixed and the compressed operand byte-correct, the end-to-end
swmmac_i32_16x16x{32,64}_iu4_w32result is still wrong (~280–530 abs err). Empirically, the hardware reads the two per-group index fields swapped relative to the survivor they govern, with an XOR-1 transform:Found via a position-pair sweep; the formula predicts all test cases and explains an earlier "always position 3" single-survivor anomaly (default field 2, XOR 1 = 3). This appears to be undocumented gfx1201
swmmac_iu4behavior; CK's compress currently emits the index in the un-swapped/un-XORed convention.Fix: swap the scan-slot→field mapping and XOR the written index value with 1 in the iu4 packed-nibble branch.
Repro / patch
Minimal standalone probes (warp-tile correctness vs CPU reference, single-nonzero and position-pair sweeps, K32/K64) reproduce each bug. A local patch against
b6759456fixing all three is available; happy to open a PR if useful — let me know the preferred form.Bugs 1 and 2 are unambiguous library defects (phantom values / OOB write). Bug 3 is a hardware-convention mismatch worth documenting regardless, since it currently blocks iu4 2:4 bring-up on RDNA4.