test: register e2e coverage, fix bug-1-encoding expectation, add pk4 e2e; fix pk4 K>=128 type check

Review-round changes, all four from doplxyz's independent verification
(github.com/doplxyz/ck3759-gfx1201-verification):

- SparseTransformsTest.SingleNonZeroPerGroup expected {V, V}: the second
  compressed slot leaking a_vec[slot3] through the pre-fix {a[2], a[3]}
  default -- bug 1 recorded as the expected result. Corrected to {V, 0}
  with a HISTORY note; this was the only registered case supplying a
  group with fewer than two survivors, which is exactly where bug 1
  bites.

- The standalone repro moves from test/ck_tile/gfx1201_sparse_swmmac/
  (unregistered, no CMakeLists) into test_amdgcn_sparse_mma.cpp as
  SparseSwmmacE2E.{AdversarialGeneratedTile,CanonicalPatternControl};
  the compile-time USE_CANONICAL_PATTERN toggle becomes a runtime
  parameter so the control runs in the same binary. Runtime-skipped on
  non-gfx12 devices.

- New SparsePk4E2E.AdversarialInt4AllShapes (K=32/64/128/256): first
  in-tree numerical coverage for the packed-nibble path (bugs 2/3).
  Ported with attribution from doplxyz's MIT-licensed pk4_e2e harness:
  dense int4 CPU oracle over logical values (independent of the
  transform under test), guard-banded host fill via CK's own register
  maps. K=128/256 also lock bug 2 and the type-check fix below at
  compile time -- idx word counts only diverge at FragsK > 1.

- checkATransformResult (sparse_mma_pipeline.hpp) re-derived the
  expected transform type via execExtVec with the defaulted
  LogicalADataType, resolving PackedSize to 1 while the real exec call
  passes ADataType and gets 2 -- so the pk4 pipeline failed to compile
  at K>=128. Fix (explicit template arguments) by doplxyz, applied
  verbatim.

- ARCH SCOPE comment extended: one-hot sweeps with bare
  v_swmmac_*_iu4_w32 builtins show the gfx1201 idx law is a plain
  identity in raw nibble coordinates; SWAP + XOR-1 encode
  CK_TILE_USE_PK4_LAYOUT_SHUFFLE's high-nibble-first convention, not a
  hardware quirk, and config.hpp pins that macro on (the pk_int4.hpp
  #else branches are unreachable). Convention dependency documented
  alongside the existing per-architecture re-measurement instruction.

Verified on gfx1201 (Radeon AI PRO R9700, ROCm 7.14): all 21 tests in
test_amdgcn_sparse_mma pass, including the four pk4 shapes (max_abs_err
= 0 exact), the adversarial e2e that fails on the unfixed tree, and the
corrected single-survivor case.

CCA

Author: jmonk

include/ck_tile/core/arch/mma/sparse/sparse_mma_pipeline.hpp

@@ -309,7 +309,8 @@ struct SparseMmaPipeline : public MmaPipelineBase<SparseMmaPipeline<ADataType_,
         using ExternalAvecRef = std::add_lvalue_reference_t<AVecType>;
         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");
 
         using CompressedVecType =

include/ck_tile/core/arch/mma/sparse/sparse_transforms.hpp

@@ -174,6 +174,20 @@ static CK_TILE_DEVICE auto compress_a_impl(AVec& a_vec)
         // before this path is trusted there. (The scalar PackedSize==1 branch
         // above is architecture-generic: its fix corrects a plain
         // out-of-spec default and changes behavior identically everywhere.)
+        //
+        // CONVENTION SCOPE: independently measured one-hot sweeps (bare
+        // v_swmmac_*_iu4_w32 builtins, no CK code in the binary) show the
+        // gfx1201 hardware idx law is a plain IDENTITY in raw nibble
+        // coordinates: idx field i pairs compressed nibble i with the B
+        // nibble at raw offset equal to the field value. The SWAP and XOR-1
+        // below are therefore not a hardware quirk -- they are the
+        // coordinate change from CK's packed convention
+        // (CK_TILE_USE_PK4_LAYOUT_SHUFFLE: logical element 0 = HIGH nibble)
+        // into raw nibble order. config.hpp currently defines that macro
+        // unconditionally and pk_int4.hpp tests it with #ifdef, so its
+        // #else branches are unreachable dead code; if the convention ever
+        // becomes switchable, these constants must change WITH it -- a
+        // source-convention dependency, independent of architecture.
         // Packed-nibble path (pk_int4_t): each iteration consumes 2
         // PHYSICAL input bytes (= 4 LOGICAL nibbles, one real 2:4 group) and
         // produces 1 PHYSICAL output byte (the <=2 survivor nibbles,