perf(GFX1100-TG200): raise Qwen3.5-4B Q4_K_M gfx1100 decode to 76.6 tok/s - #1936
Draft
ghazni101 wants to merge 77 commits into
Draft
perf(GFX1100-TG200): raise Qwen3.5-4B Q4_K_M gfx1100 decode to 76.6 tok/s#1936ghazni101 wants to merge 77 commits into
ghazni101 wants to merge 77 commits into
Conversation
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]
…kinny gate (VT_SKINNY_BF16=1) Focused ROCm gate for the arm the 15.1 attribution calls for: bf16-in/ f32-out MatmulBT at M<=4, the ProjectGdnBA population that today rides rocBLAS's MT128x32x16 tile at ~73.7us/call. Two cases: a nine-shape sweep anchored on the exact engine shape (m=1, N=32, K=2560) with gate-boundary edges in both directions (odd N, K%8!=0, m past the skinny range, N at the feature floor), asserting the sibling 1e-6 NMSE band vs a CPU oracle on BOTH arms plus cross-arm agreement; and a routing-witness case through new host-side dispatch counters, including the F1-convention TRUE-unset leg (EnvGuard(false) writes "0" and can never witness default-OFF inertness). Red-first proven: before the seam exists the binary fails to link (SkinnyF32RouteCountsForTesting undefined); once counters exist without the env arm the routing legs go red behaviorally. Registration beside test_rocm_quant_dot under VLLM_CPP_HIP. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
…ecode-skinny arm Implements the arm the 15.1 attribution calls for, behind an opt-in env (default OFF; the default path is byte-unchanged). rocm_skinny_gemm.hip gains an output-store polymorph of the adopted wvSplitKSml geometry -- same split-K/LDS/DPP-reduction tree as the bf16 arm, only the store type differs (no rounding on the f32 arm) -- with WvSplitKBT refactored onto a shared templated launcher and a new WvSplitKBTToF32 entry. rocm_matmul_hipblaslt.hip dispatches bf16-in/f32-out M<=4 shapes to it under the same donor guards as the bf16 arm (N>8, N%2==0, K%8==0, LDS fit, wave32 arch) when VT_SKINNY_BF16=1, read per call per the cuda_quant_dot convention. This is NOT bit-exact by construction (the reduction order differs from rocBLAS); correctness is gated by the new focused suite's 1e-6 NMSE band vs the CPU oracle plus engine coherence, and routing is witnessed through host-side counters bumped on exactly the branch taken. Target population: the Qwen3.5 GDN BA projections (N=32, K=2560), ~3.54 ms/tok under graph replay on rocBLAS's large-M tile. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
…reamble lost in the red-first edit The second TEST_CASE lost its DeviceAvailable guard, queue bring-up and device-buffer setup when the stale helper blocks were cut during the red-first commit, so the focused target compiled neither there nor in the first container build (scope errors at the lambda, caught before any GPU run). Restores the guard, allocations and the F1-convention comment; drops the now-unused Cpu() helper that -Werror flagged. No assertion changes. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
…ations to per-arm deltas
The first green run red-flagged my own expectation arithmetic, not the
arm: over a RunBothArms sweep the OFF dispatch always bumps the BLAS
counter once, ON bumps exactly one branch, and dispatches outside the
counted population (m>4) bump neither -- so served shapes expect
{blas:1, skinny:1}, gate-excluded shapes {blas:2, skinny:0}, and m=5
{0,0}. The first run's failures were exactly this table inverted; no
kernel or dispatch change. Focused suite now 2/2 cases, 51/51
assertions, exit 0 under the gpu-ctl lock.
FOLLOWING_AGENTS_PROTOCOL
Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:ox-alpha [omp]
…ut decode-skinny arm at +8.7% median Closes the lever with numbers: red-first proven at the link seam (undefined route-counter symbols), focused gate 2/2 cases 51/51 assertions green under the gpu-ctl lock, sibling screens untouched (test_rocm_quant_dot 752/752, test_ops_matmul 16/16), both mutations (inert knob, N-floor drift) caught with byte-equal restores, and an interleaved same-window engine A/B with ALL FIVE pairs won by ON -- medians 35.616 (OFF) vs 38.731 (VT_SKINNY_BF16=1) tok/s, +8.7%. One unique coherence md5 per arm (OFF matches the canonical adopted stream 2b29ad66...; ON fe771fb7... is self-consistent, sane prose, no loops); the cross-arm divergence at an early near-tie is the expected numerics class and its adjudication stays owed separately -- the flag ships default-OFF. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
… shows the starved Cijk tile fully displaced Post-adoption rocprofv3 capture (63 decode steps, opt-in flag on): zero MT128x32x16 launches remain at the BA decode signature; the only ones left are the single prefill pass (M=89, out of the skinny scope by design), and wvSplitKSml<1>'s call count closes exactly as 5 projections x 24 GDN layers x 63 steps. First capture attempt hit co-tenant VRAM pressure and was rerun clean, recorded per the incident convention. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
… observe getenv()==NULL
Reviewer finding F-1 (B2Review, HIGH): run_window always constructed
EnvGuard, whose constructor setenvs "0"/"1", so the claimed true-unset
window exercised getenv()=="0" and its two CHECKs pinned nothing --
reviewer mutation M-A (default return inverted) passed 51/51 green.
run_window now takes an explicit WindowEnv mode and constructs no guard
in the kTrueUnset mode; the dead never-called EnvGuard::Unset() is
removed. Re-pin under M-A: the unset-window counts flip to {blas:0,
skinny:1} and the case fails; restored byte-equal (md5 04f2a15e...) it
is green again at 51/51.
FOLLOWING_AGENTS_PROTOCOL
Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:ox-alpha [omp]
…ss against reviewer mutation M-A
Evidence 15.7 for the F-1 repair: baseline green (51/51), M-A red on the
fixed window with exactly the directed count flip ({blas:0, skinny:1})
and only those two assertions failing, byte-equal restore (md5
04f2a15e...), post-restore green, and the sibling screen unchanged
(test_rocm_quant_dot 752/752). All GPU runs under gpu-ctl; build recipe
per section 1 with /work/b2fix-src + /work/build-b2fix.
FOLLOWING_AGENTS_PROTOCOL
Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:ox-alpha [omp]
…aunches to producers Pre-code attribution artifact for the producer-fused Q8_K norm epilogue lever, per the Lever C contract step 1. From the bdb445f9ac06 rocprofv3 capture: 97 standalone single-block QuantizeQ8KK launches per decode token (~49 us each); 57/tok are fed by RmsNormRowKernel outputs (FFN gate_up x32, attn q/k/v x24 re-quantizing the same normalized row, lm_head x1) and are fusable via a producer epilogue; 40/tok (o_proj, down_proj) are not and stay owed. RmsNormGatedK finding recorded: zero quant consumers in this model, extension deferred with reason. The assignment's quoted 43/tok is honestly reconciled against the measured 97/tok. Fusion-seam gate scope finding stated. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
…orm-quant epilogue
Three focused cases ahead of any implementation: (1) routing witness --
VT_NORM_QUANT_FUSED=1 must advance the producer counter, skip the
standalone QuantizeQ8KK for matching consumers including a second
consumer of the same normalized row (the attn q/k/v pattern), and stay
byte-identical to the OFF arm; flag unset keeps the standalone route.
(2) scratch byte-equality vs the standalone quantizer over random,
tied-amax (first-occurrence tie-break adversarial), and all-zero rows
at nsb {1,3,10} and m {1,3}. (3) stale-token guard -- a non-matching
K-quant consumer takes the standalone quant and invalidates the token.
References NormQuantCountsForTesting/NormQuantResetForTesting/
NormQuantLastScratchForTesting, which do not exist yet: link-red per
the T4a convention.
FOLLOWING_AGENTS_PROTOCOL
Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:ox-alpha [omp]
… arms T10 (VT_GDN_POSTCONV_COOP=1): the chunked postconv donor hands each decode item to one thread walking dk serially twice — 27.9us/call against a sub-microsecond floor. The arm gives each item a warp with lane-strided walks and shfl sumsq trees; the association changes, so the flag is opt-in with adjudication owed. T11 (VT_GDN_SCAN_SPLIT=1, requires SCAN_COOP): the cooperative scan runs 32 blocks at decode on a 96-CU board. State rows are independent given the shared q/k/v scalars, so rows split across RS=4 blocks per head and each lane caches its row segment in registers between the two passes. Per-row arithmetic is unchanged: engine outputs are BYTE-IDENTICAL across all five A/B pairs through 24 layers. Focused gate: full suite 15/15 cases, 826 assertions including the new T10 NMSE + inertness case. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
… pointer mismatch documented The lever-C producer-token extension to the gated norm was implemented and passed the full focused gate at 16/16 cases / 850 assertions — the fused epilogue scratch is byte-identical to the standalone quantizer and the CPU oracle, and an in-process probe shows consumers_fused advancing with a same-pointer matvec. At engine level it measured inert: rocpd shows the standalone quantizer still launching 40 times per token because the FFN matvec activation does not match the recorded producer output pointer. Reverted byte-restored per the non-winner precedent; this record carries the diagnosis (trace qwen3_5.cpp FFN call sites for the buffer/view mismatch) so a successor starts from plumbing, not from the kernel again. Failed-attempt ledger: 3 of 10. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
Argmax-delimited step analysis over the campaign-config capture puts non-kernel time at ~1.83 ms/step under trace, with the largest stalls clustered on amd_copyBuffer pairs flanking Embedding — the per-token argmax-result D2H round trip. On-device token feedback (the LAGUNA path's on-device sampling) is the named candidate; an untraced paired measurement is owed before any lever claim. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
…uts found post-hoc Body inspection after landing showed both arms of the T10/T11 acceptance window produced degenerate token loops, not coherent prose. Root cause: GdnPostConvWarpK computed the conv row stride as key_dim+value_dim instead of 2*key_dim+value_dim ([q|k|v] layout) — decode rows masked it, prefill rows read wrong memory and poisoned generation from step one. The stride is fixed and the focused gate stays green (15/15 x 826), but the engine numbers in this file are retracted until a clean-window re-measurement lands; the recorded verified position returns to the T9 number (77.7 tok/s median). The coherence-check rule exists for exactly this failure and was violated by claiming prose coherence without reading the bodies. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
GdnPostConvWarpEnabled read VT_GDN_POSTCONV_COOP once per process, so the focused T10 case's ON arm silently reused the donor kernel: both arms compared donor-vs-donor, nmse was vacuously zero, and the stride bug landed behind a green suite. The toggle now reads per call (matching the T8/T9 arm convention), and the case's assertion set is restored from an earlier editing loss: g/beta must be byte-identical across arms (no reassociation touches them) and q/k/v sit inside the 1e-6 NMSE band. Mutation-verified: reintroducing the key_dim+value_dim row stride makes the case fail at nmse=1.299; the fixed kernel passes. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
…re-measurement The retraction left the restore-or-close choice to whoever reads the re-run output. The rules are now written down: coherence probe plus manual spot-check voids a window, T11 pairs must be byte-identical or the split arm dies, T10 divergences are recorded tie flips, and adoption requires the ON median to win at least four of five pairs — otherwise both arms revert byte-restored per precedent. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
…rgets Memory-copy tables show zero D2H inside decode windows — the sampled-id handoff rides rocclr_copyBuffer kernel entries. The per-step anatomy at an Argmax boundary prices three concrete items: the ~290us sampling round trip (host wakes on the id copy), ArgmaxK itself taking 154us to reduce a single 993KB logits row (~70x over its memory floor, the same single-block latency class T8-T11 attacked), and a long tail of 2-6us state-update copies behind 3-5us host gaps worth ~0.3-0.5ms/tok together. Each replaces the coarse "~1.83ms dispatch gap" with an addressable target. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
The donor argmax launches one block per row: batch-1 decode walks the full 248,320-element vocab with ~970 serial load+compare iterations per lane — 153.96us against a ~2-3us memory floor. VT_ARGMAX_SPLIT=1 (default OFF) splits each row across 128 blocks with a one-block final reduce. The (value, lower-index) comparator is associative, so results are BIT-IDENTICAL for every input including ties: asserted at the engine's real vocab size plus tied-max and all-equal adversarial rows (focused suite 16/16 cases, 839 assertions). Engagement capture shows both phase kernels at 34.2+10.5us = 44.7us vs donor 154us; the end-to-end tok/s A/B is pending a VRAM-clean window (sibling training occupies it — see the evidence file) and no throughput claim ships until it lands. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
…dot4 hardware Disassembly of an -O3 gfx1100 compile shows the repo's scalar Dp4a fallback auto-forms v_dot4_i32_iu8 with proper signedness negation — the explicit sdot4 intrinsic would additionally require the dot1-insts feature flag. Every K-quant kernel therefore already executes the hardware dot instruction; the quant families have no instruction-selection headroom on this silicon, closing the last speculative lever there at the ISA level. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
Collaborator
|
Approved the workflow runs on the current head ( Two things worth knowing before you read the results, so you don't chase them:
Leaving this as a draft since you're still pushing to it. Ping me when you want a review — the bench-evidence set and the split-scan/postconv work look substantial and I'd rather review the finished shape than a moving one. |
…d at +4.7%/+3.2% The stride-fixed arms were re-measured in a VRAM-probed window whose arm bodies passed the coherence probe: T10 OFF/ON medians 82.42/86.31 tok/s (ON wins all five pairs, +4.7%) and stacked T11 82.29/84.95 (+3.2%, all pairs byte-identical — bit-exactness held through 24 layers). Full-config steady state reads 92.9/92.7 tok/s on the canonical 70-token prompt. Prompt-length caveat recorded: tonight's pairs used a shorter prompt than older windows, so paired deltas are the verified quantity and the formal long-prompt gate rerun stays owed for absolute cross-era claims. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
…onical, graph replay confirmed Clean-GPU verification of the complete eleven-arm config: coherent output, graph replay engaged across all arms added this session (126 replays captured), steady state 92.9/92.7 tok/s on the canonical 70-token prompt and 99.9/101.1 on a shorter probe. Fresh rocpd budget at this config: 8.89 ms/tok kernel busy with the three streaming families at their audited near-peak rates; residual non-kernel time ~1.9 ms/step decomposes into the ~290us sampling round trip plus per-op launch gaps, scoped as the next session's target via the async-serving engine path. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
…d — sync-loop deferral scoped With real events, VT_ASYNC_RUNNER now resolves sched_supported=1 and the server serves coherent output on AsyncScheduler mcb=2 — the capability enablement is correct end to end. But the paired A/B through the OpenAI endpoint is a wash (55.9 sync vs 55.7 async medians) because the server path itself runs ~40% slower than the CLI on identical flags, drowning any scheduler-overlap gain; dual-server interleaving is also unavailable (two full engines cannot share the VRAM). The contained successor is a one-step-deferred sampled-id D2H inside LLMEngine::step — overlap without touching the scheduler. 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]
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
2 participants
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: ghazni101#5 (the fork campaign record; this PR is its upstream candidate).Before starting
mainon either remote; every GPU command through the operator lock); the spec committed FIRST as its own commit; the predecessor ladders (GFX1100-TG150,ROCM-QUANT-GEMM-BW) 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]