SaaS Metrics
Gemini Executive Synthesis
Compatibility and stability of w8a16 (GPTQ int8) quantized models with various linear kernels (Marlin, Exllama, Conch) and speculative decoding (DFlash2) on NVIDIA CMP 170HX, specifically addressing CUDA illegal memory access (IMA) and dtype mismatches.
Technical Positioning
Achieving robust, high-performance inference for GPTQ int8 quantized Qwen3.8-27B models on specialized hardware (CMP 170HX) by ensuring compatibility across different linear kernels and resolving critical dtype-related failures and memory corruption issues.
SaaS Insight & Market Implications
This report exposes severe stability and compatibility issues when deploying w8a16 GPTQ-quantized models on NVIDIA CMP 170HX hardware. Critical failures include CUDA illegal memory access (IMA) that wedges the GPU, and dtype mismatches blocking fp16 execution due to hardcoded bf16-only kernels. While a workaround for the IMA (sm80 VMM corruption) was found, the pervasive nature of these errors across different linear kernels (Marlin, Exllama) and the drastic performance degradation with Conch (23 tok/s vs 133.7 tok/s) highlight significant technical debt. This directly impacts the viability of deploying highly optimized, quantized models on diverse hardware, increasing development overhead and limiting market reach for performance-critical B2B applications.
Proprietary Technical Taxonomy
Raw Developer Origin & Technical Request
GitHub Issue
Aug 23, 2026
Repo: syv-ai/qwen38-27b-rtx3090
w8a16 (int8 GPTQ) target: Marlin repack CUDA IMA (wedges card until reboot); every fallback route fails — Conch 23 tok/s, fp16 blocked by bf16-only spec-decode kernels
## Environment
- This repo @ `2ae239f` (vLLM 0.27.1 fork), launched via `single-user/start_qwen.sh`
- GPU: **NVIDIA CMP 170HX 40GB** (GA100, sm80, compute capability 8.0), driver 610.43.03
- Target model: Qwen3.8-27B quantized to **GPTQ int8 (w8a16)**: `group_size=128, sym=True, desc_act=False`, config `dtype: float16`, fp16 lm_head/MTP, `--language-model-only`
- Drafter: the shipped `Qwen3.8-27B-DFlash2-W4A16`, `SPEC=dflash2 CTX=fast`
- Base command: `MODEL= SPEC=dflash2 CTX=fast GPU_UTIL=0.96 bash single-user/start_qwen.sh`, with the per-row kernel/dtype overrides below (`VLLM_DISABLED_KERNELS`, `EXTRA_ARGS`, `SPEC_ATTN=0`)
I realize w8a16 targets aren't what this repo ships — but two of the failure modes below look like real bugs independent of which checkpoint triggers them, and the crash mode (card wedge until reboot) is nasty enough to be worth a guard at minimum.
## Result matrix (5 attempts, one card)
| # | Linear kernel | `--dtype` | KV cache | Result |
|---|---|---|---|---|
| 1 | Marlin (default choice) | *(none → fp16)* | bfloat16 | **CUDA illegal memory access in `marlin_permute_scales` during `process_weights_after_loading`**; card wedged until reboot — see below |
| 2 | Exllama (Marlin disabled) | *(none → fp16)* | bfloat16 | FA2 init: `mha_varlen_fwd ... query and key must have the same dtype` |
| 3 | Conch (`conch-triton-kernels==1.3`) | bfloat16 | bfloat16 | **Boots and serves correctly — but 23 tok/s** (vs 133.7 tok/s for the shipped W4A...
Developer Debate & Comments
This is an unusually good report — five attempts with the blame isolated by controls, and a working-vs-broken vLLM version to bisect against. Taking your two asks in reverse order, because I can act on the second one and not the first. ## Ask 2: fail at startup. Done, in `7db2d1f`. You were right that this is a real bug independent of the checkpoint, and I can name the line. **The split-KV verify attention this repo adds is bf16-only, hardcoded, in two places** — `patches/spec-decode-attn.patch`: ```python qs = (q * scale).to(tl.bfloat16) # line 217 acc = acc * alpha[:, None] + tl.dot(p.to(tl.bfloat16), v).to(...) # line 236 ``` With `--dtype float16`, `v` is fp16 and the left operand of that `tl.dot` is bf16, which is your row 4 exactly: `AssertionError: Both operands must be same dtype. Got bf16 and fp16` at `spec_decode_attn.py:220`, surfaced as a `CompilationError`. The patch header does say *"bf16 KV cache only"*, but that reads as a KV-d...
**Correcting myself on row 5.** I said I had taken your rejection-sampler assert at face value as a second bf16 assumption and had not read that kernel. I have now, and **it is not one.** `rejection_sampler_utils.py` contains **zero** `bfloat16`/`float16` literals. Every load promotes to `tl.float32`, the index arithmetic is `tl.int64`, and the intermediates are allocated off `target_logits.new_empty(..., dtype=torch.float32)`. The sampler is dtype-agnostic — it cannot care what your activations are. So the assert at `rejection_sampler_utils.py:931` has some other cause. A device-side assert is the signature of an out-of-bounds index rather than a dtype mismatch (the bounds-checked loads in there are masked on `vocab_size`, and the per-request loops run to `num_draft_tokens = end_idx - start_idx - 1`), and your row 5 was Exllama on a w8a16 target whose GEMM path had already failed three other ways in rows 1-4. **Garbage logits or a bad token id out of a broken linear path would produ...
**Root cause found and worked around on our 170HX — the IMA is sm80 VMM corruption under allocation churn, not a kernel pointer bug.** Full story below; patch linked at the end. ## What actually faults With your `7db2d1f` fail-fast and `expandable_segments:False`, the w8a16 target + W4A16 DFlash2 drafter still Xid-31'd during `gptq_marlin_repack` — but the faulting frame **drifted between runs** (once `empty_cache`, once a `.to(int32)` cast, once mid-repack), always with `Xid 31 FAULT_INFO_TYPE_REGION_VIOLATION … VIRT_WRITE` at a different VA. Classic async-fault signature: the Python traceback is a bystander. Controlled experiments (5 reboots, sadly the card wedges per attempt): | Control | Result | |---|---| | Bit-exact torch reimplementation of the repack, real weights, **sm86 (3090)**, `compute-sanitizer` | **0 errors** | | Full engine-faithful drafter load on **sm80**, 10 GB allocator pressure | clean | | Same chain on **sm80**, production occupancy (~27 GB resident) | **Xid 3...
This is the best piece of debugging this tracker has seen — five reboots on a card that wedges per attempt, a bit-exact CPU reimplementation as the control, and the discipline to notice the faulting frame *drifting* and conclude the code in the traceback is a bystander. "When vanilla elementwise kernels fault, the page mappings under them are guilty" is the sentence the next GA100 owner needs, and it is now **gotcha 41** on main (`60daef8`), with your patch linked and the `expandable_segments` independent-trigger warning kept next to it. Your 170HX numbers are in the README's new "Results from other hardware" section, method-flagged (133.7 on the shipped target — the first sm80 datapoint, level with the 3090 — and 97.8 on your w8a16 with the workaround). ## Taking you up on the candidate patch Your in-tree diagnosis points at a fix I can actually build and half-test from here: keep the repack's Python-side intermediates small — chunked unpack/pack over bounded column slices, int32 st...
@ahnguyen17 The candidate patch I owed you is on main: [`patches/marlin-repack-staged-sm80.patch`](https://github.com/syv-ai/qwen38-27b-rtx3090/blob/main/patches/marlin-repack-staged-sm80.patch) (commit `3cc6afc`). ## Design — "keep the intermediates small", taken literally Rather than chunking the repack math (the Marlin tile layout makes column-slab independence a property I'd have to prove about the kernel, not assume), it removes the allocation churn itself: the `.contiguous()` copy and the pad — the two per-layer transients around `gptq_marlin_repack` — now land in **one persistent, grow-only staging buffer per device**, and the source tensor is dropped before the repack output allocates. Per-layer allocation traffic collapses from two transients + one output to the output alone; the staging buffer reallocates only when a larger shape class first appears (a handful of times per load, vs hundreds of alloc/free pairs). The repack kernel itself is untouched and receives identical b...
Adjacent Repository Pain Points
Other highly discussed features and pain points extracted from syv-ai/qwen38-27b-rtx3090.
Extracted Positioning
Bug tracking and resolution for KVarN, DFlash2, and CTX=huge configurations, specifically addressing output corruption, engine crashes, and KV cache memory limitations.
Ensuring correctness and stability for Qwen3.8-27B inference with advanced speculative decoding (MTP, DFlash2) and extended context lengths (CTX=huge) on RTX 3090, by systematically identifying and fixing critical bugs related to prefill handling, sliding window guards, and KV cache allocation.
Top Replies
Two follow-ups to [my comment on ](https://github.com/syv-ai/qwen38-27b-rtx3090/pull/13#issuecomment-5379687913), moved here since that PR is closed: the MTP result re-run to your fresh-server stan...
**Update — two of the three open items are closed, both verified on hardware.** Pushed in `82bd62d`. ## `DFLASH_TOKENS=15` does not boot at 240k → **fixed** My first fix was aimed at the wrong thin...
Our comments crossed — mine went up six minutes before yours, so this is about one line in your update rather than the conclusion. You landed on `mtp` keeping PIECEWISE, which is where I would land...
Extracted Positioning
Performance scaling of Qwen3.8-27B with vLLM using dual RTX 3090s (Tensor Parallelism TP=2) compared to a single card, specifically addressing the impact of KV_MEM pinning and overall throughput gains.
Optimizing multi-GPU inference performance for Qwen3.8-27B on cost-effective hardware (dual RTX 3090s) to achieve significant throughput improvements (+16-35%) across various concurrency levels and speculative decoding methods (MTP, DFlash2), while resolving memory allocation inefficiencies.
Top Replies
This is really useful information. I'm also running a very similar setup and had planned to provide some numbers this weekend but the information you've provided is so much more in depth. Also, you...
I've changed my configuration to roughly match your suggestions and (warm) benched at 250W and 275W. My numbers are significantly better than they were though a little lower than yours at C1. I sus...
This is the best-run field report this repo has received — a genuine controlled A/B, spread stated, negative results kept, and the two findings that touch deliberate calibrations flagged instead of...
Extracted Positioning
Functionality and performance of OffloadingConnector (KV eviction to RAM/disk) for extending KV cache capacity on the Qwen3.8-27B stack.
Enabling deployment of Qwen3.8-27B with extended context lengths beyond GPU VRAM limits by leveraging system RAM/disk for KV cache, even in virtualized environments (Proxmox VM with GPU passthrough).
Top Replies
Two follow-ups from running this in anger. **Edge 6 escalated from ops-hygiene to a production outage.** The engine died once (host OOM, most likely — nothing in the engine log), and then `restart:...
Answers to all three open questions, one shipped fix for the outage class, and one debt I am keeping. Your follow-up comment upgraded this report from useful to load-bearing — a 70-restart ghost lo...
The traced cell I owed you is done — your "store path works, load path never hits" is **reproduced on this card, mechanism found, and it is neither unliftable-by-design nor a hash miss**. It's an a...
Extracted Positioning
Stability of vLLM EngineCore with flashinfer attention backend and MTP speculative decoding, specifically concerning CUDA illegal memory access errors during long context processing.
Ensuring robust and error-free inference for Qwen3.8-27B (W4A16 AutoRound) on RTX 3090, particularly with advanced features like flashinfer and MTP, even with long contexts (~28k tokens).
Top Replies
Good report — the dump_input capture and the dmesg line are exactly what makes this analyzable from a distance. One red herring to clear first, then where the real suspicion points, then what I wou...
## Follow-up: second occurrence on 2026-08-27, same fault address The same crash recurred today (2026-08-27 14:03:48 local, after ~40h uptime, docker restart policy auto-recovered). **New data poin...
## Test result: option 1 (`--attention-backend FLASH_ATTN`) is hard-invalid on the fp8-KV tier As suggested, I added `--attention-backend FLASH_ATTN` to the `CTX=long` launch (everything else uncha...
Extracted Positioning
System RAM requirements for loading and running Qwen3.8-27B with vLLM, specifically addressing limitations with 16GB main system RAM.
Enabling deployment of Qwen3.8-27B on systems with constrained main memory, potentially through optimized loading strategies or reduced memory footprint.
Top Replies
Honest framing first: 16 GB of host RAM (10 GiB actually free) is below anything this stack has been run on, so what follows is diagnosis plus the two levers that exist — not a tested recipe. It ma...
> system ram is 16GB, using iGPU Well, your iGPU is using system RAM; what OS & have you tried to reduce background processes? Close as many web browser tabs as you can, & you might have to use thi...
Frequently Asked Questions
Market intelligence mapped to Compatibility and stability of w8a16 (GPTQ int8) quantized models with various linear kernels (Marlin, Exllama, Conch) and speculative decoding (DFlash2) on NVIDIA CMP 170HX, specifically addressing CUDA illegal memory access (IMA) and dtype mismatches..
How is Compatibility and stability of w8a16 (GPTQ int8) quantized models with various linear kernels (Marlin, Exllama, Conch) and speculative decoding (DFlash2) on NVIDIA CMP 170HX, specifically addressing CUDA illegal memory access (IMA) and dtype mismatches. positioned in the market?
Based on our AI analysis of the original developer request, its primary technical positioning is: Achieving robust, high-performance inference for GPTQ int8 quantized Qwen3.8-27B models on specialized hardware (CMP 170HX) by ensuring compatibility across different linear kernels and resolving critical dtype-related failures and memory corruption issues.
Are engineers actively discussing Compatibility and stability of w8a16 (GPTQ int8) quantized models with various linear kernels (Marlin, Exllama, Conch) and speculative decoding (DFlash2) on NVIDIA CMP 170HX, specifically addressing CUDA illegal memory access (IMA) and dtype mismatches.?
Yes, we have tracked 7 direct responses and active debates regarding this specific topic originating from GitHub Issue.
What are the foundational technologies related to Compatibility and stability of w8a16 (GPTQ int8) quantized models with various linear kernels (Marlin, Exllama, Conch) and speculative decoding (DFlash2) on NVIDIA CMP 170HX, specifically addressing CUDA illegal memory access (IMA) and dtype mismatches.?
Our proprietary extraction maps Compatibility and stability of w8a16 (GPTQ int8) quantized models with various linear kernels (Marlin, Exllama, Conch) and speculative decoding (DFlash2) on NVIDIA CMP 170HX, specifically addressing CUDA illegal memory access (IMA) and dtype mismatches. to adjacent architectural concepts including w8a16 (int8 GPTQ) target, Marlin repack CUDA IMA, card wedged until reboot, Conch 23 tok/s.
Engagement Signals
Cross-Market Term Frequency
Quantifies the cross-market adoption of foundational terms like w8a16 (int8 GPTQ) target and Marlin repack CUDA IMA by tracking occurrence frequency across active SaaS architectures and enterprise developer debates.