ROCm IQ4_XS keep-quant GEMM is 8.8x behind llama.cpp's WMMA path at the same CPU/GPU split

Row: -

Measured

Same file, same host, same 15-of-40-layer CPU-MoE-offload split, idle
isravale (RX 9060 XT, gfx1200, ROCm 7.2.3), `Nail-Qwen3.6-35B-A3B-MTP-IQ4_XS.gguf`:

Decode is ~8.8x behind at the identical split. This is a personal-box check
(no clock pinning, 1 rep on our side vs 3 on llama.cpp's), not the project's
formal gated benchmark, but a gap this size at tight llama.cpp reproducibility
(±0.47 t/s) is not noise.

Cause, read from source rather than assumed

`KERNEL-QUANT-CIQ-GEMM-ROCM-IQUANT` (this repo, row branch
`row/KERNEL-QUANT-CIQ-GEMM-ROCM-IQUANT-w1`) just landed `DotIQ4XS`/`DotIQ3XXS`
on ROCm, but scalar-only — a warp-shuffle reduction, no tensor-core tile. Its
own spec named this as owed, not attempted.

llama.cpp already runs IQ4_XS through its MMA (tensor-core) path on this exact
architecture (`ggml/src/ggml-cuda/mmq.cuh:824-826`, dispatching to
`ggml_cuda_mmq_load_tiles_iq4_xs` + `ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma`,
pin `c589f0ed1`). The loader (`ggml-cuda/mmq-load-tiles.cuh:1420-1481`)
unpacks each block's 4-bit nibbles into int8 via `get_int_from_table_16(...,
kvalues_iq4nl)` into a shared-memory tile, then the SAME generic MMA
multiply every other MMA-tiled format uses consumes it -- the codebook lookup
and the tensor-core multiply are separate steps, which is why an int8 lookup
codebook doesn't block a WMMA tile the way it looked like it might.

What already exists to port from

`KERNEL-QUANT-CIQ-GEMM-ROCM` (issue #2109) already landed a WMMA int8 tile
for Q4_K/Q6_K on this same gfx1200 hardware
(`src/vt/rocm/rocm_grouped_gemm.hip`, `KQuantGemmKWmmaQ4K`/
`KQuantGemmKWmmaQ6K`), keeping our own per-(i,j)-output launch geometry
(unlike llama.cpp's whole-tile MMQ scheme) and proving the mechanism
(`rocwmma` `fragment<matrix_a,16,16,16,int8_t>`) works on this card. Its
pattern -- a per-16-element `DequantXTile16(...) -> int8_t[16]` unpack
function feeding a WMMA fragment, with the block's scale applied once per
super-block after integer accumulation -- is the direct template for IQ4_XS
(nibble-via-codebook unpack, already signed, no recentering) and IQ3_XXS
(grid-index + sign-flip unpack, same shape as the existing scalar `DotIQ3XXS`
inner loop).

Scope

A WMMA tile for IQ4_XS and IQ3_XXS on ROCm, mirroring `KERNEL-QUANT-CIQ-GEMM-ROCM`'s
scope discipline: gfx1200/gfx1201 (RDNA4) only, scalar tier stays the
fallback everywhere else. Needs a spec before code per this repo's protocol.

Start with issue #2109 and src/vt/rocm/rocm_grouped_gemm.hip, especially KQuantGemmKWmmaQ4K, KQuantGemmKWmmaQ6K, and the rocwmma fragment pattern. Define the required spec before implementation, then compare the existing scalar DotIQ4XS/DotIQ3XXS paths with the proposed DequantXTile16 approach. Done means WMMA tiles work for IQ4_XS and IQ3_XXS on gfx1200/gfx1201, scalar fallback remains elsewhere, and the same-split benchmark improves.