Making a Ampere Optimized Quant of Qwen3.6-27b - INT4 Weights + INT8 Activations

[chriskerley78910]

An earlier experiment showed significantly improvements in TTFT with 8-bit activation's, but the decode speed was still very slow compared to the INT4 weighted model [1].

I am thinking of making a quantized model of Qwen3.6-27b that is made specifically for the Amper Arch. i.e. INT4 Weights INT8 Activations.

It would have the decode speed benefit of the current default vllm/dual setup but with the TTFT speed benefit of INT8 activations ~40% faster on prefill [2].

What do you guys think?

[noonghunna]

I addressed the decode speed bug today on dual-max, you might wanna give it a shot again before you begin your work. It's working much better now.

The problem was with int8 tph kv, the requests were hitting the older Triton kernel in vllm, which was degrading decode tps. Changing it to fp8/e5m3 makes it go to the new FlashInfer kernel which seems to have addressed the issue with decode tps.

dual-max was affected by this and now throws better tps than dual-fast (W4A8).

[tomByrer]

How far did you get in making your qant @chriskerley78910 ?

[chriskerley78910]

VLLM is the problem, it won't work with INT4 W4A8 apparently. There is a pull request to fix the bug, but there is ~3k pull requests waiting for review on that repo, ( and only a hand full of reviewers ), so it might be a long time before VLLM has the fix for it merged.

Someone made a fork of VLLM 0.23 that has the fix, and there is already a quant made that will work with it as @noonghunna mentioned in another thread of this discussion. So it looks like to run it, using that fork is the only option at the moment.

The bugfix looks small, and it wouldn't be tough to implement, the issue is trying to integrate it with club-3090's existing stuff. But given this repo is directed specifically at 3090 owners, I'll bet @noonghunna would be open to merging the bugfix (possibly a forked vllm) into this repo as a Amper specific vLLM setup, (more likely if I bench it and show the benefits first). The potential gains from it are hard to ignore. If INT4 W4A8 works as hypothesised then it'll basically be the best possible config dual Amper cards can run. (See the figure below).

Currently I am rocking the INT8 one (W8A8) because I do a lot of agentic coding work, so there is a lot of back and forth chatter while holding a long context. i.e. Prefill speeds matter more than decode speeds.

[3dluvr]

We have had a mixed W4A16 Qwen3.6-27B quant from Minachist for 2 months already. I'm sure you could just take his quantize.py script and modify it to make W4A8 variant in 30 min-1 hr total on dual 3090.

I don't see any barriers in doing so, matter a fact stay with AutoRound and make sure you specify format="auto_round:auto_gptq" in ar.quantize_and_save to avoid any vLLM issues.

[3dluvr]

LOL, I've been trying to make a working mixed Gemma-4 quant for sometime now, I don't want a distraction. ;)

[3dluvr]

Ooh interesting...AutoRound does not support W4A8 scheme:

['W4A16', 'W2A16', 'W3A16', 'W5A16', 'W6A16', 'W8A16', 'MXFP4', 'MXFP6', 'MXFP4_RCEIL', 'MXFP8', 'MXFP8_RCEIL', 'NVFP4', 'FPW8A16', 'W2A16G64', 'W2A16G32', 'FP8_STATIC', 'BF16', 'W4A16_MIXED', 'INT4', 'INT8', 'INT8_W8A8', 'FP8_BLOCK', 'MXINT4', 'GGUF:Q4_0', 'GGUF:Q4_1', 'GGUF:Q5_0', 'GGUF:Q5_1', 'GGUF:Q2_K_S', 'GGUF:Q3_K_S', 'GGUF:Q3_K_M', 'GGUF:Q3_K_L', 'GGUF:Q4_K_S', 'GGUF:Q4_K_M', 'GGUF:Q5_K_S', 'GGUF:Q5_K_M', 'GGUF:Q6_K', 'GGUF:Q8_0', 'GGUF:Q2_K_MIXED']

[chriskerley78910]

That's strange, maybe that's why Minachist didn't already create it.

[3dluvr]

I stand corrected about the time needed, it's about ~11 hrs with 3x RTX 3090 even.

Might stop it as I don't know if my parameters are good or not. It would waste ~6 kWh ($0.60) of electricity though.

[chriskerley78910]

11 hours isn't too bad. Getting a model quant built specifically for my 3090s would likely payoff quickly.

[Ezgera]

Are you aware of this project: https://github.com/avesed/vllm-ampere-optimized ? Seems to be relevant to this thread. Essentially, default vLLM gates W4A8 to Hopper+, but it can be made to work on Ampere.

[noonghunna]

thanks for sharing this, i wasn't aware of this one. Will check it out.

[chriskerley78910]

Good find dude. thanks for sharing that :) RTX 3090 owners will thank you!

W4A8 defaults to Hopper+ ins stock VLLM, but I can't find W4A8 integer quants. I am wondering why no one has created it given it seems to be the default one the repo was designed for.

[chriskerley78910]

After reading into this, it looks like it hasn't been done because autoround simply does not support that format. For that algorithm to work right it needs 16 bit activations. Got to find another method it seems.

[noonghunna]

Followed up after Ezgera's pointer (and re-read the whole thread) — and I think it resolves the exact wall you all hit, because everyone here has been trying to make a W4A8 checkpoint, and that's the wrong frame.

@3dluvr was right (that scheme-list comment): AutoRound genuinely has no W4A8 scheme — his pasted list is correct, and that's why the "run Minachist's script to emit W4A8" plan dead-ended, and why chris couldn't find any W4A8 quants. But here's the twist from avesed's release notes: W4A8 isn't a scheme you quantize to at all. It's a runtime flag — VLLM_MARLIN_INPUT_DTYPE=int8 on a normal W4A16 checkpoint. Marlin quantizes the activations to int8 dynamically at serve time and runs them on the INT8 tensor cores; the quantizer never emits "A8." So AutoRound not listing W4A8 is a non-issue — you keep your W4A16 int4 weights and flip a serve-time toggle.

So nobody needs the 11-hour build. The candidates already exist as W4A16:

• Minachist's Qwen3.6-27B-Mixed-AutoRound (the one @3dluvr linked), or
• avesed's own Avesed/Qwen3.6-27B-INT4-W4A16 (his v0.3 notes call it coherent, GSM8K-identical).

One real requirement, though — and it's the catch the thread hasn't hit yet: the int4 weights must be symmetric. avesed's v0.3 note is explicit: "W4A8 asym-weight (AWQ uint4+zp)… output is garbage."

• AWQ is inherently asymmetric → out.
• AutoRound defaults to asymmetric, so Minachist's mixed-AutoRound quant may hit the garbage path — worth testing, but not a given.
• If it does, the fix is small and 3dluvr's own tip is basically it: re-quant symmetric — AutoRound sym=True exported as auto_round:auto_gptq (or GPTQ-sym / llm-compressor int4-sym). That's a normal int4 quant, ~30–60 min like he said, not an 11-hour W4A8-specific build. Or start from a confirmed-symmetric W4A16 and skip quantizing — but check the candidate first, incl. avesed's own (his 9B is AWQ/asym, so symmetry isn't automatic).

Why it works on Ampere at all (the useful mental model): W4A8-via-Marlin only needs INT8 tensor cores, which Ampere has had since sm_80 — vLLM's Hopper gate is conservative, not silicon. That's the opposite of FA3/FA4 or native FP8/FP4, which are genuinely silicon-bound (no TMA/wgmma/tcgen05 on Ampere) and can't be un-gated. Integer low precision (INT8/INT4) runs native on Ampere; float low precision (FP8/FP4) doesn't.

And our own numbers back the "best of both" intuition. We've benched the INT8 W8A8 cousin (Avesed/Qwen3.6-27B-INT8-W8A8) on dual 3090 — the "8-bit prefill corner": +17–51% prefill vs fp8 (2062 vs 1364 tok/s @10k), decode −8%. W4A8 should be strictly better for the goal — int4 weights recover the decode bandwidth W8A8 gives up, int8 activations keep the prefill win. avesed's 9B numbers point the same way (W4A16→W4A8: prefill +49%, decode ~flat) — but that's 9B / README, not 27B, and he's already reverted one prefill opt (int8-QK attention, net-negative on decode), so the 27B magnitude is a bench-it question, not a known.

And you don't have to wait for the W4A8 eval to get a prefill win — W8A8 is ready today. Unlike W4A8, Avesed/Qwen3.6-27B-INT8-W8A8 is Ampere-native INT8 (compressed-tensors) — it loads on stock vLLM, no fork, and its int8 weights carry no symmetric-int4 gate. The only cost is decode: 8-bit weights are ~2× the memory bandwidth of int4 and decode is memory-bound, so it decodes slower than the int4 default. That makes W8A8 a prefill/TTFT corner, not a balanced-default swap — but for prefill-heavy work (long prompts, RAG, agent tool-call bursts) it's a real win you can run tonight.

Config	Prefill	Decode	Ready?
W4A16 (int4 wt) — today's default	baseline	fast (int4)	✅ shipped
W8A8 (int8 wt + act) — Avesed	+17–51% vs fp8	−8% (int8 wt bandwidth)	✅ stock vLLM, no fork
W4A8 (int4 wt + int8 act) — the goal	~+40–49%	fast (int4)	⏳ fork + symmetric-int4 + unmeasured on 27B

Why prefill moves but decode doesn't (the mechanism behind that table): the activation precision decides which tensor-core datapath the GEMM lands on. Dequant-to-bf16 (W4A16, and FP4-on-Ampere) → the bf16 cores; dequant-to-int8 (W4A8 / W8A8) → Ampere's INT8 cores at ~2× the bf16 rate — that's the prefill win, nothing more exotic. Decode is batch=1 and memory-bound on the 4-bit weight read, so the tensor-core rate is irrelevant there → decode barely moves (exactly why W4A8 keeps int4's decode while gaining int8's prefill). int4 activations would be faster still, but that's the accuracy cliff — so int8 is the floor.

W4A8 is exactly "W8A8's prefill without W8A8's decode penalty" — that's the whole reason to chase it over W8A8. Until it validates, W8A8 is the ready fallback. (One hard limit on both int8 paths: W8A8 INT8 compute is sm 7.5–9.x — not on Blackwell/5090, so this is Ampere/Hopper only.)

Caveat before anyone calls it production: the fork is vLLM v0.23.0 + FlashInfer 0.6.12 + CUDA 13 / torch 2.11 — around our vllm-stable v0.24.0 pin, so it's a fork-reconcile, not a stock flag; and there are no published 27B accuracy/throughput deltas yet.

@chriskerley78910 — want to take the first crack? You kicked this off, you've got the 3090s, and it's now a ~30-min job, not the 11-hour build. The whole experiment: grab a symmetric W4A16 (confirm avesed's, or make one — AutoRound sym=True → auto_round:auto_gptq per @3dluvr's tip), then serve it twice on avesed's fork — once plain (your W4A16 baseline) and once with VLLM_MARLIN_INPUT_DTYPE=int8 (that's W4A8). Same checkpoint both runs → the flag is the only variable → a clean A/B on prefill / decode + a quick GSM8K. That's it.

We'll run the same on the reference 2×3090 (verify / stress / soak / 8-pack + GSM8K), so whoever posts numbers first, the other corroborates. If W4A8 lands the prefill win without denting decode or long-context recall, it's a genuinely strong Ampere serving quant and we'll wire it up as an option.

Great find @Ezgera; @3dluvr your scheme-list catch was the key clue; @chriskerley78910 your instinct was right all along — it's "symmetric int4 + a flag," not a new quant. 🙏

[chriskerley78910]

"Marlin quantizes the activations to int8 dynamically at serve time and runs them on the INT8 tensor cores;"

OMFG.

I got the W4A16 weights now, just got to try the fork on it and see if it tastes good.

[noonghunna]

Since W4A4 came up earlier (@chriskerley78910's single-card idea) and the main comment stopped at W4A8, worth closing that loop — W4A8 is the floor for activations on Ampere, not a waypoint to W4A4. And "but NVFP4 is 4-bit, why not W4A4?" is the natural objection, so let me hit it head-on.

The two 4-bit formats aren't the same thing, and they live on different silicon:

4-bit format	Native compute on Ampere?	On a 3090 you get
INT4 (integer)	✅ yes — Ampere/Turing have int4 tensor cores (Hopper+ dropped them)	native int4 math — but int4 activations are too lossy
FP4 (NVFP4/MXFP4, float)	❌ no FP4 datapath (Blackwell-only)	4-bit storage only → dequant to bf16 → bf16 compute

So the surprising bit: W4A4-integer is physically runnable on a 3090 (it has int4 tensor cores — a thing Hopper actually removed). The blocker there isn't hardware, it's:

1. Accuracy cliff. Int8 activations survive LLM outlier features; 4-bit int activations (16 evenly-spaced levels) shred them → naive W4A4 output is garbage. Rotation methods (QuaRot / SpinQuant / QServe) rescue it but still cost real accuracy where W4A8/W8A8 cost ~none.
2. No decode gain anyway. Decode is batch=1, weight-bandwidth-bound — activations are tiny. So W4A4 decode ≈ W4A16 ≈ W4A8. Activation bits only touch prefill, so you'd eat the accuracy hit for a prefill-only gain W4A8 already captures.
3. No kernel + the field already voted. No mainstream Ampere W4A4 GEMM, and QServe (SOTA low-bit serving) deliberately picked W4A8KV4 over W4A4 — too much loss for too little.

And "NVFP4 is 4-bit" isn't a counterexample — because on Ampere, NVFP4 doesn't compute in 4 bits. Its 4-bit weights are unpacked to bf16 and the matmul runs in bf16 (the Marlin path): 4-bit storage, 16-bit compute — memory win, zero compute win. Native FP4 math is a Blackwell 5th-gen-tensor-core instruction that isn't in the sm_86 ISA — you can store FP4 on a 3090, you can't compute it (same wall as FA4/wgmma).

Why FP4 is the real 4-bit-activation format and INT4 isn't: FP4 (E2M1) is floating-point — an exponent gives 4 bits a wide dynamic range, and NVFP4/MXFP4 add per-block microscaling, so it captures the activation outliers uniform INT4 can't. A usable "W4A4" therefore = FP4 activations on Blackwell (5090/GB10) — which our 3090s have neither the FP4 hardware for, nor an INT4 format good enough to substitute.

On the single-card angle: A4 doesn't help there either — at batch=1 the weights are already int4 in both W4A16 and W4A4, and activations are a tiny memory slice, so A4 buys no size or decode win, just accuracy loss. Single-card levers are weight bits (int4→int3/int2 GGUF) and KV-cache quant, not activation bits.

Bottom line: W4A8 is the activation floor on Ampere — int8 is the smallest activation precision that both (a) runs native on Ampere tensor cores and (b) survives the outliers. The good 4-bit-activation format (FP4) needs Blackwell silicon; the available one (INT4) is too lossy. So this thread's target — int4 weights + int8 activations — is already sitting on the right spot.

[noonghunna]

Update — I ran the full exploration today, and the story flips one more time: you don't need the fork for the W4A8 serve-time toggle at all. It's been in stock vLLM the whole time, hiding behind two real bugs — both found, root-caused, patched (pure Python), and measured.

What's actually stock

VLLM_MARLIN_INPUT_DTYPE=int8 — the env, the per-token int8 activation quant, and the Marlin int8 kernel (upstream #24722, merged 2025-11) — has been in stock vLLM since ≤v0.22 for compressed-tensors W4A16 checkpoints, and v0.24.0 added the same hook to the GPTQ/AWQ paths. @chriskerley78910 — the "PR stuck in the 3k queue" (vllm#38066) is only needed for native W4A8-INT checkpoints (the int-quantized format nobody publishes anyway). The W4A16-checkpoint-plus-env route — the one this thread actually wants — doesn't need it.

Why it still fails for many checkpoints — two upstream bugs

1. auto-round checkpoints silently ignore the env. vLLM forces quant_method: auto-round onto its INC path, which builds the Marlin method without threading input_dtype — so your A/B shows "no change" because the flag literally never reaches the kernel. 3-line pure-Python fix.
2. Negative weight scales → garbage. The int8 path quantizes group scales to int16, but the kernel reads them back unsigned (reinterpret_cast<uint16_t*>). AutoRound legitimately emits negative scales — our shipped 27B AutoRound checkpoint has 49.7% negative (perfectly valid for W4A16, where scales stay fp16) — so half the weight groups explode ~4 orders of magnitude and you get ratoratorato…. GPTQ/AWQ/llm-compressor recipes emit positive-only scales, which is exactly why avesed's ecosystem (and upstream CI) never saw this. Minimal toy repro: identical GEMM with 50% negative scales → W4A16 rel-err 0.0003, W4A8 rel-err 16.1.

So the requirement from my earlier comment gets sharper: not just symmetric int4 — POSITIVE symmetric int4.

Measured (single RTX 3090, bf16, cudagraphs, batch 1)

checkpoint	format → route	patches needed	prefill (6K prompt)	decode
Avesed/Qwen3.6-27B-INT4-W4A16	compressed-tensors (AWQ-smoothed, sym g32)	ZERO — unmodified vllm/vllm-openai:v0.24.0	1,343 → 2,467 tok/s (+84%)	44.1 → 42.5 (−3.6%)
Gemma-4-12B QAT W4A16 (compressed-tensors, sym g32)	compressed-tensors	ZERO (on its usual serving image)	2,550 → 4,180 tok/s (+64%)	71.2 → 67.2 (−5.6%)
Qwen3.6-27B AutoRound-int4 (our shipped quant, g128)	auto-round → INC	2 pure-Python patches (input_dtype hook + scale-sign fold)	1,350 → 2,814 tok/s (+108%)	47.0 → 45.0 (−4%)

All three coherent after the fixes (chat + logprob sanity — "capital of Italy is → Rome at −0.13"; prompt logprobs healthy). Verified negative-scale counts first: Avesed 0/760M, Gemma QAT 0/340M, AutoRound 95.8M/193M.

Three honesty flags on that table:

• Decode numbers are spec-decode-OFF baselines — one variable at a time. Our shipped configs run MTP/drafters on top; W4A8 × MTP composition is untested (next on the list).
• The uniform −4–6% decode is the one thing that IS still fork-only: avesed's int8 8-row decode tile (his patch 0002, CUDA) — stock falls back to the 16-row tile at batch 1.
• These are speed + coherence numbers, not quality numbers. The full 8-pack + GSM8K on the W4A8 configs (and the committed dual-2×3090 run) is still owed — especially for the AutoRound path, where the sign-fold clamps 0.6% of weights.

@chriskerley78910 — decision rule for the weights you just got

# 1) format: config.json → quantization_config.quant_method
#    "compressed-tensors"          → stock v0.24.0+ works TODAY, zero patches
#    "auto-round" / gptq variants  → needs the small patches (coming) or a ct-format quant
# 2) scale sign — the garbage predictor:
from safetensors import safe_open
neg = tot = 0
with safe_open("model.safetensors", framework="pt") as f:
    for n in f.keys():
        if "scale" in n:
            s = f.get_tensor(n); neg += (s < 0).sum().item(); tot += s.numel()
print(f"{100*neg/tot:.2f}% negative scales")   # anything > 0% → garbage under int8, for now

Then serve with VLLM_MARLIN_INPUT_DTYPE=int8 and --dtype bfloat16 (int8-act dequant can overflow fp16 — avesed documented exactly this failure on Gemma), and A/B a long prompt. Fastest way to taste it tonight: Avesed/Qwen3.6-27B-INT4-W4A16 on stock vllm/vllm-openai:v0.24.0 — works right now, no fork, no patches.

What's next from our side

• File both bugs upstream with the minimal repro (the correct kernel fix is a signed int16 scale read; our load-time sign-fold is the interim).
• Quality gate: 8-pack + GSM8K on the W4A8 configs, W4A8 × MTP composition, and the committed dual-card numbers for this thread.
• Catalog: likely a one-flag compose knob carrying the two patches, so any of our vLLM Qwen configs can flip to W4A8.

Credit where due — this one had three authors before I typed a line:

• @chriskerley78910 called the shot and carried it. He's the one who read the perf-vs-bits figure and named W4A8 as "basically the best possible config dual Ampere cards can run," then worked the problem through every dead end — no W4A8 quants exist → AutoRound has no such scheme → the vLLM wall → the upstream PR → the fork — and publicly bet this repo would take the fix. That bet is why this exploration happened today instead of someday, and the +64–108% prefill column above says his read was right. Chris, your bench offer absolutely stands — you now have a much shorter path to it, and independent numbers from your dual rig are exactly what the table is missing.
• @Ezgera's fork pointer cracked the mechanism open.
• avesed's patch series + release notes were the map — his fork remains the only route for the decode tile, MoE int8-act, and native-W4A8 checkpoints.

But for what chris set out to get — int4 weights + int8 activations on 3090s — stock vLLM + the right checkpoint does it today, and our two small patches extend it to the AutoRound quants everyone already has.

[noonghunna]

Quality gate results for the W4A8 numbers above — it ties the baseline on both reasoning legs.

Full 8-pack on the dual-fast config (TP=2, MTP n=3, 262K, both patches): 110/150 thinking-off · 111/150 thinking-on vs the W4A16 fast tier's 109 — inside the ±5–7 noise band, with the thinking leg actually strongest (instruction-follow 15/15, bugfind 14/15).

Speed — corrected with a clean single-variable A/B (both arms fp16, both through the shipped compose, only VLLM_MARLIN_INPUT_DTYPE differs): prefill +50% (1,674 → 2,513 tok/s @ 6K prompt) · decode ~neutral (67.2 → 68.9, inside run noise) · MTP acceptance unchanged. My earlier "+8% decode" was an fp16-vs-bf16 confound between two differently-configured boots — the honest number is that int8 activations buy prefill, and leave decode alone (decode is batch-1 weight-bandwidth-bound, so activation precision barely touches it — as expected). Prefill is the whole win, and it's a big one.

So on this stack's dual config: +50% prefill, flat decode, zero detectable quality cost — including the sign-fold path for AutoRound checkpoints.

Catalog wiring is next: the patches ship with the composes (inert until enabled) and W4A8 becomes a one-env serve-time knob. The interim recipe stays as posted above.

@chriskerley78910 — the config you called "basically the best possible for dual Ampere" now has quality receipts (and an honest prefill number).

[chriskerley78910]

Very excited about this!

So a slight increase of decode over the current 4-bit INT config (the fastest decode so far) AND a ~50% increase in prefill speeds!!!