From gibberish to 61 tok/s: building llmgo on an RTX 5060 Ti

The short version

llmgo now decodes Qwen3.8-27B at 56–61 tokens a second at ~10k tokens of context on a 16 GB RTX 5060 Ti, and 44–53 tokens a second at 60–76k. It started 28 runs earlier answering in gibberish.

llmgo is a dependency-free Go inference engine for GGUF models, and it now serves Marvin through its UI and gateway router. It serves Qwen3.8-27B-UD-Q2_K_XL: 64 layers, 48 of them Gated DeltaNet (linear attention) and 16 full attention, head dim 256, a 248k vocabulary, and one MTP (multi-token prediction) layer that drafts the next tokens. The weights are 9 GB of mostly 2- and 3-bit IQ quants; the K/V cache is q4_0, with 130k tokens of context available.

The card is an RTX 5060 Ti 16 GB (Blackwell, sm_120). Its VRAM reads at 426 GB/s measured, so reading every weight once costs ~22 ms. Every speed-up below is ultimately about how close a decode step gets to that number.

Measure	llmgo today
Decode, ~10k context	56–61 tok/s
Decode, 50–76k context	44–53 tok/s
Prefill, 9.6k prompt	11.4–11.7 s

The steps that moved the ~10k number

Most of the ~10k speed came from a handful of early steps; the chart further down shows them as a curve.

1. All 64 layers on the GPU, and gibberish (s01–s04). The first deploy answered in nonsense. Putting 63 layers on the GPU (s02) or turning MTP off (s05) gave correct answers at ~28 tok/s. So the fault was in the MTP path.
2. The MTP verify fix (s06): 41 tok/s. With the output head on the GPU, the verify step recomputed every position's logits on the CPU from stale activations and overwrote the GPU's. Fixed, 88% of drafts were taken and each forward produced 2.58 tokens.
3. The production config (s09): 34 tok/s. Moving to the production setup (130k available context, q4_0 K/V cache, 1024-token prefill chunks, prompt cache) cost speed, but from here every number is what Marvin actually runs.
4. MTP drafting on the GPU, small chunks on the matvec (s10–s11): 45–46 tok/s. Priming the draft dropped from ~300 ms to 40 ms, and 8-token chunks stopped taking the dequantise-and-GEMM path.
5. Seven drafts, top-k on the GPU, the whole model kept on the GPU (s12): 49 tok/s. Up to 7 drafts at p-min 0.8; the output head's top-k on the device cut logits per forward from 11.6 to 4.7 ms.
6. k_mmv, verify on int8 tensor cores (s16): 52–54 tok/s. A verify forward checks up to 8 positions. The matvec did each position's dot products in turn; k_mmv decodes each weight once and multiplies all positions with one tensor-core mma.
7. The matvec reads 2 or 4 rows per lane (s18): 54 tok/s. Each position's activations are loaded once for several rows.

Why tokens per second lies a little

From s16 on, a forward at ~10k has cost 57–64 ms, and the tok/s spread around it is mostly how well each answer drafted.

With speculative decoding, tok/s = tokens per forward ÷ time per forward. Tokens per forward depends on how predictable the text is: 2.9 on one answer, 3.4 on the next. The same build (s18, restarted twice as s19 and s20) gave 54.2, 50.9 and 49.6 tok/s. A forward that verifies more positions also takes longer, so a good-drafting answer raises both numbers.

So from s21 on, a change only counted if the milliseconds per forward moved. Runs that didn't move it still went into the log, the CSV and the write-up, but not onto the progress chart. That rule cut a lot of false hope: several builds that looked 5 tok/s faster were the same kernel on a luckier answer.

Measuring instead of guessing

The middle stretch (s16–s24) produced few gains, and the lesson was to stop guessing and instrument.

• Guesses that didn't pay. An L2 prefetch hint did nothing in three runs and was removed. A cp.async weight ring for k_mmv doubled its shared memory; the theory was lost occupancy, but the probe later showed 6 blocks per SM without the ring were slower than 3 with it.
• Removal probes, and their limit. Startup probes timed each kernel with a part removed: no weight load, no mma, no activation staging. They showed loads and compute adding up instead of overlapping. But removing one part changes how the others overlap, so after a while they stopped being informative.
• Every weight product by role and position count. The profile's CUDA events were tagged with the weight's role (ffn_down, dn_qkv, attn_k, the output head, MTP), the forward's position count and the kernel used. This showed the int8 matvec does ~80% of decode weight time, and that each extra draft position costs ~3.5 ms.
• In-kernel counters. clock64 reads at phase boundaries, on one block in eight, show where warps spend their cycles. k_mmv: 41% unpacking weights, 28% mma, only 4–5% waiting for memory.
• Memory vs arithmetic. Each weight shape timed three ways: its bytes read alone, the product as decode runs it, and the product with the weights served from cache. At 2 positions the arithmetic alone (26.1 ms) already exceeds the memory floor (21.9 ms): the extra positions cost instructions, not bandwidth.
• A Grafana dashboard over the same counters, so a real working session shows the picture, not just a benchmark.

The last stretch: what the counters found

Once the counters were in, the fixes were small and obvious, and they took ~10 ms off a forward at 60k (66.1 → 56.6 ms).

• Attention block shapes (s25). Attention by position count showed the commonest verify (2 positions) running at half its warps: 12 query rows fill one 16-row tile, so only 2 of 4 warps worked. 6–8 positions ran 3 warps. Reshaping the blocks made attention flat across positions: at 33.7k, n=2 went 11.4 → 6.5 ms.
• Quantise each input once (s25). q/k/v, gate/up and DeltaNet's four projections each share an input, but every product re-quantised it. Doing it once: 1.9 → 1.05 ms per forward.
• Merge tiny products (s25). DeltaNet's beta and alpha are 48 rows each and ran at 10% of the roof. They sit back to back in VRAM, so they now run as one product.
• Tune the output head (s27). The 248k-row head had never been tuned and ran on the default matvec layout at 52% of the roof. Tuned: 3.2 → 2.1 ms per forward, at its floor for one position.
• Long-context attention layouts (s28). A real coding session ran to 105k, where attention reached 16–20 ms of a ~70 ms forward and read the K/V cache at only 20–25% of the roof. At head dim 256 each block needed 76–93 KB of shared memory, so one block per SM, and its loads waited for its maths. Loading K, then V, through one buffer halves that, two blocks fit, and one's loads run under the other's maths. Three layouts are timed at startup and the fastest kept per position count: 35% faster for 1–5 positions.

Context	Forward, s23	Forward, s28	Attention, s26	Attention, s28
~33k	61.2 ms	52.1 ms	7.1 ms	5.1 ms
~60k	71.4 ms	56.6 ms	11.5–12.0 ms	8.5–8.9 ms
~75k	—	59–64 ms	~14 ms	10.7–11.5 ms

Where the ceiling is now

At ~10k a forward is ~58 ms, and the weights are the biggest part: 29 ms at 2 positions and 36 ms at 4, against a 22 ms floor for reading them.

• The weights are arithmetic-bound past one position. Unpacking 2- and 3-bit IQ weights and the per-position dot products take longer than the memory read. The worst offenders are the most heavily compressed types: IQ1_S runs at twice its floor. Faster loads won't help; fewer instructions per weight will.
• MTP drafting costs ~5–7 ms a forward, the output head ~2 ms, attention ~3 ms at 10k.
• At long context, attention for 2 positions is 10.3 ms at 96k against a 4.2 ms cache-read floor, and 6–8 positions were still on the old layout. A fourth layout and attention phase counters are being tested now.

Next up: cheaper weight unpacking at 2–4 positions (up to ~7 ms a forward), and closing the long-context attention gap.

Every run, in one picture

The big jumps are early (MTP working, drafting on the GPU, k_mmv); the zig-zag from s16 on is drafting luck around a forward time that moved by 1–2 ms per change at this context. Every run's log and numbers are in the repo under docs/baselines.