volta-kernel-bakeoff

Finished 2026-10-23 and archived. What came out of it:

• Our own kernels: at one token our 4-bit kernel beats cuBLAS 2.6-3.2x and roughly matches llama.cpp. At 8 tokens and for prompts it loses: the tensor-core prompt kernel reached 57% of cuBLAS at 512 tokens, under its 70% bar (STAGE3.md), and the 8-token gap closed only to about a tie (issue #2).
• Fusing llama.cpp's small kernels (FUSION.md): five fusions, about +6% faster decode on Qwen3.6-35B-A3B, with output identical to stock in every check (text, perplexity, KL divergence). That figure is an estimate: +4.97% measured for the first four together, times +1.14% for the fifth. Patches in patches/.
• On a V100 box, the CPU's power saving costs llama.cpp about 1.7% of decode speed: its cores sleep between the host's steps.

Our own quantized matrix-multiply kernels for the Tesla V100 (sm_70, 2017), raced against cuBLAS (NVIDIA's fp16 tensor-core GEMM) and llama.cpp's own CUDA kernels, at the real table shapes of the models we run.

Numbers below are medians over 30 rounds of every contender -- ours, cuBLAS and llama.cpp -- in one session on 2026-10-20 with the SM clock locked; every round is in results/2026-10-20, medians in medians.csv (results/summarize.py).

Why the V100 is interesting here

The V100 has fp16 tensor cores but no int8 / int4 ones (Turing added those), so a quantized model's weights must be unpacked on the ordinary CUDA cores. When a model writes one token at a time, the work is reading the weights from memory, so the fastest kernel is the one that reads the fewest bytes and unpacks them cheaply where they are used.

What's in the race

contender	what it does
cublas_fp16	cuBLAS GemmEx on fp16 weights (tensor cores; 4x the bytes of 4-bit)
dq4+cublas	unpack 4-bit to fp16, then cuBLAS
fused_q4, fused_q4_v2, fused_q4_v2w	ours: one warp per output row, unpacks in registers, fp32 sums; v2 = fp16 bit trick (0x6400 | n is the half 1024 + n), v2w = a whole 32-value group per lane step
fused_q8	ours, int8 weights
moe_*	the MoE case: 8-10 experts, one token each, in one launch vs one cuBLAS call each vs cuBLAS batched
q2w_alu_v3	ours, plain 2-bit (codes -2..1, one fp16 scale per 32 = 2.5 bits per weight)
q2_lut, q2w_lut_*	the "times table": precompute the 16 possible sums per pair of activations, look them up instead of multiplying
tc_q4	ours, tensor cores for prompts: unpack a tile to shared memory, WMMA 16x16x16
llama_*	llama.cpp's kernels via ggml_mul_mat (Q4_0, Q8_0, Q2_K, IQ2_XS, IQ2_XXS, F16), bench/llamacpp_mm.cpp

Shapes: Qwen3.8-Flash-Next Coder experts 640x2560 (10 per token), its attention 12288x2560 and 2560x6144; Qwen3.6-35B-A3B experts 512x2048 (8 per token) and attention 8192x2048; Qwen3-1.7B MLP 6144x2048 / 2048x6144. Tokens per call: 1 to 512.

Making it fair

• Real memory reads: the V100's L2 is 6 MB, larger than one expert, so every contender cycles through enough copies of its weights (>= 96 MB) that each call reads them from HBM, as in a real model.
• Correct before fast: ./bakeoff --check compares every kernel with a CPU double-precision reference of the same quantized weights, at 1, 2, 4, 8, 64 and 512 tokens and for the ten-experts-in-one-launch kernel (each expert with its own weights and token, so a wrong expert index fails too): 440 checks (check.csv; both check files come from a checks-only pass of the runner right after the timing session, with the build that added the 512-token checks -- the in-session check in run.log was 404 ok and 310 sabotage failures without them). BAKEOFF_SABOTAGE=1 breaks the packed weights on purpose and the check must fail: every kernel with packed 2-, 4- or 8-bit weights fails every one of its checks (328 failures, check-sabotage.csv); the two cuBLAS fp16 contenders have no packing to break and pass, as a control.
• Same clock for everyone: the runner locks the V100's SM clock at 1,380 MHz (nvidia-smi -lgc), stops if the lock doesn't take, and resets it at the end; all 30 rounds read 1,380 MHz at 35-48 C (gpu.log).
• Each round takes the median of 5 x 50 calls after warm-up; the table reports the median round. Fastest and slowest round differ by 0.9% typically.
• Under load (issue #7, results/2026-10-20-load): a 30 + 30-round session sampled every 100 ms -- every round busy, ECC 0; on the heaviest sweep shapes the card touches its 250 W power cap for a moment (1% of busy samples, at worst 1,357 MHz). Its 964 medians match the unsampled session's to 0.04% typically (summary).
• dq4+cublas (unpack to fp16, then cuBLAS) cycles its fp16 copies too, but each multiply still reads the copy it has just written, which on an expert-sized table is still in L2 -- as it would be in a real model.

Results (one token, microseconds, lower is better)

table	cuBLAS fp16	llama.cpp Q4_0	ours 4-bit	llama.cpp Q2_K	llama.cpp Q2_0	ours 2-bit
Flash-Next Coder attention 12288x2560	101.1	32.4	30.2	31.0	18.3	23.5
Qwen3.6 attention 8192x2048	48.8	20.4	17.8	19.3	12.1	14.3
Qwen3-1.7B MLP up 6144x2048	43.7	16.4	13.9	16.0	10.2	12.1
Flash-Next Coder expert 640x2560	10.7	6.8	6.3	8.9	7.5	6.2
Qwen3.6 expert 512x2048	9.7	7.2	5.2	8.2	8.2	5.5
Qwen3.6 expert down 2048x512	6.9	6.6	4.7	7.8	7.3	4.0

Ours 4-bit = fused_q4_v2w, ours 2-bit = q2w_alu_v3 (one fp16 scale per 32, 2.5 bits per weight). llama.cpp's Q2_0 (one fp16 scale per 64, 2.25 bits) is the same kind of plain 2-bit format. llama.cpp's times include converting the activations to 8 bits for each multiply, as in a real model; our times in this table do not -- they start from activations already in fp16. That step costs ~2 us per call (mostly its launch), so this table flatters ours, most on the small tables: the like-for-like numbers are in the bullets below (issue #2). Earlier runs (2026-10-19, clock unlocked; 2026-10-20-llama30) agree with these medians to 0.09% typically; the exceptions are dq4+cublas on small tables (up to 40% slower now that its copies are cycled) and llama.cpp's IQ2_XS (~10% faster with the clock locked).

• 4-bit, one token, like for like (fused_q4_v3_fp32in: fp32 activations in, the conversion counted; 30 clock-locked rounds, results/2026-10-20-gap8-30r): at or below llama.cpp's Q4_0 on 8 of 9 -- 1-9% less time on the big tables, 9-32% on the expert-sized ones. The ninth is Qwen3-1.7B's gate + up + SwiGLU, which llama.cpp runs at one token as one fused call: 27.65 us against 27.20 (2% slower). (Without the conversion counted, the table above, ours looked 7-30% faster on 7 of 9.)
• 2-bit, one token: ours takes 12-49% less time than llama.cpp's Q2_K, but Q2_K is a more complex format (a scale and offset per 16 weights, 256 per block). Against llama.cpp's own plain 2-bit, Q2_0, ours loses on the big tables (18-29% slower) and wins on the expert-sized ones (17-46% less time). Nsight Compute on the Qwen3.6 attention table: ours keeps the GPU's arithmetic 60% busy and its memory only 37% -- it is limited by unpacking -- while Q2_0 converts the activations to 8-bit integers once and multiplies four weights per instruction (dp4a).
• 2-bit quality, measured (Qwen3.6-35B-A3B, WikiText-2, llama.cpp; from the gsq-gguf converter, not yet published): a plain 2-bit format with one scale per group, filled with ISTA-DASLab's GSQ weights (about the best such a format can hold), scores perplexity 7.93; Unsloth's UD-Q2_K_XL (llama.cpp's K and IQ formats, ~2.6 bits on the experts, same file size) 7.26; its 4-bit UD-Q4_K_XL 6.73 (lower is better). The plain format's speed comes with a measurable quality cost.
• Ten experts in one launch: 22 us against 56 for cuBLAS batched and 109 for ten separate cuBLAS calls (Qwen3.6's eight: 12 us against 31 and 79).
• A few tokens at once (issue #2), like for like: fused_q4_v3 (src/stage5_gap8.cuh) unpacks two codes per instruction into fp16 pairs, multiplies in half2, lets one activation load serve several rows and splits short tables across warps; fused_q4_v3_fp32in adds the fp32 -> fp16 conversion by one rule (in the kernel at 1 token and on expert-sized tables up to 4 tokens, a separate step otherwise). At 8 tokens (30 rounds) it is at or below llama.cpp's Q4_0 on 4 of 9 (Qwen3-1.7B's gate + up + SwiGLU counted as one, in place of its MLP-up table, at every token count): three of the four expert tables (5-13% less time) and Qwen3-1.7B's gate + up + SwiGLU (49.0 us vs 54.1 -- ours stacks gate and up into one table and converts once; llama.cpp runs them as two calls at 2-8 tokens). The big tables take 2-18% longer (Flash-Next attention 54.2 us vs 53.1; it was 83 before the rebuild); the 6,144-long rows (attention out, MLP down) are the furthest behind. At 2 and 4 tokens: 6 and 4 of 9. Without the ~2 us conversion step Flash-Next attention and Qwen3-1.7B MLP up would be ahead and Qwen3.6 attention level -- the conversion is the gap there, and folding it into the step that makes the activations is issue #3's work. Precision: products and short partial sums are rounded to fp16, so it is checked like cuBLAS fp16 (2e-3), also with large activations planted (+-20000..60000, as real down_proj inputs) and tiny ones; weights are unpacked as (code - 8) / 32 (/ 64 where a lane sums two words), so no fp16 activation can overflow a partial sum (the first version, / 8, failed the large case). Limit: activations above 65,504 in fp32 do not fit fp16 at all and come out as infinity -- llama.cpp's 8-bit conversion scales each block, so it has no such limit. A model whose activations go that high (some down_proj inputs do) needs them rescaled first, as the FP16 fixer does. Clock: in this run our side touched the V100's 250 W power cap in 24 of its 30 rounds (33 of 4,379 busy 100 ms samples, 0.75%, lowest 1,342 MHz); llama.cpp's side never did (0 of 16,974). Which of our kernels was running during each dip isn't recorded; in issue #7's run the same kind of dips moved results by 0.04%.
• Where we lose: prompts. The rebuilt tensor-core kernel (tc_q4_v2) runs at 51-62% of cuBLAS's speed at 512 tokens on the big tables (see STAGE3.md); it beats cuBLAS up to ~128 tokens.
• The times table loses on this card. Against a badly built multiply kernel it looked 13-17% faster; against a good one it is slower everywhere -- the V100 multiplies cheaply, building and reading the table costs more.
• Found on the way: a 2-bit kernel that read its activations from shared memory at a 16-float stride per lane was slower than 4-bit (bank conflicts); reading them from global memory made it 1.9x faster. Even so, our best 2-bit run moves about 420 GB/s of the card's ~900. A single expert is too small to keep the card busy: Nsight Compute shows every kernel, ours and llama.cpp's, at 5-19% of memory speed on a 512x2048 expert, with ~90% of the card idle.

Why llama.cpp's 2-bit can't do every table

llama.cpp's 2-bit formats (Q2_K, IQ2_*) pack 256 weights per block, so a table whose rows are not a multiple of 256 long can't use them -- our harness skips those (K is not a multiple of 256). Flash-Next's expert "down" tables have rows 640 long (2.5 blocks). The published Strata Coder file (ISTA-DASLab Qwen3.8-Flash-Next-GSQ-RCO-Coder-GGUF) shows the consequence: its expert gate/up tables (rows 2,560 long, 10 blocks) are stored at 2-3 bits (IQ2_S, IQ3_XXS, IQ3_S), while its expert down tables are IQ4_NL (4.5 bits, blocks of 32) in 39 of 48 layers and Q2_0 (2.25 bits, blocks of 64) in the other 9. Formats with 32-, 64- or 128-weight groups -- ours, and ISTA's GSQ (groups of 128; 640 = 5 x 128) -- fit those rows.

Stage 3: what a whole model spends its time on

Before putting our kernels into a real model, we profiled one: Qwen3.6-35B-A3B in llama.cpp on the V100 (Nsight Systems for whole tokens, Nsight Compute for single kernels). Full notes and tables: STAGE3.md.

• llama.cpp already does what stage 1 set out to show. On Volta it runs all of a layer's experts in one launch and fuses gate + up + SwiGLU (routed and shared experts together), the router's softmax + top-k, and norm + RoPE.
• A written token is ~1,526 kernels, and the GPU is busy the whole time (10.5 ms of kernels in a 10.4 ms token, measured without the profiler). The multiplies are about half: the 8-bit attention and output-head tables 38%, the 2-bit experts 15%. The other half is ~1,200 small kernels of 3-5 us each (norms, element-wise steps, copies, converting activations to 8 bits, routing).
• The weights are not the bottleneck. A token reads 2.4 GB, about 3 ms at the card's speed. Halving the experts' size (GSQ 2-bit vs Unsloth 4-bit) did not change writing speed (96 vs 95 tokens a second).
• Reading a long prompt: the linear attention's gated-delta-net kernel takes 15% of a 4,096-token prompt, one serial kernel per layer -- llama.cpp's source marks a chunked version as a TODO.
• Drafting is the biggest writing-speed lever on this card: llama.cpp's MTP drafting gives +33-37% (93 -> 128 tokens a second), more than any kernel here could.

What we'd do next, in order:

1. Fuse the small steps (about half of a token): norm + 8-bit conversion, the element-wise chains around the linear attention, router + reduction -- in a private llama.cpp build, measured end to end.
2. A chunked gated-delta-net kernel for prompt reading (Gated DeltaNet's chunkwise form), filling llama.cpp's own TODO; it matters most for long prompts (an agent's 27K-token requests).
3. A Marlin-style prompt kernel for Volta (hardest of the list): unpack quantized weights inside a tensor-core multiply, instead of llama.cpp's dequantize-to-fp16-then-cuBLAS, which is 39% of reading a 4,096-token prompt. Marlin itself needs Ampere (cp.async, ldmatrix, the m16n8k16 MMA); on the V100 the asynchronous copy would be replaced by register-staged double buffering (as GEMMs did before Ampere) and the MMA by Volta's own. Our first tensor-core kernel (tc_q4) takes 1.6-2.3x cuBLAS's time at 512 tokens, so this starts well behind (go / no-go bar: STAGE3.md, issues #1 and #5); the unpacking itself (the 0x6400 trick, from FasterTransformer) is the part we already have.
4. Our 8-bit kernel in place of llama.cpp's Q8_0 at one token (76% vs 68% of memory speed on the table family that is 38% of a token): roughly 4% end to end, cheap to try.
5. Not worth it: our 2-bit kernel in a model (llama.cpp's Q2_0 is faster on big tables, and experts are 15% of a token), and a new engine in Python around our kernels (it would add launch cost to every one of ~1,500 kernels a token, which llama.cpp's CUDA graphs avoid).
6. Still to measure: Strata, which writes Qwen3.6 faster than llama.cpp on this card (~110 vs ~95 tokens a second), profiled the same way -- to see what it does differently.

Run it

make                                   # CUDA 12.9, sm_70 (nvcc; does not run the binary)
python3 bench/test_load_flag.py        # gpu-load sample checks, no GPU
./bakeoff --check                      # correctness (exit 1 on a failure; needs the V100)
./bakeoff > sweep.csv                  # the sweep
bench/overnight.sh                     # check, sabotage check, ROUNDS sweeps, LLAMA_ROUNDS llama.cpp sweeps

GitHub Actions (.github/workflows/ci.yml) runs the sample checks, ruff, and shellcheck, and compiles with CUDA 12.9's nvcc and cuBLAS packages for sm_70. That job does not run the binary.

Hardware: Tesla V100-PCIe-32GB, i7-13700KF, 16 GB RAM, Linux, CUDA 12.9, driver 580.173.02.

Made with AI assistance (Claude); the idea of the bake-off came from an AI research assistant (Muse). MIT licence.