sixi-asi 0.9: 45 tok/s from a 27B hybrid model on Jetson Thor, and the kernel switch point that was hiding it

Radek (Sixi) · 10 November 2026 · source, patches, benchmark scripts and raw results

Two days ago sixi-asi 0.8 decoded Qwen3.8-27B at 23.2 tokens per second on an NVIDIA Jetson AGX Thor. Version 0.9 does 45.3 on the same weights, 3.8 times plain autoregressive decoding, and answers a follow-up question about a 63,000-token document in one second instead of three and a half minutes. Most of the gain came from finding out that an explanation I published in the 0.8 write-up was wrong.

What I got wrong in 0.8

Speculative decoding drafts k tokens and verifies them in one batched pass of the target model. In 0.8 I measured that every extra row in that pass cost about 17 ms on Thor, attributed it to the 48 Gated DeltaNet layers scanning the batch sequentially, and tuned the server to short drafts (k = 3).

A profile of one forward pass says otherwise. Quantized matrix products are 83 to 92% of GPU time. The recurrence is about 3 ms. The per-row cost came from llama.cpp's kernel selection: up to 8 rows it uses a matrix-vector kernel that decodes the 4-bit weights again for every row, and only from 9 rows on does it switch to a matrix-matrix kernel. Upstream keeps a tuned switch point for the RTX 4090, RTX 5090, DGX Spark and Jetson Orin. Thor (compute capability 11.0) had none and inherited the default.

The patch: 106 lines, two launch-shape changes

• Switch point. On Thor, k-quants go to the matrix kernel from 4 rows. An 8-row pass drops from 163 to 101 ms.
• Occupancy. Nsight Compute showed the matrix kernel running one 256-thread block per SM for these skinny batches: 2 active warps per scheduler, 62% of cycles with nothing to issue, 29% of DRAM bandwidth. It was waiting on memory latency, not on bandwidth. Running 128 threads with 4 blocks per SM takes the 8-row pass to 88 ms and the 16-row pass from 116 to 94 ms.

Neither change touches the weights or the arithmetic; test-backend-ops passes every MUL_MAT and MUL_MAT_ID case. One candidate from the sweep was faster than physics allows: 57 ms, below the roughly 58 ms it takes just to read 15.8 GB of weights at 273 GB/s. It failed 69 correctness cases and was dropped. A benchmark that beats the roofline is a bug report.

Drafters

With verification nearly flat from 4 to 16 rows, long drafts are free, so block drafters win. 0.9 drives three of them behind one interface: DSpine (16-token block), DSpark (Markov and confidence heads) and DFlash2. Two orchestration bugs were worth more than any tuning:

• The context n-gram drafter replaced the neural draft on any 4-token match. A trace against llama-server showed token-identical DSpine drafts in both engines; sixi-asi simply was not asking DSpine in those steps, and accepted 6 tokens where DSpine got 11.
• For 7-token drafters the adaptive policy backed off to 2 drafts on poor acceptance, which lands on the one batch size that is slower than the full block.

Numbers

NVIDIA Jetson AGX Thor developer kit (122.9 GiB unified LPDDR5X, about 273 GB/s, CUDA 13.2, MAXN). Qwen3.8-27B, greedy, batch 1, identical prompts for every engine, 256 generated tokens, thinking on, median of 3 runs per task. Every sixi-asi number was re-measured on the published build.

NVIDIA's TensorRT Edge-LLM is the reference to beat on this box: a TensorRT engine with native NVFP4 weights on the Blackwell tensor cores. sixi-asi matches it on plain Q4_K_M GGUF weights, with no 15-minute engine build per configuration, and passes it on the workload that looks like real use:

All three miss the same three problems. I first ran this gate with a 1024-token cap and briefly believed NVFP4 was worse at code (15/20 against 18/20). It was not: the cap was cutting the model off mid-reasoning on three tasks, differently per run. With an 8192-token budget the scores are identical. Edge-LLM's server also defaults to thinking off, which makes it look 20% faster until both engines are asked to do the same work.

Long documents

A hybrid model's recurrent state cannot be rolled back further than its snapshot ring, and sixi-asi took its only snapshot at the end of the prompt, after the question. So a second question about the same document paid the whole prefill again. 0.9 leaves two device-side checkpoints near the end of every long prefill, 512 and 64 tokens before the end, each about a millisecond to take.

Decode holds 19 to 29 tok/s with 63k tokens in context and 30 to 48 tok/s at 4k. The KV cache costs about 70 KB per token because only 17 layers are attention layers, so on a 122 GiB machine context length is a time budget, not a memory budget.

What the numbers do not say

• 45.3 against 45.4 is a tie, on different weight formats. Batch 8 and multi-stream decoding are not measured yet, and that is where a TensorRT engine should pull ahead.
• Plain autoregressive decoding is unchanged at 11.9 tok/s, about 70% of the bandwidth roofline. A 16-row verify is 94 ms against a floor near 68 ms. Both gaps are the next kernel work.
• Block-scaled FP4 tensor-core instructions exist only on sm_120; on Thor they are reachable through tcgen05 alone. NVFP4 weights were not better than a good 4-bit k-quant for this model (KL divergence against BF16: UD-Q4_K_XL 0.009, NVFP4 0.024, Q4_K_M 0.025).

Every claim ships with its script (scripts/ab_bench.py, scripts/intel_eval.py, scripts/longctx_bench.py), the kernel patch (patches/llama.cpp/0003) and the raw JSON under configs/benchmarks/2026-11-10/. If you can refute a number on your own Thor, please open an issue.

Source, Apache-2.0: github.com/radeksixi/sixi-asi. Previous write-up: sixi-asi 0.8.