feat(rocm): grouped + non-grouped keep-quant expert GEMM — the MoE-path terminus (issue #41)

[VikashLoomba]

Row

BACKEND-ROCM — the last kernel blocker for MoE-bearing models on discrete ROCm. Issue #41. Stacked on #509 (MoE combine/gate chain).

What changed

NEW src/vt/rocm/rocm_grouped_gemm.hip — ports the keep-quant expert GEMM family from cuda_quant_dot.cu 1:1:

• kMatmulBTQuant (op 74, non-grouped) + kMatmulBTQuantGrouped (op 75, grouped over expert_ids)
• Q8_0 / Q4_K / Q5_K / Q6_K weight formats (the ones the target GDN-MoE GGUFs carry); Q8_0 + Q8_K activation quantizers
• DotQ8_0/DotQ4K/DotQ5K/DotQ6K superblocks, __dp4a → portable Dp4a (bit-identical integer core), __shfl_down_sync reduction
• Registering the non-grouped op is what flips GgufQuantComputeAvailable() on ROCm — the grouped op alone leaves the loader dequantizing experts to bf16.

Cross-device case: all four formats vs the CPU keep-quant oracle, valid random blocks + real activations.

Evidence (4× gfx1100, ROCm 7.14, Release)

• grouped-quant case: 16/16 assertions across Q8_0/Q4_K/Q5_K/Q6_K, NMSE ≤ 5e-4 vs the CPU keep-quant reference
• ctest -R 'rocm|cross_device': 4/4; full ctest: only the 5 pre-existing host/lane failures (zero new)
• E2E: Qwen3.6-35B-A3B Q4_K_M (21GB GDN-MoE GGUF) runs end to end on one gfx1100 with keep-quant active — op 74 AND op 75 resolve vt-native, zero CPU-ref fallback, correct output. Requires --max-num-seqs 1 to fit one 24GB card (the GDN state pool otherwise pushes past — a residency note, not a kernel defect; confirmed via a backend-Alloc instrumentation run showing genuine cumulative ~23.8 GiB).
• preflight --staged + trailers green

Speed claims

• This PR makes NO speed claim.

Honest gaps

• The K-quant formats Q2_K/Q3_K and the IQ2/IQ3 family are NOT ported (throw loudly); the target models don't use them. They follow the same skeleton.
• The per-call hipMalloc/hipFree activation scratch is correctness-grade; a queue-owned grow-only pool is a perf lever (decode-step churn).
• The 35B fits only with --max-num-seqs 1 on 24GB; multi-GPU expert sharding or host streaming is the follow-on for bigger MoE / longer context.

[localai-bot]

Reviewed as part of a sweep over the open external PRs. The GEMM port is the strongest part of this — I diffed DotQ4K/DotQ5K/DotQ6K against the donor line by line and DotQ6K is byte-for-byte; all four cited anchors in cuda_quant_dot.cu resolve. Registration goes through the existing seam.

One blocker, and it is in how the op is switched on rather than in the kernel.

Registering kMatmulBTQuant flips keep-quant loader-wide for formats this kernel cannot execute.

GgufQuantComputeAvailable() (gguf_keep_quant.cpp:74) is a boolean: OpRegistered(kMatmulBTQuant, current_device). FromEnv() sets p.keep_quant from it, and KeepQuantDType admits anything vt::cpu::HasQuantDotKernel accepts — which per cpu_quant_traits.cpp:30-98 is Q4_0, Q8_0, Q2_K, Q3_K, Q4_K, Q5_K, Q6_K, IQ2_XXS, IQ3_XXS, IQ2_S and MXFP4. Twelve formats.

rocm_grouped_gemm.hip:519,592 implements four and throws for the rest. CUDA survives the same boolean because cuda_quant_dot.cu:1841-1846 falls back to the CPU kernel on an unsupported dtype. This PR's own docs/USAGE.md edit states the opposite for ROCm: "On a discrete card there is no CPU fallback tier."

So a Q4_0 or Q2_K or IQ2 GGUF that loads and generates correctly on ROCm today would now keep its blocks quantized at load and throw vt rocm: matmul_bt_quant: unsupported weight dtype on the first forward pass — after the whole model is resident on the card. That is a functional regression of a working path, discovered at the worst possible moment, and it is not mentioned in the PR body.

Same flip has collateral: expand_nk → keep_f16 turns ON, so an F16 file weight stays kF16 and MatmulBTKernelRocm (rocm_matmul_hipblaslt.hip:446-450) accepts only bf16/bf16 or f32/f32 → vt rocm: matmul_bt: unsupported.

The fix is a per-dtype availability query rather than a boolean, or a refusal at load time rather than forward time. Related: the refusal message at :592-594 lists Q5_K on both the ported and the not-yet sides, and does not name Q4_0 or MXFP4 — the two formats that actually drive this.

Second thing, worth knowing before you invest further. The per-call hipMalloc/hipFree + hipStreamSynchronize (:478,488-489,500,515-516,545,558-559,572,587-588) is not only the decode-step churn you name honestly in the body. Both are illegal under hipGraph stream capture, which hard-blocks the ROCm decode-graph work in #473/#332 — the row worth ~3x decode. The donor uses EnsureScratch(), a per-stream grow-only pool, and never synchronizes; rocm_device_bind.h:21 shows the codebase already cares about capture-safety. (Also: qact leaks if Check() throws between the malloc and the free.)

Smaller: Dp4a at :168-172 unrolls into four scalar MACs where HIP exposes __dp4a natively (lowers to v_dot4_i32_i8 on gfx9/gfx10+) — that is the innermost loop of every expert GEMM, so roughly 4x the instruction count on the hottest path for a one-line change. And there is no spec: grep -rn 'ROCM-GG-ON-MOE' .agents/ docs/ returns nothing, so the scope, gates and stop conditions for this work are not written down anywhere, and the unported arms are not recorded as owed.

Note also that the non-grouped op — the one that carries this PR's headline mechanism — has no test at all; only kMatmulBTQuantGrouped is exercised. And both new cases guard with if (!OpAvailable(op, dt)) continue;, so deleting the RegisterOp line makes them skip silently and stay green, which means nothing currently proves the registration.

On CI: the four red checks are all infrastructure, none yours. pr-size and agent-record both die on base must be an ancestor of head (unrebased fork) — and that is now fixed on our side in #619, so a rebase should clear them. Windows is the known-broken arm.

No AMD hardware here, so your 16/16, the ctest results and the 35B e2e run could not be reproduced and I am not disputing them; the finding above is read from the loader chain end to end.

[VikashLoomba]

The blocker and the rest accepted and reworked (commits 2f46038 + 29aed2f + 22ed7b9, rebased onto current main):

The blocker: keep-quant residency is now gated per-format by the running device's actual kernel set (DeviceKeepQuantSupported in the routing decision, not in the dtype-validity query). ROCm keeps {Q8_0, Q4_K, Q5_K, Q6_K} quantized; everything else keeps its pre-existing expand_bf16 residency — so a Q4_0/Q2_K/IQ2 model loads and generates exactly as before, and nothing reaches forward-time throw. The keep_f16 collateral is gated the same way (MatmulBTKernelRocm accepts bf16/f32 only). The refusal messages now name the actually-owed formats.

Capture safety: the per-call hipMalloc/hipFree/hipStreamSynchronize is replaced by a grow-only per-stream pool on hipMallocAsync, retire-never-free — mirroring the donor's EnsureScratch + graph_safe_scratch.h discipline. The qact throw-leak is gone with the frees. This also unblocks the hipGraph lane (#473).

Teeth: the non-grouped kMatmulBTQuant gains its own cross-device case (it had none), and both quant cases now REQUIRE the ROCm registration instead of skipping — which immediately caught a real regression during the rebase (the merge had silently dropped the registrations; restored, and the teeth are why it was caught). test_gguf_keep_quant gains the device-gated routing case (Q4_0/Q2_K expand, Q8_0/Q4_K/Q5_K/Q6_K keep, keep_f16 off on ROCm) plus a device-aware expectation in the TOTAL routing-table case.

On __dp4a: I verified rather than accepted — ROCm 7.14's gfx1100 toolchain has no __dp4a intrinsic, and __builtin_amdgcn_sdot4 errors with "needs target feature dot1-insts" (RDNA3 dropped the int8 dot4 hardware). The portable helper stands; the probe evidence is recorded in the spec.

The spec debt: .agents/specs/rocm-gg-keep-quant.md (the rework) + the original spike spec committed belatedly with the process miss named in its commit message.

Gates (gfx1100, flock): test_gguf_keep_quant 38/38, test_backend_cross_device 22/22, 35B Q4_K_M e2e correct on one card.

[bakon11]

Coordinating to avoid overlap: I'm scoping a (deferred) WMMA-fused FP8 expert GEMM — our rocm_fp8_channel_gemv.hip expert path currently runs scalar on gfx1201. Does #523 cover the FP8 E4M3 expert path, or only the GGUF keep-quant families (Q8_0/Q4_K/Q5_K/Q6_K)? Want to settle FP8-vs-keep-quant ownership so we don't duplicate the expert-GEMM work.

Also FYI for landing order: my #837 (GetBlas dual-slot TLS) edits rocm_matmul_hipblaslt.hip, which your #506 also touches (skinny-GEMM dispatch) — no semantic overlap with GetBlas itself, just flagging the shared file.

[bakon11]

@VikashLoomba — also read #523 against our ROCm lane, and it's a clean split: it touches none of our files, and semantically it's a different quant family — GGUF Q8_0/Q4_K/Q5_K/Q6_K, activations to Q8_0/Q8_K, Dp4a integer-dot per output wave.

We're separately looking at an FP8 E4M3 + per-channel BF16 fused-dequant WMMA expert GEMM (Fireworks-style) for the Gemma-4 prefill path — a different quant ABI, loader route, and math core — so there's no overlap and #523 doesn't subsume it. We'll keep those as explicitly separate lanes rather than fold FP8 work into this GGUF PR.

Where we can help: gfx1201 (dual R9700, wave32) correctness/portability validation across all four formats, grouped + non-grouped, M/P = 1 and > 1. Wave32 matches and your Dp4a path is software-portable, so this is validation/perf characterization rather than joint kernel design. Ping us if you'd like RDNA4 numbers before landing.