gfx900 (Vega10/V340) inference support + decode perf (dense 8.2x, MoE 3.2x)#69
Open
larkinwc wants to merge 30 commits into
Open
gfx900 (Vega10/V340) inference support + decode perf (dense 8.2x, MoE 3.2x)#69larkinwc wants to merge 30 commits into
larkinwc wants to merge 30 commits into
Conversation
- rocm.py: add _ON_GFX900 / on_gfx900() as a separate tier (not in _ON_GFX9) so all MFMA codepaths skip automatically on Vega10 - rocm.py: force TRITON_ATTN backend on gfx900 (no MFMA paged-attn kernel) - CMakeLists.txt: add gfx900 to HIP_SUPPORTED_ARCHS (MFMA kernels guard out via __HIP__GFX9__) - GFX900_SETUP.md: document the gfx900 dependency stack (rocBLAS/RCCL/amdsmi), topology/bandwidth findings, and the remaining Triton blocker Validated: ROCm 7.2.4, torch 2.12.0+rocm7.2, custom gfx900 RCCL all-reduce (2/8/16 GPU), gfx900 rocBLAS sgemm, vLLM csrc builds for gfx900. Blocker: prebuilt Triton rejects gfx900 codegen.
Document the triton-gfx900 fork (dedicated GFX900 ISA family on triton v3.7.0) that unblocks vLLM. Validated: opt-125m inference on Radeon Pro V340 via TRITON_ATTN backend.
Qwen3.5-9B (Gated DeltaNet + full attention hybrid) runs across 8x gfx900 GPUs text-only. FLA GDN Triton kernels (chunk_gated_delta_rule, causal_conv1d, l2norm) work on Vega10 with no extra fixes beyond the GFX900 ISA family. Config: TP=8, dtype=float16 (no native BF16), enforce_eager, language-model-only.
8-cell grid (TP{4,8} x concurrency{1,4} x {synthetic,coding}) via vllm bench
serve, mirroring the MI100 quant-bench methodology. FP16 (no MFMA on Vega10 so
quant kernels N/A). Best: TP8 c4 coding 20.5 out tok/s; TPOT floor ~120ms/token
reflects VALU FP16 GEMM (no matrix cores). Full TTFT/TPOT p50/p99 included.
8 more cells comparing TP2xPP2 (4 GPU) and TP2xPP4 (8 GPU), single-socket, vs the pure-TP grid. PP wins on 7/8 cells (4-GPU: +48-90%; 8-GPU: +17-25% with concurrency) because it swaps per-layer all-reduce for cheap point-to-point activation hand-offs - a big deal on a box with no XGMI and slow PCIe/QPI. Required fix: gfx900 RCCL chokes on the V1 async-scheduling PP sampled-token broadcast (EngineDeadError); run server with --no-async-scheduling.
Measures raw FP16 GEMM TFLOPS + HBM bandwidth per-GPU and aggregated: - Per-GPU sustained ~4.5 TFLOPS FP16 (21% of peak; rocBLAS uses no packed-FP16), ~365 GB/s HBM (76% of peak). - Aggregate scales linearly: 8 GPU = 38 TFLOPS/2.9 TB/s, 16 GPU = 76 TFLOPS/5.8 TB/s. - Power scan 110->50W: cards are compute-bound (draw only ~85W at stock), so capping to 80W costs zero throughput at best TFLOPS/W; knee ~70W. - Key finding: passive V340 cooling, not power, is the real limiter - a heat-soaked GPU throttled to 300MHz (-78%). Recommend 80W cap + airflow.
Back out effective compute/bandwidth from the BENCH_GFX900 serving runs: - Prefill (compute-bound) realizes ~50-76% of the parallel-compute ceiling (~24 TFLOPS/socket usable); TP-friendly. - Decode (nominally bandwidth-bound) realizes only ~5-9% of the HBM roofline - it is actually latency-bound (eager kernel launches + per-layer PCIe all-reduce, no MFMA/packed-FP16/HIP-graphs), so HBM sits >90% idle. - Realistic planning: ~24 TFLOPS/socket prefill, decode in tens of tok/s aggregate, NOT the 76 TFLOPS / 5.8 TB/s linear figure. PP beats TP for decode.
Tried to remove the ~120ms/token decode floor by enabling HIP graphs: - torch.compile/Inductor crashes on ROCm (is_current_stream_capturing returns non-Tensor) -> mode=NONE mandatory. - Pure HIP graph capture (FULL_DECODE_ONLY) at TP8 hard-hung a GPU; driver reset FAILED (stale Vega10 SMU firmware 0x1 vs 0xe) -> GPU dead until host reboot. - Post-reboot, safest retry (TP4, capture_sizes=[1]) no longer hangs but never finishes: infinite Triton autotune during capture warmup (600s timeout). Conclusion: enforce_eager is correct/required on gfx900; decode floor is intrinsic. Levers remain PP-over-TP and batching (per BENCH_GFX900.md).
…code speedup Corrects the earlier (wrong) conclusion that graphs are unusable. Key findings (A/B on TP4 Qwen3.5-9B FP16, bs=1 decode, 128 tok): - Eager: 6.69 tok/s, 149.4 ms/tok - Graph (mode=NONE + FULL_DECODE_ONLY, capture_sizes=[1]): 6.55-6.64 tok/s, 150.7-152.6 ms/tok -> within noise / slightly slower. - Capture confirmed active in logs (Capturing CUDA graphs decode FULL 1/1, finished in 1s) with identical correct output. So the ~150 ms/token decode floor is NOT kernel-launch overhead; removing it via graphs changes nothing. The floor is intrinsic hardware cost (no MFMA, no packed-FP16, per-layer PCIe all-reduce). enforce_eager stays the right default. Three blockers documented + fixed to get graphs running: - A: torch.compile/Inductor crashes on ROCm (is_current_stream_capturing) -> use mode=NONE; Inductor is the wrong target anyway. - B: capture hang + destructive BACO auto-reset (lost VRAM, wedged SMU). Root cause is NOT stale firmware (boot is clean). Fix: amdgpu.reset_method=2 (mode1, PSP-assisted) via /etc/modprobe.d -> reliable recovery. mode2/value 3 is Arcturus-only, not Vega10. - C: one-time GDN/FLA Triton autotune during capture warmup is slow (cold 647s, warm 63s, capture itself 1s). Slow, not infinite; result persists. Adds bench_scripts/ab_decode.py and bench_scripts/graph_safe.py.
…lp on gfx900 Answers "are we using all the HBM bandwidth?" - at bs=1, no (~9%), because decode is per-step overhead-bound, not bandwidth-bound. Lever 1 (#53) batching: TP4 FP16 decode scales 6.8 tok/s (bs=1) -> ~300-500 tok/s aggregate at B=64-96 (50-70x). Per-step time stays ~150ms (reads weights once per step). Even at peak we are ~10x above the HBM roofline -> never bandwidth-bound; the RCCL all-reduce + GDN Triton kernels are the ceiling. Lever 2 (#54) INT4 AWQ: QuantTrio/Qwen3.5-9B-AWQ LOADS + RUNS correctly on gfx900 via TritonW4A16LinearKernel (TP2). But it is SLOWER than FP16: - bs=1: 200ms/tok (AWQ TP2) vs 147ms/tok (FP16 TP4), +36% - peak: ~56 tok/s (plateaus B=32) vs ~300-500 tok/s (FP16), ~6-9x worse Reason: gfx900 has no native INT4 units, so Triton must dequant INT4->FP16 in software every forward = pure added compute. Quant only helps when bandwidth-bound; gfx900 decode is overhead-bound, so it adds cost with no benefit. Remaining gains are in #55 (PP-over-TP cuts all-reduce) and #56 (profile RCCL vs GDN kernel split). Adds bench_scripts batch_sweep/awq_sweep/awq_test.
Hand-written M=1 GEMV that sidesteps tl.dot (no MFMA on gfx900). Register-level reverse-AWQ int4 unpack, group scales/zeros loaded once per group, split-K to parallelize long-K/small-N shapes. Microbench (real Qwen3.5-9B MLP shapes, M=1, correctness rel err 0.0003): - gate/up K=4096 N=24576: 0.725ms FP16 -> 0.306ms int4 splitK = 2.37x (164 GB/s) - down K=12288 N=4096: 0.568ms FP16 -> 0.163ms int4 splitK = 3.48x (154 GB/s) Flips the earlier result (#54: stock AWQ via padded tl.dot was SLOWER than FP16). The win is bandwidth: int4 reads 4x fewer weight bytes and GEMV is bandwidth-bound, so no MFMA needed. Next: wire into vLLM AWQ decode path + gfx900 branch, end-to-end.
Wider autotune sweep (BLOCK_N, num_warps, split_k up to full n_groups, num_stages): - gate/up K=4096 N=24576: 0.726ms FP16 -> 0.257ms = 2.82x (196 GB/s) [BN=256 w=2 sk=32 st=2] - down K=12288 N=4096: 0.576ms FP16 -> 0.143ms = 4.02x (176 GB/s) [BN=128 w=1 sk=96 st=3] ~54% of the 365 GB/s roofline. Weighted MLP GEMM decode speedup ~3.1x. Remaining gap to roofline is the tl.sum cross-lane reduction + int4 unpack ALU; would need a different reduction structure for more. Good enough to wire into vLLM next.
…16 on half the GPUs (#57) Adds gfx900_w4a16_gemv.py (M<=8 int4 decode GEMV: register-level GPTQ-sequential unpack, per-group scales/zeros, split-K, per-shape tuned configs) and dispatches to it from triton_w4a16_gemm when on_gfx900() and M<=8 (prefill M>8 keeps tl.dot). Validation (Qwen3.5-9B-AWQ): - kernel correctness vs existing triton_w4a16_gemm: rel err <=0.001 for M=1/4/8 - microbench vs stock tl.dot path it replaces: gate/up 10.1x, down 15.6x - microbench vs FP16 GEMV: gate/up 2.82x, down 4.02x - end-to-end TP2 bs=1 decode: 95.6 ms/tok (10.46 tok/s) vs stock AWQ tl.dot 200 ms/tok (4.98) = 2.1x vs FP16 TP4 147 ms/tok (6.80) = 1.54x, on HALF the GPUs Flips INT4 from the worst decode option (#54) to the best for latency / low- concurrency / VRAM-constrained serving. High-batch still favors FP16 (decode becomes step-overhead-bound), handled by the M<=8 gate. Reversible: dispatch is gated on on_gfx900(); other arches unaffected.
In-tree TurboQuant kernels are pure Triton (no __shfl_sync/tl.dot), so the wave64 blocker that stops the llama.cpp HIP port does not apply on gfx900. Add TURBOQUANT to the gfx900 attn-backend list (opt-in via --kv-cache-dtype turboquant_*; default stays TRITON_ATTN). Validated on Qwen3.5-9B TP4: coherent output, GPU KV cache 194k->454k tokens (2.34x), decode +7-13%. Use turboquant_k8v4 (FP8 keys) since Qwen3 is key-quant-sensitive. Refs #62
AWQ int4 model (#57 GEMV) + --kv-cache-dtype turboquant_k8v4 run together: both paths active, coherent output, decode 9.97 tok/s ~= standalone AWQ (10.46), so TQ adds ~2.34x KV capacity at near-zero decode cost. Note 8GB per-die ceiling OOMs AWQ-TP2+16k-KV at batch>1; use TP4 or shorter ctx. Refs #57 #62
Prescriptive guide: AWQ int4 weights (#57 GEMV) + turboquant_k8v4 KV (#62) on TP4/one-socket is the recommended path for Qwen3.5-9B on this box. Captures dtype/eager/preset/layout rationale, layout-by-workload table, 8GB VRAM ceiling, and repro scripts. Also remove duplicated TurboQuant section from PERF_GFX900.md.
torch-profiler 127-step decode window, TP4 FP16 Qwen3.5-9B. Reality: GEMM (M=1 padded to 64x64 macrotile) = 43% wall / 88% GPU-busy, at only 20% of bandwidth roofline (padding waste). GPU-idle dispatch gap = 51%. RCCL only 2%, GDN 0.4ms, full-attn 0.2ms. So the old RCCL+GDN-bound assumption was wrong. Re-ranks candidates: lm_head GEMV (#59) and proj GEMV (#61) attack the 43%; full-attn kernel (#60) demoted. Refs #56
FP16 9.8885 vs turboquant_k8v4 9.8879 (-0.006%) on Qwen3.5-9B TP4. Confirms FP8-key preset avoids quirky-K catastrophe at zero measurable quality cost. Add bench_scripts/ppl.py. Refs #62
Root cause (from #56 profile): gfx900 is excluded from on_gfx9(), so the existing hand-written LLMM1/wvSplitK skinny GEMVs were skipped and every M=1 decode matmul fell through to rocBLAS Cijk_ (64x64 macrotile, pads a 1-row GEMV to 64 rows -> ran at 20%% of bandwidth roofline). Microbench: LLMM1 is correct on wave64 (rel err 5e-4) and 3.5x faster than rocBLAS for the lm_head shape (vocab 248320 x 4096). wvSplitK (n>1) device- asserts on gfx900, so gate to the LLMM1 n==1 branch only. Enabling it routes lm_head + the M=1 FP16 projections (qkv/o/gate_up where k<=8192, m%%4==0) through the real GEMV. E2E FP16 Qwen3.5-9B TP4 decode: 6.65 -> 14.85 tok/s (2.2x), 150 -> 67 ms/step, coherent. Re-profile: GEMM time 65.9 -> 6.5 ms/step (10x). down_proj (k=12288) stays on rocBLAS. Refs #59 #56
2.2x FP16 decode by enabling the existing LLMM1 skinny GEMV for gfx900 (was excluded from on_gfx9()). Add PERF_LMHEAD_GFX900.md and update the recommended-path table. Compounds with AWQ+TQ (9.97 -> 14.12 tok/s).
Revisited the old enforce_eager conclusion. It was valid when GPU-busy was 74.6 ms/step, but LLMM1 (#59) cut GPU-busy to 27.9 ms/step, making the host dispatch gap 59%% of the step. Graphs eliminate it: FP16 TP4 eager+LLMM1 = 67.0 ms/tok (14.93 tok/s) FP16 TP4 graph+LLMM1 = 18.4 ms/tok (54.41 tok/s) -> 3.6x, coherent, 2x reproduced Full stack (AWQ int4 + TurboQuant + LLMM1 + graphs) TP2 = 37.71 tok/s. Total decode vs original FP16 baseline: 6.65 -> 54.4 tok/s = 8.2x. Config: mode=NONE + cudagraph_mode=FULL_DECODE_ONLY + capture_sizes=[1]. Gotcha: hybrid Mamba/GDN capture needs max_num_seqs <= Mamba cache blocks. Flip recommended-path guidance to graphs ON. Add graph_ab/graph_combo scripts. Refs #63 #59 #56
Re-benchmarked PP vs TP now that LLMM1 (#59) made decode GEMMs real GEMVs and CUDA graphs (#63) removed the dispatch gap. 4-GPU head-to-head, LLMM1+graphs, Qwen3.5-9B: c=1 : TP4 52.3 vs TP2xPP2 32.5 (TP +61%) c=8 : TP4 48.9 vs TP2xPP2 54.8 (PP +12%) c=32: TP4 220 vs TP2xPP2 176 (TP +25%) Old eager data had PP winning every cell (+48-90%); the things that made pure TP slow (padded GEMM slices, per-step dispatch) are exactly what LLMM1+graphs fixed, so TP no longer needs PP. Recommend TP4 for decode now. 8-GPU TP2xPP4 crashes on offline path (eager+graph) -> not pursued. Refs #55 #59 #63
…ocket unusable (#55) Swept TP4/TP8/TP16 with LLMM1+graphs (Qwen3.5-9B FP16, offline, out tok/s): TP4 (1 socket): c1 52.3 c8 48.9 c32 220.0 <- best TP8 (1 socket): c1 14.7 c8 12.7 c32 61.8 (~3.5x worse than TP4) TP16 (cross-sok): decode stalls 20min+ in cross-QPI lm_head logits all-gather (vocab 248320 over 0.26 GB/s QPI). Killed. More TP than 4 hurts: 8-way all-reduce + splitting an M=1 GEMV 8 ways makes each slice too small to amortize launch. Cross-socket pure TP is unusable -- the 0.26 GB/s QPI link is the hard wall for any 16-GPU layout. TP8xPP2 (cross-socket PP) inconclusive: started decode then hit a collective TimeoutError while the host was overloaded; flagged for a clean retry on a quiet box (not recorded as a negative). The teardown also dropped two socket-1 dies off the PCIe bus (SMU non-responsive, needs cold power cycle). Recommend TP4 one-socket for decode; run independent per-socket engines for multi-tenant. Add topo_bench.py + run_topo.sh. Refs #55 #59 #63
…#65) MoE model runs coherently on gfx900 TP4xPP2 (one socket, GPUs 0-7): eager 8.5 tok/s -> CUDA graphs 19.3 tok/s (2.3x), bs=1 decode. Decode profile + branch instrumentation findings: - W4A16 g32 MoE correctly falls back (is_rocm) to Triton fused_moe_kernel_gptq_awq; healthy on wave64, not the bottleneck. - Every dense Linear projection already routes to LLMM1 (#59) -- the missed-gate hypothesis was disproven by direct measurement. - Remaining padded rocBLAS (MT64x64x8, ~7.6 ms/step) is the GDN linear-attn recurrence (fla ops, direct torch.matmul), not the Linear dispatch -- the real next kernel target, harder than a gate tweak. Refs #65 #59 #63
…decode (#65) Qwen3-Coder-Next shared_expert_gate is a Linear [hidden->1]; at decode its M=1 GEMV has m=1 (not %4) so it failed the LLMM1 gate and fell to padded rocBLAS, which computes a wasteful 64x64 macrotile for a 1-row output. Microbench [n=1,k=2048] m=1: rocBLAS 304 us vs fused mul-reduce 21 us (14x), rel-err 1.1e-3. Trace pinned it: aten::mm <- matmul <- linear <- rocm_unquantized_gemm, ~24 calls/step (one per layer in the PP stage). Fix: narrow gfx900 branch for n==1, m<=8, m%4!=0, bias=None -> (x*w).sum(-1). - MoE Qwen3-Coder-Next TP4xPP2 graphs: 19.3 -> 26.8 tok/s (+39%, 51.7->37.4 ms/tok), coherent. - Dense Qwen3.5-9B TP4 graphs unchanged: 54.4 tok/s (branch is narrow, does not touch it). Add gate_bench.py + trace_mm.py (trace parent-op finder). Refs #65 #59
…, graphs, TurboQuant) (#65)
…ins concurrency) (#65)
…ream) fused_moe_kernel_gptq_awq accumulates the expert GEMM with tl.dot. On gfx900 (Vega10, no MFMA) a 1-token decode is padded into a 16-row tl.dot (~16x waste), which the graph-mode profile shows is ~46% of the decode step. Add a GEMV_MODE path (tl.sum reduction over K instead of the padded matmul) and route gfx900 int4 decode (M<=2) to it with BLOCK_SIZE_M=1 and a swept-optimal BLOCK_SIZE_N=8/BLOCK_SIZE_K=128 (small N maximizes CU occupancy for the M=1 GEMV). Every other path (batched decode M>2, prefill, non-gfx900) is unchanged (GEMV_MODE defaults False -> tl.dot). Qwen3.5-35B-A3B-GPTQ-Int4, TP4, FULL_DECODE_ONLY graphs: c=1 decode 33.5 -> 47.1 tok/s (+40.6%, tpot 29.9 -> 21.2ms). c=8/c=32 unchanged (M>2 keeps tl.dot). Correctness: decode A/B vs the baseline tl.dot path (greedy, 128 tokens) is a 128/128 exact token match, chosen-token logprob delta mean 0.0019 / max 0.069 nats -- numerically equivalent (only fp accumulation order differs). Guarded by a VLLM_GFX900_MOE_GEMV kill-switch (default on; =0 forces tl.dot). The MoE analog of the dense int4 decode GEMV (#57). See BENCH_MOE_GEMV_GFX900.md for the profiling that justified it, the config sweep, and the correctness A/B. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EBCC2Amg37SsLxZ6G7GzkV
📝 WalkthroughWalkthroughAdds AMD gfx900 (Vega10) platform support: architecture detection and backend routing in ROCm platform code, a new int4 W4A16 GEMV kernel with dispatch integration, MoE fused-kernel GEMV mode, skinny GEMM fast paths, a CMake build update, extensive benchmark/profiling scripts, and multiple documentation files. Changesgfx900 Support
Estimated code review effort: 4 (Complex) | ~75 minutes Sequence Diagram(s)sequenceDiagram
participant TritonW4A16Gemm
participant OnGfx900
participant Gfx900GemvKernel
TritonW4A16Gemm->>OnGfx900: check architecture and shape
OnGfx900-->>TritonW4A16Gemm: gfx900 true, group_size ok, M<=8
TritonW4A16Gemm->>Gfx900GemvKernel: w4a16_gemv_gfx900(a, b_q, scales, qzeros)
Gfx900GemvKernel-->>TritonW4A16Gemm: fp16/bf16 output tensor
Possibly related issues
🚥 Pre-merge checks | ✅ 5✅ Passed checks (5 passed)
✨ Finishing Touches📝 Generate docstrings
🧪 Generate unit tests (beta)
⚔️ Resolve merge conflicts
Thanks for using CodeRabbit! It's free for OSS, and your support helps us grow. If you like it, consider giving us a shout-out. Comment |
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
There was a problem hiding this comment.
Actionable comments posted: 18
Caution
Some comments are outside the diff and can’t be posted inline due to platform limitations.
⚠️ Outside diff range comments (1)
vllm/platforms/rocm.py (1)
388-412: 🎯 Functional Correctness | 🟠 Major | 🏗️ Heavy liftHandle gfx900 before the MLA/sparse early returns.
With the new branch placed after
use_sparseanduse_mla, gfx900 can still selectROCM_AITER_MLA_SPARSEorROCM_AITER_MLAbefore reaching the portable Triton/TurboQuant fallback. Move or duplicate the gfx900 routing above those early returns, preserving supported MLA/sparse fallbacks or failing explicitly when unsupported.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@vllm/platforms/rocm.py` around lines 388 - 412, The gfx900 fallback in the ROCm backend selection is currently bypassed by the early use_sparse and use_mla returns. Update the backend routing in the attention selection logic so _ON_GFX900 is checked before those early exits, or explicitly gate the MLA/sparse paths for gfx900, ensuring the portable TRITON_ATTN/TURBOQUANT fallback is chosen instead of ROCM_AITER_MLA_SPARSE or ROCM_AITER_MLA when unsupported.
🧹 Nitpick comments (18)
bench_scripts/gemv_bench.py (1)
53-53: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winSplit multi-statement lines for readability and line-length compliance.
Ruff flags E702 on these
interleavechains (lines 53, 62, 104, 111), and each of these semicolon-joined lines also exceeds the repo's 88-character limit. Splitting onto separate statements fixes both.As per coding guidelines, "The line length limit for Python code is 88 characters."🔧 Example fix (applies to all 4 occurrences)
- ze = tl.interleave(qzv, qzv); ze = tl.interleave(ze, ze); ze = tl.interleave(ze, ze) + ze = tl.interleave(qzv, qzv) + ze = tl.interleave(ze, ze) + ze = tl.interleave(ze, ze)Also applies to: 62-62, 104-104, 111-111
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@bench_scripts/gemv_bench.py` at line 53, The chained tl.interleave statements in the gemv benchmark use semicolon-separated multi-statement lines, which triggers Ruff E702 and exceeds the 88-character limit. In the relevant benchmark code, split each interleave chain into separate statements on individual lines for all affected occurrences in the benchmark function(s) so the assignments remain readable and compliant. Use the existing variables like ze and any similar intermediates in the same pattern wherever these chained lines appear.Sources: Coding guidelines, Linters/SAST tools
bench_scripts/skinny_bench.py (1)
10-14: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valueCompressed one-liners reduce readability.
Ruff flags the semicolon/colon-compressed statements in
bench. Minor, benchmark-only script.♻️ Suggested cleanup
def bench(fn,*a,iters=50): - for _ in range(5): fn(*a) - torch.cuda.synchronize(); t=time.time() - for _ in range(iters): r=fn(*a) - torch.cuda.synchronize(); return (time.time()-t)/iters*1000, r + for _ in range(5): + fn(*a) + torch.cuda.synchronize() + t = time.time() + for _ in range(iters): + r = fn(*a) + torch.cuda.synchronize() + return (time.time() - t) / iters * 1000, r🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@bench_scripts/skinny_bench.py` around lines 10 - 14, The bench helper is using compressed one-line statements with semicolons, which Ruff flags and hurts readability. Refactor the bench function to use standard multi-line statements for the warmup loop, timing setup, benchmark loop, synchronization, and return value while keeping the same behavior; use the bench function name as the target for the cleanup.Source: Linters/SAST tools
bench_scripts/power_bench.py (1)
54-61: 🎯 Functional Correctness | 🔵 Trivial | 💤 Low valueZero-watt power reading treated as missing.
mean_w/peak_wuseif mean_w else None, so a legitimate0.0W reading (falsy) is treated the same as a missing sample, silently droppingtflops_per_w. Edge case only, unlikely for a running GPU under load.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@bench_scripts/power_bench.py` around lines 54 - 61, The power summary in power_bench.py treats 0.0 as missing because it uses truthiness checks on mean_w and peak_w. Update the JSON construction in the power reporting block to test explicitly for None rather than relying on if mean_w / if peak_w, so legitimate zero-watt readings are preserved and tflops_per_w is still computed when mean_w is 0.0. Use the existing mean_w and peak_w variables in the print(json.dumps(...)) section to apply the fix.bench_scripts/power_scan.sh (1)
9-9: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valueUnquoted glob/variable expansions flagged by shellcheck.
Lines 9, 15, and 28 have unquoted variable/glob expansions (
$GPU,$CAP) that could break with unexpected values.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@bench_scripts/power_scan.sh` at line 9, The power_scan.sh script has unquoted variable and glob expansions that shellcheck flags, especially around the PWR lookup and the GPU/CAP uses. Update the relevant expansions in the script to use proper quoting and safe parameter handling in the affected commands and conditionals, keeping the existing logic in place. Focus on the assignments and checks involving PWR, GPU, and CAP so unexpected values cannot break the globbing or word splitting behavior.Source: Linters/SAST tools
bench_scripts/parse_trace.py (1)
7-23: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valueRuff-flagged style issues (semicolons, colons).
Multiple E701/E702 violations across the
bucket()dispatch chain and the accumulation loop.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@bench_scripts/parse_trace.py` around lines 7 - 23, The trace parsing logic in bucket() and the event accumulation loop is using multiple statements on single lines, which triggers the Ruff E701/E702 style violations. Split the chained if/return branches and the cat/cnt/tot updates in the loop into separate lines with standard Python block syntax, preserving the existing bucket names and event filtering behavior in parse_trace.py.Source: Linters/SAST tools
bench_scripts/ab_decode.py (2)
21-25: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valueRuff-flagged style issues (semicolons, unused loop var).
Static analysis flags multiple-statements-per-line (E702) on lines 21, 23, 24 and an unused loop variable
i(B007) on line 22.🧹 Suggested cleanup
- for i in range(N): - t0=time.time(); o=llm.generate([prompt], sp); dt=time.time()-t0 - n=len(o[0].outputs[0].token_ids); tot+=dt; toks+=n + for _ in range(N): + t0 = time.time() + o = llm.generate([prompt], sp) + dt = time.time() - t0 + n = len(o[0].outputs[0].token_ids) + tot += dt + toks += n🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@bench_scripts/ab_decode.py` around lines 21 - 25, The benchmark loop in ab_decode.py has Ruff style issues from chaining multiple statements on one line and an unused loop index in the N-iteration block. Split the combined statements in the top-level timing loop and final print into separate lines, and change the for-loop in the decode benchmark so it does not bind an unused variable. Keep the existing behavior in the main timing logic around llm.generate, tot, toks, and the summary print.Source: Linters/SAST tools
7-7: 🎯 Functional Correctness | 🔵 Trivial | 💤 Low valueNo validation on required CLI arg.
sys.argv[1]is accessed unconditionally; running the script without an argument raises an unhandledIndexErrorinstead of a clear usage message.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@bench_scripts/ab_decode.py` at line 7, The CLI argument handling in the script is missing validation because MODE is assigned directly from sys.argv[1]. Add an explicit check before reading the argument, and if it is absent, print a clear usage/help message and exit cleanly instead of allowing an IndexError. Update the top-level argument parsing around MODE so the script fails gracefully when invoked without the required mode.bench_scripts/trace_mm.py (1)
16-16: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valueUnused loop variable
tid.Ruff B007: rename to
_tidsince it's unused within the loop body.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@bench_scripts/trace_mm.py` at line 16, The loop in trace_mm.py uses tid in the by_tid.items() iteration but never references it in the body, so update the loop variable to a throwaway name like _tid to satisfy Ruff B007. Keep the change localized to the loop that iterates over by_tid.items() and preserve the existing behavior of the loop body.Source: Linters/SAST tools
bench_scripts/awq_e2e.py (1)
16-19: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valueRuff-flagged style issues (semicolons, colon).
E702 on line 16, E701 on line 19.
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@bench_scripts/awq_e2e.py` around lines 16 - 19, The benchmark loop in main() has Ruff style violations from multiple statements on one line and an inline if __name__ guard. Split the chained statements in the main() body into separate lines, and rewrite the if __name__=="__main__": main() entrypoint as a standard multi-line block so the code in awq_e2e.py complies with E702 and E701.Source: Linters/SAST tools
bench_scripts/awq_sweep.py (1)
15-18: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valueRuff-flagged style issues (semicolons, colon).
E702 on line 15, E701 on line 18.
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@bench_scripts/awq_sweep.py` around lines 15 - 18, The AWQ sweep script has Ruff style violations from multiple statements on one line and a cramped main guard. In awq_sweep.py, split the combined timing/generation statements in main() into separate lines, and format the __name__ == "__main__" guard with standard spacing and line structure so it no longer triggers E702/E701 while keeping the existing logic in main() and llm.generate() unchanged.Source: Linters/SAST tools
bench_scripts/batch_sweep.py (2)
1-21: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚖️ Poor tradeoffBoilerplate duplicated across ~20 bench scripts.
The env-setdefault → LLM construction → warmup → timed generate pattern is repeated nearly verbatim across
batch_sweep.py,coder_smoke.py,coder_tput.py,combo_dec.py,graph_ab.py, and others in this cohort. Extracting a small shared harness (e.g.,bench_scripts/_common.pywith abuild_llm/warmup_and_timehelper) would reduce drift risk as configs evolve.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@bench_scripts/batch_sweep.py` around lines 1 - 21, The benchmark script repeats the same environment setup, LLM construction, warmup, and timed generation flow as other scripts in this cohort; factor that shared sequence into a small helper in a common module such as bench_scripts/_common.py. Move the repeated logic from main() into reusable functions like build_llm and warmup_and_time, then update batch_sweep.py to call those helpers while keeping only its batch-specific prompt and reporting loop.
18-20: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valueMisleading comment: no actual prefill subtraction happens.
The comment claims prefill time is subtracted via "step rate: total decode tok / wall", but
agg_tok_sandper_step_msare computed directly from wall-clockdtwithout isolating decode-only time. For a short prompt this bias is negligible, but the comment could mislead future readers into trusting per-step numbers as pure decode latency.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@bench_scripts/batch_sweep.py` around lines 18 - 20, The batch sweep summary in the print block is mislabeled: the comment claims prefill is being subtracted, but the metrics in the print statement are still based on total wall-clock time. Update the comment to accurately describe that `agg_tok_s` and `per_step_ms` are derived directly from `dt`, or adjust the calculation in the batch sweep logic if you intend to report decode-only latency. Use the existing `B`, `outtok`, `dt`, `agg_tok_s`, and `per_step_ms` symbols to keep the output and comment consistent.bench_scripts/moe35_prof.py (1)
8-9: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winHardcoded developer-specific path breaks portability.
/home/larkinwc/moe35_profonly works on the original author's machine;os.makedirswill fail for any other user whose home directory differs.♻️ Proposed fix
-PROF = "/home/larkinwc/moe35_prof" +PROF = os.environ.get("VLLM_PROF_DIR", os.path.join(tempfile.gettempdir(), "moe35_prof"))(remember to
import tempfile)🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@bench_scripts/moe35_prof.py` around lines 8 - 9, The PROF directory path is hardcoded to a developer-specific home folder, which breaks portability. Update the initialization in moe35_prof.py to derive the directory from a portable location instead of a fixed /home/... path, using tempfile as suggested, and keep the os.makedirs(PROF, exist_ok=True) call working with the new path.bench_scripts/pp_bench.py (1)
1-30: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winEntire file uses semicolon/colon-packed one-liners; ruff flags multiple E701/E702 violations.
Lines 19, 26, and 30 combine multiple statements per line, all flagged by ruff. This also pushes several lines well past the 88-character guideline. Given this repo enforces pre-commit/ruff, this will likely fail lint.
As per coding guidelines, "The line length limit for Python code is 88 characters."
🔧 Proposed fix (representative lines)
- t0=time.time(); llm=LLM(**kw); print("INIT[%s graphs=%s] %.0fs"%(layout,graphs,time.time()-t0),flush=True) + t0 = time.time() + llm = LLM(**kw) + print("INIT[%s graphs=%s] %.0fs" % (layout, graphs, time.time() - t0), flush=True) ... - t0=time.time(); o=llm.generate(prompts, sp); dt=time.time()-t0 + t0 = time.time() + o = llm.generate(prompts, sp) + dt = time.time() - t0 ... -if __name__=="__main__": main() +if __name__ == "__main__": + main()🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@bench_scripts/pp_bench.py` around lines 1 - 30, The benchmark script is packed with multiple statements on single lines and long inline expressions, which triggers ruff E701/E702 and violates the 88-character limit. Refactor the `main` function in `pp_bench.py` to split chained statements into separate lines, break long dictionary/LLM/SamplingParams constructions across lines, and keep each assignment or call readable without semicolons so the file passes lint and pre-commit.Sources: Coding guidelines, Linters/SAST tools
bench_scripts/moe35_quick.py (1)
10-27: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winDense one-liner style triggers ruff errors; also exceeds 88-char line limit.
Line 22 packs three statements with semicolons (
t0 = time.time(); out = ...; dt = ...), flagged by ruff (E702). Lines 10-14 also exceed the repo's 88-character line-length guideline. Since this repo runs pre-commit/ruff as part of its Python workflow, this will likely fail lint checks.As per coding guidelines, "The line length limit for Python code is 88 characters."
🔧 Proposed fix
- for _ in range(2): - t0 = time.time(); out = llm.generate(p, sp); dt = time.time() - t0 - n = len(out[0].outputs[0].token_ids) + for _ in range(2): + t0 = time.time() + out = llm.generate(p, sp) + dt = time.time() - t0 + n = len(out[0].outputs[0].token_ids)🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@bench_scripts/moe35_quick.py` around lines 10 - 27, The benchmark script uses dense one-liners that violate Ruff and the 88-character limit; refactor the setup and timing logic in moe35_quick.py to use one statement per line and wrap long argument lists/format strings. In particular, split the semicolon-packed timing line in the main loop into separate assignments around llm.generate, and break up the LLM and SamplingParams construction plus the cfg/print formatting so they stay within the project’s line-length rule.Sources: Coding guidelines, Linters/SAST tools
bench_scripts/ppl.py (1)
1-34: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winDense one-liner style triggers 7 ruff E701/E702 findings.
Lines 12, 16, 20, 23, 30, 34 combine multiple statements/conditionals per line. Given the repo's pre-commit/ruff usage, splitting these will avoid lint failures and improve readability.
As per coding guidelines, "The line length limit for Python code is 88 characters."
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@bench_scripts/ppl.py` around lines 1 - 34, The dense one-line style in main() is triggering ruff E701/E702 and may also exceed the 88-character limit, so rewrite the setup and loop logic into separate statements and multiline conditionals for clarity. Split the combined imports/environment setup, the preset handling around kw, the W/stride/N assignment, the prompt-building loop, the token accumulation loop, and the final print/entrypoint check into standard multi-line Python. Keep the behavior unchanged while making the main function and __name__ == "__main__" block lint-clean.Sources: Coding guidelines, Linters/SAST tools
gfx900_bench_pp.sh (1)
62-65: 🩺 Stability & Availability | 🔵 Trivial | 💤 Low valueBroad
pkill -f "vllm serve"could kill unrelated processes.On a shared multi-GPU host, this could terminate other users' or scripts'
vllm serveinstances rather than just this run's server.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@gfx900_bench_pp.sh` around lines 62 - 65, The shutdown logic in gfx900_bench_pp.sh is too broad because pkill -f "vllm serve" can terminate unrelated vllm instances on the same host. Update the cleanup flow around SERVER_PID and the server launch so it only targets the process started by this script, using the stored PID or a more specific process match tied to this run instead of a generic command-line pattern.bench_scripts/coder_prof.py (1)
3-4: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valueHardcoded user-specific path reduces portability.
/home/larkinwc/coder_profties the script to one developer's home directory; same pattern repeats inprof_decode.py/parse_any.pydefaults. Consider sourcing this from an env var or CLI arg with a fallback.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@bench_scripts/coder_prof.py` around lines 3 - 4, The PROF default is hardcoded to a single developer-specific home directory, which makes the bench scripts less portable. Update the default path setup in coder_prof.py and the matching defaults in prof_decode.py and parse_any.py to derive the location from an environment variable or CLI argument, with a sensible fallback when none is provided. Keep the existing PROF symbol as the central path value so the scripts share the same configurable source.
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
Inline comments:
In `@bench_scripts/awq_e2e.py`:
- Around line 13-18: The benchmark print in awq_e2e.py hardcodes the token count
instead of using the SamplingParams configuration. Update the calculation in the
loop that uses sp so ms_per_tok derives from sp.max_tokens (or an equivalent
shared constant) rather than a literal 128, matching the pattern used in
awq_sweep.py and keeping the metric consistent if max_tokens changes.
In `@bench_scripts/combo_dec.py`:
- Around line 5-10: The benchmark configuration in combo_dec.py does not match
the documented deployment path for this model/KV setup. Update the LLM
initialization in the combo benchmark to use the same recommended knobs as
GFX900_RECOMMENDED.md, including the documented tensor parallel size,
max_model_len, and max_num_seqs settings, or clearly separate it as a smaller
smoke benchmark. Keep the changes centered on the LLM(...) construction and the
preset handling so the benchmark result reflects the production recommendation.
In `@bench_scripts/corr_gen.py`:
- Line 25: The default output path in corr_gen.py is hardcoded to a personal
home directory, which should be replaced with a portable fallback. Update the
logic around the OUT environment lookup in the script’s main output assignment
to use a repo-relative location or a tempfile-based default instead of the
developer-specific /home/... path, so the benchmark script works across machines
without personal paths baked in.
In `@bench_scripts/moe35_prof.py`:
- Around line 36-38: The decode benchmark config is only capturing one cudagraph
size, so the runs for the other sweep values will fall back to eager execution.
Update the `CompilationConfig` in `moe35_prof.py` so `cudagraph_capture_sizes`
matches the benchmark sweep used by `bench`, and make sure the
`FULL_DECODE_ONLY` setup in that block includes all decode sizes being measured,
not just `1`.
In `@bench_scripts/parse_any.py`:
- Around line 3-5: The trace selection logic in parse_any.py can crash with an
IndexError when the glob finds no files. Update the d_dir/f lookup flow to check
the glob result before indexing, and in the parse_any script emit a clear error
or exit message when no matching *.pt.trace.json.gz file is found so the failure
is explicit and easy to diagnose.
In `@bench_scripts/parse_trace.py`:
- Around line 3-4: The trace parsing script currently hardcodes a user-specific
absolute glob path and assumes at least one match exists, so update the entry
logic in parse_trace.py to accept the trace file path from input like
trace_mm.py does via sys.argv[1] instead of using the fixed glob. Add a small
guard before loading the file so the script fails gracefully when no trace path
is provided or when the expected file cannot be found, and keep the loading
logic in the existing json.load/gzip.open flow.
In `@bench_scripts/power_scan.sh`:
- Around line 17-19: The power benchmark invocation in power_scan.sh is using a
predictable /tmp/power_bench.py path, which is vulnerable to symlink/TOCTOU
manipulation. Update the script’s benchmark launch to avoid /tmp entirely by
referencing the checked-in power_bench.py from the repo or by copying it to a
mktemp-generated private location before calling the python command. Keep the
fix localized to the shell logic around the benchmark run so the existing
HIP_VISIBLE_DEVICES, --power-path, and output redirection behavior stays
unchanged.
In `@bench_scripts/ppl.py`:
- Around line 16-21: The inline comment in the prompt-window setup is stale and
no longer matches the values used in the `W`, `stride`, and `N` loop. Update the
comment in `bench_scripts/ppl.py` near the `prompts` construction to accurately
describe the actual window size, stride, count, and total evaluated tokens, or
adjust the constants if the intended benchmark is still 20 windows of 2048.
- Line 7: The dataset load in the module-level `ds = pd.read_parquet(...)` uses
a hardcoded personal absolute path, which breaks portability. Replace the fixed
path with a configurable input for `ppl.py`, such as a CLI argument or
environment variable, and have the existing `pd.read_parquet` call read from
that value so the script works on other machines.
In `@bench_scripts/run_topo.sh`:
- Around line 7-9: The exit marker in run_topo.sh is reporting the grep status
instead of the benchmark status because the timeout-to-grep pipeline masks the
real exit code. Update the script so the pipeline’s exit status reflects the
timeout/python command, using the existing run_topo.sh flow around the
topo_bench.py invocation and the final echo marker. Keep the grep filtering
behavior, but ensure the status printed by “### EXIT $? ###” comes from the
benchmark process, not from grep.
In `@GFX900_RECOMMENDED.md`:
- Around line 33-44: Clarify the benchmark context for the “9.97 tok/s ≈
standalone AWQ 10.46” claim by tying it to the exact settings used in
bench_scripts/combo_dec.py, since its tensor_parallel_size and max_model_len
differ from the recommended TP4/16384 path. Update the markdown around the
comparison to explicitly name the config behind both numbers, or adjust
combo_dec.py to run the same TP4/16384 settings so the “compose cleanly”
statement matches the recommended setup. Refer to bench_scripts/combo_dec.py and
the decode comparison text in GFX900_RECOMMENDED.md when making the change.
In `@PERF_GFX900.md`:
- Around line 9-12: The microbenchmark references in the PERF_GFX900 doc are
stale and still point to /tmp script locations. Update the description to use
the actual shipped paths under bench_scripts for gpu_peak.py, power_bench.py,
and power_scan.sh, matching the naming used elsewhere in the doc (for example
the Repro sections for realistic.py, graph_ab.py, and pp_bench.py).
- Around line 130-133: The Reproduce block still points to stale /tmp/ script
locations; update the commands to use the bench_scripts/ versions instead.
Adjust the repro examples in PERF_GFX900.md so they reference the actual script
names/paths introduced in this PR, matching the locations used by run_peak.sh
and power_scan.sh.
In `@vllm/model_executor/kernels/linear/mixed_precision/gfx900_w4a16_gemv.py`:
- Around line 61-71: Ruff is flagging chained statements in the GEMV kernel: the
assignments to ze and bt in gfx900_w4a16_gemv should be split into separate
lines instead of using multiple statements per line. Update the relevant logic
in the mixed-precision GEMV path so each tl.interleave step is assigned
independently, while keeping the Python line length within the 88-character
guideline and preserving the existing variable names ze and bt for easy
location.
In `@vllm/model_executor/kernels/linear/mixed_precision/triton_w4a16.py`:
- Around line 223-236: The gfx900 GEMV dispatch in w4a16 should be split from
the MI100 kill switch so `_mi100_w4a16_disabled()` does not block `on_gfx900()`
routing in `vllm/model_executor/kernels/linear/mixed_precision/triton_w4a16.py`.
Update the ROCm branch around the `on_gfx900()` and `on_mi100()` checks so
gfx900 small-M AWQ decode can still return `w4a16_gemv_gfx900` when its own
conditions match, while only the MI100-specific path remains guarded by
`_mi100_w4a16_disabled()`. Use the existing symbols
`current_platform.is_rocm()`, `on_gfx900()`, `on_mi100()`, and
`_mi100_w4a16_disabled()` to keep the dispatch logic explicit.
In `@vllm/model_executor/layers/fused_moe/fused_moe.py`:
- Around line 1256-1270: The ROCm helper import inside get_default_config is
unconditional and can run on non-ROCm paths; move the import of on_gfx900 behind
a current_platform.is_rocm() check so the M <= 20 branch only touches
vllm.platforms.rocm when the platform is actually ROCm. Keep the existing
gfx900-specific config selection logic intact, but gate the ROCm-only import and
environment-based GEMV override together in that ROCm branch.
In `@vllm/model_executor/layers/utils.py`:
- Around line 181-207: The gfx900 skinny-GEMM fast paths in the utils function
should only run when the input and weight dtypes already match, since the
regular fallback handles mismatches by casting x first. Add a dtype equality
guard alongside the existing checks in both branches before the manual reduction
path and the ops.LLMM1 call, using the same x and weight symbols in this helper,
so mixed-dtype cases continue through the safe fallback instead of hitting
unsupported promotion behavior.
In `@vllm/platforms/rocm.py`:
- Around line 309-312: The on_gfx900 predicate docstring is not using the
required Google-style format. Update the docstring in on_gfx900 to use
Google-style sections and include a Returns: section that describes the boolean
result, keeping the existing detection summary concise. Ensure the formatting
matches the style used across vllm/platforms/rocm.py and other Python
docstrings.
---
Outside diff comments:
In `@vllm/platforms/rocm.py`:
- Around line 388-412: The gfx900 fallback in the ROCm backend selection is
currently bypassed by the early use_sparse and use_mla returns. Update the
backend routing in the attention selection logic so _ON_GFX900 is checked before
those early exits, or explicitly gate the MLA/sparse paths for gfx900, ensuring
the portable TRITON_ATTN/TURBOQUANT fallback is chosen instead of
ROCM_AITER_MLA_SPARSE or ROCM_AITER_MLA when unsupported.
---
Nitpick comments:
In `@bench_scripts/ab_decode.py`:
- Around line 21-25: The benchmark loop in ab_decode.py has Ruff style issues
from chaining multiple statements on one line and an unused loop index in the
N-iteration block. Split the combined statements in the top-level timing loop
and final print into separate lines, and change the for-loop in the decode
benchmark so it does not bind an unused variable. Keep the existing behavior in
the main timing logic around llm.generate, tot, toks, and the summary print.
- Line 7: The CLI argument handling in the script is missing validation because
MODE is assigned directly from sys.argv[1]. Add an explicit check before reading
the argument, and if it is absent, print a clear usage/help message and exit
cleanly instead of allowing an IndexError. Update the top-level argument parsing
around MODE so the script fails gracefully when invoked without the required
mode.
In `@bench_scripts/awq_e2e.py`:
- Around line 16-19: The benchmark loop in main() has Ruff style violations from
multiple statements on one line and an inline if __name__ guard. Split the
chained statements in the main() body into separate lines, and rewrite the if
__name__=="__main__": main() entrypoint as a standard multi-line block so the
code in awq_e2e.py complies with E702 and E701.
In `@bench_scripts/awq_sweep.py`:
- Around line 15-18: The AWQ sweep script has Ruff style violations from
multiple statements on one line and a cramped main guard. In awq_sweep.py, split
the combined timing/generation statements in main() into separate lines, and
format the __name__ == "__main__" guard with standard spacing and line structure
so it no longer triggers E702/E701 while keeping the existing logic in main()
and llm.generate() unchanged.
In `@bench_scripts/batch_sweep.py`:
- Around line 1-21: The benchmark script repeats the same environment setup, LLM
construction, warmup, and timed generation flow as other scripts in this cohort;
factor that shared sequence into a small helper in a common module such as
bench_scripts/_common.py. Move the repeated logic from main() into reusable
functions like build_llm and warmup_and_time, then update batch_sweep.py to call
those helpers while keeping only its batch-specific prompt and reporting loop.
- Around line 18-20: The batch sweep summary in the print block is mislabeled:
the comment claims prefill is being subtracted, but the metrics in the print
statement are still based on total wall-clock time. Update the comment to
accurately describe that `agg_tok_s` and `per_step_ms` are derived directly from
`dt`, or adjust the calculation in the batch sweep logic if you intend to report
decode-only latency. Use the existing `B`, `outtok`, `dt`, `agg_tok_s`, and
`per_step_ms` symbols to keep the output and comment consistent.
In `@bench_scripts/coder_prof.py`:
- Around line 3-4: The PROF default is hardcoded to a single developer-specific
home directory, which makes the bench scripts less portable. Update the default
path setup in coder_prof.py and the matching defaults in prof_decode.py and
parse_any.py to derive the location from an environment variable or CLI
argument, with a sensible fallback when none is provided. Keep the existing PROF
symbol as the central path value so the scripts share the same configurable
source.
In `@bench_scripts/gemv_bench.py`:
- Line 53: The chained tl.interleave statements in the gemv benchmark use
semicolon-separated multi-statement lines, which triggers Ruff E702 and exceeds
the 88-character limit. In the relevant benchmark code, split each interleave
chain into separate statements on individual lines for all affected occurrences
in the benchmark function(s) so the assignments remain readable and compliant.
Use the existing variables like ze and any similar intermediates in the same
pattern wherever these chained lines appear.
In `@bench_scripts/moe35_prof.py`:
- Around line 8-9: The PROF directory path is hardcoded to a developer-specific
home folder, which breaks portability. Update the initialization in
moe35_prof.py to derive the directory from a portable location instead of a
fixed /home/... path, using tempfile as suggested, and keep the
os.makedirs(PROF, exist_ok=True) call working with the new path.
In `@bench_scripts/moe35_quick.py`:
- Around line 10-27: The benchmark script uses dense one-liners that violate
Ruff and the 88-character limit; refactor the setup and timing logic in
moe35_quick.py to use one statement per line and wrap long argument lists/format
strings. In particular, split the semicolon-packed timing line in the main loop
into separate assignments around llm.generate, and break up the LLM and
SamplingParams construction plus the cfg/print formatting so they stay within
the project’s line-length rule.
In `@bench_scripts/parse_trace.py`:
- Around line 7-23: The trace parsing logic in bucket() and the event
accumulation loop is using multiple statements on single lines, which triggers
the Ruff E701/E702 style violations. Split the chained if/return branches and
the cat/cnt/tot updates in the loop into separate lines with standard Python
block syntax, preserving the existing bucket names and event filtering behavior
in parse_trace.py.
In `@bench_scripts/power_bench.py`:
- Around line 54-61: The power summary in power_bench.py treats 0.0 as missing
because it uses truthiness checks on mean_w and peak_w. Update the JSON
construction in the power reporting block to test explicitly for None rather
than relying on if mean_w / if peak_w, so legitimate zero-watt readings are
preserved and tflops_per_w is still computed when mean_w is 0.0. Use the
existing mean_w and peak_w variables in the print(json.dumps(...)) section to
apply the fix.
In `@bench_scripts/power_scan.sh`:
- Line 9: The power_scan.sh script has unquoted variable and glob expansions
that shellcheck flags, especially around the PWR lookup and the GPU/CAP uses.
Update the relevant expansions in the script to use proper quoting and safe
parameter handling in the affected commands and conditionals, keeping the
existing logic in place. Focus on the assignments and checks involving PWR, GPU,
and CAP so unexpected values cannot break the globbing or word splitting
behavior.
In `@bench_scripts/pp_bench.py`:
- Around line 1-30: The benchmark script is packed with multiple statements on
single lines and long inline expressions, which triggers ruff E701/E702 and
violates the 88-character limit. Refactor the `main` function in `pp_bench.py`
to split chained statements into separate lines, break long
dictionary/LLM/SamplingParams constructions across lines, and keep each
assignment or call readable without semicolons so the file passes lint and
pre-commit.
In `@bench_scripts/ppl.py`:
- Around line 1-34: The dense one-line style in main() is triggering ruff
E701/E702 and may also exceed the 88-character limit, so rewrite the setup and
loop logic into separate statements and multiline conditionals for clarity.
Split the combined imports/environment setup, the preset handling around kw, the
W/stride/N assignment, the prompt-building loop, the token accumulation loop,
and the final print/entrypoint check into standard multi-line Python. Keep the
behavior unchanged while making the main function and __name__ == "__main__"
block lint-clean.
In `@bench_scripts/skinny_bench.py`:
- Around line 10-14: The bench helper is using compressed one-line statements
with semicolons, which Ruff flags and hurts readability. Refactor the bench
function to use standard multi-line statements for the warmup loop, timing
setup, benchmark loop, synchronization, and return value while keeping the same
behavior; use the bench function name as the target for the cleanup.
In `@bench_scripts/trace_mm.py`:
- Line 16: The loop in trace_mm.py uses tid in the by_tid.items() iteration but
never references it in the body, so update the loop variable to a throwaway name
like _tid to satisfy Ruff B007. Keep the change localized to the loop that
iterates over by_tid.items() and preserve the existing behavior of the loop
body.
In `@gfx900_bench_pp.sh`:
- Around line 62-65: The shutdown logic in gfx900_bench_pp.sh is too broad
because pkill -f "vllm serve" can terminate unrelated vllm instances on the same
host. Update the cleanup flow around SERVER_PID and the server launch so it only
targets the process started by this script, using the stored PID or a more
specific process match tied to this run instead of a generic command-line
pattern.
🪄 Autofix (Beta)
Fix all unresolved CodeRabbit comments on this PR:
- Push a commit to this branch (recommended)
- Create a new PR with the fixes
ℹ️ Review info
⚙️ Run configuration
Configuration used: defaults
Review profile: CHILL
Plan: Pro Plus
Run ID: c96fac53-56e2-4362-98a4-a8476596544b
📒 Files selected for processing (53)
BENCH_DECODE_PROFILE.mdBENCH_GFX900.mdBENCH_MOE_GEMV_GFX900.mdCMakeLists.txtGFX900_RECOMMENDED.mdGFX900_SETUP.mdPERF_GFX900.mdPERF_LMHEAD_GFX900.mdbench_scripts/ab_decode.pybench_scripts/awq_e2e.pybench_scripts/awq_sweep.pybench_scripts/awq_test.pybench_scripts/batch_sweep.pybench_scripts/coder_prof.pybench_scripts/coder_smoke.pybench_scripts/coder_tput.pybench_scripts/combo_dec.pybench_scripts/corr_cmp.pybench_scripts/corr_gen.pybench_scripts/gate_bench.pybench_scripts/gemv_bench.pybench_scripts/gemv_sweep.shbench_scripts/gpu_peak.pybench_scripts/graph_ab.pybench_scripts/graph_combo.pybench_scripts/graph_safe.pybench_scripts/lmhead_e2e.pybench_scripts/mb_expert_gemv.pybench_scripts/moe35_prof.pybench_scripts/moe35_quick.pybench_scripts/moe_topo.pybench_scripts/parse_any.pybench_scripts/parse_trace.pybench_scripts/power_bench.pybench_scripts/power_scan.shbench_scripts/pp_bench.pybench_scripts/ppl.pybench_scripts/prof_decode.pybench_scripts/realistic.pybench_scripts/run_peak.shbench_scripts/run_topo.shbench_scripts/skinny_bench.pybench_scripts/test_gfx900_gemv.pybench_scripts/topo_bench.pybench_scripts/tq_measure.pybench_scripts/tq_smoke.pybench_scripts/trace_mm.pygfx900_bench_pp.shvllm/model_executor/kernels/linear/mixed_precision/gfx900_w4a16_gemv.pyvllm/model_executor/kernels/linear/mixed_precision/triton_w4a16.pyvllm/model_executor/layers/fused_moe/fused_moe.pyvllm/model_executor/layers/utils.pyvllm/platforms/rocm.py
Comment on lines
+13
to
+18
| sp=SamplingParams(max_tokens=128,temperature=0,ignore_eos=True) | ||
| llm.generate([prompt], SamplingParams(max_tokens=8,temperature=0,ignore_eos=True)) | ||
| for B in [1,8,32]: | ||
| t0=time.time(); oo=llm.generate([prompt]*B, sp); dt=time.time()-t0 | ||
| ot=sum(len(x.outputs[0].token_ids) for x in oo) | ||
| print("B=%2d agg_tok_s=%.2f ms_per_tok=%.1f"%(B, ot/dt, 1000*dt/128), flush=True) |
There was a problem hiding this comment.
🎯 Functional Correctness | 🟡 Minor | ⚡ Quick win
Hardcoded 128 duplicates sp's max_tokens.
ms_per_tok is computed as 1000*dt/128, hardcoding a value that's already defined via sp=SamplingParams(max_tokens=128,...) on line 13. If max_tokens is changed later without updating the literal, the printed metric silently becomes wrong. awq_sweep.py avoids this by referencing the GEN variable directly in its analogous print statement.
🔧 Suggested fix
-sp=SamplingParams(max_tokens=128,temperature=0,ignore_eos=True)
+GEN=128
+sp=SamplingParams(max_tokens=GEN,temperature=0,ignore_eos=True)
llm.generate([prompt], SamplingParams(max_tokens=8,temperature=0,ignore_eos=True))
for B in [1,8,32]:
t0=time.time(); oo=llm.generate([prompt]*B, sp); dt=time.time()-t0
ot=sum(len(x.outputs[0].token_ids) for x in oo)
- print("B=%2d agg_tok_s=%.2f ms_per_tok=%.1f"%(B, ot/dt, 1000*dt/128), flush=True)
+ print("B=%2d agg_tok_s=%.2f ms_per_tok=%.1f"%(B, ot/dt, 1000*dt/GEN), flush=True)📝 Committable suggestion
‼️ IMPORTANT
Carefully review the code before committing. Ensure that it accurately replaces the highlighted code, contains no missing lines, and has no issues with indentation. Thoroughly test & benchmark the code to ensure it meets the requirements.
Suggested change
| sp=SamplingParams(max_tokens=128,temperature=0,ignore_eos=True) | |
| llm.generate([prompt], SamplingParams(max_tokens=8,temperature=0,ignore_eos=True)) | |
| for B in [1,8,32]: | |
| t0=time.time(); oo=llm.generate([prompt]*B, sp); dt=time.time()-t0 | |
| ot=sum(len(x.outputs[0].token_ids) for x in oo) | |
| print("B=%2d agg_tok_s=%.2f ms_per_tok=%.1f"%(B, ot/dt, 1000*dt/128), flush=True) | |
| GEN=128 | |
| sp=SamplingParams(max_tokens=GEN,temperature=0,ignore_eos=True) | |
| llm.generate([prompt], SamplingParams(max_tokens=8,temperature=0,ignore_eos=True)) | |
| for B in [1,8,32]: | |
| t0=time.time(); oo=llm.generate([prompt]*B, sp); dt=time.time()-t0 | |
| ot=sum(len(x.outputs[0].token_ids) for x in oo) | |
| print("B=%2d agg_tok_s=%.2f ms_per_tok=%.1f"%(B, ot/dt, 1000*dt/GEN), flush=True) |
🧰 Tools
🪛 Ruff (0.15.20)
[error] 16-16: Multiple statements on one line (semicolon)
(E702)
[error] 16-16: Multiple statements on one line (semicolon)
(E702)
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
In `@bench_scripts/awq_e2e.py` around lines 13 - 18, The benchmark print in
awq_e2e.py hardcodes the token count instead of using the SamplingParams
configuration. Update the calculation in the loop that uses sp so ms_per_tok
derives from sp.max_tokens (or an equivalent shared constant) rather than a
literal 128, matching the pattern used in awq_sweep.py and keeping the metric
consistent if max_tokens changes.
Comment on lines
+5
to
+10
| preset = sys.argv[1] | ||
| kw = dict(model="QuantTrio/Qwen3.5-9B-AWQ", tensor_parallel_size=2, dtype="float16", | ||
| enforce_eager=True, gpu_memory_utilization=0.85, max_model_len=4096, | ||
| **{"language_model_only": True}) | ||
| if preset!="auto": kw["kv_cache_dtype"]=preset | ||
| llm=LLM(**kw) |
There was a problem hiding this comment.
🗄️ Data Integrity & Integration | 🟡 Minor | ⚡ Quick win
🧩 Analysis chain
🏁 Script executed:
#!/bin/bash
set -euo pipefail
printf '\n== bench_scripts/combo_dec.py ==\n'
git ls-files bench_scripts/combo_dec.py GFX900_RECOMMENDED.md
printf '\n-- combo_dec.py (relevant lines) --\n'
cat -n bench_scripts/combo_dec.py | sed -n '1,80p'
printf '\n== GFX900_RECOMMENDED.md ==\n'
cat -n GFX900_RECOMMENDED.md | sed -n '1,220p'
printf '\n== search for QuantTrio/Qwen3.5-9B-AWQ and turboquant_k8v4 ==\n'
rg -n "QuantTrio/Qwen3\.5-9B-AWQ|turboquant_k8v4|tensor_parallel_size|max_model_len|max_num_seqs|9\.97 tok/s|10\.46 tok/s" bench_scripts GFX900_RECOMMENDED.mdRepository: larkinwc/vllm-gfx908
Length of output: 21846
🏁 Script executed:
#!/bin/bash
set -euo pipefail
sed -n '220,240p' GFX900_RECOMMENDED.md | cat -n
sed -n '1,80p' bench_scripts/graph_combo.py | cat -n
sed -n '1,80p' bench_scripts/awq_e2e.py | cat -nRepository: larkinwc/vllm-gfx908
Length of output: 3646
Align the benchmark with the documented deployment path. bench_scripts/combo_dec.py uses TP2/4k and omits max_num_seqs, while GFX900_RECOMMENDED.md recommends TP4/16k/32 for this exact model+KV combo. If the 9.97 tok/s result is meant to support the production recommendation, the harness should use the same knobs; otherwise call it a smaller smoke benchmark.
🧰 Tools
🪛 Ruff (0.15.20)
[error] 9-9: Multiple statements on one line (colon)
(E701)
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
In `@bench_scripts/combo_dec.py` around lines 5 - 10, The benchmark configuration
in combo_dec.py does not match the documented deployment path for this model/KV
setup. Update the LLM initialization in the combo benchmark to use the same
recommended knobs as GFX900_RECOMMENDED.md, including the documented tensor
parallel size, max_model_len, and max_num_seqs settings, or clearly separate it
as a smaller smoke benchmark. Keep the changes centered on the LLM(...)
construction and the preset handling so the benchmark result reflects the
production recommendation.
| for tid, lp in zip(o.token_ids, o.logprobs): | ||
| # lp is {token_id: Logprob}; take the chosen token's logprob | ||
| seq.append([int(tid), float(lp[tid].logprob)]) | ||
| out = os.environ.get("OUT", "/home/larkinwc/corr_out.json") |
There was a problem hiding this comment.
🔒 Security & Privacy | 🟡 Minor | ⚡ Quick win
Hardcoded developer home-directory path as default output.
The fallback path embeds a specific developer's username (/home/larkinwc/...). Even for a dev-only benchmark script, avoid hardcoding a personal path into the repo; prefer a repo-relative or tempfile-based default.
🛡️ Suggested fix
- out = os.environ.get("OUT", "/home/larkinwc/corr_out.json")
+ out = os.environ.get("OUT", os.path.join(os.getcwd(), "corr_out.json"))📝 Committable suggestion
‼️ IMPORTANT
Carefully review the code before committing. Ensure that it accurately replaces the highlighted code, contains no missing lines, and has no issues with indentation. Thoroughly test & benchmark the code to ensure it meets the requirements.
Suggested change
| out = os.environ.get("OUT", "/home/larkinwc/corr_out.json") | |
| out = os.environ.get("OUT", os.path.join(os.getcwd(), "corr_out.json")) |
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
In `@bench_scripts/corr_gen.py` at line 25, The default output path in corr_gen.py
is hardcoded to a personal home directory, which should be replaced with a
portable fallback. Update the logic around the OUT environment lookup in the
script’s main output assignment to use a repo-relative location or a
tempfile-based default instead of the developer-specific /home/... path, so the
benchmark script works across machines without personal paths baked in.
Comment on lines
+36
to
+38
| compilation_config=CompilationConfig(mode=CompilationMode.NONE, | ||
| cudagraph_mode=CUDAGraphMode.FULL_DECODE_ONLY, cudagraph_capture_sizes=[1]), | ||
| **common) |
There was a problem hiding this comment.
🎯 Functional Correctness | 🟡 Minor | ⚡ Quick win
🧩 Analysis chain
🌐 Web query:
vllm CUDAGraphMode FULL_DECODE_ONLY behavior for batch sizes not in cudagraph_capture_sizes
💡 Result:
In vLLM's FULL_DECODE_ONLY CUDA graph mode, the behavior for batch sizes not explicitly present in cudagraph_capture_sizes depends on whether the batch size is smaller or larger than the largest configured capture size [1]. If the current batch size is smaller than or equal to one of the cudagraph_capture_sizes, vLLM will pad the batch to the closest (smallest sufficient) captured CUDA graph size and execute the graph [1][2]. If the current batch size is larger than the largest configured cudagraph_capture_size, the system cannot utilize a pre-captured graph and will fall back to eager mode execution [1][3]. Because FULL_DECODE_ONLY specifically captures graphs for uniform-decode batches (and none for prefill or mixed batches) [4][5], failure to fall into a captured size range results in a loss of the performance benefits provided by CUDA graphs for those decode operations [3]. Users should ensure their cudagraph_capture_sizes configuration accounts for their expected workload requirements to avoid this fallback, as vLLM will not dynamically capture new graphs at runtime for sizes outside the initial configuration [1][3].
Citations:
- 1: https://docs.vllm.ai/en/stable/api/vllm/config/vllm/
- 2: https://discuss.vllm.ai/t/why-is-cuda-graph-capture-sizes-limited-by-max-num-seqs/954
- 3: [Bugfix] Fix max_num_batched_token not captured in cuda graph vllm-project/vllm#40734
- 4: https://docs.vllm.ai/en/stable/design/cuda_graphs/
- 5: [Core] Allow full cudagraph with separate attention routines and orthogonal to compilation, add support for FA2 and FlashInfer vllm-project/vllm#20059
🏁 Script executed:
#!/bin/bash
set -euo pipefail
# Inspect the relevant benchmark file and the comparison target.
git ls-files bench_scripts/moe35_prof.py bench_scripts/moe_topo.py
echo '--- moe35_prof.py ---'
cat -n bench_scripts/moe35_prof.py | sed -n '1,140p'
echo '--- moe_topo.py ---'
cat -n bench_scripts/moe_topo.py | sed -n '1,180p'Repository: larkinwc/vllm-gfx908
Length of output: 5134
🏁 Script executed:
#!/bin/bash
set -euo pipefail
rg -n "cudagraph_capture_sizes|FULL_DECODE_ONLY|c=1/8/32|c in \[1,8\]" bench_scriptsRepository: larkinwc/vllm-gfx908
Length of output: 2184
Include the sweep sizes in cudagraph_capture_sizes.
bench runs decode at c=1/8/32 under FULL_DECODE_ONLY, but this config only captures size 1. c=8 and c=32 will fall back to eager decode, so those throughput numbers won’t reflect graph-enabled runs. Match the sweep with [1, 8, 32] here.
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
In `@bench_scripts/moe35_prof.py` around lines 36 - 38, The decode benchmark
config is only capturing one cudagraph size, so the runs for the other sweep
values will fall back to eager execution. Update the `CompilationConfig` in
`moe35_prof.py` so `cudagraph_capture_sizes` matches the benchmark sweep used by
`bench`, and make sure the `FULL_DECODE_ONLY` setup in that block includes all
decode sizes being measured, not just `1`.
Comment on lines
+3
to
+5
| d_dir = sys.argv[1] if len(sys.argv) > 1 else "/home/larkinwc/moe35_prof" | ||
| f = sorted(glob.glob(d_dir + "/*rank0.*.pt.trace.json.gz"))[0] | ||
| d = json.load(gzip.open(f)) |
There was a problem hiding this comment.
🩺 Stability & Availability | 🟡 Minor | ⚡ Quick win
Unhandled IndexError when no trace file is found.
If the glob pattern matches nothing (wrong d_dir, profiler not run, or filename mismatch), this crashes with an opaque IndexError instead of a clear message.
🛠️ Proposed fix
-f = sorted(glob.glob(d_dir + "/*rank0.*.pt.trace.json.gz"))[0]
+matches = sorted(glob.glob(d_dir + "/*rank0.*.pt.trace.json.gz"))
+if not matches:
+ sys.exit(f"No trace files found matching pattern in {d_dir!r}")
+f = matches[0]📝 Committable suggestion
‼️ IMPORTANT
Carefully review the code before committing. Ensure that it accurately replaces the highlighted code, contains no missing lines, and has no issues with indentation. Thoroughly test & benchmark the code to ensure it meets the requirements.
Suggested change
| d_dir = sys.argv[1] if len(sys.argv) > 1 else "/home/larkinwc/moe35_prof" | |
| f = sorted(glob.glob(d_dir + "/*rank0.*.pt.trace.json.gz"))[0] | |
| d = json.load(gzip.open(f)) | |
| d_dir = sys.argv[1] if len(sys.argv) > 1 else "/home/larkinwc/moe35_prof" | |
| matches = sorted(glob.glob(d_dir + "/*rank0.*.pt.trace.json.gz")) | |
| if not matches: | |
| sys.exit(f"No trace files found matching pattern in {d_dir!r}") | |
| f = matches[0] | |
| d = json.load(gzip.open(f)) |
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
In `@bench_scripts/parse_any.py` around lines 3 - 5, The trace selection logic in
parse_any.py can crash with an IndexError when the glob finds no files. Update
the d_dir/f lookup flow to check the glob result before indexing, and in the
parse_any script emit a clear error or exit message when no matching
*.pt.trace.json.gz file is found so the failure is explicit and easy to
diagnose.
Comment on lines
+61
to
+71
| ze = tl.interleave(qzv, qzv); ze = tl.interleave(ze, ze); ze = tl.interleave(ze, ze) | ||
| zf = ((ze >> shifts) & 0xF).to(tl.float32) | ||
| else: | ||
| zf = tl.full((BLOCK_N,), ZP_BIAS, dtype=tl.float32) | ||
|
|
||
| offs_k = g * G + tl.arange(0, G) | ||
| mask_k = offs_k < K | ||
| qw_ptrs = qw_ptr + offs_k[:, None] * (N // 8) + offs_n8[None, :] | ||
| m_w = (mask_k[:, None] & in_range) & mask_n8[None, :] | ||
| qwt = tl.load(qw_ptrs, mask=m_w, other=0) | ||
| bt = tl.interleave(qwt, qwt); bt = tl.interleave(bt, bt); bt = tl.interleave(bt, bt) |
There was a problem hiding this comment.
📐 Maintainability & Code Quality | 🟡 Minor | ⚡ Quick win
Split the chained assignments flagged by Ruff.
Lines 61 and 71 use multiple statements per line, which will fail E702.
🧹 Proposed lint fix
- ze = tl.interleave(qzv, qzv); ze = tl.interleave(ze, ze); ze = tl.interleave(ze, ze)
+ ze = tl.interleave(qzv, qzv)
+ ze = tl.interleave(ze, ze)
+ ze = tl.interleave(ze, ze)
@@
- bt = tl.interleave(qwt, qwt); bt = tl.interleave(bt, bt); bt = tl.interleave(bt, bt)
+ bt = tl.interleave(qwt, qwt)
+ bt = tl.interleave(bt, bt)
+ bt = tl.interleave(bt, bt)As per coding guidelines, **/*.py should keep Python lines within 88 characters.
📝 Committable suggestion
‼️ IMPORTANT
Carefully review the code before committing. Ensure that it accurately replaces the highlighted code, contains no missing lines, and has no issues with indentation. Thoroughly test & benchmark the code to ensure it meets the requirements.
Suggested change
| ze = tl.interleave(qzv, qzv); ze = tl.interleave(ze, ze); ze = tl.interleave(ze, ze) | |
| zf = ((ze >> shifts) & 0xF).to(tl.float32) | |
| else: | |
| zf = tl.full((BLOCK_N,), ZP_BIAS, dtype=tl.float32) | |
| offs_k = g * G + tl.arange(0, G) | |
| mask_k = offs_k < K | |
| qw_ptrs = qw_ptr + offs_k[:, None] * (N // 8) + offs_n8[None, :] | |
| m_w = (mask_k[:, None] & in_range) & mask_n8[None, :] | |
| qwt = tl.load(qw_ptrs, mask=m_w, other=0) | |
| bt = tl.interleave(qwt, qwt); bt = tl.interleave(bt, bt); bt = tl.interleave(bt, bt) | |
| ze = tl.interleave(qzv, qzv) | |
| ze = tl.interleave(ze, ze) | |
| ze = tl.interleave(ze, ze) | |
| zf = ((ze >> shifts) & 0xF).to(tl.float32) | |
| else: | |
| zf = tl.full((BLOCK_N,), ZP_BIAS, dtype=tl.float32) | |
| offs_k = g * G + tl.arange(0, G) | |
| mask_k = offs_k < K | |
| qw_ptrs = qw_ptr + offs_k[:, None] * (N // 8) + offs_n8[None, :] | |
| m_w = (mask_k[:, None] & in_range) & mask_n8[None, :] | |
| qwt = tl.load(qw_ptrs, mask=m_w, other=0) | |
| bt = tl.interleave(qwt, qwt) | |
| bt = tl.interleave(bt, bt) | |
| bt = tl.interleave(bt, bt) |
🧰 Tools
🪛 Ruff (0.15.20)
[error] 61-61: Multiple statements on one line (semicolon)
(E702)
[error] 61-61: Multiple statements on one line (semicolon)
(E702)
[error] 71-71: Multiple statements on one line (semicolon)
(E702)
[error] 71-71: Multiple statements on one line (semicolon)
(E702)
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
In `@vllm/model_executor/kernels/linear/mixed_precision/gfx900_w4a16_gemv.py`
around lines 61 - 71, Ruff is flagging chained statements in the GEMV kernel:
the assignments to ze and bt in gfx900_w4a16_gemv should be split into separate
lines instead of using multiple statements per line. Update the relevant logic
in the mixed-precision GEMV path so each tl.interleave step is assigned
independently, while keeping the Python line length within the 88-character
guideline and preserving the existing variable names ze and bt for easy
location.
Sources: Coding guidelines, Linters/SAST tools
Comment on lines
223
to
236
| if current_platform.is_rocm() and not _mi100_w4a16_disabled(): | ||
| from vllm.platforms.rocm import on_mi100 | ||
| from vllm.platforms.rocm import on_gfx900, on_mi100 | ||
| # gfx900 (Vega10) has no MFMA -> tl.dot is a padded FP32 GEMM and is | ||
| # *slower* than FP16 at small M. For the decode hot path (M<=8) route to | ||
| # the dedicated bandwidth-bound int4 GEMV (10-15x vs the tl.dot path). | ||
| if on_gfx900() and group_size in (32, 64, 128) and M <= 8: | ||
| from vllm.model_executor.kernels.linear.mixed_precision.gfx900_w4a16_gemv import ( | ||
| w4a16_gemv_gfx900 as _gfx900_gemv, | ||
| ) | ||
| return _gfx900_gemv( | ||
| a=a, b_q=b_q, scales=scales, | ||
| qzeros=qzeros, group_size=group_size, zp_bias=zp_bias, | ||
| ) | ||
| if on_mi100() and group_size in (32, 128): |
There was a problem hiding this comment.
🚀 Performance & Scalability | 🟡 Minor | ⚡ Quick win
Do not gate gfx900 dispatch on the MI100 kill switch.
_mi100_w4a16_disabled() now disables the gfx900 GEMV path too, causing gfx900 small-M AWQ decode to fall back to the slower generic Triton path when only the MI100 path was meant to be disabled.
⚙️ Proposed dispatch split
- if current_platform.is_rocm() and not _mi100_w4a16_disabled():
+ if current_platform.is_rocm():
from vllm.platforms.rocm import on_gfx900, on_mi100
@@
if on_gfx900() and group_size in (32, 64, 128) and M <= 8:
- from vllm.model_executor.kernels.linear.mixed_precision.gfx900_w4a16_gemv import (
+ from .gfx900_w4a16_gemv import (
w4a16_gemv_gfx900 as _gfx900_gemv,
)
@@
- if on_mi100() and group_size in (32, 128):
+ if (
+ not _mi100_w4a16_disabled()
+ and on_mi100()
+ and group_size in (32, 128)
+ ):📝 Committable suggestion
‼️ IMPORTANT
Carefully review the code before committing. Ensure that it accurately replaces the highlighted code, contains no missing lines, and has no issues with indentation. Thoroughly test & benchmark the code to ensure it meets the requirements.
Suggested change
| if current_platform.is_rocm() and not _mi100_w4a16_disabled(): | |
| from vllm.platforms.rocm import on_mi100 | |
| from vllm.platforms.rocm import on_gfx900, on_mi100 | |
| # gfx900 (Vega10) has no MFMA -> tl.dot is a padded FP32 GEMM and is | |
| # *slower* than FP16 at small M. For the decode hot path (M<=8) route to | |
| # the dedicated bandwidth-bound int4 GEMV (10-15x vs the tl.dot path). | |
| if on_gfx900() and group_size in (32, 64, 128) and M <= 8: | |
| from vllm.model_executor.kernels.linear.mixed_precision.gfx900_w4a16_gemv import ( | |
| w4a16_gemv_gfx900 as _gfx900_gemv, | |
| ) | |
| return _gfx900_gemv( | |
| a=a, b_q=b_q, scales=scales, | |
| qzeros=qzeros, group_size=group_size, zp_bias=zp_bias, | |
| ) | |
| if on_mi100() and group_size in (32, 128): | |
| if current_platform.is_rocm(): | |
| from vllm.platforms.rocm import on_gfx900, on_mi100 | |
| # gfx900 (Vega10) has no MFMA -> tl.dot is a padded FP32 GEMM and is | |
| # *slower* than FP16 at small M. For the decode hot path (M<=8) route to | |
| # the dedicated bandwidth-bound int4 GEMV (10-15x vs the tl.dot path). | |
| if on_gfx900() and group_size in (32, 64, 128) and M <= 8: | |
| from .gfx900_w4a16_gemv import ( | |
| w4a16_gemv_gfx900 as _gfx900_gemv, | |
| ) | |
| return _gfx900_gemv( | |
| a=a, b_q=b_q, scales=scales, | |
| qzeros=qzeros, group_size=group_size, zp_bias=zp_bias, | |
| ) | |
| if ( | |
| not _mi100_w4a16_disabled() | |
| and on_mi100() | |
| and group_size in (32, 128) | |
| ): |
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
In `@vllm/model_executor/kernels/linear/mixed_precision/triton_w4a16.py` around
lines 223 - 236, The gfx900 GEMV dispatch in w4a16 should be split from the
MI100 kill switch so `_mi100_w4a16_disabled()` does not block `on_gfx900()`
routing in `vllm/model_executor/kernels/linear/mixed_precision/triton_w4a16.py`.
Update the ROCm branch around the `on_gfx900()` and `on_mi100()` checks so
gfx900 small-M AWQ decode can still return `w4a16_gemv_gfx900` when its own
conditions match, while only the MI100-specific path remains guarded by
`_mi100_w4a16_disabled()`. Use the existing symbols
`current_platform.is_rocm()`, `on_gfx900()`, `on_mi100()`, and
`_mi100_w4a16_disabled()` to keep the dispatch logic explicit.
Comment on lines
+1256
to
+1270
| from vllm.platforms.rocm import on_gfx900 | ||
| import os as _os | ||
| if on_gfx900() and M <= 2 and _os.environ.get("VLLM_GFX900_MOE_GEMV", "1") != "0": | ||
| # gfx900 has no MFMA: a 16-row tl.dot wastes ~16x on a 1-token | ||
| # decode. Use a real per-token GEMV (BLOCK_SIZE_M=1 + tl.sum). | ||
| # Small BLOCK_SIZE_N maximizes CU occupancy for the M=1 GEMV | ||
| # (swept: N=8,K=128 optimal on Vega10, 56 CUs). | ||
| config = { | ||
| "BLOCK_SIZE_M": 1, "GROUP_SIZE_M": 1, "SPLIT_K": 1, | ||
| "GEMV_MODE": True, | ||
| "BLOCK_SIZE_N": 8, "BLOCK_SIZE_K": 128, | ||
| "num_warps": 4, "num_stages": 2, | ||
| } | ||
| else: | ||
| config = {"BLOCK_SIZE_M": 16, "GROUP_SIZE_M": 1, "SPLIT_K": 1} |
There was a problem hiding this comment.
🩺 Stability & Availability | 🟠 Major | ⚡ Quick win
Guard the ROCm-only import by platform.
get_default_config() can run during non-ROCm config selection, but this branch imports vllm.platforms.rocm unconditionally once the M <= 20 path is reached. Keep the ROCm helper import behind current_platform.is_rocm().
🛡️ Proposed guard
- from vllm.platforms.rocm import on_gfx900
import os as _os
- if on_gfx900() and M <= 2 and _os.environ.get("VLLM_GFX900_MOE_GEMV", "1") != "0":
+ is_gfx900 = False
+ if current_platform.is_rocm():
+ from vllm.platforms.rocm import on_gfx900
+
+ is_gfx900 = on_gfx900()
+ gfx900_moe_gemv_enabled = (
+ _os.environ.get("VLLM_GFX900_MOE_GEMV", "1") != "0"
+ )
+ if is_gfx900 and M <= 2 and gfx900_moe_gemv_enabled:📝 Committable suggestion
‼️ IMPORTANT
Carefully review the code before committing. Ensure that it accurately replaces the highlighted code, contains no missing lines, and has no issues with indentation. Thoroughly test & benchmark the code to ensure it meets the requirements.
Suggested change
| from vllm.platforms.rocm import on_gfx900 | |
| import os as _os | |
| if on_gfx900() and M <= 2 and _os.environ.get("VLLM_GFX900_MOE_GEMV", "1") != "0": | |
| # gfx900 has no MFMA: a 16-row tl.dot wastes ~16x on a 1-token | |
| # decode. Use a real per-token GEMV (BLOCK_SIZE_M=1 + tl.sum). | |
| # Small BLOCK_SIZE_N maximizes CU occupancy for the M=1 GEMV | |
| # (swept: N=8,K=128 optimal on Vega10, 56 CUs). | |
| config = { | |
| "BLOCK_SIZE_M": 1, "GROUP_SIZE_M": 1, "SPLIT_K": 1, | |
| "GEMV_MODE": True, | |
| "BLOCK_SIZE_N": 8, "BLOCK_SIZE_K": 128, | |
| "num_warps": 4, "num_stages": 2, | |
| } | |
| else: | |
| config = {"BLOCK_SIZE_M": 16, "GROUP_SIZE_M": 1, "SPLIT_K": 1} | |
| import os as _os | |
| is_gfx900 = False | |
| if current_platform.is_rocm(): | |
| from vllm.platforms.rocm import on_gfx900 | |
| is_gfx900 = on_gfx900() | |
| gfx900_moe_gemv_enabled = ( | |
| _os.environ.get("VLLM_GFX900_MOE_GEMV", "1") != "0" | |
| ) | |
| if is_gfx900 and M <= 2 and gfx900_moe_gemv_enabled: | |
| # gfx900 has no MFMA: a 16-row tl.dot wastes ~16x on a 1-token | |
| # decode. Use a real per-token GEMV (BLOCK_SIZE_M=1 + tl.sum). | |
| # Small BLOCK_SIZE_N maximizes CU occupancy for the M=1 GEMV | |
| # (swept: N=8,K=128 optimal on Vega10, 56 CUs). | |
| config = { | |
| "BLOCK_SIZE_M": 1, "GROUP_SIZE_M": 1, "SPLIT_K": 1, | |
| "GEMV_MODE": True, | |
| "BLOCK_SIZE_N": 8, "BLOCK_SIZE_K": 128, | |
| "num_warps": 4, "num_stages": 2, | |
| } | |
| else: | |
| config = {"BLOCK_SIZE_M": 16, "GROUP_SIZE_M": 1, "SPLIT_K": 1} |
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
In `@vllm/model_executor/layers/fused_moe/fused_moe.py` around lines 1256 - 1270,
The ROCm helper import inside get_default_config is unconditional and can run on
non-ROCm paths; move the import of on_gfx900 behind a current_platform.is_rocm()
check so the M <= 20 branch only touches vllm.platforms.rocm when the platform
is actually ROCm. Keep the existing gfx900-specific config selection logic
intact, but gate the ROCm-only import and environment-based GEMV override
together in that ROCm branch.
Comment on lines
+181
to
+207
| if ( | ||
| envs.VLLM_ROCM_USE_SKINNY_GEMM | ||
| and on_gfx900() | ||
| and x.dtype in [torch.float16, torch.bfloat16] | ||
| and n == 1 | ||
| and m <= 8 | ||
| and m % 4 != 0 | ||
| and bias is None | ||
| ): | ||
| # weight is [m, k]; broadcast x [1, k] over rows and reduce over K -> [m] | ||
| x_view = x.reshape(-1, x.size(-1)) | ||
| out = (x_view * weight).sum(dim=-1) | ||
| return out.reshape(*x.shape[:-1], m) | ||
|
|
||
| if ( | ||
| envs.VLLM_ROCM_USE_SKINNY_GEMM | ||
| and on_gfx900() | ||
| and x.dtype in [torch.float16, torch.bfloat16] | ||
| and k % 8 == 0 | ||
| and m % 4 == 0 | ||
| and n == 1 | ||
| and k <= 8192 | ||
| and bias is None | ||
| ): | ||
| x_view = x.reshape(-1, x.size(-1)) | ||
| out = ops.LLMM1(weight, x_view, 4) | ||
| return out.reshape(*x.shape[:-1], weight.shape[0]) |
There was a problem hiding this comment.
🩺 Stability & Availability | 🟠 Major | ⚡ Quick win
Require matching weight dtype before taking gfx900 fast paths.
The fallback below casts x when x.dtype != weight.dtype, but these new branches bypass that safeguard. The manual path can return a promoted dtype, and ops.LLMM1 may receive unsupported mixed dtypes.
🔧 Proposed guard
and on_gfx900()
and x.dtype in [torch.float16, torch.bfloat16]
+ and weight.dtype == x.dtype
and n == 1
@@
and on_gfx900()
and x.dtype in [torch.float16, torch.bfloat16]
+ and weight.dtype == x.dtype
and k % 8 == 0📝 Committable suggestion
‼️ IMPORTANT
Carefully review the code before committing. Ensure that it accurately replaces the highlighted code, contains no missing lines, and has no issues with indentation. Thoroughly test & benchmark the code to ensure it meets the requirements.
Suggested change
| if ( | |
| envs.VLLM_ROCM_USE_SKINNY_GEMM | |
| and on_gfx900() | |
| and x.dtype in [torch.float16, torch.bfloat16] | |
| and n == 1 | |
| and m <= 8 | |
| and m % 4 != 0 | |
| and bias is None | |
| ): | |
| # weight is [m, k]; broadcast x [1, k] over rows and reduce over K -> [m] | |
| x_view = x.reshape(-1, x.size(-1)) | |
| out = (x_view * weight).sum(dim=-1) | |
| return out.reshape(*x.shape[:-1], m) | |
| if ( | |
| envs.VLLM_ROCM_USE_SKINNY_GEMM | |
| and on_gfx900() | |
| and x.dtype in [torch.float16, torch.bfloat16] | |
| and k % 8 == 0 | |
| and m % 4 == 0 | |
| and n == 1 | |
| and k <= 8192 | |
| and bias is None | |
| ): | |
| x_view = x.reshape(-1, x.size(-1)) | |
| out = ops.LLMM1(weight, x_view, 4) | |
| return out.reshape(*x.shape[:-1], weight.shape[0]) | |
| if ( | |
| envs.VLLM_ROCM_USE_SKINNY_GEMM | |
| and on_gfx900() | |
| and x.dtype in [torch.float16, torch.bfloat16] | |
| and weight.dtype == x.dtype | |
| and n == 1 | |
| and m <= 8 | |
| and m % 4 != 0 | |
| and bias is None | |
| ): | |
| # weight is [m, k]; broadcast x [1, k] over rows and reduce over K -> [m] | |
| x_view = x.reshape(-1, x.size(-1)) | |
| out = (x_view * weight).sum(dim=-1) | |
| return out.reshape(*x.shape[:-1], m) | |
| if ( | |
| envs.VLLM_ROCM_USE_SKINNY_GEMM | |
| and on_gfx900() | |
| and x.dtype in [torch.float16, torch.bfloat16] | |
| and weight.dtype == x.dtype | |
| and k % 8 == 0 | |
| and m % 4 == 0 | |
| and n == 1 | |
| and k <= 8192 | |
| and bias is None | |
| ): | |
| x_view = x.reshape(-1, x.size(-1)) | |
| out = ops.LLMM1(weight, x_view, 4) | |
| return out.reshape(*x.shape[:-1], weight.shape[0]) |
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
In `@vllm/model_executor/layers/utils.py` around lines 181 - 207, The gfx900
skinny-GEMM fast paths in the utils function should only run when the input and
weight dtypes already match, since the regular fallback handles mismatches by
casting x first. Add a dtype equality guard alongside the existing checks in
both branches before the manual reduction path and the ops.LLMM1 call, using the
same x and weight symbols in this helper, so mixed-dtype cases continue through
the safe fallback instead of hitting unsupported promotion behavior.
Comment on lines
+309
to
+312
| def on_gfx900() -> bool: | ||
| """Detect gfx900 (Vega10): GFX9 ISA but lacks MFMA, FP8, and XGMI. | ||
| Must route around MFMA paged-attn/skinny-GEMM and XGMI custom all-reduce.""" | ||
| return _ON_GFX900 |
There was a problem hiding this comment.
📐 Maintainability & Code Quality | 🟡 Minor | ⚡ Quick win
Use the requested Google-style docstring format.
The new predicate docstring should include at least a Returns: section.
📝 Proposed docstring update
def on_gfx900() -> bool:
- """Detect gfx900 (Vega10): GFX9 ISA but lacks MFMA, FP8, and XGMI.
- Must route around MFMA paged-attn/skinny-GEMM and XGMI custom all-reduce."""
+ """Detect gfx900 (Vega10).
+
+ gfx900 is GFX9 ISA but lacks MFMA, FP8, and XGMI.
+
+ Returns:
+ Whether the current ROCm device architecture is gfx900.
+ """
return _ON_GFX900As per coding guidelines, **/*.py docstrings should use Google-style sections such as Args:, Returns:, and Raises:.
📝 Committable suggestion
‼️ IMPORTANT
Carefully review the code before committing. Ensure that it accurately replaces the highlighted code, contains no missing lines, and has no issues with indentation. Thoroughly test & benchmark the code to ensure it meets the requirements.
Suggested change
| def on_gfx900() -> bool: | |
| """Detect gfx900 (Vega10): GFX9 ISA but lacks MFMA, FP8, and XGMI. | |
| Must route around MFMA paged-attn/skinny-GEMM and XGMI custom all-reduce.""" | |
| return _ON_GFX900 | |
| def on_gfx900() -> bool: | |
| """Detect gfx900 (Vega10). | |
| gfx900 is GFX9 ISA but lacks MFMA, FP8, and XGMI. | |
| Returns: | |
| Whether the current ROCm device architecture is gfx900. | |
| """ | |
| return _ON_GFX900 |
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
In `@vllm/platforms/rocm.py` around lines 309 - 312, The on_gfx900 predicate
docstring is not using the required Google-style format. Update the docstring in
on_gfx900 to use Google-style sections and include a Returns: section that
describes the boolean result, keeping the existing detection summary concise.
Ensure the formatting matches the style used across vllm/platforms/rocm.py and
other Python docstrings.
Source: Coding guidelines
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.
gfx900 (Vega10 / Radeon Pro V340) inference support + decode perf
Adds a gfx900 support tier to this ROCm vLLM fork and brings up fast LLM inference on an 8× V340 (16 die) box. Validated end-to-end on two models: dense Qwen3.5-9B and sparse-MoE Qwen3-Coder-Next-AWQ-4bit.
Headline results
What's in the platform/kernel changes
vllm/platforms/rocm.py:on_gfx900()detection + capability(9,0), gfx900-appropriate attention backend selection (Triton/TurboQuant), 56 CUs / wave64 / no-MFMA characterization.vllm/model_executor/layers/utils.py:shared_expert_gate, n==1, m≤8, m%4≠0) that fails the LLMM1m%4==0gate and otherwise falls into a wasteful rocBLAS macrotile — +39% MoE decode ([gfx900] Bring up Qwen3-Coder-Next-AWQ-4bit (MoE) + tune MoE expert-GEMM kernels #65).gfx900_w4a16_gemv.py: int4 AWQ weight GEMV decode kernel (split-K tuned) ([gfx900] Hand-written int4 AWQ GEMV decode kernel (no tl.dot) to make quant win #57).FULL_DECODE_ONLY, inductor off) — flips an earlier "no speedup" result into 3.6× after LLMM1 ([gfx900] HIP CUDA graphs NOW give 3.6x decode (revisit after LLMM1) #63).turboquant_k8v4) validated lossless (WikiText-2 PPL within noise); 2.3–2.5× KV capacity ([gfx900] Enable + validate in-tree TurboQuant KV-cache quant (Qwen3.5-9B, quirky-K) #62).Key findings (documented in
GFX900_RECOMMENDED.md+ linked issues)Docs added
GFX900_RECOMMENDED.md(best validated path, both models),PERF_GFX900.md,BENCH_DECODE_PROFILE.md,BENCH_GFX900.md,GFX900_SETUP.md, plus benchmark scripts underbench_scripts/.Testing
All numbers measured on the 16× gfx900 box (ROCm 7.2.4, torch 2.12.0+rocm7.2, custom gfx900 RCCL + Triton fork). gfx900-specific code paths are gated behind
on_gfx900(); other archs unaffected. Tracked across issues #55–#65.Summary by CodeRabbit
New Features
Bug Fixes
Documentation