perf(ROCM-QUANT-GEMM-BW): 622 GB/s effective weight streaming on gfx1100 (65% of peak) by ghazni101 · Pull Request #2 · ghazni101/vllm.cpp

ghazni101 · 2026-08-22T00:57:29Z

What changed

One optimization to the W1 keep-quant GEMM (src/vt/rocm/rocm_quant_dot.hip), zero numeric change:

The tail path of QuantDotGemmKernel splits each super-block across all 32 lanes (DotSuperblockWarp, 8 elements per lane) instead of one super-block per lane, which idled 22 of 32 lanes at decode shapes (nsb = k/256 = 10..36).
The dmin/bsums correction terms apply once per block on lane 0; integer partials warp-reduce via shuffle. Two's-complement int32 addition is associative, so the lane split is bit-exact.
DotQ2K / DotQ4K / DotQ5K lose their int8_t aux8[256] stacks to fused inline decode: QuantDotGemmKernel<Q4_K> drops from VGPR 192 to VGPR 64.

Measured result

Qwen3.5-4B Q4_K_M greedy decode on RX 7900 XTX (gfx1100, ROCm 7.14.0 container), rocprofv3 GPU-busy per steady token:

	baseline	this PR
GPU-busy per token	8.16 ms	4.41 ms
Effective weight-streaming	~336 GB/s	~622 GB/s (65% of ~960 peak)
Warm wall clock	9.8 tok/s	17.8 tok/s

Goal was >= 576 GB/s (60%). Met and confirmed on a second capture.

Gates

test_rocm_quant_dot: 132,094 assertions green, test file UNCHANGED (integer core bit-exact vs CPU reference, NMSE <= 1e-6).
Focused sweep ctest -R 'rocm|cross_device|quant': 20/21, the one failure the pre-existing MoeSiluMul bf16 case (Characterize Qwen3.5-0.8B CPU against ROCm numerics on gfx1100 mudler/vllm.cpp#1588).
End-to-end smoke deterministic across runs; output coherent.

Closes mudler#1587

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:ox-alpha [omp]

…tion call site (mudler#1273) Closes the "Owed" item from mudler#1065: CheckKvCacheShape's production call site in GPUModelRunner::initialize_kv_cache now has a test. Building the runner with a non-multiple-of-16 block size (kBlockSize = 8) asserts the throw at construction, from the FLASH_ATTN backend's own get_kv_cache_shape — the executable statement of the contract the server's --block-size validation and the bench rounding exist to prevent at the entry points. Validated on gfx1151 (Strix Halo, ROCm 7.2.3) as part of the mudler#41 M3 battery: test_runner 20/20 cases, 544 assertions, on the current-main build with mudler#1056 + mudler#1065 merged. Issue: mudler#41 FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: DeepSeekV4:Flash [Freebuff]

The ROCm backend registers no quantized-weight GEMM provider, so every GGUF k-quant weight on an AMD card computes off device. The spec scopes W1 as kMatmulBTQuant and kMatmulBTQuantGrouped providers mirroring the CUDA sibling, which the GGUF loader reaches automatically once they are registered, and owes the upstream csrc/rocm W4A16 family behind a loader consumer that does not exist yet. The issue index gains the three 2026-08-21 campaign issues: mudler#1586 toolchain adoption and optimization, mudler#1587 this row, mudler#1588 the Qwen3.5-0.8B numerics characterization. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]

… kROCM The GGUF loader routes a block-typed weight to MatmulBTQuant whenever the running device has the provider, so registering these two ops lights up keep-quant compute on every ROCm board with no model-path change: the dense and grouped MoE towers stage once through ResidentWeight and dispatch to the new device GEMM. Coverage mirrors the CUDA sibling exactly — the ten Q8_K-family encodings plus a native Q8_0 arm. The integer dots are the portable scalar forms of the CPU reference bodies in the CPU accumulation order, because gfx1100 exposes no signed byte dot (v_dot4_i32_iu8 is unsigned-only; sdot4 needs a feature this target does not offer), and the gate is bit-exactness against the CPU tier at NMSE 1e-6 with the f64 dequant band at 5e-4. Unsupported dtypes throw naming the dtype instead of silently falling back to a host kernel that cannot follow device pointers; VT_GGUF_KEEP_QUANT=0 restores load-time expansion. Gates on gfx1100 / ROCm 7.14.0: test_rocm_quant_dot 132,094 assertions green across all ten encodings (decode through prefill shapes, broadcast and per-row grouped arms over a poisoned output buffer), focused ctest 'rocm|cross_device|quant' 20/21 with only the pre-existing MoeSiluMul bf16 exactness failure (mudler#1588) remaining, and an end-to-end Qwen3.5-0.8B Q4_K_M decode that is deterministic on device. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]

QuantDotGemm is 48.3% of GPU busy in the 97,721-dispatch rocprofv3 capture of Qwen3.5-4B Q4_K_M decode at 6236e9e, and it streams weights one byte per lane, so the model decodes at ~163 GB/s effective weight streaming, about 17% of the RX 7900 XTX peak. The spec scopes an attempt ladder that vectorizes the weight loads to 16 bytes per lane and reshapes the wave geometry without touching the accumulation order, because test_rocm_quant_dot pins the integer core bit-exact against the CPU reference and stays unchanged. The issue index gains the campaign row under mudler#1586. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]

The keep-quant GEMM assigned one super-block per lane (sb = lane), so a decode-shape row with nsb = k/256 in {10..36} idled 22 of 32 lanes on every block and serialized each output on one lane. The rocprofv3 capture priced the result at 392.8 us for a [1,2560]x[9216] Q4_K GEMM (13.3 MB of weights, ~34 GB/s) and 1438.5 us for the [1,9216]x[2560] Q6_K arm (~13 GB/s). The tail path now splits every super-block across all 32 lanes (8 elements per lane, DotSuperblockWarp) and warp-reduces the integer partials; whole super-blocks beyond the first multiple of 32 keep the scalar path. Two's-complement int32 addition is associative, so the lane split of the reference element order is bit-exact, and the float scales still apply only after the integer sums. The dmin/bsums correction terms apply once per block on lane 0. The IQ families keep the scalar body inside the warp variant (grid lookups dominate there, and their nsb is large), and DotQ2K/DotQ4K/DotQ5K lose their int8_t aux8[256] stacks to fused decode, dropping QuantDotGemmKernel<Q4_K> from VGPR 192 to 64. test_rocm_quant_dot: 132094 assertions green, unchanged. Steady decode of Qwen3.5-4B Q4_K_M on gfx1100 drops from 8.16 ms GPU-busy per token (~336 GB/s effective weight streaming) to 4.41 ms (~622 GB/s, 65% of peak); warm wall clock goes from 9.8 to 17.8 tok/s. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]

…ttn arm The S1 rocprofv3 re-take on the acceptance workload (512-token prompt, 256 generated tokens, gfx1100) attributed 24.1 of the 54.8 ms/token decode budget to QuantDotGemmKernel<Q6_K> at n=2560/K=9216: 16 calls per token at ~1507us each, streaming the 19.35 MB of weights at 12.8 GB/s against the ~296 GB/s its Q4_K sibling achieves. The hot path for nsb >= 32 shapes is the SCALAR DotQ6K — with full_sbs = 32, each lane decodes one whole super-block through the old switch-on-quadrant body with size_t address arithmetic per element. A standalone HIP microbench proved the dot bodies are equivalent in isolation; the cost is the serial per-element decode on the lane-split path. DotQ6K is restructured to unroll the four quadrants per 128-chunk with pure-int addressing and no switch: same element order, same scale-per-group application point, so the integer core stays bit-exact (test_rocm_quant_dot 132,094 assertions unchanged). Measured on the gate workload: ffn_down falls to ~509us/call and end-to-end warm tg rises from a 13.1 to an 18.5 tok/s median (+41%). A microbench of the new body shows it at parity or better versus both old forms. Also lands the first S2 attention lever behind VT_ATTN_DECODE_GQA4=1 (default OFF): PagedAttnDecodeGqaBf16<4,8,8> for the qg==4/d==256 geometry, which currently falls to the generic PagedAttnOnline (818us/call, 6.1 ms/token across the 8 full-attention layers). NOT YET EFFECTIVE for this model: the GGUF dense path passes an f32 query while DecodeGqa is bf16-only, so the arm is recorded as owed the f32-query extension plus the near-tie adjudication before any default flip. The focused cross_device gate keeps its documented pre-existing MoeSiluMul bf16 failure (mudler#1586/mudler#1588); all other cases green. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]

The S2 budget after the branch-free Q6K fix still showed ffn_down at ~510us per call: 2560 warps each serially walking 36 super-blocks is latency-bound, not bandwidth-bound — lm_head proves the memory system delivers 598 GB/s on the same kernel family. This adds a split-K variant for decode shapes (m == 1, nsb >= 8, n >= 512, non-f32 out): K is divided across 4 warps per output, each warp writes its scaled float partial to a dedicated buffer, and a small reduce kernel folds the partials and converts to bf16. The integer core is untouched; only the float scale-sum reassociates, inside the NMSE band the file's gate already accepts. Two defects caught during bring-up, both by end-to-end A/B rather than the op-level gate: the first A/B showed no delta because the arm gated on f32 out while every model-path GEMM emits bf16; the second produced garbage tokens because the partials buffer was taken from EnsureScratch — the same per-stream buffer that backs the quantized activation, so the split-K kernel clobbered the very activations it was reading. Partials now live in their own grow-only allocation. Measured on the acceptance workload (512-token prompt, 256 generated, greedy, idle box, gpu-coord lock held): 18.5 -> 22.7 tok/s median (+23%). Cumulative campaign position: 13.1 -> 22.7 tok/s (+73%) today. Output text matches the plain path token-for-token for roughly the first 40 tokens and then drifts — the expected float-reassociation cascade in greedy decode; the op gate stays green (132,094 assertions) and the near-tie adjudication before any default-flip decision is recorded as owed. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]

…A4=1 The S2 re-profile at the split-K head showed PagedAttnOnline as the top remaining lever: 6.05 ms/token across the model's 8 full-attention layers, 803us per call. The generic kernel handles one (token, head) pair per CTA and walks the KV sequence with a block-wide __syncthreads() online softmax per context token — a sync storm at decode shapes, with 16 CTAs on a 96-CU part. The DecodeGqa geometry fixes both: QG q-heads fused into one CTA read each K/V row once, and the walk is warp-strided with only end-of-loop warp reduces. The existing DecodeGqaBf16 kernel is bf16-in/bf16-out, while this model's GGUF dense path runs "Phase 1" attention numerics — an f32 query and f32 output over the bf16 KV cache — so it never qualified. This adds the mechanical sibling: PagedAttnDecodeGqaF32Q<QG,EPL,NWARPS> (f32 query, bf16 K/V, f32 out) via templated LoadRowEplF32/StoreRowEplF32 helpers, plus a dispatch arm gated on VT_ATTN_DECODE_GQA4=1 for the exact dtype/geometry combination (f32 query/out, bf16 cache, d=256, hq=16/kv=4). Measured on the acceptance workload: attention calls drop from ~803us to sub-100us; end-to-end rises from a 22.7 to a 25.7 tok/s median (+13%). Cumulative campaign position: 13.1 -> 25.7 tok/s (+96%) today. Output is coherent text; the reduction-order change means greedy anchors can move at exact ties versus the generic path, same policy class as the d128 arm — the near-tie adjudication before any default flip stays owed. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]

…GEMV Two levers from the decode kernel-sequence dump at the split-K head. KEEP-QUANT UNDER THE V-ROW REORDER (qwen3_5_gguf_weights.cpp). The GDN in_proj_qkv/attn_gate weights are Q5_K, but the V-head row reorder made them kTransformedWeight, forcing bf16 expansion at load and hipBLASLt decode GEMMs. A row permutation cannot cut a ggml K-block — each row is whole blocks — so OwnGgufQuantBlocksRowPermuted now applies the reorder to the quantized rows directly. 48 tensors stay packed (~0.9 GB less host+device memory); the resident weight encodes exactly what dequantize→reorder→bf16 encoded. End-to-end neutral on throughput: these GEMMs were already bandwidth-fine on hipBLASLt; the win is memory and quant-path coverage. TINY-N DECODE GEMV (rocm_matmul_hipblaslt.hip). The per-layer kernel sequence showed two hipBLASLt calls plus epilogues between in_proj and the GDN conv — the ssm_alpha/ssm_beta projections (N=32, K=2560), ~78us of fixed library launch cost for a 164 KB weight read, ~4.1 ms/token across 24 layers when both their bf16-out and f32-out forms are counted. Both forms now route to a row-per-block GEMV (bf16-out via the existing Bf16GemvBT, f32-out through a new GemvBTF32OutKernel templated on the activation type), forced for N<=64 independent of VT_ROCM_GEMV. Measured on the acceptance workload with the attention arm enabled: 25.0 -> 27.6 tok/s median (+10%). Cumulative campaign position: 13.1 -> 27.6 tok/s today (+110%). Also adds an env-gated VT_MM_TRACE shape log to MatmulBTKernelRocm that made the caller attribution possible; it stays as a diagnostic. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]

…ive, garbage-fast instance #2 The VT_GEMV_MMVQ=1 arm is bit-exact vs the CPU oracle at the op seam under every constructed condition (331 assertions: bf16/f16 activations, bf16 outputs, tied-amax quant byte-equality, ON-vs-OFF sweep identity at model shapes) yet in-engine it degrades the acceptance stream to fluent number-loops at 5-11 tok/s vs the OFF arm's coherent 34.81 tok/s median. Interleaved same-window medians recorded with uptimes; graph replay ruled out; one ON-arm rep wrapped in rocprofv3 for operator attribution (raw db at /work/t4a-prof-on). Verdict: CLOSED NEGATIVE, flag stays default-OFF, defect correctable, next hypotheses named (per-layer dispatch trace, geometry-vs- fold split arms, near-tie adjudication). Also recorded: mutation log incl. one NOT-caught mutation with the gate gap named; the gpu-ctl exclusion incident timestamps verbatim; and the campaign-level lesson that garbage-fast now has two independent instances (T3a LDS, T4a dtype coverage) -- op parity never substitutes for engine token-coherence. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]

…69.8 tok/s median T5a re-attribution at a5bfddb (GPU busy 15.21 ms/tok) showed the #2 item was not kernel waste but a routing hole: f32-query decode fell to the generic PagedAttnOnline (276us/call, serial per-key __syncthreads walk) because the T3a f32-Q DecodeGqa arm required d == 256. T5b adds the EPL=4 instantiation behind VT_ATTN_DECODE_GQA4=1: acceptance A/B interleaved x5 gives median 69.780 vs 61.468 tok/s (+13.5%, ON wins all five pairs), and the ON arm's gate-prompt output is byte-identical to the original pre-campaign baseline output — zero tie flips recorded; the teacher-forced logprob-band ceremony stays owed before any default flip of the flag. wvSplitKSml recorded near-roofline (408 GB/s of ~598 peak) with no ceiling declared. The test_gguf_keep_quant and one test_backend_cross_device failure reproduce on the pristine head and are owned separately from this lever. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]

…ive, garbage-fast instance #2 The VT_GEMV_MMVQ=1 arm is bit-exact vs the CPU oracle at the op seam under every constructed condition (331 assertions: bf16/f16 activations, bf16 outputs, tied-amax quant byte-equality, ON-vs-OFF sweep identity at model shapes) yet in-engine it degrades the acceptance stream to fluent number-loops at 5-11 tok/s vs the OFF arm's coherent 34.81 tok/s median. Interleaved same-window medians recorded with uptimes; graph replay ruled out; one ON-arm rep wrapped in rocprofv3 for operator attribution (raw db at /work/t4a-prof-on). Verdict: CLOSED NEGATIVE, flag stays default-OFF, defect correctable, next hypotheses named (per-layer dispatch trace, geometry-vs- fold split arms, near-tie adjudication). Also recorded: mutation log incl. one NOT-caught mutation with the gate gap named; the gpu-ctl exclusion incident timestamps verbatim; and the campaign-level lesson that garbage-fast now has two independent instances (T3a LDS, T4a dtype coverage) -- op parity never substitutes for engine token-coherence. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]

…69.8 tok/s median T5a re-attribution at a5bfddb (GPU busy 15.21 ms/tok) showed the #2 item was not kernel waste but a routing hole: f32-query decode fell to the generic PagedAttnOnline (276us/call, serial per-key __syncthreads walk) because the T3a f32-Q DecodeGqa arm required d == 256. T5b adds the EPL=4 instantiation behind VT_ATTN_DECODE_GQA4=1: acceptance A/B interleaved x5 gives median 69.780 vs 61.468 tok/s (+13.5%, ON wins all five pairs), and the ON arm's gate-prompt output is byte-identical to the original pre-campaign baseline output — zero tie flips recorded; the teacher-forced logprob-band ceremony stays owed before any default flip of the flag. wvSplitKSml recorded near-roofline (408 GB/s of ~598 peak) with no ceiling declared. The test_gguf_keep_quant and one test_backend_cross_device failure reproduce on the pristine head and are owned separately from this lever. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]

tbrasser and others added 9 commits August 21, 2026 12:54

Navigation Menu

perf(ROCM-QUANT-GEMM-BW): 622 GB/s effective weight streaming on gfx1100 (65% of peak) - #2

perf(ROCM-QUANT-GEMM-BW): 622 GB/s effective weight streaming on gfx1100 (65% of peak)#2
ghazni101 wants to merge 9 commits into
mainfrom
row/ROCM-QUANT-GEMM-BW

ghazni101 commented Aug 22, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Conversation

ghazni101 commented Aug 22, 2026

What changed

Measured result

Gates

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants