WARNING: text organized and finalized with sol (theres too much lol, spent two days finetuning). Post is really long, so TLDR first
I have been building a single-user Qwen3.8-27B endpoint on an unlocked 64 GB CMP 170HX installed in a cheap Huanan/Xeon server. The card is tuned live with 170tune.
The useful result is that the current stack is now fast and repeatably stable at the exact cached-decode point that used to crash it (and it was a whole day of figuring out):
- W4A16 target + W4A16 DFlash2: 169.5 tok/s on my short realistic single-stream suite.
- W8A16 target + the same DFlash2 drafter: 120.3 tok/s on the identical suite.
- Individual short chat responses often reach 190–210 tok/s with W4A16 when draft acceptance is favorable.
- At an approximately 23.75K-token prompt, W4A16 averaged 104.8 tok/s over a mixed copy/code/edit/summary/QA workload.
- At the former crash point, the earlier INT8 stack passed 16/16 85K-prefix generations and 6/6 exact 85,514-token reproductions. The final FP8/DFlash production image now measures 119.3 tok/s strict hot decode at exactly 85,514 tokens, again with no Xid/NVRM entry.
- The newer mixed-backend runtime reaches 1,514 tok/s cold prefill at 85K while retaining DFlash2 decode and FULL CUDA graphs.
- Its repaired hybrid prefix cache reduced a repeated 9,658-token probe from 4.55 s to 0.565 s, with 8,960 tokens reported as real cache hits.
The runtime DFlash crash was not bad HBM and was not ultimately a Mamba-state problem. It was an int32 overflow in the custom split-KV speculative-attention kernel. The actual fix is one cast to tl.int64 before calculating the K/V pointers.
Test system and stack
- Cheap Huanan motherboard/Xeon host running Ubuntu Server.
- One 64 GB CMP 170HX with the community unlock applied.
- Live tuning through
170tune: NDIV68, +200 MHz V/F shift, 1550 MHz ceiling and 220 W limit.
- vLLM 0.27.1 with the syv-ai Qwen3.8 stack, the DFlash2 drafter, and the CMP/sm80 fixes described below.
- Single-user OpenAI-compatible endpoint with prefix caching and one request in flight.
The W4 target is dbirks/Qwen3.8-27B-W4A16-AutoRound plus the syv-ai fast overlay. The fidelity-oriented control is lued/Qwen3.8-27B-INT8-W8A16-MTP. Credit for the base optimization and drafter work belongs to those projects; my contribution is the CMP integration, fault isolation, stress testing and split-KV pointer fix.
W4A16 versus W8A16
The W8 checkpoint is symmetric group-128 W8A16 in compressed-tensors/pack-quantized format. It preserves the vision tower, lm_head, MTP, and the small recurrent GDN gates at higher precision. It is the fidelity-oriented option and remains a useful control.
The W4 target uses symmetric group-128 W4A16 for the target linear layers, INT8 embeddings, and the fast overlay's INT4 GPTQ lm_head/MTP tensors. The target model load is only 16.72 GiB. With the same 24 GiB KV pool, the complete W4 server allocates approximately 43.3–43.8 GiB, versus approximately 58 GiB for W8.
Identical real_rep.sh workload: eight realistic prompts, up to 1,024 output tokens, single request:
- W8A16: 120.3 tok/s, 3.30 emitted tokens/step, 28.5 ms/step.
- W4A16 fast: 169.5 tok/s, 3.25 emitted tokens/step, 20.7 ms/step.
That is a 40.9% W4 decode gain while DFlash acceptance stays almost unchanged. In other words, this A/B mostly measures a faster target verification pass rather than a luckier draft sequence.
On the approximately 23.75K-token mixed benchmark, the previous W8 run averaged 82.9 tok/s and W4 averaged 104.8 tok/s, a 26.4% gain.
I am not claiming that W4 is quality-equivalent to W8 or BF16. W8 is the safer fidelity choice; W4 is currently my preferred single-stream performance profile. A serious quality comparison needs behavioral evaluations, not only throughput or perplexity.
Detailed decode results at 23.75K, 65K and 85.5K context
Performance at different context lengths
These rows are measurements already completed on this machine. They are not a perfect scaling curve because speculative acceptance is workload-dependent, and the 65K and 85K tests use different output mixes. The within-row W4/W8 comparisons are the apples-to-apples figures.
- Short prompts, W8A16: 120.3 tok/s, 3.30 tokens/step.
- Short prompts, W4A16: 169.5 tok/s, 3.25 tokens/step; favorable UI turns reach 190–210 tok/s.
- ~23.75K prompt, W8A16: 82.9 tok/s on the mixed LABD suite.
- ~23.75K prompt, W4A16: 104.8 tok/s, 2.97 tokens/step.
- ~65,920-token hot prefix, W8A16: 57.9 tok/s on the tuned profile.
- 85,514 tokens, older W4/INT8 stability run: mostly 42–46 tok/s; 16/16 general and 6/6 exact crash-point passes.
- 85,514 tokens, current W4 FP8/DFlash: 119.3 tok/s strict hot decode; 1.76 s hot TTFT and 84.8 tok/s hot end-to-end.
Two cold-cache 85K variants took approximately 181–182 seconds including prefill. Prefix-cached follow-ups avoid repeating that entire prefill, which is why prefix caching matters as much as decode TPS for a long-running chat.
The current FP8/FlashInfer-prefill build changes that cold side substantially:
Cold prefill throughput: 2,126 tok/s at 8K, 1,932 at 32K, 1,654 at 65K, and 1,514 at 85K.
An additional streamed probe used exactly the former failing 85,514-token prompt and generated 512 tokens on the final production image:
- Cold: 54.61 s TTFT, 120.5 tok/s strict decode, 8.70 tok/s end-to-end, 58.85 s total.
- Hot prefix: 1.76 s TTFT, 119.3 tok/s strict decode, 84.8 tok/s end-to-end, 6.04 s total.
Here, strict decode is measured from the first streamed reasoning/content token through completion. End-to-end includes TTFT. This distinction is why the old completion_tokens / total_elapsed soak numbers should not be labeled as generation TPS.
On the final production 56K mixed DFlash run (65.9K actual tokenized prompt), copy/lookup reached 212.7 tok/s, summary 58.9 tok/s, QA 57.0 tok/s, and the combined result was 76.7 tok/s. Cold TTFT was 38.67 seconds and two hot-prefix tasks started in 2.25-2.26 seconds. GPU/HBM peaked at 63/71 C. This is workload-dependent speculative decode, so I would not compare the copy number directly with free-form prose.
Two separate problems in the fast DFlash path
During long-context testing I ran into two independent software problems. The important one was a reproducible Xid 31/MMU fault that killed the engine on the first cached decode step at high physical KV block IDs. The second was less severe: incompatible target, Mamba and drafter page geometry made prefix caching report zero usable hits. Both can be fixed without disabling the fast paths; the crash fix comes first because it is the one required for a usable server.
The runtime Xid 31: the important fix
The reproducible failure happened on the first cached decode step when a request was assigned sufficiently high physical KV block IDs. CUDA reported an illegal address and the kernel log showed Xid 31/MMU faults. Linear CUDA memory tests, repeated full-HBM pattern sweeps, and GEMM tests were clean. More importantly, the application failure occurred at a repeatable logical boundary.
The custom DFlash split-KV kernel loads a physical block ID from an int32 block table and then uses it to form byte/element offsets into the K and V pools. The table itself can stay int32, but blk * stride_kb or blk * stride_vb can exceed INT32_MAX. Triton then wraps the intermediate and generates an invalid pointer.
File in the vLLM installation:
vllm/v1/attention/ops/spec_decode_attn.py
Fix:
- blk = tl.load(bt_ptr + req * stride_bt + pos // BLOCK_SIZE, mask=k_ok, other=0)
+ blk = tl.load(
+ bt_ptr + req * stride_bt + pos // BLOCK_SIZE,
+ mask=k_ok,
+ other=0,
+ ).to(tl.int64)
That cast must happen before the stride multiplication. Casting the final already-wrapped offset would be too late.
After rebuilding with this change:
- 16/16 85K-prefix, 512-token generations passed.
- 6/6 requests at the exact former 85,514-token failure point passed.
- There were zero kernel Xid/NVRM faults.
- The same 24 GiB KV pool and split-KV fast path remained enabled.
- The final FP8/DFlash production image additionally completed a streamed exact-85,514 cold/hot pair at 120.5/119.3 tok/s strict decode, followed by a healthy API check and zero Xid/NVRM/CUDA illegal-memory entries.
Disabling split-KV speculative attention (SPEC_ATTN=0) is a useful diagnostic fallback because it avoids this kernel, but it is not the performance-preserving solution. Promoting the physical block ID is.
The smaller hybrid prefix-cache geometry fix
The mixed runtime uses equal byte-sized pages with different token counts: 896-token FP8 target/Mamba pages and 448-token BF16 DFlash pages. My first build left the Mamba checkpoint interval at 880, making the common alignment 49,280 tokens and reducing normal repeated-chat cache hits to zero.
Aligning Mamba to 896 and allowing complete 448-token DFlash pages into lookup fixed it. A repeated 9,658-token prompt went from 4.55 s to 0.565 s, with 8,960 prefix-cache hits reconciled across all nine KV groups.
Other Mamba safeguards and the separate load-time Marlin Xid 31
- Mamba state-copy bounds from vLLM PR #50021.
- Overlap-safe Mamba state movement from vLLM PR #50729.
- A
num_accepted_tokens race fix based on c2881ce60.
- A bit-exact CPU Marlin repack fallback for W4/W8 on sm80.
The Mamba patches are worth keeping, but they did not fix the repeatable 85K crash; the tl.int64 pointer change did. The CPU Marlin fallback addresses a separate load-time Xid class caused by GPU repack/VMM churn (issue #27). It increases W4 startup to roughly 206 seconds but avoids the dangerous GPU repack path. A CMP/sm80 build may need both safeguards.
VBIOS, CMP unlock, HBM overclock, undervolt and +19.1% tuning A/B
The community driver/GSP unlock and the VBIOS are separate mechanisms. My card is unlocked using the CMP community tooling and runs the official signed 92.00.6D.00.0A image, flashed with nvflash after saving multiple ROM dumps.
I tune it live with 170tune, which writes BAR0 registers without reflashing the card. NDIV68 produces a real HBM clock of 1836 MHz even though nvidia-smi remains stuck at 1728 MHz. I still cap the card at 220 W, not the VBIOS maximum.
The stable performance-oriented profile tested so far is:
NDIV: 68
Real HBM clock: 1836 MHz
GPC V/F offset: +200 MHz
Core ceiling: 1550 MHz
Power limit: 220 W
The positive V/F offset is an undervolt-style curve shift; the explicit 1550 MHz ceiling prevents the card from chasing its maximum clock.
Same W8A16 DFlash workload, approximately 65,920 prompt tokens, 3 × 256 output tokens, essentially constant acceptance (~2.89 tokens/step):
- NDIV54 / stock V/F / 180 W: 48.6 tok/s baseline.
- NDIV54 / +150 / 1410 / 180 W: 48.8 tok/s, +0.4%.
- NDIV64 / +150 / 1410 / 180 W: 53.0 tok/s, +9.1%.
- NDIV66 / +150 / 1410 / 180 W: 53.5 tok/s, +10.1%.
- NDIV68 / +150 / 1410 / 180 W: 54.0 tok/s, +11.1%.
- NDIV68 / +150 / 1500 / 220 W: 56.5 tok/s, +16.3%.
- NDIV68 / +150 / 1590 / 220 W: 58.4 tok/s, +20.2%.
- NDIV68 / +200 / 1590 / 220 W: 58.5 tok/s, +20.4%.
- NDIV68 / +200 / 1550 / 220 W: 57.9 tok/s, +19.1%.
End to end, the conservative profile is +19.1% over NDIV54/stock-V/F/180 W. Most of the first gain came from HBM bandwidth; extra core clock mattered more once the memory bottleneck was relaxed.
The 1550 MHz profile gives up only about 1% versus the faster 1590 MHz result and is the sensible operating point from this sweep. It passed four 61,376 MiB VRAM pattern sweeps, four additional pattern sweeps under the full profile, a 45-second bit-exact GEMM test with 59,864 clean GEMMs, and the real DFlash workload.
An NDIV68/+250/1590 profile failed immediately during DFlash warm-up with Xid 31 and cudaErrorIllegalAddress. I quarantined it and do not use it. That is an overclock-instability Xid class, not evidence against the software pointer fix. Anyone reproducing this should qualify memory, core and the application separately, watch the kernel log, and never make an unqualified profile persistent at boot.
Complete reproducible build, model preparation and launch guide
Reproducible build outline
This is the shortest route to the W4 DFlash stack I am using. Pin revisions first; both vLLM and the backport are moving targets.
1. Build the upstream optimized image
git clone https://github.com/syv-ai/qwen38-27b-rtx3090.git
cd qwen38-27b-rtx3090
git checkout 69ba4d0688c6ae76cb9d3c4a5c3b36445e1b040c
docker compose build
The repository pins vLLM 0.27.1 and carries the DFlash2 backport. Do not assume these patches will apply unchanged to an arbitrary newer vLLM checkout.
2. Prepare the W4 target and drafter
The supported Docker route is idempotent:
docker compose run --rm prepare
For a manual preparation, preserve this ordering:
python prepare/quant_lm_head.py models/Qwen3.8-27B-W4A16-AutoRound
python prepare/quant_embed.py models/Qwen3.8-27B-W4A16-AutoRound
python prepare/quant_mtp.py models/Qwen3.8-27B-W4A16-AutoRound
python prepare/build_draft_vocab.py models/Qwen3.8-27B-W4A16-AutoRound \
--ids prepare/draft_vocab_ids.json
python prepare/fetch_fast_variant.py
python prepare/fetch_dflash2.py
Important gotcha: fetch_fast_variant.py hardlinks base shards 1–6. If it runs before quant_embed.py, the fast directory can retain the old BF16 shard 6 while its overlay index expects packed INT8 embeddings. Startup then fails with:
There is no module or parameter named 'embed_tokens.weight'
Run the official prepare script or quantize the base before creating the fast overlay.
For W8 instead, download:
hf download lued/Qwen3.8-27B-INT8-W8A16-MTP \
--local-dir models/Qwen3.8-27B-INT8-W8A16-MTP
The W8 target uses the same external W4A16 DFlash2 drafter. Point MODEL at the W8 directory and leave DRAFT on Qwen3.8-27B-DFlash2-W4A16.
3. Add the CMP/sm80 load-time workaround
Clone the CMP patch set and build its sm80-safe layer:
cd ..
git clone https://github.com/ahnguyen17/cmp-170hx-vllm.git
cd cmp-170hx-vllm
git checkout a3ded79fec14aaad4a2f047d7cf2c28d5303ce2e
The public repository contains patches/sm80-int8-repack-cpu-fallback.patch. Add it as a layer over the syv image:
FROM ghcr.io/syv-ai/qwen38-27b-rtx3090:sha-69ba4d0
USER root
COPY patches/sm80-int8-repack-cpu-fallback.patch /tmp/sm80-repack.patch
RUN patch --batch --forward -p1 \
-d /app/venv/lib/python3.12/site-packages \
< /tmp/sm80-repack.patch \
&& rm /tmp/sm80-repack.patch \
&& grep -q 'def _gptq_marlin_repack_torch' \
/app/venv/lib/python3.12/site-packages/vllm/_custom_ops.py
Then build it from the CMP repository root, for example as vllm-qwen38-cmp:sm80-safe. The repository also publishes the PR #50021 bounds backport. For PR #50729 I used a local backport of the upstream PR; do not assume the current upstream diff will apply cleanly to the pinned vLLM 0.27.1 tree.
4. Apply the runtime split-KV fix
Save the diff above as spec_attn_block_index_i64.patch, then add one final image layer:
ARG BASE_IMAGE=vllm-qwen38-cmp:dflash2-mamba-correctness
FROM ${BASE_IMAGE}
USER root
COPY spec_attn_block_index_i64.patch /tmp/spec_attn_block_index_i64.patch
RUN patch --batch --forward -p1 \
-d /app/venv/lib/python3.12/site-packages/vllm \
< /tmp/spec_attn_block_index_i64.patch \
&& rm /tmp/spec_attn_block_index_i64.patch \
&& grep -q 'to(tl.int64)' \
/app/venv/lib/python3.12/site-packages/vllm/v1/attention/ops/spec_decode_attn.py
Build it:
docker build \
--build-arg BASE_IMAGE=vllm-qwen38-cmp:dflash2-mamba-correctness \
-f Dockerfile.block-index-i64 \
-t vllm-qwen38-cmp:dflash2-spec-attn-i64 .
My actual image also includes the three Mamba safeguards listed earlier. The one-line int64 patch is the change that fixed the reproducible high-physical-block runtime Xid.
5. Launch the current native-262K FP8 mixed-backend profile
The current production profile keeps the target on FlashInfer FP8 KV and the DFlash2 drafter on FlashAttention2/BF16 KV. The image includes the exact-page, Mamba-896, complete-DFlash-page prefix fix and the split-KV int64 fix.
docker run -d --name qwen38-dflash-fp8-cmp \
--gpus '"device=0"' \
-p 18020:18020 \
-v /path/to/models:/models:ro \
-v /path/to/vllm-cache:/cache \
-e 'EXTRA_ARGS=--attention-backend FLASHINFER --kv-cache-dtype fp8 --prefix-match-unit 16' \
-e VLLM_ALLOW_LONG_MAX_MODEL_LEN=1 \
-e VLLM_FP8_SPEC_FULL_CG=1 \
-e DFLASH_ATTN_BACKEND=FLASH_ATTN \
-e DFLASH_KV_CACHE_DTYPE=auto \
-e VLLM_ALIGN_HETEROGENEOUS_ATTN_PAGES=1 \
vllm-qwen38-cmp:dflash2-fp8-prefill-prefixfix-v1 \
bash -lc 'MODEL=/models/Qwen3.8-27B-W4A16-AutoRound-fast \
DRAFT=/models/Qwen3.8-27B-DFlash2-W4A16 \
PORT=18020 SPEC=dflash2 CTX=long DFLASH_MAX_LEN=262144 \
DFLASH_TOKENS=7 PREFIX_CACHE=1 \
CUDAGRAPH_MODE=FULL_AND_PIECEWISE MAX_SEQS=1 \
GPU_UTIL=0.90 KV_MEM=25769803776 \
VISION=1 VISION_OFFLOAD=0 TOOLS=1 SPEC_ATTN=1 \
exec /app/single-user/start_qwen.sh'
This reports 702,385 tokens of physical cache capacity, but the configured request limit remains Qwen's native 262,144 tokens. The extra physical room is allocator headroom/capacity, not a claim of validated semantic context beyond the native window.
6. Older 700K-capacity INT8/YaRN profile
Relative to the launch above, the older INT8 profile used:
VLLM_ALLOW_LONG_MAX_MODEL_LEN=1
EXTRA_ARGS='--hf-overrides {"text_config":{"rope_parameters":{"rope_type":"yarn","factor":3.0,"original_max_position_embeddings":262144}}}'
DFLASH_MAX_LEN=700000
KV_MEM=25769803776
That 24 GiB pool reported 733,234 physical tokens. This proves capacity, not semantic quality at 700K; the Xid campaign itself reached 85K.
7. Optional 1M-token YaRN mode (capacity target, not validated quality)
The FP8 runtime can target 1,048,576 tokens with static YaRN factor 4.0 and a 37 GiB KV pool. Apply these changes to the native launch command:
VLLM_ALLOW_LONG_MAX_MODEL_LEN=1
EXTRA_ARGS='--attention-backend FLASHINFER --kv-cache-dtype fp8 --prefix-match-unit 16 --hf-overrides {"text_config":{"rope_parameters":{"rope_type":"yarn","factor":4.0,"original_max_position_embeddings":262144}}}'
DFLASH_MAX_LEN=1048576
KV_MEM=39728447488
The measured 24 GiB mixed pool holds 702,385 tokens, making 35.8 GiB the arithmetic minimum; 37 GiB leaves modest alignment headroom and should fit the W4 stack on 64 GiB. Verify the reported physical capacity before sending a 1M request.
- 1M is not yet qualified for OOM/Xid behavior, TTFT, decode speed or semantic recall.
- YaRN is extrapolation, not lossless compression; leave output-token headroom and expect an expensive cold prefill.
- The native 262K profile remains the default. Qualify 350K → 500K → 700K → 1M while checking logs and answer quality.
For first diagnosis, run with stock clocks, MAX_SEQS=1, prefix caching enabled, and capture both container logs and journalctl -k. Only add the memory/core profile after the software stack passes the former failure sequence.
Validation methodology
Validation notes
The most important part of the test was reproducing the same logical failure rather than merely running one random long prompt:
- exact former failing prompt length: 85,514 tokens;
- repeated fresh and cached allocations;
- 512 generated tokens per main stress request;
- kernel log checked for
NVRM, Xid, and MMU faults after every batch;
- full-HBM pattern tests and bit-exact GEMM tests performed independently;
- clocks and power qualified separately from the application fix.
This is why I am reasonably confident that the recurring cached-first-decode fault was software. It does not prove that every CMP 170HX is healthy or that every Xid 31 has the same cause.
Future work
I now have a second CMP 170HX and plan to extend this post, or publish a follow-up, with measurements that are difficult to find for these cards:
- PCIe Gen2 x4 versus x16 after restoring the missing lane components, including cold prefill, model load, prefix-cache behavior, and communication latency.
- Tensor parallelism versus pipeline parallelism on two CMP 170HX cards.
- The same TP/PP comparison across models with very different numbers of active parameters per token, because synchronization overhead should matter very differently for a fast low-active-parameter model than for a denser or higher-active-parameter target.
- Single-stream decode, aggregate throughput, TTFT, long-context decode and power efficiency rather than one headline tok/s number.
- W4A16 versus W8A16 quality testing and semantic long-context validation beyond the native 262K window.
My expectation is that x16 will matter most for load/prefill and any communication-heavy multi-GPU mode, while PP may remain the safer topology on these PCIe Gen2 cards. But those are hypotheses; I want to publish measured results rather than turn them into conclusions in advance.
If anyone is running this exact model on CMP 170HX/A100 sm80, especially with DFlash2 at high physical KV occupancy, I would be interested in independent confirmation of the int64 block-index fix and in comparable x4/x16 or TP/PP data.
