ROCm KQuantGemmKWmmaQ4K/Q6K: redesign around a cooperative shared tile, matching llama.cpp's mul_mat_q mechanism

Row: KERNEL-QUANT-CIQ-GEMM-ROCM

Row

KERNEL-QUANT-CIQ-GEMM-ROCM — already ACTIVE. Follow-on to #3032/#3033
(wider-block A/B, measured and rejected, geomean -4.3%). This issue is the
sharper hypothesis that measurement pointed at.

Confirmed mechanism (read, not inferred)

KQuantGemmKWmmaQ4K/Q6K give each WARP an independent 16x16 output tile:
its own shared-memory slice, its own load_matrix_sync calls for both the
weight and activation fragments, zero sharing with any other warp in the
same block — even when two warps share the same M-tile (it) and are
reading the identical 16 activation rows.

llama.cpp's mul_mat_q does not do this. At this row's shapes
(ggml/src/ggml-cuda/mmq-config-rdna4.cuh:127,
GGML_TYPE_Q4_K, 256 threads, I=128, J=128), ONE BLOCK owns a 128x128
output region:

• ggml_cuda_mmq_load_tiles_q4_K (ggml/src/ggml-cuda/mmq-load-tiles.cuh:703-741)
  cooperatively loads and dequantizes the I=128 weight rows ONCE, striped
  across every warp's threads (i0 += nrows*nwarps) — no warp repeats work
  another warp already did.
• mmq_get_nbytes_shared (ggml/src/ggml-cuda/mmq.cuh:1379-1383) sizes the
  activation-tile shared allocation by J alone, not J*nwarps — one
  shared copy of the activation tile serves the whole block too.
• ggml_cuda_mmq_write_back_mma (ggml/src/ggml-cuda/mmq.cuh:476-500)
  splits the I rows across warps (rows_per_warp = I/nwarps); each warp
  loops the full J range doing 16x16 MMA against that one shared load.

Our kernel's per-warp-independent design is the actual gap the rejected
wider-block experiment could not reach, because widening only added more
independent warps to the same per-warp-load pattern — it never introduced
sharing.

The work

Redesign KQuantGemmKWmmaQ4K/KQuantGemmKWmmaQ6K so one BLOCK cooperatively
loads and dequantizes a bigger shared I x J tile (weight rows x activation
rows), with warps splitting the I dimension for compute and looping the
full J range per warp — matching the mechanism above, not necessarily
llama.cpp's exact bucket sizes or its multi-bucket (64x16 .. 128x128)
dispatch table on the first pass. A spec proposing a concrete, scoped tile
size (fitting this card's 64 KiB LDS budget, reusing the existing
per-superblock staging/epilogue shape where possible) lands before any
kernel code, per this row's own recorded PR-shape preference.

Named risk: llama.cpp's own multi-bucket dispatch exists because a single
big I/J tile wastes work on shapes smaller than it (our M is prompt length,
essentially never aligned to a large tile; our benchmarked N shapes happen
to already be multiples of 128, so N is likely fine with one bucket, M is
the harder axis). The M/N tail-fill mechanism this row already has absorbs
the remainder either way; whether one bucket size is enough or a second,
smaller bucket is needed is a measured decision, not assumed up front.

Confirmed capped at RDNA4 (gfx1200/gfx1201) by the same two existing gates
this row's kernels already use — #if defined(__gfx1200__) || defined(__gfx1201__) at compile time, Gfx12QuantWmmaHostOk at runtime —
so this redesign has no reach into CDNA, RDNA3, or RDNA3.5 as long as it
stays inside those same guards, which is also the only place rocwmma.hpp
is includable at all.

Start with the existing KQuantGemmKWmmaQ4K/Q6K implementations and their RDNA4 compile-time and runtime gates. Read the referenced llama.cpp paths—mmq-config-rdna4.cuh, mmq-load-tiles.cuh, and mmq.cuh—then write the required scoped tile specification within the 64 KiB LDS budget. Done means a measured cooperative-loading redesign, including any needed M-tail bucket decision, while remaining inside the existing guards.