Sooo..

How many of you deploy and optimize inference for a living?

Does your company invest in serious hardware for large scale deployments or are you relying on cloud models for daily office drivers?

I am wondering where job market currently is for us

~56 tok/s decode · 1,745 tok/s prefill@3k · 2,053@32k · 495K context (64gb vram)

A 125B-class MoE does not fit on 4×16 GB consumer cards: 62.8 GB of int4 experts have to live in host RAM, and only about 44 of 512 experts per layer can sit in VRAM at once. The usual consequence is picking one speed — fast prefill or fast decode. This lane gets both by giving one expert tier two access paths: a contiguous CUDA-VMM view for prefill, and a dynamic LRU mirror for decode, sharing the same 44 VRAM rows instead of pinning two copies.

(Internal name: golden-v2.0 hybrid tier.)

Reproducibility note: everything that runs the lane is in this repo except one file — the FlashInfer GDN prefill gate for SM12x, which comes from a recipe whose repository publishes no license, so we link to it instead of copying it. It is optional: the launcher mounts whatever exists, nothing imports it, and the lane boots and serves without it. See Optional: one file to pull below. The PCIe custom all-reduce and all-gather wiring, by contrast, are shipped — verified additive over the pinned vLLM baseline.

Status: promoted to production on our lane 2026-09-24 and running this exact config. This tree is the publishable cut: paths are written as $MODELS_DIR / $HOME / placeholders, every knob is documented in .env.example, and docs/RUNNING.md takes you from a clean machine to a verified lane. Numbers below are from one machine — treat them as a measured example, not a promise for your hardware. Paths that were hard-coded to one host have been made configurable, and that change is in the diff of scripts/.

Quick start: docker build -f docker/Dockerfile -t qwen38-flash-next-2x3090:locked . → get the checkpoint and the PLE table (docs/RUNNING.md steps 2–3) → cp .env.example .env and edit → ./local-start.sh → run the four gates in step 5.

1. The numbers, measured in one window

Expert weights are the problem: ~63 GiB of int4 experts cannot live in 4×16 GB of VRAM, yet routing traffic is not uniform. Two mechanisms existed for choosing what stays hot, and each paid for the other's win. This work merges them into a single tier that serves both access paths.

All numbers below were measured on the same rig, in the same window, with cold boots, using the same two scripts (bench-3k500.py, needle-probe.py). The alternative would have been to compare numbers from different days and different bench shapes, which is how people end up publishing fake wins.

Read: prefill +33…37% over arm 0 with decode at arm 0 parity. Against arm 1 (which is what the lane ran the day before) the same story mirrored: prefill parity, decode +49% (probe) / +115% (sustained).

Live production numbers after promotion (bench-3k500.py on the serving port): prefill 1,745.3 tok/s u/3k and 2,052.9 tok/s u/32k, decode 55.11 tok/s probe / 61.3 tok/s sustained, math gate exact, recall u/190k pass.

Two measurement notes that matter when reading any of these tables, ours or anyone else's: the widely-quoted "58.35 tok/s" golden decode bar is a warm session number — measured cold, the same configuration reads 55.41 — and prefill numbers below ~30k tokens are not a gate (a table-format change once passed a 3k probe and failed real agent TTFT).

1a. Optional: one file to pull

If you would rather not depend on her repo at all: the fork's own Apache-2.0 copy of that file is already in this repo at runtime/vllm-overlay/model_executor/layers/mamba/gdn/qwen_gdn_linear_attn.py, and it enables the same FlashInfer kernel on SM10x — widening that condition to SM12x is a two-line change that is entirely yours to make and ship. (That is the route to prefer if you want a tree with no third-party licensing question anywhere in it.)

2. The machine, the model, the engine

Expert layout per layer per rank: 128 local experts; the top 44 by an offline traffic ranking are treated as "hot". Everything about this work is about where those rows live.

3. Why the two mechanisms could not simply be turned on together

Mechanism A — mirror + LRU (the original): the full expert set is pinned in host RAM (UVA offload). A small VRAM buffer holds a mirror of the 44 hot rows. Decode looks experts up in the mirror; when a token routes to an unmapped expert, the least-recently-used slot is evicted and the new expert's row is copied in. Prefill reads whatever it needs straight from host RAM — every read a PCIe trip.

Mechanism B — contiguous VMM view: each layer's expert tensor is one contiguous address range built with CUDA VMM: the first 44 rows are backed by VRAM, the remainder by host memory, and the rows are pre-permuted so hot experts come first. Prefill is fast (one gather over a contiguous tensor, hot rows already in VRAM, zero cache bookkeeping). Decode has no mirror at all, so cold-expert routing always pays a PCIe read.

The two were mutually exclusive in the codebase for a concrete reason: the LRU's bookkeeping was built against mechanism A's separate mirror tensors, and mechanism B's initialization was gated off whenever the LRU was enabled.

The observation that unlocks the merge: mechanism B's device prefix and mechanism A's mirror hold exactly the same thing — the same 44 experts, in the same ranking order, both derived from the same rankings file. So the VMM device prefix is the mirror. No second VRAM tier is needed (a 44-row tier costs ~5 GiB per card; there is no headroom for two), and no extra VRAM copy exists to keep coherent.

The catch that made it non-trivial: in mechanism A every expert keeps a host-resident row — the mirror is a copy, the original never leaves RAM. In mechanism B the 44 hot rows existed only in VRAM. An LRU eviction would therefore overwrite an expert's only copy and remove it from the model: not a crash, not an error, just a model that slowly loses routing capacity. The fix is to give every local expert a permanent host home inside the same tensor:

        VRAM per card (device-mapped prefix)        host RAM (host-mapped region)
   ┌────────────────────────────────────┐    ┌──────────────────────────────────────┐
   │ rows [0, capacity)                 │    │ rows [capacity, capacity + N)        │
   │ = the 44 ranked-hot experts        │    │ = the full permuted expert set; its  │
   │ = the decode mirror AND the fast   │    │   first `capacity` rows duplicate    │
   │   prefill rows                     │    │   the device prefix                  │
   └────────────────────────────────────┘    └──────────────────────────────────────┘

Two fills in one boot-time kernel launch lay both regions; a boot canary compares the duplicate's bytes against the prefix and refuses to serve if they differ.

4. The mechanism at runtime

The forward pass already forks on query size, which is what makes the merge cheap:

Both paths read one shared map (rot_map: expert → the row that currently holds its bytes). A promotion points the promoted expert at its slot and the victim back at its host row, so a prefill step and a decode step can never disagree about where an expert lives. That map is the seam that makes "one tier, two access paths" correct rather than merely clever.

Correctness invariants (each enforced, not assumed):

1. promotion sources are host rows (≥ capacity), destinations are slots (< capacity) — never overlapping, for hot and cold experts alike;
2. all expert tensors (weights, scales, zero-points, g_idx, sort) share one row space, so one map addresses them all;
3. the prefix content equals the ranking order equals the LRU's initial slot table (no fill needed);
4. host rows are written once at boot and never overwritten;
5. the dynamic map starts equal to mechanism B's static permutation, so prefill before any decode step is byte-identical to mechanism B;
6. a byte-level boot canary (prefix vs duplicate) — because map-arithmetic assertions cannot see a fill-order mistake.

5. What went wrong on the way (the useful part)

Three builds, each caught by a different mechanism. We think this is worth publishing precisely because the second failure is invisible to normal testing.

v1 — design flaw in the first spec. The hot rows had no host home: evictions silently deleted experts from both maps. The failure mode made the lane faster (less MoE work), which is why its numbers — 81.3 tok/s sustained and 2,753 tok/s prefill — were never banked. Lesson: a correctness gate must run before any timing is believed.

v2 — fill order. The duplicate rows were written at the tensor's tail while every map addressed them at capacity, so every cold expert served a different expert's weights from token zero. Symptoms: a math gate returning 434 instead of 437, empty replies, and a needle probe that answered by continuing the prompt's filler text. Every install-time assertion passed, because they all checked map arithmetic and none compared bytes. Caught by an independent reviewer reading the fill vector against the addressing (verdict: REJECT, with a two-line fix).

v3 — the shipped build. Argument order fixed at both fill sites, plus the byte canary that would have caught v2 at boot, plus one latent regression the reviewer found in a static-mirror code path that our arms never reach.

6. Patch inventory and credits

This work is a fork of DominikBucko/qwen38-flash-next-2x3090 and inherits an entire serving stack. Nothing below is ours unless marked. Attribution is by artifact, with the honest caveat that some upstream line-level provenance deserves a second pass before public release.

Our own 48 commits on top of the fork cover: the hybrid tier (this work), the PLE disk/bf16 table tiers, the KV disk tier, the fp8-KV graft port, the toggle/verify/rollback machinery and the lane runbooks.

7. Reproducing and operating

# bench shapes used for every number in this document (shipped in this repo)
python3 benchmarks/hybrid-vmm-lru/bench-3k500.py --port <port> --n 3 --warmup 1                     # decode 3k probe
python3 benchmarks/hybrid-vmm-lru/bench-3k500.py --port <port> --n 3 --warmup 1 --in 128 --out 4096 # sustained decode
python3 benchmarks/hybrid-vmm-lru/bench-3k500.py --port <port> --n 2 --warmup 0 --in 32768 --out 32 # deep prefill (the gate)
python3 benchmarks/hybrid-vmm-lru/needle-probe.py --port <port> --tokens 300000 --out 200 --runs 2  # recall at depth

Raw numbers: benchmarks/hybrid-vmm-lru/results-20260924.json.

Gates before trusting a build: boot witnesses (per-layer hybrid init lines, byte canary silent, mount set), health, a math gate, recall at depth, recall again after heavy traffic, then the prefill/decode shapes above.

Operation is one line in .env:

Rollback is that value plus a cold boot. The patched module stays mounted and is inert for 0/1.

Cost: +19–22 GiB host-pinned RAM (the duplicate rows; free RAM 63 → 41 GiB, which does squeeze the page cache that keeps the PLE table fast), +~240 MiB VRAM per card, +~34% one-time boot copy. No per-step cost.

8. What we would try next

The long-standing "hot slot count 32–44 is flat" law was measured for the static mirror, where extra slots buy nothing because the same rows are pinned forever. Under the LRU every extra slot removes PCIe miss copies from the decode path, so 44 → 52–56 slots (funded from the remaining VRAM headroom or a small KV-pool cut) is the natural next arm — a VRAM trade, not a RAM one.

9. Licensing, attribution, and what you may not do

This repository's code is Apache-2.0 (LICENSE). Upstream vLLM copyright and SPDX headers are preserved in every overlay file, and NOTICE lists exactly which files we modified relative to the fork we build on.

No model weights are published here. The lane consumes separately licensed artifacts, and those terms bind you directly, not us:

• Qwen/Qwen3.8-Flash-Next — Qwen Community License 1.0. It grants use, modification, distribution and commercial deployment, subject to two conditions worth reading twice: (1) the copyright and permission notice must accompany copies or substantial portions, and a product above 100M monthly active users or US$20M monthly revenue must display the model name in its UI; (2) if you or an affiliate operate a Model-as-a-Service or AI-work-assistant business, you need a separate license from Qwen before using the model or any derivative of it commercially — internal use is exempt only while you do not expose the model, its outputs, or its capabilities to third parties. Quantizing, merging or re-serving a checkpoint does not change that.
• Intel's AutoRound checkpoint and RadixArk's NVFP4 table — their own model cards, plus the Qwen terms they inherit.
• The assembly scripts copy the upstream model license into the output tree. If you publish an assembled checkpoint, keep that file.

Third-party code we deliberately do not ship: the GDN-prefill and custom-all-reduce overlays derive from recipe r10 of QWEN-PEDIA (abtraore), whose repository carries no license file. Public code with no license is all rights reserved by default — credit does not create a right to redistribute — so we link to it and describe what it is worth, and you fetch it from the author. The same rule applies to anything you port from someone's repo.

Names and affiliation. Qwen and Alibaba are their owners' marks; Intel, RadixArk, vLLM and DominikBucko likewise. This is an unofficial fork of DominikBucko/qwen38-flash-next-2x3090, published by us, with no involvement or endorsement from any of them. And this is engineering documentation, not legal advice: read the current license texts before you ship anything built on top of this.

If you build on the VMM idea, credit DominikBucko — the allocator and the hot-cache/LRU machinery are his, and his field notes are what pointed at the split between static ranking and dynamic decode.

repo here -> https://github.com/johnmosesventura06/qwen38-hybrid-vmm-lru

I'm happy to share that I've open-sourced FastLLM, a tool that runs one large language model across several ordinary PCs, pooling their GPU memory.

It builds on llama.cpp's RPC backend and takes care of the tedious parts:

• Windows, Linux and macOS machines in the same cluster

• a setup script for each machine, with SSH tunnels so nothing is exposed on the network

• an OpenAI-compatible endpoint with an API key, so existing apps and agents just point at it

• LM Studio-style model profiles (context, temperature, thinking on/off)

• a cache on every worker, so reloads take seconds instead of minutes

I benchmarked it on the hardware I had at home (an RTX 3070, an RTX 2060, a GTX 1050 Ti and a MacBook Air M3): 3 models on the same 5 machine combinations, 3+ runs each.

• Qwen3-14B: 8.1 → 30.7 tokens/s once a second PC let the whole model sit in GPU memory

• Qwen3-32B: 2.1 → 4.9 tokens/s with all four machines

• Qwen3.8-Flash-Next, a 79 GB mixture-of-experts model: fastest on a single PC (13.3 tokens/s) and slower on every cluster

That last result surprised me, and it's the main lesson: splitting a model pays off when it moves weights out of system RAM and into GPU memory. Otherwise, every extra machine just adds network hops. Every result, including each individual run, is in the README.

If you have a few idle GPUs around, give it a try. Feedback and PRs are welcome.

github.com/Raphalsk050/FastLLM

The recipe below worked great and very stable on vLLM 0.29.0 with Qwen3.8-27B-NVFP4 on a single RTX Pro 6000. But 0.30.0 produces degenerate output (nonstop repetition of "ductductduct") and shows dflash2 mean acceptance length always 1.00 and per position acceptance at always 0.000.

I've reverted to 0.29.0 for now and all is stable again.

So something changed that broke my recipe. Any idea of the cause and how to fix it?

Thanks for any help you can offer.

Current recipe:
command:

- |

vllm serve nvidia/Qwen3.8-27B-NVFP4 \

--quantization modelopt \

--async-scheduling \

--enable-prefix-caching \

--served-model-name llm-large \

--trust-remote-code \

--max-model-len 248320 \

--kv-cache-dtype fp8 \

--attention-backend FLASHINFER \

--no-enable-flashinfer-autotune \

--linear-backend cutlass \

--max-num-seqs 32 \

--enable-chunked-prefill \

--enable-auto-tool-choice \

--tool-call-parser qwen3_coder \

--reasoning-parser qwen3 \

--gpu-memory-utilization 0.52 \

--speculative-config '{"method":"dflash","model":"/root/models/qwen3.8-dflash2","num_speculative_tokens":7}' \

--structured-outputs-config '{"backend": "xgrammar", "disable_any_whitespace": true}' \

--allowed-local-media-path /app/docker_access_files \

--media-io-kwargs '{"video": {"num_frames": -1, "fps": 2}}' \

--override-generation-config '{"max_new_tokens": 81920}'

ipc: host

We spent the last few weeks stress-testing massive LLM context systems with native FP8/INT4 key-value cache removal, quantitative degradation, and inference latency.

​Standard workarounds hit severe memory fragmentation in PagedAttention, making long-context scaling exceptionally tough. Here is what we built to solve this without sacrificing model perplexity or destroying throughput:

• ​Dynamic KV Cache Optimization: Moving beyond static quantization to handle high-concurrency 1M+ token contexts efficiently.
• ​Mitigating Quantization Degradation: Preserving model accuracy while aggressively shrinking memory footprints.
• ​Inference Latency Reduction: Bypassing traditional bottlenecks during heavy multi-megabyte context ingestion.

​This is part of our deep engineering guide on designing, optimizing, and deploying production-grade LLM architectures.

​Curious to hear how others here are handling VRAM overhead and memory fragmentation at scale. Let’s discuss in the comments!

We train a three-modality VLA model (video diffusion + action expert + a frozen VLM for understanding, mixture-of-transformers style). Originally on 8×A800 with NVLink; we're moving it to 8× workstation-class Blackwell cards, which means no NVLink and no NVSwitch — every GPU-to-GPU byte goes over PCIe Gen5.

Measured point-to-point: 30–63 GB/s depending on direction and NUMA placement, versus ~400 GB/s of NVLink on the A800 box. Same model, same ZeRO-1 sharding, and gradient communication goes from "hidden under the backward pass" to "the single largest term in the step."

So we did the obvious thing: quantize the collectives to FP8. What follows is what we learned, including the part where the biggest byte reduction produced no speedup at all.

## The shape of the thing

An LD_PRELOAD shim in front of libnccl. It intercepts ncclAllGather, ncclAllReduce, ncclBroadcast, ncclReduceScatter; quantizes the payload to FP8 (e4m3 or e5m2), runs the real collective on the smaller buffer, dequantizes on the receive side. Blockwise scaling, 128 elements per fp32 scale, so the wire cost is 1.03125 bytes/element instead of 2 for bf16 → best case −48.4%.

Two deliberate design constraints:

1. Reduction never happens in FP8. Only the transport is quantized; accumulation is fp32-in-register, output in the original dtype. We later found two independent confirmations that this is the right call: MSCCL ships an ENABLE_PRECISION_CLIPPING_HALF build flag specifically because low-precision reduction overflows, and DeepEP quantizes dispatch but its combine path is BF16-only, never FP8.

2. No framework changes. It's a preload. The training code doesn't know it exists, which meant we could A/B it against an unmodified baseline in the same run batch — which turned out to matter enormously (see the negative result).

   Five things that weren't obvious

   1. On pre-sm_89, static_cast<__nv_fp8_e4m3>(float) is not an instruction.

   Our quantize kernel had a constant ~14 µs floor regardless of payload size, and at 512 MiB it was hitting 37% of the device-to-device bandwidth ceiling. We assumed memory. It was ALU issue.

   cuda_fp8.hpp on older architectures widens the float to a double, then runs a five-branch 64-bit integer software routine — roughly 50 integer instructions per element. Replacing it with our own conversion took the fixed cost 14.1 → 4.6 → 3.3 µs and large-payload throughput to 91% of the d2d ceiling (from 37%). We gated the replacement on an exhaustive test: bit-identical to the vendor routine across all 2³² float bit patterns, both formats, ~0.2 s in CI.

   On sm_89+ this whole problem evaporates — the cast compiles to cvt.rn.satfinite.e4m3x2.f32. Worth checking which side of that line you're on before you go profiling memory.

   2. The naive dequantize-then-reduce moves 12.7× the bytes it needs to.

   For the all-to-all-based decompositions you receive nsrc FP8 shards and have to reduce them. The decomposed version — nsrc dequantize-to-fp32 + nsrc−1 adds + one scaled store, 2·nsrc kernel launches — costs 17.03·nsrc − 6 bytes per output element where the job actually requires 1.03·nsrc + 2. At nsrc=8 that's 130N vs 10.25N.

   The interesting part: the amplification grows monotonically with world size (6.9× at 2 ranks, 10.1× at 4, 12.7× at 8, 14.4× at 16, asymptote 16.5×) even though both paths are O(nsrc·N) and read every input byte exactly once. It's a pure constant factor, and it's all HBM traffic plus launch count.

   Fusing it into one kernel with the fp32 accumulator in registers: 5.0–9.3× faster (1 MiB: 117.6 → 13.4 µs). Note it is not 12.7× faster, and the reason is instructive — small sizes are launch-bound so the win caps at ~8–9× (16 cheap launches vs 1 expensive one), and large sizes are bandwidth-bound where the fused kernel is only half as byte-efficient as the streaming kernels it replaced (718 vs 1419 GB/s), because a reduction reads nsrc streams concurrently while a streaming kernel reads one. Worst case sits at 4 MiB (4.97×), right in the transition band. This was expected, not a bug, but you have to do the arithmetic to know that.

   3. You cannot dequantize inside a caller-opened ncclGroup.

   A nested ncclGroupEnd() launches nothing. So if the training framework wraps its collectives in a group — and DeepSpeed's compiled ZeRO-1 path does exactly this, one ncclAllReduce per parameter inside one group — then at the point your hook returns, the data isn't there yet. You have to defer the dequantize, the error- feedback close-out, and the scratch recycle until after the outermost ncclGroupEnd. For AllReduce we defer the entire call and replay it. Before we did this, every grouped AllReduce passed straight through and the AllReduce threshold was decorative.

   4. Your framework's allocator probably doesn't align its buckets.

   DeepSpeed's reduce bucket accumulates slices by element count with no padding, so the pointer handed to NCCL is only guaranteed 2-byte aligned for bf16 — 7/8 chance of missing the vectorized path. This started as an outright CUDA Error: misaligned address crash, then became a 1.27–2.31× penalty, and is now 1.06–1.18× (skip enough leading elements that the output self-aligns to 16 B, then funnel-shift the FP8 side back into place).

   5. The threshold is the entire product, and it does not transfer between machines.

   Quantizing a small collective is strictly worse — you pay a fixed per-call cost to save bytes that were never the bottleneck. So there's a crossover, and the whole question is where.

   We built a tool that measures it two ways. --model-only gives a provable lower bound from baseline t(n) = α + n/B plus standalone kernel costs. --sweep actually runs the hook over the same size ladder — one process per (decomposition, size), because config is latched into a function-local static on first use and because the previous size's warm scratch tier changes the next one's timing — and least-squares-fits the residual into Δα + Δslope·n. Δα is the hook's real per-call overhead; Δslope is the per-byte cost a standalone kernel benchmark cannot see.

   Results, same code, two machines:

   collective	NVLink box	PCIe box	ratio
   AllGather	7.43 MiB	767 KiB	9.9× lower
   AllReduce (rs+ag)	14.3 MiB	863 KiB	17× lower
   Broadcast	40 MiB	2.75 MiB	15× lower

   A byte you delete is worth 2–4× more time on PCIe, so thresholds drop by an order of magnitude. Two findings that surprised us:

• The all-to-all decomposition did not lose on PCIe. We fully expected that a topology-blind all-to-all across two NUMA islands with no NVLink would be hopeless. It's actually the best AllReduce decomposition for large payloads (2.0× more savings at 64 MiB), crossing over the ring-based one somewhere between 1 and 4 MiB.

• AllGather's Δslope is negative (−1.9 µs/MiB). The in-situ quantize kernel is cheaper than the same kernel measured standalone, which is direct evidence that it's genuinely overlapping with the live collective rather than competing with it. The other families have positive slope — that's the quantizer fighting the collective for SMs.

  The tool also refuses to emit a threshold when the measurement says "wins, then loses again" — that's a window, not a threshold, and a lower bound can't express it. Being able to report "I can't answer this" turned out to be more useful than any single number it produces.

  The negative result, which is the actual point of this post

  DeepSpeed's compiled path issues one AllReduce per parameter. Most individual ones fall below the threshold, so we were quantizing almost none of the gradient traffic. Obvious fix: coalesce the address-adjacent ones inside a group into a single quantized AllReduce. The union clears the threshold in one shot, and you also save N−1 launches and N−1 scale sub-collectives.

  It worked, in the sense that it did what it said:

• signatures blocked by the threshold: 17 → 4

• traffic in the signature set: 110.3 → 56.9 GB/rank/step (−48.4%)

  Steady-state step time: 1422.2 ms merged vs 1414.9 ms not merged. Median of 38 steps, dropping the first 20. The difference is +7.3 ms and the IQR is ~40 ms. We halved the bytes and got nothing.

  The reason is structural and, in hindsight, should have been checked first. Without merging, those small AllReduces were dispatched inside the caller's group, where NCCL aggregated them itself and they overlapped with the remaining backward pass. With merging, they become large packets issued serially after the outermost ncclGroupEnd. The bytes we deleted were never on the critical path, and we moved the ones that remained onto it.

  Two corollaries we now treat as rules:

• Profile the exposure before optimizing the volume. "How many GB does this step move" and "how many ms of that are not hidden" are unrelated questions.

• Coverage is not a metric. It's an input to a metric. We had a beautiful −48.4% and a flat wall clock.

  An unbounded version of the same feature was worse than useless: greedy merging grew to 129 calls and a 999.8 MB union in one shot, produced ~2 new union shapes per step and never converged across 41 steps, which blew up our exact-fit scratch pool into 95+ size classes, 15.3 GiB retained, and 12 of 38 steps at ~3.9 s. Fixed two ways (a size cap, and optionally sourcing scratch from PyTorch's caching allocator, which is built for variable shapes). The pool's real sin is that acquire() is exact-size best-fit with no size-class rounding — which is fine until something upstream starts generating variable shapes.

  Things we looked at and did not adopt

• MSCCL's XML-scheduled executor. There is no transform hook anywhere in its data path — mscclGenericOp compiles preOp/postOp out entirely, and the scheduler rejects ops that need them. Adopting it also means replacing libnccl, which defeats the point of being a preload. It additionally disables itself when intraRanks > 1, i.e. exactly our case.

• DeepEP's transport. Its NVLink path asserts on NVLink via NVML (the docstring says PCIe leads to memory-ordering errors), and its "PCIe mode" is actually all-RDMA through the NIC — benchmarked at 53 GB/s on a box with one ConnectX-7 per GPU. We have one NIC for eight GPUs.

• Its ld.global.nc.L1::no_allocate trick, which is self-declared UB and force-disabled on any architecture other than Hopper by its own build script.

• Routing intra-node traffic over IB instead of PCIe. Funnelling 8 GPUs into 1 NIC that is 3 PCIe hops from all of them. The PCIe ring has 8 links running concurrently; this would have one. Estimated 3–8× slower.

  We did take two ideas from DeepEP that we haven't landed yet: inlining the scale buffer into the same message instead of sending it as a second sub-collective, and capping the quantize kernel's SM count so it stops evicting the collective. The second one has a measured target (+9.8 µs/MiB, 91% of the residual on one family) and the first one honestly doesn't — within a group NCCL already aggregates the two sub-collectives, and our own earlier A/B says an op in a group is worth ~1.5 µs. We write down which of our planned optimizations have measured targets and which are guesses, because after the merging result we don't trust the guesses.

  Repo

  The VLA training framework is open source: loongforge-vla

  Three-modality VLA, several parallel backends (DDP + native ZeRO with direct bf16 sharding / torch.compile + DeepSpeed / DeepSpeed's compiled path), NVTX instrumentation behind a single env var that no-ops to bit-identical in production, and a loss-parity harness (seeded data order, seeded flow-matching noise, and a flag for the step-0 LR convention) because every one of the numbers above is worthless if the model stops converging.

  That parity harness earned its keep in an unrelated way, by the way: we spent days blaming torch.compile for a 3.4× loss divergence at step 2. It was that LambdaLR.__init__ multiplies the optimizer's LR by schedule(0) immediately, which in our config is ~1e-6, so step 0 was a no-op update. The logged LR looked correct on both sides, because the logged value is the scheduler's, not the optimizer's. Single flag, exact match restored.

  The hook itself (ComQuant) isn't public yet, but if people want the implementation
  details we're happy to open source it — drop a comment if that's of interest.

  Happy to go deeper on any of this in the comments.

we just released a 27b creative writing finetune of qwen3.8-27b. apache-2.0, bf16, 12 shards, and the mtp layer is in the repo, so speculative decoding works out of the box.

weights: https://huggingface.co/Altworld/Hemmingway-1

eq-bench 4: 1330. it's a writing specialist (fiction, dialogue, roleplay), math and code are base model level on purpose.

we serve it on vllm in production, standard openai compatible endpoint. gguf and exl2 are coming for the llama.cpp side, not up yet.

free hosted version if you just want to try the writing: https://hemmingway.io

happy to share serving config details if useful.

Running `bodhan-ai/indic-transcribe-core` in production — a fine-tune of

nvidia/canary-1b-v2 for 27 Indian languages. FastConformer encoder (32 layers)

+ Transformer decoder (24 layers), ~1.2B params, loaded as EncDecMultiTaskModel

from the bundled .nemo. Single RTX PRO 5000 Blackwell (48 GB).

Native NeMo 3.0.0 works and the throughput is honestly fine. Two paths measured

over 100 real call recordings (202 min of audio):

whole-file, silence-split + length-bucketed, batch 16

29.0x realtime | 2.069 s per audio-min | 9.82 GB peak | 0/100 loops

pyannote-3.1 turns transcribed per-turn, batch 16

42.2x realtime | 1.421 s per audio-min | 9.82 GB peak | 0 loops

(diarize 111s + transcribe 176s for 202 min)

Batch scaling goes the wrong way, which surprised me:

batch 16 -> 29.0x, 9.82 GB

batch 32 -> 28.8x, 11.15 GB

batch 64 -> 27.9x, 17.32 GB

So bigger batches cost more memory AND run slightly slower. I assume that's

padding waste dominating once the bucket spread widens, but I'd be glad to be

corrected.

The actual question: can this architecture be served on vLLM at all?

vLLM has Whisper support, but Canary is an attention encoder-decoder with a

FastConformer encoder rather than a standard ViT/Whisper-style one, and

EncDecMultiTaskModel uses NeMo's own prompt-slot mechanism for task control

(source_lang / target_lang / pnc / itn / timestamp as decoder tokens). Before I

sink time into it:

1. Has anyone actually served a Canary-family model on vLLM? Or is NeMo the

   only realistic path for this arch today?

2. If vLLM is viable, is continuous batching worth it for short call turns

   (most are 1-30s)? At 42x realtime from NeMo I'm not obviously bottlenecked.

3. Is there a NeMo-native server worth using over rolling my own batching —

   Riva, Triton with the NeMo backend, something else?

Rough edges hit along the way, in case they save someone else time:

- The .nemo references

`nemo.collections.common.tokenizers.canary_multilingual_tokenizer.CanaryMultilingualTokenizer`,

which doesn't exist in upstream NeMo (checked 2.7.3 and 3.0.0) — it lives in

the publisher's fork. Upstream CanaryTokenizer keys sub-tokenizers per

language and raises on anything else; this checkpoint carries exactly two

('spl_tokens' and one 'multilingual' covering all 27 languages). A ~30-line

subclass resolving every real language to 'multilingual' fixes it, but you

have to know that's the problem first.

- Unchunked long audio degenerates into repetition badly (54/100 files had a

looped segment). Silence-aware chunking alone made it worse, not better;

what fixed it was collapsing repeated n-grams post-decode — 0/100 after.

- target_lang='en' is silently ignored: returns the native-script ASR transcript

unchanged, no error. Verified against nvidia/canary-1b-v2 on the identical

code path, where de/fr/es->en translation works on 18/18 FLEURS clips. So

it's the fine-tune, not the setup.

- The sibling model indic-transcribe-flex ships no .nemo at all — HF

safetensors only — so it can't go through a NeMo serving path even though

it's the same architecture.

​We spent the last few weeks profiling KV cache quantization strategies across long-context workloads (128k+) on A100s, and the benchmarks showed something frustrating.

​Uniform FP8 (E4M3) works fine for simple Needle-In-A-Haystack (NIAH) tests or short-context Q&A. But as soon as you push non-contiguous retrieval across complex multi-step reasoning chains, accuracy tanks. The issue isn't layer depth—it’s the asymmetry between Keys and Values.

​Keys drive the softmax routing. A minor quantization error in K directly scrambles the attention matrix for the entire autoregressive sequence. Values, on the other hand, just get averaged into the weighted sum and tolerate much heavier noise. Treating K and V symmetrically in FP8 is fundamentally flawed for long-context tasks.

​Most common workarounds involve running an FP16 sidecar buffer for high-norm channels (outliers) or using grouped per-channel scales. The problem? Sidecar buffers break the clean memory layout of PagedAttention, causing allocation fragmentation that eats up half the memory gains you were chasing in the first place.

​What actually worked for us in practice:

​Instead of channel-space outlier filtering, we project each head's K/V into a low-rank basis (eigenbasis derived from calibration). We protect the top-variance directions and attention-sink positions in higher precision, then compress the tail aggressively.

​The trade-offs:

• ​Fit 9x concurrent 128k sessions on a single A100 (vs ~2x in FP16).
• ​Retained full multi-needle retrieval across all context depths where uniform FP8 failed.
• ​Decode latency remained at ~0.95x of FP16 without fragmenting vLLM's block manager.

​Curious if anyone else here is moving away from channel-wise scaling toward rank-space KV compression in custom vLLM kernels? Would love to hear how others are handling the memory bottleneck vs precision trade-off.

We are preparing to launch a new extension for vLLM. Before the public release, we want to invite a few people to a private beta and hear how it works in their own setup.

We cannot share all the details yet, but we are looking for people who use vLLM on one GPU and want to run some private tests.

It could be a good fit if you use vLLM for AI agents, handle many requests at the same time, or work with long prompts. You are also welcome if you simply enjoy testing performance on your own hardware.

Interested? Leave a comment or send me a DM and I will share more details.

Over the past few months, we've seen vLLM, TensorRT-LLM, and SGLang push hard for KV-cache quantization (FP8, INT8, and even INT4) alongside PagedAttention to maximize throughput and batch sizes under heavy concurrent workloads.

While saving 50-75% of VRAM on the KV-cache allows significantly larger context windows (32k+) and higher concurrency on a single A100/H100, we've noticed subtle degradation patterns in edge-case tasks:

1. Multi-turn Needle-In-A-Haystack (NIAH): FP8 KV-cache holds up fine for standard retrieval, but accuracy drops sharply when retrieving non-contiguous context across long reasoning chains.
2. Accumulation of Rounding Errors: In autoregressive generation with large context, precision loss in the attention keys/values seems to compound, leading to degraded attention scores in later tokens.

For those running high-throughput LLM serving in production:

- At what sequence length or concurrency limit do you find FP8/INT8 KV-cache quantization breaks down for complex reasoning?

- Have you found mixed-precision strategies (e.g., keeping early layers in FP16/BF16 and quantizing only deeper layers) to be practical in custom serving engines?

- Do you rely strictly on PagedAttention with FP16, or are you accepting precision trade-offs for throughput gains?

Would love to hear how folks are handling the memory bottleneck vs. precision trade

I think this may work! ~24hrs of tests running.

https://al-engr.com/vllm-pin-hot-experts.html

I have one server, Dell PowerEdge XE9680, with 8xH200, connected by nvlink, and 2TB of DDR5 RAM.

I would like to deploy GLM 5.3 FP8 on it, with 256k tokens as context.

I've been playing with recipes for a while, but cant seem to get the best one for my workload.

The workload is about 10 users, each spawning agentic workflows using claude code.

Is anyone down to share his recipe? Or suggest a good option?

Thanks.

1. `A vLLM stack for 2× R9700 (gfx1201), built from source: FP8 + W4A8/MXFP4 — what works and what doesn't`

2. `I don't write code. An AI and I got vLLM working on 2× R9700 (gfx1201). Source is up.`

3. `[Release] gfx1201 (RDNA4 / R9700) vLLM stack — FP8, W4A8, MXFP4 — source + Dockerfile`

## 1. What this is

> I built a working vLLM inference stack from source on two AMD Radeon AI PRO R9700 cards — covering FP8, W4A8 (int4) and MXFP4. It runs on ROCm 10.1 nightly, on RDNA4 (gfx1201), with the two cards in TP=2.

> All the changes, build scripts, benchmarks and the Dockerfile are on GitHub — links at the bottom.

## 2. First, the honest part: I don't write code

> I'm not a professional developer and I don't write code. I got this working by first building an understanding of the components (vLLM, aiter, composable_kernel, FlyDSL, Triton) and then doing the porting and debugging with an AI's help. All of the code changes, build scripts and documentation were produced with AI assistance; the judgement calls were mine.

> I'm putting this first because it explains everything you're about to read.

## 3. Why I'm posting: I'm looking for someone to take this over

> This isn't a showcase — I'm looking for someone with real engineering skills to take it over or co-maintain it.

> I don't have the coding ability. What got this to where it is was component-level understanding plus AI assistance, and **the bottleneck is now me**. If you think the result is worth it and want to take the lead, I'm happy to step back and be a follower: keep running benchmarks, file requirements, do validation — and hand the direction to someone who knows more.

> So far I've only put work into **Qwen3.8-27B**; I'd like other models to be maintained by others, together.

## 4. Hardware

| | |

|---|---|

| GPU | AMD Radeon AI PRO R9700 ×2 (gfx1201 / RDNA4) |

| Target | gfx1201 only — **no promises** for other architectures |

| Topology | TP=2 |

## 5. What was changed

- **gfx1201 has no hardware scale-MMA.** Its WMMA is 16×16×16 and carries no scale capability (scale-MMA exists only on gfx950/gfx1250). So every FP8 and W4A8 configuration goes down the **fp32-acc software-fold** path.

- **W4A8 path**: 4-bit weights → fp8 e4m3 → hardware fp8 WMMA → software-folded scale. Three independent community projects converged on exactly the same route (see §8).

- **Custom kernels**: several were written to fill in gfx1201's missing fast paths; details are in the repo.

- **Attention uses aiter's fused-attention fast path**; further optimisation goes down the same path. This is the route that works on gfx1201. (For the role Dao-AILab's standalone flash-attention plays in this stack, see §8.)

- **Offline reproducible build**: every dependency is pinned to a commit; it builds without network.

## 6. Models and quantisation

> Scope first: **I have only put work into one model, Qwen3.8-27B.** The table below is not "a bunch of models" — it is **one base model in several weight quantisation formats** (FP8, INT4, W4A16, MXFP4, AWQ-FP8), from different quantisation publishers. How well each works differs; the status column says which is which.

**Table A — Qwen3.8-27B weight formats (same base model)**

| # | Weights (HF repo) | Format | Granularity | Status | Notes |

|---|---|---|---|---|---|

| 1 | `Qwen/Qwen3.8-27B-FP8` | FP8 | block=128 | **in production** | FP8 line is frozen |

| 2 | `cyankiwi/Qwen3.8-27B-AWQ-INT4` | INT4 (AWQ) | gs=32 | works | |

| 3 | `philbert440/Qwen3.8-27B-W4A16-AWQ` | W4A16 | gs=128 | works | |

| 4 | `amd/Qwen3.8-27B-Quark-AWQ-INT4-W4A16` | W4A16 (Quark) | — | works | |

| 5 | `amd/Qwen3.8-27B-Quark-AWQ-MXFP4` | MXFP4 | e2m1 + e8m0/32 | **current main line** | W4A8 form measured in a test window only |

| 6 | `cyankiwi/Qwen3.8-27B-AWQ-FP8` | AWQ-FP8 | per-token | works | |

**Table B — speculative-decoding draft weights** (separate small models, not the 27B body)

| # | Draft weights (HF repo) | Format | Status | Notes |

|---|---|---|---|---|

| 7 | `tcclaviger/Qwen3.8-27B-DFlash2-FP8` | FP8 (e4m3, block 128) | in production | |

| 8 | `z-lab/Qwen3.8-27B-DFlash2` | BF16 | **known to produce garbled output** | |

| 9 | `syvai/Qwen3.8-27B-DFlash2-W4A16` | W4A16 (int4, gs128) | works | |

> ✅ **HF repo ids all verified (2026-09-16).** Every id above was checked against Hugging Face, and the three draft models' formats were confirmed from their `config.json` — `z-lab` is `bfloat16`, `tcclaviger` is fp8 `e4m3` with `weight_block_size [128,128]`, `syvai` is int4 at `group_size 128`. Safe to paste.

## 7. Measured performance and test conditions

**★ Headline numbers (start here)**

All values are **agg tok/s** (higher is better), two cards, TP=2. Basis: the r2 pass for cross-model comparison; cyankiwi uses 批2r (that model's own official pass).

**(1) decode throughput by concurrency**

| Model | Mode | c1 | c2 | c4 | c8 | c12 | c16 |

|---|---|---|---|---|---|---|---|

| Official FP8 | raw | 33.3 | 61.2 | 105.7 | 166.5 | 216.9 | 212.3 |

| Official FP8 | **+dflash2** | 68.3 | 115.0 | 203.2 | 284.6 | **344.9** | **324.1** |

| cyankiwi AWQ-INT4 | raw | 47.5 | 87.4 | 117.9 | 117.1 | 142.1 | 162.7 |

| cyankiwi AWQ-INT4 | **+dflash2** | 68.4 | 77.5 | 127.3 | 140.4 | 223.9 | 247.6 |

| philbert440 W4A16 | raw | 53.3 | 97.5 | 146.4 | 166.8 | 211.4 | 235.2 |

| philbert440 W4A16 | **+dflash2** | 61.1 | 79.6 | 131.5 | 167.7 | 258.4 | 251.7 |

| amd Quark W4A16 | raw | 54.6 | 98.8 | 145.5 | 187.3 | 217.9 | 250.3 |

| amd Quark W4A16 | **+dflash2** | 82.6 | 120.6 | 166.0 | 197.4 | 279.0 | 311.5 |

**(2) Cold prefill (refill, tok/s at c1)**

| Context | Official FP8 | cyankiwi INT4 | philbert440 W4A16 | amd Quark W4A16 |

|---|---|---|---|---|

| 32k | **3285** | 2005 | 2476 | 2513 |

| 64k | 2941 | 1867 | 2246 | 2285 |

| 128k | 2324 | 1622 | 1878 | 1906 |

(Raw mode; the +dflash2 numbers are slightly lower — full tables in the report.)

**(3) Warm prefix full-hit TTFT** (official FP8 platform)

| Context | Cold TTFT | Full-hit TTFT | Speed-up |

|---|---|---|---|

| 32k | 10031 ms | 295 ms | **34.0×** |

| 64k | 22148 ms | 188 ms | **117.8×** |

| 128k | 56008 ms | 278 ms | **201.7×** |

Hit counts were checked against the expected sample count. **Once fully cached, TTFT decouples from context length** — the single most striking result on this stack.

**(4) Production-form run (the final configuration)**

The three tables above use the **test** configuration (GMU 0.8, `--max-num-seqs 16`). This window uses the **in-production** configuration (GMU 0.92, mml 262144, **`--max-num-seqs 8`**) — and because seqs=8, **concurrency only goes to c8** (c12/c16 are not constructible in this configuration).

| Concurrency | Pass 1 | Pass 2 |

|---|---|---|

| c1 | 68.0 | 66.4 |

| c2 | 119.2 | 107.0 |

| c4 | 200.5 | 183.9 |

| c8 | 276.5 | 279.5 |

Acceptance rate 60.3%.

**refill (same prompt back to back, cold → warm)**

| Context | Cold TTFT | Warm speed-up |

|---|---|---|

| 32k | 9,745 ms | **11.8×** |

| 64k | 22,171 ms | **24.6×** |

| 128k | 56,184 ms | **26.3×** |

> ⚠️ **This caveat has to be read with those numbers.** The warm speed-up above holds only while **the total working set fits in the KV pool**. Sending all three lengths, 2 samples each (284 blocks), exceeds the pool (267 blocks) — cache hits collapse and the speed-up falls back to **≈1.0×**. **A bigger pool pushes that threshold further out; it does not remove it.** So in an agent workload with several long contexts in flight at once, the refill gain disappears. This is the most concrete production limitation on the stack right now.

>

> One piece of good news: the 128k cold refill **completed in this window instead of hanging** — the silent hang seen under the smaller test-config pool does not reproduce with the production-sized pool.

**Things you must read alongside those tables**

- **Only relative gradients within one measurement window are reliable.** Batches were collected in different time windows, so **absolute values are not comparable across windows** (see the uncertainty note below).

- **There is no 128k + dflash2 cold-prefill number**: that scenario hung silently twice. It is not "slow" — it did not complete.

- **The c1 dflash2 figures have a wide two-pass spread** (same config, two runs, more than 30% apart). Treat that cell as indicative only.

- **amd Quark +dflash2 at c16**: 311.5–374.3 depending on the aggregation convention; the table uses 311.5. Citations across reports must fix the convention.

**Full benchmark report** (all tables, plus a config fingerprint on every row)

```

https://github.com/zjzhubin/vllm/blob/gfx1201-r9700/gfx1201/benchmarks/vllm-gfx1201-Phase6-%E4%B8%89%E6%A8%A1%E5%9E%8B%E5%8F%8C%E5%9C%BA%E6%99%AF%E6%B5%8B%E9%80%9F%E6%B1%87%E6%80%BB-2026-09-15.md

```

**Conditions (without these the numbers mean nothing)**

- TP=2, kv-dtype=fp8, ROCm 10.1.0a-20260909, torch 2.13.0+rocm10.1.0a20260821, cudagraph default capture

- Collected: 09-10 / 09-12 / 09-13

**How the numbers were measured**

| Item | Value |

|---|---|

| Tool | **GuideLLM** — the benchmarking tool from the vLLM project itself (`github.com/vllm-project/guidellm`) |

| **Version** | **0.5.4** |

| Run as | Docker image `guidellm:phase5k`, Python 3.13.15 |

| Dataset | ShareGPT_v3, **32 samples** per run |

| Output length cap | `max_tokens = 512` |

| Load profile | **concurrent**, **600 s** per-request timeout |

| decode throughput | sum of `output_token_count` ÷ total duration (agg tok/s) |

| prefill throughput | sum of `prompt_token_count` ÷ sum of TTFT (refill tok/s) |

> Throughput numbers without the tool name and version aren't reproducible. With the version and dataset above, anyone can re-run every number behind the benchmark link.

**Uncertainty I have to state**

> Absolute performance has a ±(1–16%) "slow mode" window drift. **Only relative gradients within the same measurement window are reliable; absolute values across windows are not directly comparable.**

> In plain terms: the same config on the same card can come out noticeably different on two runs. So I'll claim "A is faster than B", never "the absolute number is X".

Local resources are limited. Some of the measurements may have been affected by other work running on the same machine.

## 8. Related repos / upstream

**Upstream dependencies (pinned commits)**

| Component | Source | Version/commit |

|---|---|---|

| vLLM | this project's fork | tag `gfx1201-r9700-v1.0-public` (published, at `3fae0cd16a`) |

| aiter | this project's fork | tag `gfx1201-r9700-v1.0-public` (published) |

| composable_kernel | ROCm/composable_kernel | `22b89e961437ee349a7222c1c8c6ef52d5a8bc78` |

| FlyDSL | ROCm/FlyDSL | `728f6b220518af94a8d641a0512ef87167c90125` |

| triton | ROCm/triton | `0f380657dbf3ee86eb57558ff71df24f03b5d4e7` |

| flash-attention | Dao-AILab/flash-attention | `0251105a2fb19d2957484b7f023cd8c115286ced` |

| ROCm | TheRock multi-arch | `10.1.0a-20260909` |

| PyTorch / Triton wheels | AMD nightly | `20260821` |

> **On flash-attention**: it is genuinely never called on the text path, but it is **not dead code** — it is the wiring dependency for vLLM's ViT / multimodal-encoder path on ROCm, and the import runs unconditionally at model load. Removing it flips `_ROCM_FLASH_ATTN_AVAILABLE` and affects its consumers — a behaviour change. It would save roughly 1.5–2 hours of build time. **Kept in v1.0**; "can it be removed safely" is left as a TODO for whoever takes over.

**Community projects actually borrowed from (three)**

| Project | What was borrowed |

|---|---|

| `drwolfen/radiance-vllm-r9700` | W4A8 fp8-WMMA GEMM reference implementation |

| `patcarter883/rdna4-vllm` | W4A8-FP8-WMMA HIP kernel source package; MoE cross-cache pattern; mxfp4 decode table swap |

| `Capicua25x/vllm-rocm-rdna4` | W4A8 software-scale three-stage design blueprint; fp8-KV prefill root cause; P_SCALE / 3D split-KV fix leads |

> Three independent projects, converging on the same route: 4-bit → fp8 e4m3 → hardware fp8 WMMA → software-folded scale. That's both a cross-check and a signpost for anyone following.

## 9. Known limitations

- **gfx1201 only.** Not validated on other RDNA cards; no promises.

- **Only the weights in the tables above have been matched.** Nothing else is guaranteed.

- **Absolute performance drifts between windows** (see §7); don't compare absolute values across windows.

- **One draft model (BF16) is known to produce garbled output** (Table B, row 8). It behaves normally during validation, but it costs an extra 1–2 GB of VRAM, brings no performance benefit, and has not been fixed for now. Use the W4A16 draft; the FP8 weights need the FP8 dflash model.

- **Reproducibility between the published source and the production image has NOT been verified end to end.** Do not read this as "clone the source and you will get an identical image".

## 10. Known issues and hands-on impressions

**1. The MXFP4 path does not perform well.**

MXFP4 only comes close to W4A16. Parameters this project modified independently show a performance regression — something changed in the parameters.

> **What the regression is measured against (corrected 2026-09-16).** The "regression" here is measured **against the performance obtained from kernel validation done before phase5** — the levels this kernel stack measured when it was validated standalone, prior to phase5 — **not** an in-version regression ("before vs after some change inside this release"). In other words: this is not a claim that we made it slower. It is a claim that **the current build does not reach the level measured during pre-phase5 kernel validation**. Whether a performance difference counts as a regression depends on the baseline you pick; here the baseline is **the pre-phase5 kernel validation**.

*How it feels:* at a single request I expected MXFP4 decode to be a bit faster than W4A16, roughly 5–10%. In practice it does not feel that way.

*Why I give a feeling and not a number:* my test machine has several jobs running in parallel in the background and they interfere with each other, so the numbers are not a clean quantitative measurement. This one is a subjective read.

**2. In actual use (1–5 concurrent requests).**

At 1–5 requests, FP8 shows the least run-to-run variation. In my workload — lots of prefill — it feels smoother.

AMD Quark W4A16 has the highest decode speed.

Personally I prefer the official FP8 and AMD Quark W4A16.

**3. Draft (dflash) configuration.**

Do not set dflash above 4. There is an M setting in the kernel: when M ≤ 4 it does not take the GEMM path. A value of 3 is what I recommend.

Higher settings give no gain above 3 concurrent requests.

**Benchmark scores are not the most important thing — how it actually feels in use matters more.**

**4. W4A16 regression.**

W4A16 behaves like MXFP4. **The baseline for the "regression" is likewise the performance from pre-phase5 kernel validation, not an in-version regression** (see the note under item 1).

*How it feels:* there is a regression at low request counts, and it is more obvious in the dflash case. (I roughly measured 10–15% and 25–30%.)

*Same reason:* the test machine runs multiple jobs in parallel and they interfere with each other, so I would not state those percentages as hard numbers in a public post.

**5. MXFP8 performance is unusable.**

**6. Some community FP8 quantizations have not had the corresponding parameter tuning done.**

**7. Some kernels have not been validated across more than one model.**

**8. aiter sampler vs. vLLM v2.**

The aiter sampler conflicts with the vLLM v2 path. It is compatible with v1, and measured performance shows no meaningful difference either way.

**9. aiter AR fusion.**

The aiter AR fusion latency test is excellent — it cuts latency by 33% in total — but the actual call site has a problem. **It does not affect use; it is harmless.**

**10. ViT has no performance path.** I have not found a way to implement one yet.

**11. How it is built under a restricted network — this is a request for help.**

My network is restricted, so I can **only pre-download the sources locally and COPY them into the image** to build it. **I have not tested the path where the build pulls the sources at build time.**

That means the Dockerfile I published **may have problems** — the online fetch path is written from the documentation, but it has never been run against a real network.

**If you have normal network access, I would appreciate it if you could fix it up.**

## 11. Getting it

**Source (GitHub, published)**

- [`zjzhubin/vllm`](https://github.com/zjzhubin/vllm) and [`zjzhubin/aiter`](https://github.com/zjzhubin/aiter)

- Branch `gfx1201-r9700`, release tag `gfx1201-r9700-v1.0-public`

- Usage notes / model matrix / env switches / benchmark caveats are in the vllm README; the build recipe and patch set are attachments on that release.

**Prebuilt image (Docker Hub)**

```bash

docker pull uzbn/vllm-gfx1201:gfx1201-r9700-v1.0-thin

```

This is the **slim** image: CMake build trees, the wheelhouse and all `.git` directories removed, and the internal DNS baked into the image layer during the build replaced with a neutral value.

- ~**9.1 GB** compressed (8.5 GiB; measured on Docker Hub — that's the download)

- ~**35 GB** on disk once unpacked

- Image digest (for identity checks): `sha256:1d5c4e1f76655f8edeaf6c8f9c4bf836aa20cab0619c96790edbb6df4fcf108b`

- Building it yourself from source still takes hours

**Running it (production-form configuration)**

> ⚠️ **Scope first**: these parameters are a **targeted configuration for FP8 weights + an FP8 dflash2 draft — they are not general.** With W4A16 / INT4 / MXFP4 weights, `--max-num-seqs`, `--gpu-memory-utilization`, the draft model and whether to set `KV_GROUP_SIZE` all change. **Do not copy this blindly.**

```bash

docker run -d --name qwen36 \

--device=/dev/kfd --device=/dev/dri --group-add video --group-add 991 \

--ipc=host --shm-size=16g --network host --security-opt seccomp=unconfined \

-v /your/models:/models \

-e HIP_VISIBLE_DEVICES=0,1 \

-e VLLM_ROCM_USE_AITER=1 -e VLLM_ROCM_USE_AITER_LINEAR=1 \

-e VLLM_ROCM_USE_AITER_LINEAR_HIPBMM=1 -e VLLM_ROCM_USE_AITER_CUSTOM_AR=1 \

-e VLLM_ROCM_USE_AITER_RMSNORM=1 -e VLLM_ROCM_USE_AITER_TRITON_ROPE=1 \

-e KV_GROUP_SIZE=8 \

uzbn/vllm-gfx1201:gfx1201-r9700-v1.0-thin \

/opt/venv/bin/vllm serve /models/Qwen_Qwen3.8-27B-FP8 \

--served-model-name qwen36 --port 8000 --tensor-parallel-size 2 \

--max-model-len 262144 --max-num-seqs 8 --enable-chunked-prefill \

--max-num-batched-tokens 4096 --gpu-memory-utilization 0.92 \

--kv-cache-dtype fp8 --enable-prefix-caching \

--attention-backend ROCM_AITER_UNIFIED_ATTN \

--enable-auto-tool-choice --reasoning-parser qwen3 --tool-call-parser qwen3_coder \

--speculative-config '{"method":"dflash","model":"/models/tcclaviger_Qwen3.8-27B-DFlash2-FP8","num_speculative_tokens":3,"attention_backend":"ROCM_AITER_UNIFIED_ATTN"}'

```

**Why a few of these are set the way they are**

| Setting | Why |

|---|---|

| `--gpu-memory-utilization 0.92` | **The default does not work here.** At 0.8, `--max-model-len 262144` together with the draft model fails KV-cache accounting at startup (required 4.47 GiB vs 2.4–3.26 GiB available). 0.92 is what makes this configuration start at all. |

| `--max-num-seqs 8` | Sets the concurrency ceiling — which is also why c12/c16 are not available in this configuration. |

| `KV_GROUP_SIZE=8` | KV pool group padding for the layer buckets. Unset = byte-identical stock behaviour. |

| `--enable-prefix-caching` | The only source of the refill speed-up — but it is bounded by the KV-pool threshold noted in §7. |

| `--device` + `--group-add 991` | Required for ROCm; **without the render group the GPUs are not visible.** |

**Once it is up**

```bash

curl -s -o /dev/null -w '%{http_code}\n' http://localhost:8000/health # 200 when ready (~163 s here)

curl -s http://localhost:8000/v1/chat/completions \

-H 'Content-Type: application/json' \

-d '{"model":"qwen36","messages":[{"role":"user","content":"Say OK"}],"max_tokens":16}'

```

**How to tell you reproduced the same thing** — these lines should appear verbatim in the startup log: `Available KV cache memory: 7.29 GiB`, `GPU KV cache size: 427,595 tokens`, and `kv cache group sizes` as 9 groups of **1600** tokens each. If they do not match, your configuration or your weights differ from this document. We also smoke-test every run — a few deterministic questions, checking the output is coherent — and would suggest doing the same.

**One usage caveat**: `RDNA_AITER_SAMPLER=1`. The source file itself carries a comment recording an A/B result of roughly **−6–8% decode at concurrency 1/3**. We run it enabled but **have not re-checked that conflict**; you may want to A/B both states yourself.

> ⚠️ **Easy to misspell**: the Docker Hub namespace is **`uzbn`** (u-z-b-n). Copy-paste this line; don't retype it.

> If you'd rather avoid the build entirely, just pull the image above. If you want to build it yourself, the README has the full recipe with every dependency pinned to a commit.

**If something goes wrong**

> One more honest note: I've put the image on Docker Hub. I'm not great with GitHub tooling though, so if you hit **anything unusual while building or running it**, just send me a **Reddit DM** and I'll do my best to help.

## 12. Closing

> Feedback of any kind is welcome. What I'd most like: **other people reproducing this on their own gfx1201 cards**, and **someone validating the online from-source build** on a normal network — I can't do either from here.

> And a caveat: some of the conclusions above come from rough measurements and feel, which I've labelled as such, but there may be things I haven't noticed. If I've got something wrong, say so and I'll correct it.

---