[Just for Fun — Advanced] Why Isn't Your 4-Bit Quant Faster on a 2080 Ti? I Tore Open the CUDA Backend to Find Out

Intro

Repacking a truckload of cargo from big boxes into small ones saves warehouse rent. It doesn't save labor — the movers still unpack every small box back into a big one before they carry it. You paid to shrink the boxes, and the carrying cost didn't move.

That's the trap in this article. You quantize a model. The file drops to half its size, but tok/s doesn't change. Sometimes it gets worse.

This is part 12 of the modded 2080 Ti series, picking up right where #11 left off. That one got MiniMax-H3 running on this card with INT8 and ruled out W4A4 on quality grounds — but it left one question unanswered: why is there nothing between them, no way to keep 4-bit weights without sacrificing quality? This article follows that question all the way down to the instruction level.

Turing Has No Instruction for 4-Bit Times 16-Bit

I dumped comfy-kitchen's CUDA backend, _C.abi3.so, with nm -D and c++filt, then classified every symbol by cutlass::arch::SmXX and by operand width. Look at the last two columns — they're empty on every architecture:

Arch	s4×s4	s8×s8	s4×s8	s4×f16
Sm75 (Turing, this card)	32 symbols/dtype	64 symbols/dtype	—	—
Sm80 (Ampere)	—	96 symbols/dtype	—	—
Sm89 (Ada)	56 symbols/dtype	—	—	—

The hardware instructions mma.s4.s4 and mma.s8.s8 exist. mma.s4.s8 does not. Both operands to a tensor core matrix multiply have to be the same width — you can't multiply a 4-bit weight against a 16-bit activation in one instruction, on any of these three architectures. This isn't a software gap. The instruction simply doesn't exist.

So GGUF, AWQ, GPTQ, and Marlin All Take the Same Detour

Every 4-bit format on the market routes around the missing instruction the same way: dequantize the 4-bit weights back to fp16, once per layer, then run an ordinary fp16 GEMM.

The TurboMind paper (arXiv 2508.15601v2) spells out the sequence for its own INT4 path: load the INT4 weights, convert with I2F, dequantize to FP16, then run the FP16 tensor core GEMM. That's the same dequant-to-fp16 principle GGUF's Q4 kernels use.

Quantization on this card buys VRAM headroom. It does not buy time.

The Proof: a Smaller Quant Format Ran Slower

If the win were bandwidth, a smaller file should always be faster. IQ3_M is smaller and more aggressive than Q4_K, and it measured slower on all three workloads:

Workload	Q4_K	IQ3_M	Change
struct	45.9 tok/s	37.1 tok/s	−19%
reasoning	45.3 tok/s	39.9 tok/s	−12%
create	29.3 tok/s	28.6 tok/s	−3%

Q4_K's dequant is a linear block multiply-add. IQ3_M's is a nonlinear codebook lookup — much heavier per weight. That moves the bottleneck from VRAM bandwidth to dequant compute. "Smaller file means faster" needs two things to hold: the workload has to be genuinely bandwidth-bound, and the dequant has to be cheap. Turing fails both.

The One Native 4-Bit Path Is Broken by Design

W4A4 — 4-bit weights and 4-bit activations — does have a native instruction. At the kernel level it's fast: blocks.0.attn.out_proj went from 19.256ms in eager mode to 1.373ms in CUDA, a 14.02× speedup, and the biggest matrix in the model, which OOMs in eager, completes in 3.063ms under W4A4. Caveat: that was measured on a nearly-full GPU and may be inflated. End to end, across a full fp32 pipeline, it only bought 6.3%.

Quality is what kills it. Relative error is 0.0110 for INT8 versus 0.2005 for W4A4 — 18× worse. The full autopsy, including the color tearing it produced on a video model, lives in #11's "Why W4A4 doesn't work" section.

So this card has exactly two rungs on the quantization ladder: W4A4, native but broken, and INT8, native and fine. There's no middle band.

What to Actually Run

This table covers the decision; the row underneath is the part people skip.

Situation	Use	Why
Model doesn't fit in VRAM	GGUF Q4	It runs — don't expect the dequant to be free
Fits, and you want speed	W8A8-INT8	Only natively accelerated quant path: 2.07–2.83× vs FP16, measured on Wan FFN shapes
NVFP4 / FP8	Skip	Emulated here — GGUF Q4_K_M ran 53.083s vs NVFP4's 142.434s (video-gen benchmark), GGUF 2.68× faster
Quality doesn't matter	W4A4	Fast but 18× the error of INT8

INT8 isn't a free pass either. Wan 2.2's LowNoise DiT under W8A8 dropped to SSIM 0.345642 — visible confetti noise in the output. Always eyeball a quality A/B before you ship a quant, even the one with a real hardware path.

Quantization on this card still saves the rent. The labor bill doesn't shrink.

Deep dive: the W4A8 case file

Reading this section is optional. It doesn't change what to run — that's already decided above. It's the investigation that got me there, including a hypothesis I killed with my own measurement and a claim from July I had to walk back.

Reading 36% fewer weight bytes — why didn't I get a single second back?

comfy-kitchen's W4A8 path — 4-bit weights, 8-bit activations — ships with a claim from its author of roughly 1.09× INT8's speed. My expectation: W4A8's weights come to 12.54GB against INT8's 19.53GB, 36% less to read per step. Under fixed memory bandwidth, that alone should win.

I measured both on the same H3 DiT from #11, under identical conditions: 864×480, 124 frames, 14 steps, Spectrum on, seed 42.

Format	Time	Residency
INT8	238.55s	loaded partially, 778.39MB offloaded
W4A8	454.19s	loaded completely, zero offload

W4A8 was 90% slower with a fully resident model and zero offload — the best-case residency scenario. If it were bandwidth-bound, that's exactly backwards. The slowness is compute, not memory traffic, which means tuning --reserve-vram for this comparison is pointless. My own experiment killed my own hypothesis.

A related datapoint, with an explicit caveat attached: on the H3 video DiT (2026-08-05, same box), a fully resident Q3 GGUF took 189.23s against an offloaded INT8's 110s — 72% slower. But the two arms there differ in more than quantization (Q3_K_M vs INT8-ConvRot, different residency), so on its own it only proves that a fully-resident GGUF still loses to an offloaded INT8. The W4A8-vs-INT8 pair above is the clean isolation.

The banner prints asym_w4a8_int8 — so why is it slow?

First suspicion: maybe W4A8 isn't actually dispatched as a native op — maybe it's just masquerading as one while a slow generic kernel runs underneath. The startup banner seemed to refute that; it lists asym_w4a8_int8 under Native ops:.

Reading the code showed the suspicion was right, just not in the way I expected:

# around comfy_kitchen/backends/cuda/__init__.py:2258
used = _C.cutlass_int8_dequant(...)      # returns False on Turing (guard requires major >= 8)
if not used:
    return eager_w4a8_int8_linear(...)   # ← the slow path

_cuda_device_supports_cutlass_int8_dequant requires major >= 8. Turing is major == 7, so every call falls into eager_w4a8_int8_linear. The banner only means "the runtime recognizes this op's name." It says nothing about whether that op has a fast path on your card.

That distinction generalizes, so I kept a table of it:

What you see	Why it isn't evidence of a fast path
Native ops: lists the op name	Not-in-disabled-list, not hardware acceleration
--force-fp16	Storage format only
Using sage attention in the startup banner	The flag was accepted, not that it dispatches fast
pip install sageattention==2.2.0 succeeds	Installable isn't the same as dispatchable on this card

Is the guard hiding missing software, or missing hardware?

My assumption after finding the guard: the cutlass fused-dequant template just wasn't compiled for sm_75 yet — a software gap, patchable with enough time.

I dropped below the application layer and went back to the nm -D dump, this time with output dtypes broken out:

Arch	Operand width	Output dtype	Symbols
Sm75	s4×s4	bf16	32
Sm75	s4×s4	f16	32
Sm75	s4×s4	f32	32
Sm75	s8×s8	bf16	64
Sm75	s8×s8	f16	64
Sm75	s8×s8	f32	64
Sm80	s8×s8	bf16	96
Sm80	s8×s8	f16	96
Sm80	s8×s8	f32	96
Sm89	s4×s4	bf16	56
Sm89	s4×s4	f16	56
Sm89	s4×s4	f32	56

No mixed-width symbol shows up on any architecture, at any output dtype. The absence is at the instruction-set level, not the compiled-template level. That distinction matters because the fix is different: a software gap means wait for an update; a hardware gap means change your method.

Two landmines I stepped on along the way

(a) A silent METADATA claim. comfy-kitchen 0.2.28's wheel METADATA declares no torch dependency, which reads as version-agnostic. But its na.py uses PEP-585 list[int], and torch 2.6.0's infer_schema only accepts typing.List[int]. Nothing errors until schema inference actually runs. Production was down for 4 minutes. METADATA silence isn't proof of no dependency.

(b) My own leftover script. h3_rescale.py, with SCALE=1.0, is a numeric no-op — but it still calls sc.mul_(INV) unconditionally. W4A8 stores weight_s_rel as fp8, and Turing has no fp8 units, so the call throws "mul_cuda" not implemented for 'Float8_e4m3fn'. Code written for an old config becomes a roadblock in a new one, and it never announces that it's idle.

What my July selection matrix got wrong — and what it still got right

My selection matrix from July claimed three things: AWQ/W4A16 uses native mma.sync.s4 on Turing, Marlin needs sm_80+, and AWQ beats GGUF by 1.5–2×.

Six days later, the TurboMind paper overturned the first claim — it's dequant-to-fp16, not a native mixed-width instruction. A month after that, vLLM v0.27.1's docs overturned the second — Marlin has supported sm_75 since roughly v0.14.

Both specific claims were wrong. The underlying conclusion never moved: Marlin is still a dequant kernel, just a well-optimized one, and at batch size 1, no 4-bit format on this card runs on native mixed-width MMA. The software path changed; the hardware limitation didn't.

A signal worth waiting for

The s8×s8 kernel is already in the official comfy-kitchen wheel. The only thing blocking W4A8 on this card is the one arch guard at major >= 8. If fused dequant opens up for sm_75, interpolation puts W4A8 (12.54GB) around 104s — 12% faster than INT8, roughly matching the author's ~1.09× claim.

That's interpolation, not a measurement, and even verifying it is currently walled off. My attempt to build it from source died at phase 0: the official build requires CUDA ≥ 12.8, this box has nvcc 12.4, and no .so was ever produced. For now it's a reason to wait, not a verified result.

---

Also in this series:

• #9 0xc0000409: When My AI Service Died Silently and the Log Ate the Evidence
• #10 Two Identical 2080 Tis, One 3.4× Slower — the Culprit Was a Single dtype Log Line
• #11 Wait, a 2080 Ti Can Run MiniMax-H3? 1080p With Audio on a 2018 Card