ck_tile: fix 2:4-sparse SWMMAC correctness on gfx1201/RDNA4 (3 bugs) + fail→pass repro

[The-Monk]

0. TL;DR for the maintainer (schung-amd)

Three correctness bugs in ck_tile's SPARSE MmaOpFamily path produce wrong results on gfx1201/RDNA4 (v_swmmac_*_iu4). Root-caused, fixed (two headers), and reproduced with a committed standalone test that fails on current CK and passes with the fix. We found these while building a production-grade 2:4-sparse fp8/int4 GEMM for LLM inference on RDNA4 — §3-4 give that context, which doubles as validation that the fixed path works end-to-end at scale.

1. The three bugs (files: sparse_mma_pipeline.hpp, sparse_transforms.hpp)

Bug 1 — compress_a_impl writes phantom values on <2-nonzero 2:4 groups.
The fallback ADataType nonzero_elems[2] = {a_vec[i*4+2], a_vec[i*4+3]} seeds the compressed pair from fixed input positions rather than true zero. A 2:4 group with 0 or 1 real nonzeros then reconstructs with a phantom value from a wrong lane (observed as register aliasing / UB). Fix: default nonzero_elems to true {0,0} and fill only real survivors — 0/1/2-nonzero groups all reconstruct exactly.

Bug 2 — packed sub-byte (pk_int4_t) group-of-4 scan treats a byte as one element.
The compaction's "nonzero per group of 4" scan operates on bytes, but a pk_int4_t byte holds two 4-bit values; a byte is "nonzero" if either nibble is, so the group-of-4 assumption over-counts and writes out of bounds into nonzero_elems[2]/[3]. Fix: packed-sub-byte-aware nibble counting (guarded with a static_assert that the packed path currently implements only pk_int4_t).

Bug 3 — compress_a_impl emits the two per-group idx metadata fields in an order that mismatches what the hardware reads, requiring a SWAP + XOR-1 to correct.
Within CK's own pk_int4_t packing, the metadata CK writes has the two idx fields swapped relative to which compressed survivor they govern, plus a XOR-1, versus what v_swmmac_i32_iu4 consumes on gfx1201 — so the sparse result comes out wrong (pos(HIGH-nibble survivor, found first) = idx1_field XOR 1, pos(LOW-nibble survivor, found second) = idx0_field XOR 1). Dense iu4 was confirmed correct first (11 seeds) to isolate this to the sparse metadata path. Fix: emit the idx fields with the swap+XOR so the hardware reconstructs the correct positions (Idx0 < Idx1, the documented 2:4 contract, is naturally satisfied).
Precision note (so this isn't over-stated): we verified this is specific to CK's compress/packing convention, not a universal "the hardware swaps." An independently-derived encoder in our own driver produces correct metadata without the swap — and applying the swap to it breaks a working encoding. So the fix corrects CK's own ordering to match the hardware; it is not a claim about v_swmmac_iu4 behavior for all encoders.

2. The test (committed; fails on current CK, passes with the fix)

test/ck_tile/gfx1201_sparse_swmmac/sparse_swmmac_correctness_repro.cpp drives CK's own machinery for the SPARSE MmaOpFamily — it builds Pipeline::AWarpDstrEncoding internally (no manual compression math), generates an adversarial 2:4 A tile in-source (each 4-group cycles through all six two-survivor position pairs, all four single-survivor positions, and the zero-survivor group — survivors at every position, values keyed to (row, group, position) so misplacement shows numerically rather than cancelling), runs the SWMMAC GEMM (K=32/64/128, three FragsK tile sizes), and checks against a CPU int64-accumulate reference. The legacy slots-0,2 pattern is retained as an opt-in control (-DUSE_CANONICAL_PATTERN); it cannot detect Bug 1 and passes on both trees.

Verification (2026-08-15, head ad79359d0 vs merge-base 8fc1ac24e9, identical test source on both trees, 1x R9700 gfx1201, ROCm 7.14):

• Unfixed (merge-base): FAIL — max_abs_err = 112 / 127 / 272 at K=32/64/128
• Fixed (this PR): PASS — max_abs_err = 0 on all three shapes

Reproduction (self-contained — no external files):

git clone -b gfx1201-sparse-swmmac-fixes https://github.com/The-Monk/composable_kernel ck-3759 && cd ck-3759
hipcc -std=c++17 -O2 --offload-arch=gfx1201 -I include \
  test/ck_tile/gfx1201_sparse_swmmac/sparse_swmmac_correctness_repro.cpp -o repro
HIP_VISIBLE_DEVICES=0 ./repro     # -> ALL PASS (max_abs_err=0)
# unfixed A/B: git worktree add ../ck-base $(git merge-base HEAD origin/develop)
#   then build the SAME test file with -I ../ck-base/include -> FAIL (112/127/272)

Historical note: earlier revisions of this description cited max_abs_err=352 from a real Quark-quantized weight tile (REAL_A_16x128, -DUSE_REAL_TILE); that input is superseded by the in-source generator (same root cause, now reproducible from the PR alone) — the 352 figures apply only to that tile.
(We compiled with the ROCm-devel clang toolchain directly — clang++ -x hip --offload-arch=gfx1201; the distro hipcc on this box is a stale 5.7-era wrapper unrelated to CK's target toolchain and will not build ck_tile headers. Any current ROCm/HIP clang works the same way.)

ASAN attempted, not obtained (honest note, not a blocker): we tried to get a device-sanitizer stack trace on the Bug-1/Bug-2 OOB write for extra evidence (-fsanitize=address -fgpu-sanitize). Device ASAN needs an xnack+ target-ID variant; RDNA4 (gfx1201) does not accept an xnack+/xnack- feature suffix at all (clang++: error: invalid target ID 'gfx1201:xnack+') — HIP device-side ASAN is a gfx9/CDNA-class feature (gfx90a/gfx94x with xnack+), not available on RDNA4 consumer/workstation targets in this ROCm toolchain. We did not chase this further since it's not required to reproduce the bug — the max_abs_err numbers above are a complete, deterministic fail→pass.

We'll port this into CK's gtest format for the PR; the standalone repro is included so it's runnable without the full CK test build.

3. Context — why these bugs mattered (the "whole thing")

We hit these building a library-grade 2:4-sparse GEMM for RDNA4 LLM inference (llama.cpp/ggml fork, gfx1201, 2× R9700). The journey, honestly:

• ISA ceiling is real: isolated v_swmmac microbench hits 765 TOP/s fp8-2:4 (2× dense) and 1531 TOP/s int4-2:4 (4×) at ILP≥4 — 88-100% of the R9700 spec.
• A hand-written MMQ-grade 2:4 kernel (cooperative tiles, LDS staging, ILP≥4) reaches ~93% of our native-fp8 WMMA MMQ (which we built — none exists upstream) on the whole model, and is at or ahead of the library on the three dominant GEMM shapes (+1.8% / +3.9% / +32% on gate-up / q-o / down-proj).
• The honest finding worth AMD seeing: the ISA 2×/4× sparsity ceiling does not translate to a compute-bound, well-tiled model GEMM — measured across three quality tiers, the sparsity factor is 2× (isolated ISA) → 1.25× (crude kernel) → ~1.01× (library-grade). Once the kernel is well-tiled, the GEMM is bound by the tiled system (LDS staging, occupancy), not the raw SWMMAC issue rate that sparsity accelerates. This is a real, reproducible microbench-vs-model boundary for RDNA4 structured sparsity.
• The mechanism, made precise (int4-2:4): the "4×" decomposes as 2×(int4 vs int8) × 2×(2:4 sparsity). The first 2× is real and already captured by dense wmma_i32_16x16x32_iu4 — gfx1201's dense int4 engine is already K=32-wide. The second 2× washes: swmmac_i32_16x16x32_iu4 (749,761 GOP/s) ≈ wmma_i32_16x16x32_iu4 (765,061) = 0.98×, instruction-for-instruction, because dense int4 already occupies the K=32 slot. The only sparse tensor edge is the K=64 form (1.77×) — the same "2×-K" mechanism fp8-2:4 already showed washing at library grade. So on high-arithmetic-intensity LLM GEMM shapes (compute-bound), structured 2:4 is a ~1.2× weight-bandwidth win (memory-bound regime only), not the tensor 2×/4×. (Correctness first regardless — the perf question is separate, and now answered.)
• The remaining gap to the library was pinned by direct profiling to two specific, non-magical things — a 2×-slow activation-quantize helper (LDS+sync vs shuffle) and one occupancy-starved small-N shape — both fixed [§4].

4. Closing the gap — result

The 8% was pinned (by direct rocprofv3 head-to-head) to two things, NOT the big GEMMs (we're at/ahead of our own fp8 MMQ there: +1.8% / +3.9% / +32% on gate-up / q-o / down-proj). Fixing them:

build	dense-MMQ (vs our fp8 MMQ)	2:4-MMQ (vs our fp8 MMQ)
capstone baseline	3786 (0.924×)	3823 (0.933×)
+ shuffle-based quantize (kept)	3900 (0.952–0.978×)	3944 (0.962–0.989×)

• Quantize helper rewrite (kept): replaced a 32-thread / LDS+__syncthreads() reduction with the upstream mmq.cuh quantize shape — 128 threads, float4 loads, warp-shuffle max-reduction, zero LDS/sync (halved it, 36µs→~18µs). Closed ~40% of the remaining gap; both kernels now high-90s% of our native-fp8 kernel (0.989× same-session).
• The N=1024 k/v small-shape (honest negative): both candidate levers — LDS bank-conflict padding and coarser K-per-sync — were implemented, correctness-verified, and measured worse (padding −2.5%; coarser-K a monotonic regression as it forced dropping double-buffering into an occupancy/LDS cliff). Both reverted. The outlier remains open, wants a different lever (shape-adaptive smaller tile without touching sync granularity).

Net: the hand-written 2:4 kernel now reaches ~96–99% of our native-fp8 WMMA MMQ (which we built — none exists upstream), is ahead on the FLOP-dominant GEMMs, and carries the measured ~1% sparsity edge over a comparable dense kernel — a library-competitive RDNA4 2:4 GEMM on top of the correctness fixes.

5. A related RDNA4 iu4 quirk (same class, different instruction — FYI, not part of the fix)

The Bug-3 element-ordering quirk on swmmac_iu4 is not isolated: an independent kernel-authoring effort on this box hit the same class on the dense v_dot8_i32_iu4/sudot8 path — "dots mismatched element pairs." RDNA4's iu4 instruction family appears to carry undocumented element-ordering conventions in both the sparse (SWMMAC metadata) and dense (dot8 operand pairing) paths. Documenting these in the ISA/CK would save the next implementer the multi-day reverse-engineering we did. Happy to write up the dense one too if useful.

6. Method (for reproducibility)

All numbers: gfx1201 (R9700), ROCm 7.14, GPU-isolated, warm, medians ≥3. Correctness gated by execution (CPU-reference compare), never inspection. The full capability-optimizer method (measure → grade vs the published peak → drive the lever) is what surfaced both the bugs and the microbench-vs-model boundary.

[doplxyz]

Thanks for digging this out and writing it up in this much detail — this had been sitting without a
review for three weeks, and the analysis in the description made it possible to check the claims
rather than guess at them. I have a gfx1201 box, so I ran it.

Short version: bug 1's fix holds up — I reproduced a failure on the base tree and, for what this
test measures, traced it to that one line. But the committed test can't be built from the PR alone, it passes on
the unfixed tree in the configuration that is buildable, and it never reaches the iu4 path that
bugs 2 and 3 are about. Separately, the changed code is not restricted to gfx1201, and I think that's
the thing to sort out before merge.

This is not a formal approval: I can only speak for gfx1201 and the int8 path.

Environment. AMD Radeon RX 9070 XT, gfx1201, amdgcn-amd-amdhsa--gfx1201; Linux 6.14.0-37,
amdgpu 6.19.14.31400100; container
rocm/pytorch:rocm7.14_ubuntu24.04_py3.12_pytorch_release_2.12.0, AMD clang 23.0.0git
(ROCm/llvm-project 46fcb339fb61). Trees compared: base
8fc1ac24e9bd7b431663a15e2122ce02c2979d37 — which is git merge-base of this PR's head and
develop, and base..head is exactly the three files in this PR — versus head
ac24ac28d662acf279478545cec541d4dde00f31. Identical test source and compile flags on both sides,
separate clean trees and build directories, three fresh processes per variant, identical stdout on
every repeat.

Two environment-side notes so the commands reproduce, neither of which is a problem with this PR:
that container needs --rocm-device-lib-path=<sdk>/lib/llvm/amdgcn/bitcode, and it ships only
libamdhip64.so.7, so a libamdhip64.so symlink is needed for the link step.

---

1. The repro can't be built from the PR alone

sparse_swmmac_correctness_repro.cpp:34 has

#include "real_a_tile.h"  // REAL_A_16x128: Quark int8 2:4 weight tile

outside the #ifdef USE_REAL_TILE at line 264, and real_a_tile.h isn't one of the PR's three files.
From a clean checkout the translation unit doesn't compile either way, so the max_abs_err=352
figures can't be checked by a reviewer.

I generated a substitute tile to get something running: same int8_t[16][128] shape, at most two
non-zeros per group of four along K, covering all six two-survivor position pairs, all four
single-survivor positions and the zero-survivor case, with values derived from (row, group, position) so that a misplaced survivor is more likely to show up numerically rather than cancel.
This is my input, not yours — nothing below should be compared against 352.

2. base fails, head passes

variant	K=32	K=64	K=128
base 8fc1ac2	max_abs_err=56	76	92	FAIL
head ac24ac2	0	0	0	PASS

The harness wraps hipMalloc / hipMemcpy / hipMemset / hipDeviceSynchronize / hipFree in
HIP_CHECK_ERROR, and none tripped, so the failing numbers are a real kernel result rather than a
silent launch failure.

3. The whole effect is bug 1

Two more variants — base plus only the nonzero_elems true-zero default, and head with only that
line reverted:

variant	K=32	K=64	K=128
base + bug-1 line only	0	0	0	PASS
head − bug-1 line	56	76	92	FAIL

The head − bug-1 stdout is identical to base's, byte for byte (diff on the run logs). So for the
failure this test measures, on these two commits, with this input and these three shapes, that one
line is both necessary and sufficient. I'm not claiming more than that — this doesn't establish
correctness over all valid inputs, types or architectures. The reasoning in your comment matches what
I see.

4. The test never reaches iu4, so bugs 2 and 3 get no numerical coverage from it

I disassembled the code objects of the kernels that actually launch, rather than grepping the
executable. The three SparseGemmKernel symbols (WaveTileK = 32 / 64 / 128) contain 1, 2 and 4
SWMMAC instructions respectively, zero v_wmma, zero v_mfma — and all seven are
v_swmmac_i32_16x16x32_iu8. There is no iu4 instruction anywhere in the binary.

That follows from the instantiation: the test uses SparseMmaPipeline<int8_t, int8_t, int32_t, ...>,
and int8_t picks up the generic numeric_traits::PackedSize == 1, so the if constexpr(PackedSize == 1) branch is taken, the packed-nibble path of bug 2 is compiled out, the SWAP + XOR-1 transform of
bug 3 is never reached, and TotalCompressedElems * MmaOp::APackedSize evaluates unchanged.

So the description reasons about all three bugs from the iu4 side, but the committed test only
exercises iu8. I haven't verified bugs 2 or 3 either — building an iu4 oracle independently of the
transform under test (nibble order, sign extension, logical vs packed K) is its own job, and reusing
your transform as the oracle would be circular.

What I could check is the shape side, with a probe that instantiates
SparseMmaPipeline<pk_int4_t, pk_int4_t, int32_t, ...> directly and launches it:

	base	head
TotalCompressedElems, K=32	4	8
TotalCompressedElems, K=64	8	16
TotalUncompressedElems	8 / 16	unchanged
IdxNumWords	1	unchanged at these shapes

Both trees compile, both emit v_swmmac_i32_16x16x32_iu4 and v_swmmac_i32_16x16x64_iu4, and
hipDeviceSynchronize() returns success. That's consistent with your ISA reading — 8 index values
per lane for a K=32 iu4 tile — but it's a shape and codegen check, not a correctness one.

5. The buildable configuration passes on the unfixed tree — and the reason is the one you identified

Built without -DUSE_REAL_TILE, so the #else synthetic path runs. (The substitute header from
§1 is still needed even here, since the #include is unconditional — dropping the define alone does
not make the file compile.)

	result
base 8fc1ac2, no USE_REAL_TILE	ALL PASS
head ac24ac2, no USE_REAL_TILE	ALL PASS

I expected that to be because the synthetic input has no under-filled groups, but that isn't it. I
replayed the exact host-side fill (mt19937(42), uniform_int_distribution(-8, 8), then
apply_sparse_pattern) and counted:

K	groups	0 non-zeros	1 non-zero	2 non-zeros
32	128	0	16	112
64	256	2	26	228
128	512	2	55	455

So under-filled groups are plentiful — the distribution includes 0 — and it still passes on the buggy
code. The reason is exactly the precondition you call out in the comment: apply_sparse_pattern
always zeroes slots 1 and 3, so survivors only ever sit at slots 0 and 2, and the old
{a_vec[i*4+2], a_vec[i*4+3]} default therefore always seeds slot 1 from a guaranteed zero. My tile
breaks it because survivors also sit at positions 2 and 3.

The practical consequence: once the missing header is supplied in whatever form, the configuration
that does not define USE_REAL_TILE reports a green run on the bug. Worth committing an input (or
generating one in the test) that puts survivors at position 3, and making that the default.

6. The change isn't scoped to gfx1201, and bug 1's fix demonstrably reaches CDNA

This is my main question before merge, and I don't think it's answerable from gfx1201 alone.

Neither changed hunk has an architecture predicate; both branch on PackedSize / APackedSize.
More to the point, the transforms selector in sparse_transforms.hpp:381 is specialized purely on
the op family:

struct MmaTransformsDefaultSelector<MmaOp, CompilerTarget,
                                    std::enable_if_t<MmaOp::OpFamily == MmaOpFamily::SPARSE>>
{ using SelectedTransforms = MmaDefaultTransformsSparse<MmaOp::kCompressionRatio>; };

with no enable_if_target_family_gfx*, unlike the dense gfx9 / gfx11 / gfx12 selectors right
alongside it. So every SPARSE op on every target resolves to this compress_a_impl. And
sparse/mfma/sparse_gfx9.hpp defines SPARSE ops for GFX942 / GFX950 over fp16_t, bf16_t,
int8_t, fp8_t, bf8_t — none of which specialize PackedSize, so they all take the same
PackedSize == 1 branch that bug 1's fix changes.

I checked that rather than inferring it. A compile-only probe that instantiates
SparseMmaPipeline<int8_t, int8_t, int32_t, 16, 16, 64, ..., Gfx942Target> and asserts

static_assert(std::is_same_v<
    typename MmaTransformsDefaultSelector<MmaOp942, Gfx942Target>::SelectedTransforms,
    MmaDefaultTransformsSparse<MmaOp942::kCompressionRatio>>);

compiles on both trees at --offload-arch=gfx942, and the resulting code object contains
v_smfmac_i32_16x16x64_i8. So that CDNA instantiation does route through the compress_a_impl this
PR changes, and takes the PackedSize == 1 branch.

That makes bug 1's fix a behaviour change for gfx942 as well, for inputs whose survivors don't sit
where the old default assumed. I think it's the correct fix and my gfx1201 result supports it — but
whether CDNA results actually change in practice, and whether anything downstream depended on the old
behaviour, is a runtime question I can't answer without the hardware.

The narrower half of this: I found no pk_int4_t SPARSE op for gfx9 in the tree, so bug 2's packed
branch and bug 3's SWAP + XOR-1 look gfx12-only in practice today, which limits the blast radius of
the empirically-derived metadata mapping considerably. It's still worth saying out loud that those
branches carry no architecture condition, so a future CDNA pk_int4 sparse op would silently inherit a
mapping that was measured on gfx1201.

On CI: the GitHub API reports no check suites and no commit statuses for either cd34d2f or
ac24ac2. I can't tell from outside whether anything ran elsewhere, but from the PR there's no
visible regression signal for the architectures the selector reaches. That's a project-side
infrastructure question rather than something to put on you.

7. Minor, non-blocking

• The comments carry internal process labels (Stage-17b fix (local, not upstreamed),
  Stage-17c CLOSURE fix, this Stage-17c audit's own probe) — worth rewording for upstream.
• The new static_assert(PackedSize == 2 && is_same_v<LogicalADataType, pk_int4_t>) narrows the
  packed path to pk_int4_t explicitly. I haven't tested whether any other packed A type reaches
  this code today, so I don't know if it changes anything in practice — just noting it's a scope
  change beyond the three bugs.
• You already note the gtest port is pending. For what it's worth, test/ck_tile/CMakeLists.txt
  enumerates its subdirectories with explicit add_subdirectory(...) calls rather than a glob, and
  gfx1201_sparse_swmmac isn't among them, so the port will need that line too.

What would let me say more

1. Commit real_a_tile.h, or generate an equivalent tile inside the test, so the fail→pass is
   reproducible from the PR alone — and make sure the default-built configuration is one that fails
   before the fix.
2. Add an iu4 / pk_int4_t case, since bugs 2 and 3 have no numerical coverage without one.
3. Say what architecture scope is intended, given §6.

Happy to re-run any of this on gfx1201. If it would help, I can also try to pin bug 3's mapping
empirically and independently of sparse_transforms — a minimal wave-level kernel issuing raw
v_swmmac_*_iu4 with one-hot operands, enumerating the metadata encodings and reading the
contributing positions back out of the result — and then compare what that gives against SWAP + XOR-1.
That would be a second measurement rather than a derivation from the doc, which as you note doesn't
capture the iu4-specific detail. Say the word and I'll put the generator, flags, logs and per-symbol
disassembly somewhere you can pull them.

[The-Monk]

Thanks for this — running it on your own gfx1201 and bisecting the line both directions is exactly the review this needed, and the §5 finding (the canonical pattern structurally can't detect the bug) plus §1 (missing header) are both fair hits. Fixes pushed in ad79359:

§1/§5 — repro now self-contained, and the default config fails-before/passes-after. real_a_tile.h is gone; the test now generates its adversarial tile in-source: deterministic 2:4 input cycling every group through all six two-survivor position pairs, all four single-survivor positions, and the zero-survivor group, with values derived from (row, group, position) so misplacement shows numerically — same design as your substitute tile. That's the default build now. The old slots-0,2 synthetic pattern is demoted to an opt-in control (-DUSE_CANONICAL_PATTERN) with a comment stating it cannot detect the bug. Re-verified before pushing, same procedure you used: identical test source on both trees — head ALL PASS (max_abs_err 0 at K=32/64/128), base 8fc1ac2 FAIL (max_abs_err 112/127/272). So the buildable-by-default configuration is the one that fails on the unfixed tree.

§4 — iu4 coverage. Agreed this is the real gap, and I'll take you up on the one-hot v_swmmac_*_iu4 metadata sweep — that's the independent second measurement bug 3 needs, since my SWAP+XOR-1 mapping is device-measured but through my own transform stack. Say where you drop the generator/logs/disassembly and I'll cross-check against my derivation. Related question: do you have a working HIP-level iu4 route already — a __builtin_amdgcn_swmmac_*_iu4 builtin that lowers correctly on this clang, or an inline-asm wrapper from your probe work? Everything I have reaches the instruction through the CK pipeline; a bare intrinsic path would make the one-hot sweep trivially shareable and would also give my pk4 test an oracle route that doesn't touch the transform under test. On my side I've added a pk_int4 case scaffold to the test with an honest marker: the CPU oracle is straightforward (dense int4 reference from the same uncompressed logical values — not circular), but the host-side register-map fill convention for packed tensors is unproven on hardware, so it's compiled out by default (-DENABLE_PK4_CASE) until validated rather than pretending coverage.

§6 — intended scope. Bug 1's fix is intentionally architecture-generic: the old {a[2], a[3]} default is out-of-spec for any input whose survivors don't sit at 0/2, on every target the sparse selector reaches — your gfx942 compile-probe confirms the reach, and I'd argue the CDNA behavior change is the fix working as intended. If anyone with CDNA hardware can run the (now self-contained) repro there, great, but I don't think it should block. Bugs 2/3 are gfx12-only in practice today (no gfx9 pk_int4 sparse op, as you found); I've added an explicit ARCH SCOPE comment at the packed path stating the SWAP+XOR-1 mapping is gfx1201-measured and must be re-measured before any future CDNA pk4 sparse op trusts it.

§7 — housekeeping. Internal stage labels reworded to bug-number references throughout. Noted on the static_assert scope point — it's declaring the current implementation boundary, not narrowing behavior (nothing else reaches the packed branch today), and the scope comment now says so. Thanks for the add_subdirectory catch for the eventual gtest port — on the list for that change.

Your env notes (--rocm-device-lib-path, the libamdhip64.so symlink) match what I hit as well — I'll fold them into the test's header comment with the gtest port so the next person doesn't rediscover them.

[The-Monk]

Since you have gfx1201 and clearly care about this corner of the stack — some context on where this PR came from, and the wider map it belongs to. While building a 2:4-sparse fp8/int4 inference path for RDNA4 we kept hitting the same pattern: the capability is in the ISA, but no library exercises it, so the first real user finds the bugs (or the absence). The list of gaps we personally hit and had to fill with hand-written kernels, in case any of it is useful to you or worth more upstream issues:

Matrix/dot datapaths:

• 2:4-sparse SWMMAC (v_swmmac_*, ISA §7.12.3) — this PR's territory. CK's SPARSE family was silently wrong for arbitrary-position 2:4 (bug 1), the sparse selector is un-arch-specialized as you found, and as far as we can tell no inference stack used sparse prefill on RDNA4 at all before we did (our llama.cpp fork uses it for prefill: The-Monk/llama.cpp, roc8/roc10 branches).
• iu4 packed sparse (v_swmmac_i32_16x16x{32,64}_iu4) — CK's packed path was structurally unreachable/broken (bugs 2/3), no gfx9 pk_int4 sparse op exists, the ISA's generic sparse pseudocode doesn't capture the iu4 idx semantics (hence the empirical SWAP+XOR-1), and we still don't know of a HIP builtin that lowers it correctly (the question upthread).
• dp4a family for sub-2-bit quant (v_dot4_i32_iu8) — the generic HIP lowerings for 1-bit/ternary decode were poor on AMD (measured +37% on Q2_0 from one HIP-guarded vec_dot, cuda: AMD RDNA4 Q1_0/Q2_0 — HIP-path vec_dots (+37%/+16% decode), opt-in quant dedup, opt-in hipBLASLt prefill routes PrismML-Eng/llama.cpp#116) and binary had no dp4a decode at all; we wrote bit-spread+dp4a kernels for both.
• FP8 (E4M3/E5M2) WMMA — mainline llama.cpp had no fp8 tensor types whatsoever; our fork added the types + kernels (bit-exact verified, ~96% of bandwidth roofline at 14B). MXFP8 similarly hardware-supported, zero stack coverage.

Verification gaps that compound it — validating new datapath kernels on gfx1201 is disproportionately hard because: device ASAN is unavailable (no xnack+ target-ID — exactly the tool that would have caught bug 1/2's OOB class), stochastic PC sampling is unsupported (host_trap only), and the GL2C EA size-split counters read zero even with full ppfeaturemask + profile_standard (base GL2C counters work — so it's counter-plumbing, not the perfmon clock). If you've gotten any of those three working on your 9070 XT I'd genuinely like to know how.

Happy to compare notes on any of these, or split out issues where they belong in other repos — the sparse/iu4 items are the CK-relevant ones, which is why they became this PR.

[doplxyz]

Re-ran everything on ad79359d0 on the 9070 XT. Short version: the new default
test reproduces exactly as you describe, and yes — there is a bare builtin
route to iu4, it runs on hardware, and it links no CK code.

All artifacts, logs, compile commands and disassembly for both rounds:
https://github.com/doplxyz/ck3759-gfx1201-verification @ 8cee1e3.

---

1. The bare iu4 intrinsic route you asked about — it exists

The builtins are on this toolchain. The name needs a _w32 suffix; without
it clang reports an undeclared identifier, which may be what you hit:

__builtin_amdgcn_swmmac_i32_16x16x32_iu4_w32(bool, int,  bool, int2, int8, int, bool);
__builtin_amdgcn_swmmac_i32_16x16x64_iu4_w32(bool, int2, bool, int4, int8, int, bool);

Seven arguments. I left the parameters unnamed above because I measured what
they do rather than assuming what they mean — one kernel per boolean
combination, read off the emitted asm:

(arg1, arg3, arg7)	modifiers on v_swmmac_i32_16x16x32_iu4
F, F, F	(none)
T, F, F	neg_lo:[1,0,0]
F, T, F	neg_lo:[0,1,0]
T, T, F	neg_lo:[1,1,0]
F, F, T	clamp
T, T, T	neg_lo:[1,1,0] clamp

So arg1 sets neg_lo[0], arg3 sets neg_lo[1], arg7 sets clamp, and
neg_lo[2] is not reachable from these arguments. CK passes arg1/arg3 as the A
and B signedness flags; I'm reporting the modifier mapping I measured rather
than re-asserting that interpretation. Probe source is
src/iu4_builtin_modifier_probe.cpp.

src/iu4_bare.cpp is a complete HIP program that includes nothing from
composable_kernel — no headers, no transforms, no layout helpers. It compiles,
launches, and returns data. Its device object contains exactly one iu4
instruction per kernel, zero iu8 instructions, and zero ck_tile
symbols:

v_swmmac_i32_16x16x32_iu4 v[0:7], v11, v[8:9], v12       neg_lo:[1,1,0]
v_swmmac_i32_16x16x64_iu4 v[0:7], v[12:13], v[8:11], v15 neg_lo:[1,1,0]

Per-lane operand sizes for wave32: K=32 takes A = 1 dword, B = 2 dwords,
D = 8 dwords; K=64 takes A = 2 dwords, B = 4 dwords, D = 8 dwords; idx is one
dword per lane in both. With every A and B nibble set to 1, every accumulator
dword in every lane reads 16 at K=32 and 32 at K=64 — the sparse product
counts — and a zero-A control returns all zeros. Bit-exact across fresh
processes.

Toolchain: AMD clang 23.0.0git, ROCm/llvm-project @ 46fcb339fb61, inside
rocm/pytorch:rocm7.14_ubuntu24.04_py3.12_pytorch_release_2.12.0. I have not
checked which older ROCm releases carry these builtins, so treat the _w32
names as confirmed for this snapshot only.

What this is and is not. It's a numeric route to the instruction that
doesn't pass through the transform under test. That's what makes a non-circular
second measurement possible, and it may be useful as an oracle route for your
pk4 case. It is not a measurement of idx semantics, and it is not a solved
register map: an all-ones stimulus is invariant under exactly the permutations
that matter, so it says nothing about which logical element landed where. I've
deliberately not formed a view yet on what any idx bit means. I'll pin your
current formula and its predicted table with a hash before I measure, so the
comparison can't drift after the fact.

To be precise about one thing, since you may want to reuse it: iu4_bare.cpp
doesn't avoid packed representation — the nibbles are still packed into dwords.
What it avoids is depending on CK's host-side fill helpers to do that packing,
because it writes and reads raw dwords directly. Establishing which raw nibble
position corresponds to which logical element is exactly the job of the
calibration below, not something the harness assumes.

2. ad79359d0 re-verified — your numbers reproduce exactly

I used one immutable copy of your new test for every tree, so base and head
are never compared with different test sources. Clean builds (each variant's
output directory removed first), three fresh processes each. max_abs_err:

build	tree	K=32	K=64	K=128
default (adversarial)	base 8fc1ac2	112	127	272	FAIL
default	head ad79359d0	0	0	0	PASS
-DUSE_CANONICAL_PATTERN	base 8fc1ac2	0	0	0	PASS
-DUSE_CANONICAL_PATTERN	head	0	0	0	PASS
default	head minus bug 1	112	127	272	FAIL
default	base plus bug 1 only	0	0	0	PASS

All six are run1 == run2 == run3 byte-for-byte. Your 112/127/272 match to the
digit. §1 and §5 are closed from my side: the test is self-contained, the
default configuration is the one that fails on the unfixed tree, and the
canonical control passes on base — preserving the evidence that the old pattern
could not detect this.

Two further things from the same table. The head minus bug 1 build produces
stdout byte-identical to base, and bug 1 only passes — so for this test's
inputs, bug 1's one-line change is both necessary and sufficient for the
observed difference. That's a statement about this test, not about the fix's
full effect. And I disassembled all six builds: each contains 7
v_swmmac_i32_16x16x32_iu8 and zero iu4 instructions, so this table is a
bug-1 regression result and I'm not counting any of it as evidence for bugs 2
or 3.

I also checked that your header delta from ac24ac28d is comment-only: with
comments stripped, both sparse_transforms.hpp (163 lines) and
sparse_mma_pipeline.hpp (206 lines) are byte-identical across the two commits.
So my round-1 four-way bisect carries over to the current head without
re-running it. Hashes are in EVIDENCE_round2.md.

3. Nothing in the tree references the test

test/ck_tile/gfx1201_sparse_swmmac/ contains the .cpp and nothing else — no
CMakeLists.txt — and test/ck_tile/CMakeLists.txt has no add_subdirectory
for it. I grepped the whole head tree: the string gfx1201_sparse_swmmac
appears nowhere outside that directory. No CMake file, no script, no
Jenkinsfile stanza refers to it. So as the PR stands, nothing in the project's
own tooling can find this test — every result in this thread comes from
hand-driven builds, including mine.

I raised add_subdirectory last time as a footnote for the eventual gtest port.
That was too soft: whatever form the test finally takes, something in the tree
has to point at it, or the fix has no standing guard.

I want to be careful about the CI half of this, because I was about to overstate
it. The repository has no .github/ directory at all, so it isn't using GitHub
Actions and the empty check-suites on all three commits are the expected
consequence of that, not evidence of anything. CI here runs from the top-level
Jenkinsfile, which doesn't report into GitHub Checks — meaning I can't see from
outside whether it ran on this PR, what it covered, or whether it has gfx1201
hardware. That's worth asking a maintainer directly, and it's a different
question from the registration gap, which is visible in the tree and cheap to
close.

4. A thought on splitting bug 1 out

Entirely the maintainers' call and yours — I'm a drive-by reviewer with one GPU.

Bug 1 is in good shape: the test's default configuration fails before and passes
after on real hardware, independently reproduced, and the change is a plain
out-of-spec default. I find your §6 argument for keeping it
architecture-generic convincing. My probe result there is narrower than the
claim, though, and I should have said so last round: it shows that a gfx942
compile of the sparse path succeeds and emits v_smfmac_i32_16x16x64_i8. That
establishes the selector reaches gfx942 at compile time. It is not a runtime
result and doesn't generalize across gfx9 by itself.

Bugs 2 and 3 sit differently: the pk4 case is compiled out by default behind
-DENABLE_PK4_CASE, and the executing kernels are iu8-only, so as far as
anything in this thread shows, no number yet distinguishes bug 2 or bug 3 being
present from being absent.

Merging bug 1 now and keeping bugs 2 and 3 in a follow-up until the sweep and a
CK-level pk4 test land would get the confirmed fix into the tree without waiting
on the part that still needs measurement. Splitting inside this PR works just as
well if you'd rather not open another one — the point is only that the
verified part shouldn't have to wait on the unverified part.

5. What's next, and what I can't do

In order: the one-hot metadata sweep proper — full raw idx enumeration against
individually one-hot compressed slots, low and high nibble stimulated
separately, B basis-swept so I can tell which physical B element each result came
from, all 32 lanes and all 8 accumulator dwords collected, K=32 and K=64 swept
independently rather than one extrapolated from the other, with the lane/element
map established by calibration rather than borrowed from CK. Then the same
question end to end through CK's pk4 path, with bug-2-only and bug-3-only revert
mutants, since if those two can mask each other a head-vs-base comparison won't
separate them. Everything lands in the repo above with the raw tables.

Two things I can't close from here. I have no CDNA hardware, so my gfx942 result
stays compile-level; someone with a gfx942 part running the now-self-contained
repro would turn that into a runtime answer. And I can't tell whether this
repository's maintainers want a hardware-specific test registered in the normal
test graph, gated behind an option, or kept standalone — that decision, and what
evidence they'll accept for a target they may not have in CI, has to come from
them. Worth asking explicitly rather than either of us guessing.

On your verification-tooling questions (device ASAN, stochastic PC sampling, the
GL2C EA size-split counters reading zero): I have the same card and would rather
give you measurements than guesses, so I'll try all three and report back —
after the sweep, purely because I don't want to leave that half-finished. I've
subscribed to ROCm/ROCm#6613. I can see why you grouped them with this work; the
absence of device ASAN is precisely why bug 2's class had to be found the hard
way.

[doplxyz]

Sweep done, plus an end-to-end pk4 test. Headline: your bug-3 mapping is
correct, and I can now say why it has the shape it does. The run also turned up
something worth handling before this merges: the PR makes an existing test in
this repository fail.

Everything I measured is reproducible from
https://github.com/doplxyz/ck3759-gfx1201-verification @ b42b707. (Claims
below about what does or doesn't exist in the tree are from grepping it; claims
about history or anyone else's environment aren't things I can check.)

---

1. The one-hot iu4 metadata sweep — SWAP + XOR-1 confirmed

Method, as offered: bare __builtin_amdgcn_swmmac_*_iu4_w32, no
composable_kernel in the binary at all. For each (group, field pair, compressed slot) I drove one compressed A nibble and scanned every raw B position; a
nonzero accumulator names the position the hardware paired with that slot. No
logical coordinate system is assumed in the measurement — the grouping of B
positions and their order within a group come out of the data.

I pre-registered your rule and its predicted table, hashed and committed, before
taking any data (PREREGISTERED_HYPOTHESIS.md).

Result: your transform reproduces the hardware pairing exactly. 384 cells
across K=32 and K=64 — the full 4×4 field grid, both slots, every group, all 32
lanes — zero mismatches.

The raw law underneath:

idx field i governs compressed nibble i; a field carrying raw value v
pairs its nibble with the B nibble at raw offset v within that group.

Plain identity. No swap, no XOR. Both halves of your rule are the two places
CK_TILE_USE_PK4_LAYOUT_SHUFFLE enters — element 0 of a packed byte being the
HIGH nibble:

	CK index	raw index
compressed slot	s (0 = high)	r = 1 - s → the SWAP
uncompressed position	j (0 = byte0 high)	o = j XOR 1 → the XOR

Substituting both into the identity law reproduces your code line for line, and
your corollary falls out: raw value 2 selects raw offset 2, which is CK position
3 — constant regardless of the real survivor, exactly the "always reconstructs at
position 3" you reported.

Adversarial checks, each breaking one way the first pass could have looked like
an identity without being one:

• lane uniformity — one active A lane, every other lane carrying the opposite
  idx: 1024/1024 followed the active lane's own idx. That rules out the specific
  failure of another lane's metadata being substituted; it isn't a general proof
  that no cross-lane routing exists.
• arbitrary idx words — all groups live, a different random word per lane:
  192/192.
• superposition — leave-one-out on dense random inputs: 1919 nibbles, 0
  mismatches.
• operand range — the full signed 4-bit range including −8, with a non-zero
  starting accumulator: 256/256 exact.
• identical at -O0, -O2, -O3.
• dataflow — between the operand loads and the instruction there are no
  instructions at all touching the loaded values; the only ALU work in that
  window is address arithmetic on the thread id. So the mapping isn't something
  the lowering introduced.

A correction against my own case: my pre-registered transcription of your
rule was wrong. I wrote it in raw coordinates when it is defined in CK's,
silently assuming CK slot and position indices equal raw nibble indices. Scored
literally it "matches" 32/128, and anyone comparing the raw table against my
pre-registration would wrongly conclude your fix is broken. The 128/128 figure is
a post-hoc coordinate translation, not a confirmed prior prediction. Both are
recorded separately and the pre-registration is unmodified in git history.

On the ARCH SCOPE comment

Keep the instruction to re-measure — I was going to argue against it and I was
wrong. A future CDNA sparse op could differ in field assignment, operand
numbering or lane routing even with identical CK packing, so per-architecture
re-measurement stays necessary.

What I'd add is that the constants aren't encoding a gfx1201 oddity — the
gfx1201 encoding is the identity — they're encoding CK's packing convention.
Today that convention isn't switchable: config.hpp:181 defines
CK_TILE_USE_PK4_LAYOUT_SHUFFLE whenever it isn't already defined and
pk_int4.hpp tests it with #ifdef, so setting it to 0 changes nothing and
the #else branches there are unreachable. But that also means those constants
depend on a convention that a source change could flip without touching
architecture at all, and the #else branches are already dead code that would
silently disagree with the sparse path if revived. Worth a sentence in the same
comment.

2. SparseTransformsTest.SingleNonZeroPerGroup fails on this branch

This is the one I'd act on first.

test/ck_tile/core/arch/mma/pipeline/test_amdgcn_sparse_mma.cpp is registered
via _add_mma_gtest in test/ck_tile/core/arch/mma/CMakeLists.txt:17. It exists
at base 8fc1ac2 and this PR does not modify it. One of its cases drives
compress_a_impl on device with a single survivor per group and asserts:

// Single non-zero per group of 4 (at slot 3).
// nonzero_elems initializes to {a_vec[slot2]=0, a_vec[slot3]=V}.
// Only j=3 triggers: nonzero_elems[0]=V, field0=0b11, pos becomes 1.
// nonzero_elems[1] keeps its init V. Output: {V, V}.
expected_output[g * 2]     = val;
expected_output[g * 2 + 1] = val;

That expected value is bug 1's behaviour — the {a[2], a[3]} default reaching
the second compressed slot — recorded as the expected result.

I built and ran that file, unmodified, against both include trees, varying
nothing but -I:

include tree	result
base 8fc1ac2	0 failing of 18
head ad79359d0	1 failing of 18 — SparseTransformsTest.SingleNonZeroPerGroup

EXPECT_EQ failed at test_amdgcn_sparse_mma.cpp:202   (x9, the compressed values)
compressed out, base: 5 5 6 6      <- what the test expects
compressed out, head: 5 0 6 0      <- what your fix produces
idx unchanged (0x000000bb) in both

MixedSparsityPattern, NonZerosAtSlots{1And3,0And3} and all eight
FullMatrixVerify_* cases pass on both trees; this is the only one that moves.

Caveat on method: googletest isn't installed in my container, so I supplied a
minimal gtest.h shim (TEST, EXPECT_*, ASSERT_*, GTEST_SKIP) rather than
building the CMake target. The test source itself is byte-for-byte upstream, and
the shim clearly does detect failures since it found this one — but I haven't run
the registered target through CMake, so treat this as "the upstream test's own
assertions fail on head", not "the CI job goes red" (which I can't observe
anyway).

Two things follow. The test needs updating as part of this PR, and that diff is
some of the better evidence in the change, since it shows the old behaviour
written down and corrected. And it explains why the two-survivor cases don't
move: bug 1 only bites when a group has fewer than two survivors, which
SingleNonZeroPerGroup is the only case in that file to supply. That's also the
shape of the gap new cases would fill.

It makes the registration point from my last comment concrete too: there's
already a registered home with the right idiom, so the standalone repro could
become extra cases in that file rather than a new directory — no new CMake entry
needed at all if it goes there.

3. End-to-end pk_int4 through CK

ENABLE_PK4_CASE expands to a printf and a TODO; it doesn't instantiate the
packed pipeline and asserts nothing. Grepping test/ and example/, the only
two files referencing the sparse pipeline are your standalone repro and
test_amdgcn_sparse_mma.cpp, and neither instantiates it with pk_int4_t. So as
far as the tree shows, bugs 2 and 3 have no numerical coverage in it, and the
numbers below are the first I can find evidence of for the packed path.

The fill convention you flagged as unproven, resolved. Asked CK's own types
rather than guessed:

• for pk_int4_t the register map's vector index enumerates logical 4-bit
  elements, not physical bytes — at K=32, num_vector_items is 16 while
  sizeof(AWarpTensor) is 8, so porting the int8 fill verbatim overruns the
  tensor by 2×;
• logical element v goes to byte v/2, high nibble for even v;
• the coordinate order is pinned without copying it from the code under test:
  K ≠ N, so the two components have different ranges. Measured, A and B both
  return coord[0] in [0,15] and coord[1] in [0,K-1], so B is
  B[coord[1] * N + coord[0]]. I had this backwards at first and it produced a
  plausible-looking wrong answer, which is why I went back and pinned it from the
  ranges.

Logical k landing at raw nibble offset k XOR 1 is the same XOR the sweep
found from the hardware side. Not fully independent evidence — both involve the
same nibble convention — but they're arrived at from opposite ends.

Kill matrix. One independent build and run per shape, so a compile failure at
one shape can't mask another. max_abs_err against a dense int4 CPU reference
over the same logical values; int64 accumulation, exact comparison.

Every mutant row sits on top of the §4 fix, since without it nothing builds at
K≥128 and the rows would not be comparable. The first two rows show that fix's
effect on its own.

tree	K=32	K=64	K=128	K=256
head ad79359d0, unmodified	0 PASS	0 PASS	compile fail	compile fail
head + §4 fix — the baseline for the rows below	0 PASS	0 PASS	0 PASS	0 PASS
… + bug 2 reverted	0 PASS	0 PASS	compile fail	compile fail
… + bug 3 fully reverted	311 FAIL	451 FAIL	635 FAIL	904 FAIL
… + bug 3, swap kept, XOR removed	460 FAIL	810 FAIL	1142 FAIL	2290 FAIL
… + bug 3, XOR kept, swap removed	351 FAIL	588 FAIL	887 FAIL	1397 FAIL
… + bug 1's true-zero default reverted, packed path	64 FAIL	68 FAIL	89 FAIL	132 FAIL

• Bug 3 is load-bearing, and in these configurations neither half alone
  suffices — the CK-side counterpart to §1. (It doesn't rule out some other
  equivalent formulation; only that these two partial forms are wrong.)
• Bug 1's true-zero default is load-bearing in the packed path too, not only the
  scalar path — the same defect §2's existing test records.
• Bug 2's mutant survives at K=32 and K=64 and is killed only at K=128 and
  K=256, as a compile error. The condition isn't "multi-fragment" as such but
  "shapes where the idx word count diverges": at K=32/64 both accountings round
  to a single word, so the difference is unobservable there. Whatever form a
  regression test for bug 2 takes — end-to-end or a type-level unit test — it
  needs a case where those counts differ, or it won't exercise that fix.

The test carries guard bands around every per-lane buffer, verified intact after
the fill so an off-by-a-factor is loud rather than silent; asserts stimulus
coverage rather than assuming it (all eleven at-most-two-survivor group patterns
present — the six two-survivor pairs, four single-survivor, one all-zero — and
the signed end point −8 present in both operands); and uses a B spanning the full
signed range, asymmetric in k and n so a transposed reading can't cancel out.

4. Multi-fragment packed A doesn't compile

At K=128 and K=256 the pk4 pipeline fails to build on head:

sparse_mma_pipeline.hpp:311: 'ATransformResult must match the return type of
ATransform::exec'   SparseIdxPack<2> vs SparseIdxPack<1>

checkATransformResult() re-derives the expected type as

decltype(ATransform::execExtVec(std::declval<ExternalAvecRef>()))

leaving LogicalADataType at its default. ext_vector_t<pk_int4_t, N> has
signed char as its scalar type, so the default resolves PackedSize to 1 —
while the real call path, exec(), passes ADataType explicitly and gets 2. The
check disagrees with the call it validates. It's hidden at K=32/64 because 8 and
16 two-bit fields both round to one idx word; the counts first diverge at K=128.

         static_assert(
             std::is_same_v<ATransformResult,
-                           decltype(ATransform::execExtVec(std::declval<ExternalAvecRef>()))>,
+                           decltype(ATransform::template execExtVec<AVecType, ADataType>(
+                               std::declval<ExternalAvecRef>()))>,
             "ATransformResult must match the return type of ATransform::exec");

With that, all four shapes compile and pass, and your int8 tests are unaffected
(adversarial ALL PASS, canonical control ALL PASS). Whether K≥128 is a shape the
packed path is meant to support is your and the maintainers' call — the argument
that it is comes only from your int8 test covering K=128 and the pipeline being a
template over K, which isn't conclusive for a different data type. What is
measured is that the check and the call disagree. I haven't run a compile matrix
over every type/shape combination, so read the diff as "makes the check agree
with the call", not as "verified side-effect-free".

5. What this doesn't show

Mutant kills show the test is sensitive to the code under test; they aren't proof
head is correct. If head and my harness shared a wrong layout assumption, mutants
could still die while a correlated error passed — the range-based determination
of the coordinate order and the CK-free sweep are what reduce that risk, not the
kill matrix. Everything is FragsM = FragsN = 1, gfx1201, one toolchain, wave32,
full EXEC. I still have no CDNA hardware, so bug 1's reach there stays a
compile-level result on my side.

Happy to open a PR against your branch with any of this rather than leave you to
lift it — the pk4 case and the SingleNonZeroPerGroup update being the two that
matter most for someone reviewing without a gfx1201 part.