QuantizeQ8KK is thread-per-superblock: 95.5us/call against llama.cpp's 1.6us, and 41% of the ROCm decode gap

[joral]

Row: BACKEND-ROCM

QuantizeQ8KK decomposes one thread per 256-element superblock. At decode that
leaves 16 threads on a 32-CU GPU, and it costs 95.5 us per call against
llama.cpp's 1.6 us for the same work — a 35x per-call gap that accounts for
41% of the entire ROCm decode deficit on a dense model.

Measured

rocprofv3 1.1.0 (rocprofiler-sdk 7.2.3, native), --kernel-trace --stats,
Ornith-1.5-9B-Q4_K_M.gguf (dense qwen35), RX 9060 XT (gfx1200), ROCm 7.2.3.
Decode isolated by differencing a 4-token run against a 36-token run over 32
tokens, on both sides. vllm.cpp 4b1154bc5; llama.cpp b10451 (10bf611e5)
HIP. Free VRAM asserted above 13 GiB with no resident model process.

	ms/tok	calls/tok	us/call
ours, QuantizeQ8KK	12.321	129	95.5
llama.cpp, quantize_q8_1	0.352	217	1.6

Total kernel time per decode token is 52.735 ms against llama.cpp's 23.466 ms.
The quantizer delta alone is 11.969 ms, 40.9% of that 29.269 ms gap.
llama.cpp issues more quantize calls than we do and still pays 35x less each.

Cause

src/vt/rocm/rocm_grouped_gemm.hip:146, and the comment says it outright —
Q8_K (thread-per-256-superblock):

const int64_t t = blockIdx.x * blockDim.x + threadIdx.x;   // one thread = one superblock
const int64_t elem0 = i * a_rs + sb * kQK_K;
float mx = 0.0f, amax = 0.0f;
for (int j = 0; j < kQK_K; ++j) {
    const float ax = fabsf(DLoadAct(a, adt, elem0 + j));
    if (ax > amax) { amax = ax; mx = DLoadAct(a, adt, elem0 + j); }
}

launched at rocm_grouped_gemm.hip:479:

QuantizeQ8KK<<<static_cast<unsigned>((m * nsb + 127) / 128), 128, 0, s>>>(...)

At dense decode m is 1 and nsb is K/256. For a 4096-wide hidden that is
16 threads, one workgroup, on a 32-CU part, each thread walking 256
elements serially. The MoE path at :547 has the same kernel with m =
T x top_k, which is why #1294
measured it at 35% of GPU time there.

Two defects, not one. DLoadAct is called twice per element on the
amax scan — once for the comparison and once for the value kept — so the loop
issues 512 loads where 256 would do. That is independent of the decomposition
and is a strict improvement on any backend.

The donor

ggml/src/ggml-cuda/quantize.cu:53 decomposes per element:

const int64_t i0 = (int64_t)blockDim.x*blockIdx.x + threadIdx.x;
if (i0 >= ne0) return;

launched (quantize.cu:567) as:

const int64_t block_num_x = (ne0 + CUDA_QUANTIZE_BLOCK_SIZE - 1) / CUDA_QUANTIZE_BLOCK_SIZE;
const dim3 num_blocks(block_num_x, ne1, ne2*ne3);
const dim3 block_size(CUDA_QUANTIZE_BLOCK_SIZE, 1, 1);   // 256

with a warp reduction for the scale. Thousands of threads instead of 16.

Note llama.cpp is a secondary oracle here and never the mirror source: it
supplies the shape of a fix and the floor to beat, not the behaviour. vLLM
defines behaviour, and this kernel's own header already names its vLLM donor
(cuda_quant_dot.cu QuantizeQ8KKernel). Any change must stay byte-identical to
the reference path.

Scope

• MatmulBTQuantKernelRocm (:479), the dense path — measured above.
• MatmulBTQuantGroupedKernelRocm (:547), the MoE path — the same kernel,
  measured at 35% by ROCm decode is kernel-bound: the activation quantizer is 35% of GPU time and GdnPostConvK 19%, while H2D+D2H is under 1% #1294. One fix moves both.
• The arithmetic is unchanged: this is a decomposition and a redundant-load
  removal, so every existing golden must stay byte-identical. That is the
  correctness gate, and it is stronger than a tolerance.

Not claimed

No throughput number is predicted here. Closing the quantizer to llama.cpp
parity would remove 11.969 ms of 52.735 on this one model and prompt, which is
about 1.3x on the arm measured, and says nothing about other shapes, batch
sizes, or the MoE path's own ratio. A same-binary A/B behind an environment
flag, on an idle host, is what a speed claim would owe — and note that
VT_GGUF_KEEP_QUANT=0 is not available as the OFF arm on a 16 GiB card
(#1870), so the A/B has to be
against the unmodified kernel rather than against expansion.

The residual GEMM gap of 1.54x (KQuantGemmK + wvSplitKSml 32.574 ms against
MMVQ 21.229 ms) is not in scope here and stays with
#1863.

Owned by BACKEND-ROCM.

[localai-org-maint-bot]

Added the Row: line to the body, which #2472 notes its authenticated account could not do.

Scope note, so this does not get closed by the wrong change. #2472 ("cooperate across Q8_K activation blocks") is the fix for this defect, and it is measured: median decode 30.943 -> 40.316 tok/s, quantizer time 7.8811 -> 0.3426 ms/token over 129 calls per token.

But it is architecture-scoped to gfx1100. From its own description, the cooperative arm is selected on unset only when the queue-device resolver returns gfx1100, and "unset remains legacy for gfx1200, gfx1201, unknown architectures, and resolver failure".

This issue was measured on gfx1200 (RX 9060 XT, ROCm 7.2.3), which is precisely one of the architectures that keeps the legacy thread-per-superblock arm.

So when #2472 merges it should not carry a closing keyword for this issue, and I will not add one. What it closes is the gfx1100 half. The gfx1200 measurement that opened this issue — 95.5 us/call against llama.cpp's 1.6 us, 41% of the ROCm decode deficit — still stands unaddressed on that part.

Leaving this open, owned by BACKEND-ROCM, for the gfx1200/gfx1201 arm. The mechanism is already written and reviewed; what is owed is the evidence to widen the architecture predicate, not a new kernel.