gitarist/Qwen3-8B-BPDQ-2bit

Qwen3-8B — 2-bit BPDQ (packed)

Qwen3-8B with the transformer weights stored as 2-bit bit-planes and kept that way for the whole run. They are never expanded into a dense fp16 tensor. 3.0 GB on disk, 4.3 GiB resident, against 15.3 GiB for the same weights as fp16.

Runs on one CUDA GPU or an Apple M-series Mac.

Run it

Three files sit on top of the weights: bpdq.py, demo.ipynb and this README. bpdq.py works out what it is running on and picks the kernels.

import sys; sys.path.insert(0, MODEL_DIR)
import bpdq

model, tok = bpdq.load(MODEL_DIR)                      # cuda | mps | cpu
ids = tok("The capital of France is", return_tensors="pt").to(model.device)
out, s_per_step = bpdq.decode(model, ids.input_ids, 64,
                              sampling=bpdq.Sampling.for_mode(think=False))
print(tok.decode(out[0]))

decode takes a batch; ragged prompts are left-padded (tok.padding_side = "left", pass the attention_mask). Pass a transformers streamer to stream.

On CUDA, serve through vLLM — importing bpdq registers the method:

import sys; sys.path.insert(0, MODEL_DIR)
import bpdq
from vllm import LLM

llm = LLM(model=MODEL_DIR, quantization="bpdq", dtype="float16",
          max_num_seqs=64, reasoning_parser="qwen3")

demo.ipynb detects the machine and gives one generate() either way.

python bpdq.py selftest [MODEL_DIR]   # kernels vs a dense reconstruction
python bpdq.py bench    [MODEL_DIR]   # per-matmul and end-to-end
python bpdq.py chat      MODEL_DIR    # interactive, streaming

Decoding

bpdq.Sampling.for_mode(think) carries Qwen3's published settings:

	temperature	top_p	top_k	min_p
think	0.6	0.95	20	0
nothink	0.7	0.80	20	0

Do not decode this checkpoint greedily. At 2 bits it falls into restating the same paragraph. Sampling also sets no_repeat_ngram=12 on the torch path and presence_penalty=0.3 on the vLLM path — each backend's native control for that.

It rarely stops thinking on its own. On vLLM, Sampling(think_budget=N).vllm() sets thinking_token_budget, which needs LLM(..., reasoning_parser="qwen3") so it knows where </think> is. The torch path has no equivalent; give think mode a generous max_tokens there.

Measured

RTX 4090, vLLM 0.28, 64-token decode. The fp16 column is these same weights reconstructed dense and served by the same vLLM, so only the matmul differs.

concurrent seqs	packed 2-bit	dense fp16
1	133.5 tok/s	58.3	2.29x
4	521.0	226.0	2.31x
16	1963.4	886.6	2.21x
64	3883.3	3084.0	1.26x
128	4727.2	5359.2	0.88x
weights resident	4.27 GiB	15.26 GiB

Decode is bandwidth-bound until the batch amortises the weight read, so the format wins by roughly its size ratio to ~16 sequences, still wins at 64, and loses past ~100 where both are tensor-core-bound.

Apple M4 (10-core GPU, 17 GB unified), 230-token prompt. fp16 Qwen3-8B does not fit on this machine.

batch	decode	prefill
1	20.4 tok/s	142 tok/s
4	44.9	174
8	58.1	175
16	66.6	176

Per matmul against dense fp16 of the same shape: 3.7–5.8x at batch 1, 1.2–1.7x at 16, 0.68–0.85x from 32 up.

Quality

	value
WikiText-2 PPL (seq 2048)	18.58    (fp16 Qwen3-8B: 9.72)
KL to fp16 (32 chunks)	0.864
GSM8K think, 8-shot, 250q, 1024 tok	0.588

The packed path and a dense reconstruction of the same weights agree to fp16 round-off — worst relative error 5e-4 over all 252 weight matrices, checked by bpdq.py selftest.

Format

w[r, c] = sum_i coeffs[g(c), r, i] * bit_i(r, c) + coeffs[g(c), r, msbits]

Two bit-planes per weight, group_size = 256, one fp16 coefficient set per (group, output row) plus a per-group offset — about 2.2 bits per weight. embed_tokens and lm_head ride as 8-bit GPTQ-v2 integer sets and are the only tensors expanded to fp16 at load; they are roughly half the 4.27 GiB.

At load the packed tensors are repacked, not decoded:

codes   int32  [in_features / 16, out_features]    16 columns per word, 2 bits each
coeffs  fp32   [groups, 3, out_features]

out_features is innermost so a warp (or SIMD group) covering 32 rows reads one contiguous line, and a column's two bits are interleaved into one code so a weight is peeled with code = word & 3; word >>= 2.

Kernels

CUDA is one Triton kernel for every batch size: each program owns a [BT, BM] output tile, walks K, rebuilds a weight slab in registers from the code words and hands it to tl.dot. The weight is read once per tile whatever the batch. K is split across programs when a decode-shaped launch would not fill the card.

Apple is two Metal kernels. Decode reduces in registers against a 4-entry value table per (row, group) held in registers for a whole group. Prefill feeds simdgroup_half8x8 matrices from a weight slab rebuilt in threadgroup memory. Neither writes anything dense to device memory.

Limits

• TP = 1. No tensor parallelism.
• Past ~100 concurrent sequences on a 4090, dense fp16 is faster.
• On Apple, batches of 32 tokens and up run at 0.68–1.0x of a dense fp16 matmul. The win is at decode batch sizes.
• At 2 bits the model invents confident detail about anything obscure, and past roughly 250 tokens long open-ended answers drift. Use it for short exchanges and bounded tasks.
• The torch decode loop does not capture CUDA graphs; transformers bakes part of the attention mask into a capture and the replayed step drifts. vLLM's own capture is fine.
• Greedy output is not bit-stable across cache implementations — this loop uses a StaticCache, model.generate a dynamic one. Both are faithful to the weights.
• Launch shapes were measured on an RTX 4090 and an M4; other chips run correctly but may not be at their own optimum.

License

Apache-2.0, inheriting Qwen/Qwen3-8B.