v0.1.2

Highlights

• RTX 5090 (sm_120) support. v0.1.1 crashed there. On a card with that many multiprocessors the decode attention kernel's merge tile passed the per-block shared memory ceiling, and the driver refused the kernel. It now falls back when its splits outgrow the card. Decode on sm_89 and sm_120 was also up to 10x slow, from an L2 prefetch at a 32-bit-truncated address; prefetch now goes only with 64-bit indices. Throughput on the 5090 is not tuned yet: on Qwen3.5-4B Q4_K_M it measures 0.88x llama.cpp at pp128, 0.73x at pp512 and 0.82x at tg128.
• Prompt processing at or above llama.cpp on the RTX 4090 and RTX 4060 for Qwen3.5-4B Q4_K_M. Token generation stays within 1 to 3 percent of llama.cpp.
• GLM-4.6V-Flash (glm4) runs, with a native Q5_0 path.

Qwen3.5-4B Q4_K_M	pp128	pp512	tg128
RTX 4090	9783 vs 8077 (1.21x)	15367 vs 14536 (1.06x)	197.9 vs 203.5 (0.97x)
RTX 4060	3268 vs 2872 (1.14x)	3480 vs 3343 (1.04x)	74.6 vs 75.4 (0.99x)
RTX 5090	4977 vs 5642 (0.88x)	7552 vs 10415 (0.73x)	174.4 vs 211.4 (0.82x)

Tokens per second, phobos-bench against llama-bench (llama.cpp b11497, CUDA), interleaved in one session per card with no other process on it: medians of six rounds for pp, where v0.1.2 led every round on both cards, and of four for tg. Rented cards vary by host; across four 4090 hosts pp512 ranged from 1.03x to 1.07x.

What got prefill there:

• The gated delta rule's prompt scan holds its state in registers (delta_scan_t), with lanes split into row and column groups so few-SM cards are not bound by shuffles.
• Attention on the tensor cores for prompts, after fixing the mma.sync B fragment that dot_t addressed wrongly on sm_80+.
• qgemm: the epilogue's minimum term folded and cvt conversion on sm_80+, split-K on a starved grid, paired 8-byte output stores, and a split reduce that adds into the residual.
• Prompt passes no longer allocate and free their quantized activations every pass, and issue their launches as they record them.
• The gated norm in registers (rms_norm_gated_q_t), the delta convolution computed once per key head, the alpha and beta gates as one stacked projection, and no f32 SwiGLU output when the down projection reads only its quantized copy.

What's Changed

• Glm4 by @joa in #13
• gguf: Decode attention falls back when its splits outgrow the card by @joa in #14
• lang: Prefetch only with 64-bit indices by @joa in #15
• vastai: Smoke test, benchmark and profiling for rented cards by @joa in #16
• gguf: Split-K on the Q8_0 prompt projection by default by @joa in #17
• gguf: Split q8_qmma's starved grids harder on cards past 48 SMs by @joa in #18
• lang: Reach the sm_80 tensor-core path, and stage i8 at 16 bytes by @joa in #19
• vastai: qg.c checks a kernel against a reference by @joa in #20
• vastai: qg.c shapes and split-K mode, renting gotchas by @joa in #23
• lang: Fold qgemm's minimum term, and convert with cvt from sm_80 by @joa in #21
• gguf: Split-K for the staged projections on a starved grid by @joa in #22
• lang: Address dot_t's B fragment along k on sm_80 and later by @joa in #24
• gguf: Prompt attention on the tensor cores by @joa in #29
• gguf: The chunked delta scan owns 32 state columns a program by @joa in #26
• gguf: Quantize the FFN's SwiGLU on the way for a prompt by @joa in #27
• gguf: Project a prompt's q/k/v straight into the conv history by @joa in #28
• gguf: Move the rotary angle table into layers/rope.rs by @joa in #30
• lang, gguf: The gated delta rule with its state in registers by @joa in #31
• gguf: Cut a prompt pass's host time by @joa in #32
• lang, gguf: delta_scan_t with lanes split into row and column groups by @joa in #33
• lang: qgemm_t stores a lane's two adjacent outputs as one 8-byte store by @joa in #34
• lang, gguf: rms_norm_gated_q_t, the gated norm in registers by @joa in #35
• gguf: Fewer launches around a prompt's small projections by @joa in #36
• gguf: delta_conv computes each query and key head once by @joa in #37
• release: 0.1.2 by @joa in #38

Full Changelog: v0.1.1...v0.1.2