perf(GFX1100-TG200): raise Qwen3.5-4B Q4_K_M gfx1100 decode to 76.6 tok/s - #6
Open
ghazni101 wants to merge 98 commits into
Open
perf(GFX1100-TG200): raise Qwen3.5-4B Q4_K_M gfx1100 decode to 76.6 tok/s#6ghazni101 wants to merge 98 commits into
ghazni101 wants to merge 98 commits into
Conversation
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]
The developer set a goal on 2026-08-22: serve Qwen3.5-4B Q4_K_M on the RX 7900 XTX at 150 tok/s text generation, pure autoregressive greedy decode with no speculative path, on an fp8-e4m3 KV cache. Issue mudler#1651 records the goal and this spec turns it into a gated, staged campaign: S1 re-prices the attribution on the exact acceptance workload before any lever is chosen, S2-S5 take the kernel and dispatch levers in S1's order, S6 extends the landed KV-FP8 CPU brick to a ROCm store and paged-attention read behind the cache_dtype=fp8 surface that rocm_paged_attn.hip refuses by name today, and S7 runs the median-of-five idle-host acceptance gate. The index gains the campaign row; mudler#1586 stays scoped to the quant-GEMM bandwidth ladder. The wall-vs-GPU-busy gap at current head (17.8 tok/s wall against 4.41 ms GPU busy per token) is recorded as an inference to verify, not a finding: one rocprofv3 capture at a different prompt length priced it, and S1 owns the re-take. 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]
ghazni101
force-pushed
the
row/GFX1100-TG200
branch
from
August 22, 2026 23:03
6e89e5b to
bb8545d
Compare
The developer set a goal on 2026-08-22: serve Qwen3.5-4B Q4_K_M on the RX 7900 XTX at 200 tok/s single-stream text generation, pure autoregressive greedy decode, no speculative path. Issue #5 (ghazni101/vllm.cpp) records the goal and this spec turns it into a gated, staged campaign. Feasibility is recorded as settled -- llama.cpp sustains ~200 tok/s on this exact checkpoint and GPU with a q8 KV cache, our own lm_head streams 598 GB/s on this board, and the ceiling arithmetic puts the target at ~47% of peak -- so no stage relitigates it. The base is pinned at upstream tip 019f66c rather than the TG150-era base, because main has since landed three levers inside exactly the budget TG150 measured remaining (GdnPostConvK single-thread value_dim copy, VT_ATTN_DECODE_D128 default-on for ROCm, wvSplitK decode-skinny GEMM routing); T1 exists to re-price the tip on the exact acceptance workload before any new lever is chosen. Stages T1-T6: attribution re-take, dispatch-collapse (HIP graph / FusedChain), GDN family decode levers, residual quant-GEMM arms continuing mudler#1586's ladder, hipBLASLt/wvSplitK arms, acceptance gate + landing. The index gains the campaign row; mudler#1586's scope is unchanged. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
…usy 8.64 + gap 2.08 ms/tok T1a wall clock on the exact acceptance workload at base 019f66c: runs 40.639/40.671/40.712/40.594 tok/s after one warmup, median 40.65 tok/s. The pre-merge position of 27.6 tok/s is stale: the three upstream levers that landed inside this budget (GdnPostConvK value_dim fix, VT_ATTN_DECODE_D128 default-on, wvSplitK skinny GEMMs) bought about +13 tok/s before any new campaign work. T1b rocprofv3 capture, steady-state window over 511 tokens: wall 10.72 ms/token = GPU busy 8.64 ms + host dispatch gap 2.08 ms (inter- dispatch idle). Top items: the SECOND GdnPostConvK instantiation still runs grid=1-block at 183 us per call, 10.8 calls/token = 1.98 ms/tok -- the f4ccabb repair covered the K-variant single-thread copy and this sibling kept the pathology; Q6K QuantDotGemmKernel 1.24 ms/tok; bf16 PagedAttnOnline 1.07 ms/tok at grid=1 with a block-wide sync per context token; hipBLASLt Cijk 0.70; GdnScan 0.51; one straggler SplitK launch (grid=124160) 0.44. Target arithmetic recorded in the evidence file: 200 tok/s = 5.00 ms/tok, so the campaign needs busy ~8.64 -> ~3.2 AND gap ~2.08 -> ~0.5. The lever order T1 outputs is: second GdnPostConvK geometry, HIP-graph dispatch collapse, PagedAttnOnline coverage, then Q6K quant-GEMM bandwidth. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
…=1 pathology was a profiler artifact T1b's budget table priced GdnPostConvK at grid=1, 183 us per launch, 1.98 ms/token -- the largest GPU item. That row was an artifact: the GDN decode step dispatches kGdnPostConv twice per layer with two different host-side argument tuples, and both collapsed into ONE kernel symbol in the rocprofv3 capture, averaging their durations into a misleading entry. A measurement build (NOT landed) registered a byte-for-byte duplicate TU of rocm_gdn_postconv.hip with renamed symbols as a second provider named "vt-native2" for (kGdnPostConv, kROCM), which wins the priority-0 tie by name order and routes every dispatch through it. RegisterOpProvider rejects duplicate provider names, so a second "vt-native" registration is not an option; the duplicate must carry its own name. With symbols separated the same work reads 10.8 launches/token at ~28 us = 0.30 ms/token: the "grid=1 pathology" was the collision's average, not a real defect. The measurement build stays out of the tree -- landing it would touch CMakeLists.txt (a user_usage/landing_page file) and owe a docs/USAGE.md edit for zero user-visible change. Wall clock was throughput-neutral on three 5-rep sets under co-tenancy: 33.3 / 33.07 / 33.22 tok/s median (host load 12-20; two earlier attempts died with hipMalloc OOM when a co-tenant model was resident) against the T1a idle-host baseline of 40.65. The re-taken capture's top item is now QuantizeQ8KK: 129 launches/token at ~59 us, grid=128 each -- about 16K super-blocks quantized per launch where the decode shapes (m=1, K=2560 -> nsb=10) need grid=1. That ~3.4 ms/token of activation quantization is the next lever, not postconv. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
…ured budget Two-source study: vLLM via web (subagent), SGLang via a local shallow clone read directly (sglang-src at the shared checkout root; kernels under python/sglang/kernels/aot/csrc/quantization/gguf/). Findings, ranked by expected ms/token recovered on our T2a capture: 1. SGLang GGUF MMVQ (fused_mul_mat_gguf -> ggml_mul_mat_vec_a8, mmvq.cuh): for batch <= mmvq_safe the activation is quantized ONCE to q8_1 by a single small kernel (quantize_row_q8_1_cuda, one warp per 512-padded row, wave reduce), then mul_mat_vec_q runs ONE WARP PER OUTPUT ROW with q4_K/q5_K/q6_K blocks dequantized in registers (vec_dot_*_q8_1). Our engine instead runs 129 QuantizeQ8KK launches/token at grid=128 (~16K super-blocks each, ~59 us) = ~3.4 ms/token of activation quantization where the decode shapes (m=1, K=2560 -> nsb=10) need grid=1. Porting the MMVQ geometry to rocm_grouped_gemm.hip is the top lever; it also removes the Q8_K scratch round-trip. The q8_1 vs Q8_K convention changes reduction order, so the arm lands behind the ratified near-tie adjudication, not a bit-exactness claim. 2. vLLM W4A16 (gptq_marlin/triton): activations stay bf16; weight dequant happens inside the GEMM registers. Same destination as (1). 3. Graph capture: both engines capture the whole uniform-decode step as ONE replay (vLLM FULL_AND_PIECEWISE; SGLang "full" backend default). Our seam exists (BACKEND-ROCM W1 hipGraph + BreakableGraph; platforms/rocm.cpp keeps support_static_graph_mode false pending the model-path W2). Flipping it for this model attacks the 2.08 ms/token dispatch gap. 4. Merged gate_up: vLLM MergedColumnParallelLinear packs gate+up into one GEMM. Our bf16 arm already has gate_up_proj; the GGUF keep-quant arm loads separate gate/up owners and runs two quant GEMMs. Merging halves those sites (rides layers::MlpGateUpMethodBase). 5. RDNA3: no first-party gfx1100 tuning upstream (AMD CI targets CDNA); autotune locally. LDS 64KB/workgroup cap; no MFMA (WMMA only); hipBLASLt Cijk tiles are large-batch tuned -- at m=1 a custom N-major skinny GEMM usually beats them. Checkpoint facts verified by direct GGUF header parse: Qwen3.5-4B H=2560, 32 layers (24 GDN + 8 full-attn), ffn 9216; every GEMM weight is Q4_K/Q5_K/Q6_K keep-quant (attn_qkv + ssm_out Q5_K, rest Q4_K/Q6_K mix), so all 129 quant launches/token come from the keep-quant GEMM path. Record-only: no production code changes in this commit. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
…tor recycling breaks it Follow-up to the T2a instrument repair. The 129 QuantizeQ8KK launches per token looked like pure launch-overhead waste, so a per-stream cache keyed on the activation pointer was tried: first GEMM over an activation quantizes, later same-key GEMMs reuse the scratch. Two cuts, both REJECTED for correctness: - Pointer-only key: throughput rose to ~45 tok/s median but output degenerated into repeated garbage. Root cause: the DevicePool recycles activation blocks ACROSS steps -- the same pointer carried different content and stale quantized data was served. - Epoch keying (vt::BumpQuantEpoch/CurrentQuantEpoch bumped once per model forward): still degenerate. Within ONE step the pool hands the SAME address to DIFFERENT activations (a DBuf is freed and its block re-allocated mid-forward), so pointer identity does not imply content identity even inside a single step. Both cuts reverted before landing; revert verified by coherent acceptance- workload output. Evidence file records both cuts with numbers. What this fixes in the campaign's aim, not just documents: the sound levers for that budget are structural -- merged gate+up keep-quant GEMM (halves the sites), MMVQ-style dequant-in-register decode GEMV following SGLang's mmvq.cuh (removes the separate quant kernel entirely), and ROCm decode-graph capture (removes the ~59us/launch queue overhead that makes each tiny kernel expensive). The probe instrumentation also stays out; it served its purpose. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
…MLP instead of two The T2a probe census showed the dense SwiGLU MLP issuing TWO kMatmulBTQuant calls per layer (gate, then up) over the SAME activation, each paying its own QuantizeQ8KK launch. Following vLLM's MergedColumnParallelLinear topology, this loads ffn_gate + ffn_up rows-stacked into ONE [2I, H] nk=true block owner and runs ONE quant GEMM into [T, 2I], feeding the unchanged SiluAndMul + down. Byte-exactness: each output row's integer dot runs over its own whole K- blocks; stacking only concatenates output rows, so every row's result is identical to the split call (verified: coherent acceptance-workload output, token stream matching the pre-change build's greedy run). Implementation: - OwnGgufKeptStacked (gguf_weights.cpp): copies both keep-quant halves into one contiguous owned buffer; two disjoint file spans cannot be borrowed as one mapping span, so this arm always copies. - LoadMergedKeptGateUp: returns empty when encodings or K differ -- caller falls back to the split pair. - DenseMlpBlock: the existing gate_up_proj branch now serves both the bf16 and the stacked keep-quant owners; dispatch keys on emptiness as before. A/B on the acceptance workload: both arms ~35.8 tok/s median in a co-tenancy window (the same-window split baseline read 32.2-35.9). No regression; the win is one fewer QuantizeQ8KK + one fewer GEMM launch per layer per token, which compounds with the decode-graph stage where per-launch cost dominates. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
…de-graph capture engaged The last false predicate in the dense decode-graph gate chain falls: support_static_graph_mode() returns true on ROCm. With W1's hipGraph capture/replay seam already implemented (rocm_backend.hip, with the mutate-src-then-replay assertion) and the keep-quant scratch pool already capture-safe (hipMallocAsync, stream-ordered, never freed), the flip engages Qwen3_5DenseDecodeGraph for uniform decode steps. Verified live: "[DenseDecodeGraph] captured ... padded size S=1" then "14 total replays across 1 captured size(s)" on a 16-token run; output coherent. A/B (acceptance workload, 256 tok x 5): 36.4 tok/s median in a co-tenancy window vs 35.8 same-window split-arm baseline -- neutral-to-slightly-positive under contention, as expected since a shared GPU queue partially hides the dispatch gap. The full ~2.08 ms/tok gap removal shows only in an idle-host window (projected ~46+ from the 40.65 baseline); the definitive idle-host capture is queued as campaign follow-up. Evidence: docs/bench-evidence/gfx1100-tg200-t2b-20260823.md FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
…indsight 500s) Records-only: appends the T2b session's hindsight note to docs/bench-evidence/gfx1100-tg200-t2b-20260823.md so the next session resumes from the recorded position instead of re-deriving it. No source, test, or measurement change. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
… -- and fixes its LDS underallocation PagedAttnOnline<f32,bf16,f32> is the top GPU item at this head (8 calls per token, grid=1-block, one __syncthreads() online-softmax step PER context token -- a serial walk measuring ~593us/call in the T2c capture). The model's full-attn layers run f32 query x bf16 cache x f32 out, which excludes every bf16 decode kernel, and FA2 stays CUDA-only on ROCm. This ports PagedAttnDecodeGqaF32Q<4,8,8> from c112d88 on the TG150 ladder: QG=4 fused q-heads per KV group read each K/V row once, warp-strided sequence walk, register online softmax, no per-key block sync. Dispatch arm gated behind VT_ATTN_DECODE_GQA4=1 for the exact geometry (f32 q/out, bf16 k/v, d=256, hq=16/kv=4) so the default path is unchanged. The port landed with the TG150 dispatch's shared-memory formula -- nwarps*2*d floats -- but this kernel's layout needs nwarps*QG*d; with QG=4 the arm ran on HALF the required LDS. The op-level suite stayed green (its GPU-parity cases are HasCuda()-guarded and skip on ROCm-only boxes) while the engine degenerated into number-loops after ~20 tokens at an apparent 39.5 tok/s. Caught by the token-coherence check on the acceptance workload; the fix keys the smem formula on QG. The garbage-fast lesson is recorded in the evidence file: a throughput median without an engine-level coherence check would have shipped it. After the fix: coherent output, full 256-token streams differ only in mid-stream near-tie moves (reduction order changes vs PagedAttnOnline, same policy class as VT_ATTN_DECODE_D128). Clean-window interleaved A/B under co-tenant load noise: ON ~36.8 vs OFF ~36.2 tok/s median; the definitive idle-host capture is owed once the co-tenant compile storm clears. Near-tie adjudication before any default-ON flip remains owed; the arm ships opt-in. Evidence: docs/bench-evidence/gfx1100-tg200-t3a-20260823.md FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
Records-only: appends the T3a session's hindsight note to the t3a evidence file so the next session resumes from the measured position. No source, test, or measurement change. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
The campaign spec names tests/vt/test_rocm_quant_dot.cpp as the quant-path
lever gate, but that file never existed: every GPU-parity keep-quant case
lives in test_cuda_quant_dot.cpp behind HasCuda(), so on this ROCm-only box
the whole op-level suite silently skipped (the exact T3a blind spot). This
creates the focused ROCm kMatmulBTQuant gate, guarded on ROCM availability,
never on CUDA, registered under VLLM_CPP_HIP next to test_rocm_backend.
The gate is written RED for the T4a lever on purpose. The new MMVQ-style
decode arm behind VT_GEMV_MMVQ=1 will reproduce the CPU oracle's float
association (positional sums[] chains plus the sequential dmin chain of
VecDot{Q4,Q5,Q6}_KQ8_K), so its parity case asserts BIT-exact equality with
host vt::MatmulBTQuant -- strictly tighter than the 1e-6 NMSE band the
baseline warp-reduction kernel can claim, because its __shfl_down tree
reassociates the float sum. Until that dispatch arm exists the flag is
inert, the baseline runs, and the bit-exact compare fails.
FOLLOWING_AGENTS_PROTOCOL
Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:ox-alpha [omp]
…arm -- bit-exact vs the CPU oracle
The T4a capture prices QuantizeQ8KK (2.593 ms/tok, 43.7 launches/tok at
59.3us avg on grids of <=1 block) plus KQuantGemmK (2.943 ms/tok combined)
as the top remaining GPU family. This adds an opt-in decode arm for
kMatmulBTQuant's K-quant branch behind VT_GEMV_MMVQ=1, m==1 only; the flag
is read per call (the cuda_quant_dot.cu convention) and the default path is
byte-unchanged.
Geometry-only change: the warp still owns one output element j, but its 32
lanes walk 32-elem CHUNK units (4 super-blocks x 8 chunks per pass) instead
of lane-strided whole super-blocks, so every lane stays busy down to nsb=1
(the baseline idles 22 of 32 lanes at this model's nsb=10), the scale/min
unpack is split per chunk instead of duplicated per lane, and Q6_K drops
its aux8[256] local-array rebuild for positional in-register dequant.
Numerics are bit-exact vs the host oracle BY CONSTRUCTION, not within a
band: the integer cores are exact under any association, and the float side
reproduces VecDot{Q4,Q5,Q6}_KQ8_K's association exactly (8 positional
per-super-block sums chains accumulated sequentially in super-block order,
plus the sequential dmin*sumi chain). The new focused gate
tests/vt/test_rocm_quant_dot.cpp asserts BIT equality with vt::MatmulBTQuant
on host tensors: 54/54 parity asserts across {Q4_K,Q5_K,Q6_K} x
nsb{1,3,10} x N{1,7,129} x 2 seeds, plus the default-arm 1e-6 NMSE probe.
RED-first recorded: pre-arm build failed 47/55 (the baseline shfl tree
cannot meet bit-exactness); post-fix 55/55, exit 0.
FOLLOWING_AGENTS_PROTOCOL
Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:ox-alpha [omp]
…ologue and widens the gate to engine dtypes Operator steer from the fresh HEAD capture: QuantizeQ8KK bills ~59us per launch on grids of <=1 block (2.593 ms/tok across 43.7 launches/tok) -- the prize is deleting the standalone launch, not shaving the GEMM alone. The per-super-block body is factored into QuantQ8KSBlock; the new KQuantGemvMmvqFusedK quantizes the row into block LDS via that SHARED body (same thread-per-super-block walk as the standalone grid), barriers, then runs the unchanged GEMV row body against the LDS copy. Byte-identity is by construction and asserted directly: MmvqQuantScratchForTesting exposes both semantics and the focused gate memcmps them on pseudo-random rows plus a tied-amax adversarial row (+max first, equal-magnitude negatives later) and an all-zero row. The fold engages only when nsb*sizeof(BlockQ8_K) fits 32KiB LDS; larger rows take the standalone-quant arm. The parity case now covers the ENGINE'S actual dtypes -- bf16 and f16 activations, bf16 outputs -- after the first acceptance window showed the op-green/f32-only gate while the engine degenerated: the second instance of the garbage-fast lesson (T3a was LDS underallocation, T4a is dtype coverage). Suite: 331 assertions green including ON-vs-OFF byte identity at model-like shapes (bf16 act/out, K up to 10240, N up to 10240). 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]
…patch and makes the GEMV bit-equal to the baseline kernel
Two defects closed red-first against the round-1 arm, which was op-green
but engine-garbage and uniformly slow.
Defect 1 (the engine killer): round-1 gated only the LDS fold on m==1, so
the non-fused GEMV branch captured EVERY multi-row call -- including the
engine's observed m=39 prefill chunks -- and the GEMV kernels write row 0
only, leaving rows 1..m-1 of prefill outputs UNWRITTEN while every op-level
test (all m==1) stayed green. The m==1 restriction now gates gemv_mmvq
itself. Red-first: a new MULTI-M gate case (m in {3,39} incl. the real
18432x2560 prefill shape, canary-filled outputs) fails 4/4 at the unfixed
code and passes after the one-line gate move.
The GEMV row body is also now BIT-EQUAL TO THE BASELINE KQuantGemmK output
at every (Fmt, nsb, j), not merely oracle-exact: the octet chunk-walk keeps
the exact integer phase (dp4a word cores; Q6_K nibble bias removed exactly
in the integer domain via a constant-word dp4a), then each super-block's
float term is reconstructed as the baseline's own d*isum (- dmin*sumi)
expression and replayed under the baseline's lane ownership and
__shfl_down(16,8,4,2,1) tree. ON==OFF byte identity at every shape is now
by construction, asserted by an extended sweep over the REAL model shape
set from the checkpoint GGUF manifest plus all exact engine dispatch tuples
-- including lm_head-class N=151936 and N=248320 (max offset 0.52 GB < 2^31,
falsifying the 32-bit-overflow hypothesis; the sweep's round-1 reds were
isolated near-tie rows from the oracle-vs-baseline ULP difference).
QuantQ8KSBlock also loads each activation once instead of twice (same
values). Perf defect: per-grid timing shows the fused fold's per-block
requantization scales with n/4 and loses 2.15-3.3x at lm_head-class grids,
so the fold is now gated to n <= 512 where it measurably wins; the
non-fused arm beats KQuantGemmK at every captured grid (0.31x-0.85x).
Focused suite: 719/719 across 6 cases incl. multi-m, engine-shape sweep,
and timing evidence. Mutations M-A (ownership predicate), M-B (qh 2-bit
mask), M-C (dmin term), M-D (m-gate removal) each caught and restored
byte-equal. Engine A/B: interleaved x5+5 same-window medians OFF 35.751 vs
ON 40.508 tok/s (+13.2%), all five ON outputs byte-identical to their OFF
pairs.
FOLLOWING_AGENTS_PROTOCOL
Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:ox-alpha [omp]
…+13.2% median with byte-identical engine outputs Amends the T4a evidence file (history retained) with the fourth-session repair cycle: the operator's round-1 ON-capture parse; the red-first extended ON-vs-OFF sweep over the real model shape set incl. lm_head-class N=151936/248320, which falsifies the 32-bit-overflow hypothesis (max offset 0.52 GB < 2^31; isolated near-tie-row ULP signature instead); the TRUE defect-1 root cause -- the m-gate hole that let m=39 prefill chunks into the row-0-only GEMV branch, proven red by a canary-filled MULTI-M gate case and by an instrumented dispatch trace; the baseline-bit-equal GEMV rewrite; the per-grid timing table isolating the fused fold's block-count-scaled requant cost and the measured n<=512 hybrid crossover; mutation log additions M-A..M-D; and the post-repair acceptance A/B -- interleaved x5+5 same-window medians OFF 35.751 vs ON 40.508 tok/s (+13.2%), all five ON outputs byte-identical to their OFF pairs. Verdict updated to LEVER ADOPTED; flag stays default-OFF. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
…ers and the F1/F2 routing-witness gate cases Closes the two reviewer findings that the round-1 gate could not see: (1) no case exercised VT_GEMV_MMVQ truly unset (EnvGuard(false) writes "0") while ON==OFF are bit-equal by construction, so outputs cannot witness routing; (2) nothing detected a kMmvqFoldMaxRows crossover drift, and the reviewer's 512->4096 mutation went fully green while flipping measured per-call ratios. The fix is the operator-contracted shape: process-global HOST-side counters bumped once per MatmulBTQuantKernelRocm dispatch on exactly the branch taken (baseline / non-fused GEMV / fused fold), exposed via MmvqRouteCountsForTesting + a reset hook; during stream capture kernel launches are recorded as graph nodes and not executed, so replay multiplicity can never skew the witness. Two new cases assert true-unset routes to baseline with the ON pair asserting the reverse, and pin the fold crossover at n=256 vs n=2304. Red-first proven: both cases fail to link before the seam exists. Replay of reviewer mutation M3 (getenv default inverted) is caught by F1 and M4 (fold cap 512->4096) by F2; restores are byte-equal (md5 5419b3f91dcdbb2321db823c60063f06). Focused suite green at 731/731 assertions; spec ctest gate unchanged vs a proven HEAD baseline; 0.8B engine coherence smoke byte-identical across arms. Evidence section 13 appended. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
…hind VT_GEMV_MMVQ_FOLD_MAX The fresh capture at b80a0bd prices the n>512 shapes' standalone QuantizeQ8KK launches at 2.177 ms/tok -- the top remaining GPU item -- while reviewer-mutation M4 evidence shows the fused fold leg running 1.30x baseline at grid=576 vs 0.53x unfolded, so folding wins whenever the deleted ~50us quant launch outweighs the folded-call penalty. Whether that net win extends past the shipped 512-row crossover is an empirical, end-to-end question, so this change makes the crossover runtime-tunable for measurement: VT_GEMV_MMVQ_FOLD_MAX takes integer rows, defaults to kMmvqFoldMaxRowsDefault (512, unchanged), and falls back to the default on empty or invalid values; it is read per call like VT_GEMV_MMVQ so in-process tests and graph capture pick it up at dispatch time. Default behavior is byte-unchanged and stays pinned by the F2 constants; the new F3 witness case asserts through the existing host-side route counters that the env actually moves routing both ways (n=2304 folds at 4096, n=256 stops folding at 128), that the boundary is inclusive, and that garbage values fall back to the default. Red-first proven: before the knob exists the widening and narrowing legs fail while every default-pinning leg passes (exit 1). Mutations M-B1 (knob inert) and M-B2 (<= narrowed to <) are both caught by F3; restores are byte-equal (md5 e0841e2083c1d85e75617c0b2f248df2). Full focused suite green at 752/752 assertions across 9 cases. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
…e closed negative, 512 already optimal Section 14 records the full B1 session: the VT_GEMV_MMVQ_FOLD_MAX knob and F3 routing witness (red-first, mutations M-B1/M-B2 caught, suite green at 752/752), the interleaved three-arm acceptance-workload A/B under gpu-ctl (OFF 35.594 / ON-default 40.348 / ON-FOLD_MAX=4096 36.142 median tok/s -- on4096 loses in all five paired triads), the middle-value refinement probe (1024 ties default at 40.149 vs 40.305), byte-identical coherence across all arms (md5 2b29ad66eea3ee3a99ff0694127ce88f, same as the adopted window), four honestly-recorded protocol incidents, and the negative-close verdict with the knob kept inert-documented: the fold's per-block requant already outweighs the deleted ~50us quant launch at the first n>512 shape class, so the 2.177 ms/tok standalone-quant item needs a cooperative-quant or graph-level fusion instead. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
…k to the two f32-out GDN BA projections Per-site attribution of the rank-2 GPU item (Cijk_Alik_Bljk_BSS_BH_ MT128x32x16_SE_1LDSB0, 21.6 calls/tok amortized at ~73.6us) from the fresh capture DB, committed before any kernel code per the lever-B2 contract. Parsing the rocprofv3 dispatch stream and correlating it with the GdnBlock op order and the GGUF tensor map shows the 12288-call population closes exactly as 48 calls per decode step x 255 steps plus one prefill pass: the Qwen3.5 GDN blocks' in_proj_b/in_proj_a (ssm_beta/ssm_alpha, N=32 x K=2560). Both sites emit f32 (ProjectGdnBA, qwen3_5.cpp:3663-3664), so every decode-skinny gate in MatmulBTKernelRocm -- which requires a bf16 output -- skips them and they land on hipblasGemmEx -> rocBLAS's large-M Tensile tile: ~73.7us to stream a 164 KiB weight, ~3.54 ms/tok combined under graph replay, 100% of the arm-coverage target. The sibling bf16 projections (in_proj_ qkv/z, out_proj) already ride wvSplitKSml at bandwidth-bound times. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
…ivity quantified The rmsnorm-row LDS-epilogue attempt measured -38% (LDS bank conflicts on u16 lane-strided reads vs an L1-resident gmem re-read) and was reverted byte-restored. Separately, byte-identical code measured 92.9 tok/s at load ~3.5 but 53-58 under sibling host contention at load ~5.5: launch- bound decode tracks host scheduling. Rule added — acceptance numbers are only comparable at recorded host load; untraced absolute claims need load < ~4, paired A/Bs stay valid under matched conditions. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
…he rate picture Assigning all 72 per-step split-K calls to step positions across 505 steady steps resolves three clean sites at 46.0/23.7/26.5us — attn_qkv at 456 GB/s, attn_gate 442, ssm_out 396 — replacing the blended 700 GB/s figure and restoring real headroom (~0.5 ms/tok to a 550-600 GB/s practical target). The concrete lever is a per-shape launch-config sweep of the donor's kYtile/wvPrGrp/split-factor math for exactly these three (N,K) shapes on gfx1100. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
VT_WVSPLIT_YTILE (1|2|4) selects precompiled kYtile variants of the skinny split-K kernel and VT_WVSPLIT_PRGRP overrides the runtime work-groups-per- grouping — enabling per-shape sweeps of the donor launch math for the three hot GDN shapes on gfx1100. Sweep under host load ~5 found YTILE=4 directionally positive (+1.2% paired median) but not conclusive under contention; knobs stay env-gated with donor defaults until an idle-host re-sweep. The f32-out lever-B2 arm keeps donor geometry regardless. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
The 10:41Z watcher-run chain included the T14 arms: OFF median 82.180 vs ON 82.897 tok/s across five interleaved pairs, all byte-identical as the bit-deterministic argmax requires. The lever's pending tok/s A/B is closed; session total stands at ~92.8 tok/s canonical with every adopted lever enabled. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
…ention taskset to the free cores under load ~5.5 reads identically to unpinned: the sibling services' pressure is host memory bandwidth, which launch- bound decode cannot be shielded from by core selection. Confirms idle- host as the only valid absolute-measurement condition. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
…anchors The sync-loop deferral is smaller than assumed: EngineCore::step already supports depth-2 pipelining via sample_tokens_async, and the only blocker was QueueSupportsAsyncInputCombine — which the real event primitives make TRUE on ROCm. The plan enables VT_ASYNC_RUNNER=1 in the acceptance config, verifies the depth-2 drain engages, and A/Bs paired x5 through the CLI. Fallback is the async-serving path with its ~40% overhead attributed first. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
…tical The idle-sweep watcher fired and ran the conclusive YT4 vs baseline paired verification: ON wins all five pairs (+1.8% median, 52.95->53.91). Output is bit-identical on a separate coherence check. The pre-committed decision rule (adopt iff ON wins >=4/5) is satisfied. Default changed from YT=2 to YT=4 in WvCfg; the f32-out B2 arm keeps donor geometry via the existing cfg.yt!=2 guard. Gate 16/16, 839 assertions. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
…torm closed, roofline analysis Acceptance gate at load 1.45: median 100.46 tok/s (runs 2-6, +8.1% from YT4 adoption under unconstrained bandwidth). T13 async-runner A/B is a wash on CLI (depth-1 drain, all byte-identical) — closed for CLI path. Copy-storm attribution: per-step small copies total 0.035 ms/tok, negligible — closed as a lever. Roofline: model 2.74 GB at 800 GB/s = 3.43 ms minimum weight read; wvSplitKSml at 57% bandwidth is the top headroom target. Next: v_dot2_f32_bf16 to reduce compute bottleneck. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
…y-bound kernel The dot2 instruction replaces 5 scalar f32 ops with 1 per bf16x2 pair, but wvSplitKSml is memory-bound at 57% bandwidth — compute is already hidden behind memory latency. Paired A/B on idle host (load 0.55): OFF 88.784 vs ON 88.897 (+0.13%, WASH). All 5 pairs differ (reduction order change) but both outputs are coherent analytic prose. The runtime branch also caused a 12% regression in the OFF path from code-size pressure; reverted and 100.47 tok/s confirmed restored. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
…vqRow Replace scalar Dp4a emulation with amd_mixed_dot(char4,char4,...) in the MMVQ Q4_K/Q5_K/Q6_K inner loops. The scalar Dp4a pattern (reinterpret_cast<int8_t*> + multiply + add) does NOT auto-lower to v_dot4_i32_iu8 when the input comes from a shift+mask expression because the compiler cannot prove the values fit in 8 bits. amd_mixed_dot forces the hardware dot instruction, reducing instruction count by ~10x and VGPRs by 30-43%. ISA verification (gfx1100, -O3): Q4_K: 0->8 v_dot4, 31->0 v_mul, 62->6 v_bfe, 68->42 VGPRs Q5_K: 0->8 v_dot4, 31->0 v_mul, 46->6 v_bfe, 73->51 VGPRs Q6_K: 0->16 v_dot4, 30->0 v_mul, 46->6 v_bfe, 86->49 VGPRs Correctness: gate 16/16, 839 assertions. Output BYTE-IDENTICAL to scalar baseline (integer arithmetic is exact regardless of evaluation order). Body coherence verified — analytic prose, no degeneration. Matched-load A/B (loadavg ~1.0, 5 runs each): baseline median: 90.133 tok/s v_dot4 median: 92.570 tok/s (+2.7%, all fixed > all baseline) Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:glm-5-2 [omp]
…ted — engine wash T20 rewrote KQuantGemvMmvqRow for full-warp cooperation (32 threads per super-block, zero intermediate barriers, single warp_reduce_sum). The kernel wins 2.4-3.1x on large grids (lm_head Q6_K 248320x2560: 2134->682us) but the engine A/B is a wash: ON 92.9 vs OFF 92.8 tok/s median (5 paired reps, full campaign config). The kernel win doesn't reach the engine because the dominant Q4_K path (2.46 ms/tok, 25% of wall) has small grids (ffn_gate/up ~576 rows) where the kernel is launch-overhead-bound, not reduction-barrier-bound. The Q6_K path (1.20 ms/tok) is mostly small-grid ffn_down (22 calls/tok, 1.03x). The large-grid lm_head save is 1 call/tok — invisible when averaged. Conclusion: kernel micro-optimization is exhausted for the dominant paths. The 4.2 ms/tok overhead above the 4.38 ms/tok weight-read floor is the bottleneck. Path to 200 tok/s requires launch-overhead reduction (HIP graph capture, kernel fusion, or persistent kernels). Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:GLM-5-2 [OMP]
…ctions The GDN layers attn_qkv (Q5_K, 24 tensors [2560,8192]) and attn_gate (Q4_K, 24 tensors [4096,2560]) were expanded to bf16 at load time because the V-head row reorder classified them as kTransformedWeight. The reorder is a ROW permutation — quantization blocks are along the K (column) dimension and are self-contained per row — so it is block-safe. T21 routes these tensors as kMatmulWeight to allow keep-quant, copies the blocks via OwnGgufQuantBlocks(mmap_src=nullptr), and applies ReorderVRows to the block bytes at load time. The forward pass already dispatches quantized nk=true weights through vt::MatmulBT, so no forward-pass change was needed. A/B: +3.9% (87.4 to 90.8 tok/s median, 5/5 pairs). Gate 16/16, 839 assertions. Output coherent but not byte-identical (Q5_K integer dot product vs bf16 float MAC). VT_GDN_ROWPERM_KEEP_QUANT=0 reverts to the old bf16 expansion path for A/B isolation. The improvement is less than the projected 14% because the Q5_K GEMV kernel has lower effective bandwidth on small grids (n=2560) than assumed, and wvSplitKSml is more efficient on these grids than projected. Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:glm-5-2 [omp]
…consumers The prior bridge invalidated the producer token on any non-matching K-quant consumer query, forcing standalone QuantizeQ8KK launches when a non-matching GEMV (e.g. attn o_proj reading the attention output) queried between the norm producer and a matching consumer that shares the normed hidden state. A non-matching consumer reads a DIFFERENT activation buffer and cannot stale the token — only a new producer overwrites it. Removing the invalidation lets matching consumers reuse the fused scratch even after irrelevant queries intervene. The gate test confirms: a non-matching query no longer invalidates, and a subsequent shape-matching call on the original buffer still reuses the fused scratch (consumers_fused == 1, was 0). Byte-identical outputs (5/5 A/B pairs, 1068 bytes). Gate 16/16, 839 assertions. Correctness: the token records the PRODUCER's output pointer and scratch. A non-matching consumer reads a different buffer that cannot overwrite the producer's output or scratch. The token is overwritten only when a new producer calls NormQuantRecordProducer. Stream-ordering is unchanged (single stream, sequential dispatch). Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:GLM-5-2 [TOOL]
…pKernel The fused Q8_K quant epilogue in RmsNormRowCoopKernel re-reads the normalized output from global memory (DLoadAct on orow) after Pass 3 stores it. On gfx1100 the 5 KB bf16 row (h=2560) competes with the weight and input in the 16 KB L1, so the re-read can miss to L2. T24 stores the normalized row to dynamic shared memory during Pass 3 (when the value is already in registers) and reads from LDS in the quant epilogue, eliminating the global re-read. The LDS buffer is h * sizeof(Tout) bytes (5 KB for bf16 h=2560), well within the 64 KB per-CU limit. Env gate VT_RMSNORM_LDS_QUANT (default ON) controls the optimization: set to 0 to revert to the global re-read path for A/B isolation. The gate is read per-call so captured graphs and in-process tests pick it up at dispatch time. Byte-identity: the LDS store uses the same conversion as Store (bf16 RNE for bf16 output, exact copy for f32), and DLoadAct reads the same bytes from LDS as from global. Gate test: 16/16 cases, 839 assertions, all passed. A/B measurement pending: the co-tenant 27B model holds the GPU VRAM, blocking the acceptance workload. The A/B script is staged at agent-artifacts/tg200-t24/ab-t24.sh for when the GPU is available. Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:GLM-5-2 [OMP]
ssm_out (out_proj) is Q5_K in the GGUF checkpoint but was expanded to bf16 at load time because the V-head column reorder (ReorderVCols) cuts across Q5_K 256-element block boundaries. T25 keeps the weight in tiled Q5_K order (no ReorderVCols) and permutes the 4096-element GEMV input from grouped to tiled order at runtime instead, cutting weight bandwidth ~4x (Q5_K ~5 MB vs bf16 20 MB per call). The permutation is a simple gather of 128-element groups within each of the 4096-element rows, gated by VT_GDN_COLPERM_KEEP_QUANT=1 (default OFF). A new out_proj_tiled flag on GdnLayerWeights distinguishes the tiled Q5_K path (needs input permutation) from the gdn_expand_nk bf16 path (already column-reordered, no permutation needed) — the nk flag alone conflates both. A/B (5 interleaved pairs, --max-tokens 256 --temperature 0 --seed 0): OFF median=90.930 tok/s, ON median=91.703 tok/s, +0.85%, 5/5 ON>OFF. Output coherent but NOT byte-identical (Q5_K vs bf16 weight precision). Gate test: 16/16, 839 assertions. The improvement is modest because the permutation kernel launch overhead (~13.4 us x 24 calls = ~322 us/tok) offsets most of the weight bandwidth savings (~368 us/tok). The net gain is ~46 us/tok. Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:GLM-5-2 [OMP]
The standalone QuantizeQ8KK kernel used 1 thread per 256-element superblock, each doing a serial scan of 256 elements (~800 instructions). For decode (m=1, nsb=10) only 10 of 128 threads were active, and on wave32 each thread is its own wave, so the kernel took ~13.4 us/call = 540 us/tok (6.0% of wall time). The new QuantizeQ8KKWarpCoop kernel uses 8 threads per superblock (32 elements each). The amax scan is done per-chunk (ascending, ax > amax first-occurrence), then reduced across 8 threads via __shfl_xor_sync with lower-chunk-index tie-break — equivalent to a sequential scan of all 256 elements. The quantization (iscale = -127/mx, DNearestInt, clamp 127) and bsums are order-independent. Output is BYTE-IDENTICAL to the original QuantQ8KSBlock, asserted by the gate test (16/16, 839 assertions) under VT_QUANT_Q8K_WARP=1. For m=1, nsb=10: 1 block, 80/128 threads active (vs 10/128), 3 waves of ~100 instructions (vs 10 waves of ~800) = ~8x fewer wave-cycles. A/B on acceptance workload (Qwen3.5-4B Q4_K_M, 256 tokens, temp 0, seed 0): OFF median: 91.532 tok/s ON median: 93.417 tok/s +2.06%, 5/5 pairs ON>OFF, all 5 byte-identical (1039 bytes) Gated by VT_QUANT_Q8K_WARP (default OFF, read per-call). Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:GLM-5-2 [OMP]
The Q6_K GEMV inner loop ran two amd_mixed_dot calls per iteration: one for the value dot product (sub = v . W8) and one for the bias correction (pre = 32 . W8), then subtracted. Since (v - 0x20) . W8 equals sub - pre exactly, subtracting the bias word before the dot product eliminates the second dp4a entirely. Unsigned byte subtraction wraps to the correct signed char value (v < 32 maps to 256+v-32, which is signed v-32), so the char4 reinterpret is correct. Also sets kGemvWarps from 4 to 8 for better L2 weight sharing on gfx1100; measured neutral-to-positive across the acceptance workload. Measured: 96.0 tok/s median (5 reps, 256 tokens, all coop + GQA4 env vars), unchanged from baseline. The Q6_K path is 15% of decode GPU time and the kernel is memory-bound, so halving the dot-product count does not move the wall clock. The change is kept because it reduces ALU pressure without a correctness cost. Correctness: test_rocm_quant_dot passes. The fused dot product is bit-exact with the original two-dot formulation because unsigned subtraction modulo 256 equals signed subtraction for 8-bit values. Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:glm-5-2 [omp]
S1-S3 complete at 96.0 tok/s median (target 150). The campaign exhausted optimization attempts within the allowed edit scope (src/vt/rocm/ and tests/vt/). The Outcome section records what was measured, what was rejected, why each default has its value, and the next traceable hypothesis: porting the CUDA-only async device mirror and executor to ROCm (requires runner.cpp changes outside the allowed scope) and a fundamentally different GEMV kernel design for >53% bandwidth efficiency. Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:glm-5-2 [omp]
Brings in the GFX1100-TG150 spec (.agents/specs/gfx1100-tg150.md) and its issue-index row. Pure spec addition, no code changes.
Brings in the KERNEL-QUANT-CIQ-GEMM-ROCM W1 keep-quant providers (rocm_quant_dot.hip) and the ROCM-QUANT-GEMM-BW bandwidth work (split-K decode, branch-free Q6K, f32-query GQA, row-permuted keep-quant). Both quant-dot paths coexist: TG200's rocm_grouped_gemm.hip routing and CIQ/BW's dedicated rocm_quant_dot.hip provider. # Conflicts: # CMakeLists.txt # src/vllm/model_executor/models/qwen3_5_gguf_weights.cpp # src/vt/rocm/rocm_paged_attn.hip # tests/CMakeLists.txt # tests/vt/test_rocm_quant_dot.cpp
ghazni101
force-pushed
the
row/GFX1100-TG200
branch
from
August 27, 2026 23:27
cb5088c to
70fb405
Compare
Cherry-pick 872b53e onto row/GFX1100-TG200 to enable fp8 KV cache on the production server. Resolved tests/CMakeLists.txt conflict by keeping both the TG200 test additions and the new test_rocm_fp8_kv_cache. Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:GLM-5-2 [OMP]
) The fp8 KV cache store and correctness-grade read landed in W6 (fp8-kv-cache.md ## W6), but the fp8 read through the fast decode kernel is owed: --kv-cache-dtype fp8 forces PagedAttnOnline because every optimized decode kernel requires k_cache.dtype == kBF16. Measured 7.5x regression at 16K context vs bf16 KV on gfx1100. This spec covers widening the VT_ATTN_DECODE_GQA4 dispatch guard to accept kI8 KV and adding an fp8 dequant load path inside the GQA kernel. The dequant arithmetic already exists in LoadKv(uint8_t*, ...) at rocm_paged_attn.hip:176; the fast kernel just does not call it. Fork issue #7. Separate spec and implementation PRs (developer preference 2026-08-27). Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:GLM-5-2 [OMP]
Widen the VT_ATTN_DECODE_GQA4 dispatch guard to accept DType::kI8 KV cache when kv_cache_dtype != kAuto, and add an fp8 dequant load path inside the GQA kernel. Previously --kv-cache-dtype fp8 forced PagedAttnOnline (the reference kernel with per-key __syncthreads reduction) because every optimized decode kernel required k_cache.dtype == kBF16. Measured 7.5x regression at 16K context vs bf16 KV on gfx1100. Changes to rocm_paged_attn.hip: - Template PagedAttnDecodeGqaF32Q on TKV (default __hip_bfloat16) with k_scale/v_scale parameters - Add LoadRowEplFp8<EPL>: vectorized uint8_t loads (uint32_t for EPL=4, uint2 for EPL=8, uint4 for EPL=16) with F8E4M3ToF32Dev dequant + scale - Add LoadRowEplKv<EPL, TKV>: if constexpr dispatch to LoadRowEplBf16 or LoadRowEplFp8 - Widen dispatch guard: kBF16||kI8 with kv_cache_dtype != kAuto - Add fp8 launch paths: PagedAttnDecodeGqaF32Q<...,uint8_t> with args.k_scale/args.v_scale The dequant arithmetic (F8E4M3ToF32Dev * scale) is identical to the existing LoadKv(uint8_t*, ..., scale) at line 176, so the fp8 read values are bit-identical between the fast and fallback kernels. A/B benchmark (Qwen3.5-4B Q4_K_M, RX 7900 XTX, ROCm 7.14.0, 128-token greedy decode, 4 reps): ctx=256: 99.94 -> 140.19 tok/s (1.40x), bf16=143.15 (0.98x) ctx=1024: 56.28 -> 122.31 tok/s (2.17x), bf16=129.02 (0.95x) ctx=4096: 20.53 -> 79.53 tok/s (3.87x), bf16=92.08 (0.86x) ctx=8192: 11.08 -> 54.55 tok/s (4.93x), bf16=66.85 (0.82x) ctx=16384: 5.78 -> 33.47 tok/s (5.79x), bf16=43.16 (0.78x) Token-exact: 3/3 simple prompts match the PagedAttnOnline fallback output byte-for-byte. Divergence at thinking-tag tokens is the reduction-order risk the bf16 GQA4 path already carries. Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:GLM-5-2 [OMP]
Replaces the serial <<<n,1>>> random-sample kernel with a block-cooperative argmax reduction over Gumbel scores, mirroring the CUDA fix from mudler#1984. The RNG and argmax reduce come from include/vt/sample_common.h, shared across CPU/CUDA/ROCm so bit-identity is a build property, not a copy-sync property. The serial path is retained behind VT_FAST_RANDOM_SAMPLE=0 for same-binary A/B gating. Adds ROCm + CUDA subprocess A/B test cases that re-exec the binary with VT_FAST_RANDOM_SAMPLE=0 and =1 and assert byte-identical token ids across widths up to 248320 (Qwen3.8-27B vocab). Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:MODEL [TOOL]
…mplate fallback backend.h: RocmAttentionBackend now advertises fp8 and fp8_e4m3 in supported_kv_cache_dtypes, matching the fp8-e4m3 KV cache read path in rocm_paged_attn.hip (W6). server_main.cpp: a .gguf model has no tokenizer_config.json — its chat template lives in GGUF metadata. Falls back to LoadChatTemplateFromGguf when the config path throws and the model is a .gguf file. Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:MODEL [TOOL]
upstream/main moved 101 commits past this branch's base (a73b269): the SAMPLE-CORE parallel Gumbel draw (mudler#1984/mudler#1997), SAMPLE-PROMPT-LOGPROBS, the SPEC-DFLASH2 waves, and the LTX-2 oracle tooling. Two files conflicted: - scripts/env-doc-allowlist.txt: this branch sorted the list and dropped the orphaned comment block; upstream appended VT_TT_SLOT_TRACE, VT_DUMP_QKVZ, VT_DUMP_TRUST around VT_FAST_RANDOM_SAMPLE. Kept the sorted shape and folded the four upstream names in. - tests/vt/test_ops_sample.cpp: both sides carry the mudler#1997 sampling work (this branch cherry-picked it as 0617b3f), so the auto-merge reunited the shared primitives; kept this branch's ROCm random_sample suites verbatim where upstream added nothing beside them. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:GLM-5.3 [OMP]
…d red check-agent-record and check-env-doc both failed on the branch head, so no push of this row could chain a green gate. - issue-index: mudler#1586 was listed twice (BACKEND-ROCM umbrella row and the ROCM-QUANT-GEMM-BW row). The umbrella row keeps the link; the quant-gemm row's spec already records "Owned under issue mudler#1586", and the duplicate is what the checker refuses. Dropped the duplicate append before it lands, so no union merge can resurrect it. - issue-index: the #7 row carried `kBF16 || kI8` verbatim, and the two pipes split the table row (7 pipes, 5 expected). Reworded to "accept `kI8` KV as well". - engine-matrix: the upstream-main merge shifted qwen3_5_gguf_weights.cpp by +52/+131 lines; re-anchored LoadGgufSharedEmbedAndHeadBf16 to :1067 and LoadQwen3_5MTPFromGguf to :1556, restoring anchor rot to the upstream baseline (33). - env-doc-allowlist: the keep-quant campaign kernels read eight kernel-internal knobs (VT_ASYNC_DEBUG, VT_GDN_COLPERM_KEEP_QUANT, VT_GDN_ROWPERM_KEEP_QUANT, VT_MM_TRACE, VT_QDOT_SPLITK, VT_QDOT_TRACE, VT_QUANT_Q8K_WARP, VT_RMSNORM_LDS_QUANT) that never made the list; allowlisted in sorted position. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:GLM-5.3 [OMP]
Conflict resolutions taken from row/GFX1100-ALL, which already combined this campaign with later main: LoadChatTemplateForModel in server_main (mudler#2077/mudler#2079), keep-quant grouped GEMM plus mudler#1910 Q6_K coop, and both campaign and upstream test registrations. Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:cursor-grok-4.6 [OMP]
…code-attn Keep the campaign performance work (including bf16 PagedAttnDecodeGqaF32Q) and upstream W6 store/read. The fast fp8 GQA path, dtype advertise, and rocm-fp8-kv-decode-attn spec move to row/fp8-kv-decode-attn. FOLLOWING_AGENTS_PROTOCOL: true AI-Assisted: true Assisted-by: AGENT:glm-5-2 [OMP]
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Row
GFX1100-TG200— raise Qwen3.5-4B Q4_K_M single-stream decode throughput on RX 7900 XTX (gfx1100) toward >=200 tok/s pure-autoregressive greedy. Spec:.agents/specs/gfx1100-tg200.md. Tracking issue: #5.Before starting
row/GFX1100-TG200), spec committed FIRST as its own commit; predecessor laddersrow/GFX1100-TG150-SPEC(spec(GFX1100-TG150): commit the 150 tok/s campaign spec #4) androw/ROCM-QUANT-GEMM-BW(GFX1100-TG150: serve Qwen3.5-4B-Q4_K_M on the RX 7900 XTX at >= 150 tok/s pure autoregressive tg #3) are consumed, not duplicated.docs/bench-evidence/gfx1100-tg200-*.md.What changed
Measured position moved 40.65 -> 76.60 tok/s median (+89%; the pre-campaign position was 27.6) on the acceptance workload: canonical 109-token prompt, 256 generated tokens, greedy, batch 1, idle host, gpu-ctl lock held. Levers, in landing order:
support_static_graph_mode) — wall/token approaches GPU-busy/token.VT_GEMV_MMVQ=1(+13.2%, byte-identical outputs) plus the m-gated dispatch and routing-witness gates; lever-B1 fold crossover tunable behindVT_GEMV_MMVQ_FOLD_MAX; lever-B2 f32-out skinny-BF16 armVT_SKINNY_BF16=1(+8.7%) displacing starved hipBLASLt Cijk tiles.RmsNormRowKernelbehindVT_NORM_QUANT_FUSED=1, producer-token consumption in the MatmulBTQuant dispatch; kills standalone single-block quant launches with byte equality by construction.QuantQ8KSBlockbody vectorized (dtype unswitched once, 16-byte loads, amax scan kept in strict ascending element order so the first-occurrence lowest-index tie-break stays bit-exact) — +23% alone; every consumer (standalone quantizer, fused epilogue, MMVQ prologue) benefits.PagedAttnDecodeGqaF32Qinstantiated at head_dim 128 behindVT_ATTN_DECODE_GQA4=1— f32-query full-attention decode had fallen to the generic serial kernel at 276us/call.VT_GDN_SCAN_COOP=1(donor walk was uncoalesced, ~17 GB/s effective).VT_ATTN_PREAMBLE_COOP=1.All seven knobs are classified kernel-internal on
scripts/env-doc-allowlist.txtperdocs/ENVIRONMENT.md's own policy; none flips a default.Evidence
docs/bench-evidence/gfx1100-tg200-t1..t4a-2026*.mdandgfx1100-tg200-t5-native-baseline-20260825.md. Closing table: GPU busy 12.13 ms/tok of ~13.1 wall (dispatch gap ~1 ms after graph capture).tests/vt/test_rocm_quant_dot12/12 cases 797 assertions including tied-amax adversarial;ctest -R 'rocm|cross_device|quant'green except two failures proven PRE-EXISTING by rerun on the pristine head (test_gguf_keep_quant, one cross_device case — owned separately).python3 scripts/agent-ready.py: All gates green on this tree (native gfx1100 build,-DVLLM_CPP_HIP_ARCHITECTURES=gfx1100).Speed claims
gpu-ctl run/reserve, standing serve parked via reservation) and they are recorded with repro recipes indocs/bench-evidence/gfx1100-tg200-*.md; gate scriptstools/tg200-run-gate{,2}.shreproduce every window.Honest gaps
.agents/specs/rocm-m4-oracle.md) is OWED before any default flip ofVT_ATTN_DECODE_GQA4,VT_GDN_SCAN_COOP, orVT_ATTN_PREAMBLE_COOP; until then they ride the campaign config and stay default-OFF.docs/BENCHMARKS.md/docs/USAGE.mdrows ride the T6 landing stage, which fires when the gate is met or the campaign reports its final measured position.FOLLOWING_AGENTS_PROTOCOL
Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:ox-alpha [omp]