v0.2.0

Highlights

The language says what one thread does. v0.1.2 wrote whole-tile programs and reached for a built-in whenever a kernel needed per-lane work. v0.2.0 writes that in the language, and the built-ins are gone:

• ctx names where a kernel runs: ctx.block, ctx.thread, ctx.warp and ctx.lane. The compiler infers how far each value varies between threads, uniform, warp-uniform or per lane, with no annotation, and rejects what the rules forbid. ctx.block.x replaces program_id(0).
• Collectives and shuffles over lanes and the CTA, barrier(), elem and store for one thread's run of a shared tile, sum(e) over a thread's own slices with dp4a for int8, word, byteperm, popcount, byte masks and f16bits. rowsum and rowmax of a register tile fold with shuffles.
• A register-resident tile is a distribution type with two layouts, Blocked and Gemm. The int8 scaled dot_t accumulates in tensor-core registers, the fragment accumulator takes acc += r * c and acc -= r * c, and generate(f) is a tile computed where it is read, so a weight dequantized per fragment never touches shared memory.
• Functions, inlined at every call. for<unroll n>, @registers, @pipeline over a dynamic extent, @dynshared up to the chip's block limit, @swizzle. Bitwise, logical and compound operators, where, iota, pow, rowargmax, atomic_cas, A[:e], break, continue and return, u8, u16 and u32 storage, and views that slice and write through.

Every decode matvec, the quantizing norms, the gated delta scan, decode attention's warp partials, MoE's top-k and the raw formats' prompt projections on sm_75 are now .ph sources in phobos-gguf. Dropped with them: <fmt>_qdot_t, <fmt>_qdot_i8_t, <fmt>_qdecode_t, delta_scan_t, warp_partial, rms_norm_q_t and rms_norm_gated_q_t.

A shared-memory race check, on the host. phobos-compile FILE.ph --check-races runs every thread of one CTA through the MLIR as emitted, with no GPU, so it sees the barriers the membar pass kept and elided. A race between warps, a barrier not every thread reaches, or an access outside its buffer is an error on the source, as the compiler's own are. It found several barrier placements in the tree's own kernels that this release fixes. See RACECHECK.md for what it models and what it does not see yet.

Errors point at the source. Every error after the parser carries a source span and renders in rustc's layout: labels on the lines involved, a note on the declaration, and a suggested var or barrier() where that is the fix. Ops carry their source locations into MLIR, so race reports and traces name the line too.

phobos-lang.dev. The language guide, every ctx field and attribute with a figure, the spec, ENV.md and the benchmarks rendered in place, the inference guides, the kernel ABI, a supported-platforms page, and a kernel visualizer that steps through recorded traces of real kernels. The compiler checks the site: every whole-kernel fence compiles, every expected error fails, and the diagnostics shown are what phobos-compile prints.

Inference.

• Q8_0 experts stream on the device, so a mixture-of-experts file in Q8_0 takes the expert cache as the K-quant ones do. An expert set in an unsupported format is refused by name. The opt-in lookahead (PHOBOS_MOE_LOOKAHEAD) now spends a block's one copy on its heaviest predicted miss, ahead of the token, and an expert still copying in when its block runs is a host miss, never a wait; it stays off because the host hides those misses on the development card.
• Decode on cards with many multiprocessors: a narrow K-quant matvec over a long k splits into partials and a sum, a split projection adds into the residual in its reduce, the fused MLP splits its down projection on a wide grid, and independent fused nests start on the blocks the last left idle. A decode row's attention operands prepare in two launches, and the output gate writes the quantized copy the output projection reads. grid_barrier orders memory, not only progress.
• Prompt: the int8 projections accumulate on tensor-core registers, qgemm_t contracts with m16n8k32 at one CTA per SM on Blackwell, qgemm copies its activation tile with cp.async on sm_80 and later, and the IQ2_XXS, IQ2_XS, IQ2_S, IQ3_XXS, IQ3_S and Q5_0 prompt projections run as .ph kernels on sm_75.
• phobos-compile --set NAME=VALUE pins an @autotune parameter and --all-variants writes every variant. A batch compile lowers at most one kernel per core at a time.
• The --listen dashboard wears the website's night sky, and its throughput panel shows the mean rate beside the peak.

The kernel cache shipped with each archive was warmed under this release's compiler for sm_75, sm_80, sm_86, sm_89, sm_90 and sm_120, so the v0.1.2 cache's m8n8k16 Blackwell kernels are replaced by the m16n8k32 ones.

No head-to-head benchmark was run for these notes. The published figures are in Benchmarks, measured on the RTX 2080 SUPER, and the supported platforms page gives the rented cards' rough standings against llama.cpp, each measured in October 2026 on the tree before this tag.

What's Changed

• lang: qgemm_t contracts with m16n8k32 on Blackwell, one CTA an SM by @joa in #39
• gguf: Decode that fills a many-SM card by @joa in #40
• gguf: A fused decode convolution carries its history in place by @joa in #41
• gguf: Fewer launches around decode attention by @joa in #42
• lang: Bitwise, logical and compound operators, where, iota and the missing math by @joa in #43
• lang: An untyped elementwise var fuses into one sweep by @joa in #45
• lang: The remaining language gaps by @joa in #47
• lang: Code generation gaps by @joa in #48
• lang: @registers keeps a tile in registers across a loop by @joa in #49
• lang: A dot stages an operand it reads more than once by @joa in #50
• lang: A dot applies the update of its product as it writes it out by @joa in #51
• lang: Barrier elision no longer doubles its pending accesses at every loop by @joa in #52
• Kernels on the new language, qgemm cp.async, two compiler fixes by @joa in #53
• kernels: A batch compile lowers at most one kernel a core at a time by @joa in #54
• lang: x += y broadcasts a row or a column of x by @joa in #55
• lang: ctx names where a kernel runs; ctx.block.x replaces program_id(0) by @joa in #56
• lang: ctx.thread, ctx.warp and ctx.lane, with the rules for per-thread values by @joa in #57
• lang: Collectives and shuffles over lanes and the CTA by @joa in #58
• lang: sum(e) over one thread's own slices, with dp4a for i8 by @joa in #59
• lang: rowsum and rowmax of a register tile fold with shuffles by @joa in #60
• Lanes: moe_topk and the Q4_K decode matvec in .ph; Q8_0's .ph matvec by default by @joa in #61
• Lanes: word and dp4a; Q5_K and Q6_K decode matvecs in .ph by @joa in #62
• Lanes: every raw format's decode matvec in .ph; popcount, byteperm, shared-tile runs, var slots in registers by @joa in #63
• SPEC A[:j] row; a typed var of a partial slice allocates once by @joa in #64
• lang: barrier(), elem and store; the quantizing norms written with lanes by @joa in #65
• lang: Float multiply-add contracts to fma; the delta scan written with lanes by @joa in #66
• gguf: Decode attention's warp partials written with lanes by @joa in #67
• gguf: The IQ formats' f32 matvec and qdecode written with lanes by @joa in #68
• gguf: IQ lanes matvecs stage their grid wide, and run by default by @joa in #69
• lang: Drop _qdot_t; the row-major formats' matvec is lanes only by @joa in #70
• lang: grid_barrier orders memory, not just progress by @joa in #71
• gguf: backend_check judges quantized-activation projections by spread by @joa in #72
• gguf: IQ lanes matvecs by default below sm_80 only; K-quants and Q8_0 always lanes by @joa in #73
• lang: for; IQ lanes matvecs two blocks a turn at four CTAs an SM by @joa in #74
• gguf: IQ formats' and PTQ1_0's decode matvecs on lanes by default on every card by @joa in #75
• gguf: fused and MoE K-quant decodes on lanes by @joa in #77
• website: phobos-lang.dev by @joa in #76
• website: Mixed, like Phobos by @joa in #78
• Drop the _qdot_i8_t built-in; decode matvecs on lanes only by @joa in #79
• cli: The dashboard takes the website's night sky by @joa in #81
• lang: int8 scaled dot_t accumulates in tensor-core registers by @joa in #80
• gguf: q8_qmma as a .ph loop instead of qmma_t by @joa in #82
• lang: A barrier ahead of an op when no trailing barrier orders its conflict by @joa in #83
• lang: staged int8 operands for the register accumulator by @joa in #84
• lang: pipelined register accumulators, word stores, padded slices, fences by @joa in #85
• lang: ldmatrix at 32-bit indices; a fenced turn skips its store barrier by @joa in #86
• lang: Same-thread sweeps need the same bytes behind each element by @joa in #87
• lang: shared int8 tiles the scaled dot reads are swizzled by @joa in #90
• lang: A shared-memory race check on the host by @joa in #88
• website: the @Swizzle page by @joa in #91
• lang: A CTA atomic's slot read waits before its bytes are reused by @joa in #89
• gguf: iq2xxs_qgemm as .ph on sm_75 by @joa in #92
• lang: a stage read past a barrier is stored at it by @joa in #93
• compile: --set pins an @autotune parameter, --all-variants writes every variant by @joa in #96
• lang: Diagnostics with source spans, rustc's layout and highlighting by @joa in #99
• lang: Expressions, attributes, params and names carry source spans by @joa in #100
• lang: the int8 accumulator takes acc += r * c and acc -= r * c by @joa in #102
• gguf: iq3xxs_qgemm as .ph on sm_75 by @joa in #94
• gguf: qgemm_time covers every format with a prompt projection by @joa in #103
• gguf: IQ2_XS and IQ2_S projections as .ph on sm_75 by @joa in #104
• gguf: the K-quant block sources are re-exported only with cuda by @joa in #105
• lang: Every error after the parser points at the source by @joa in #101
• lang: the int8 accumulator's column pairs load and store as one 8-byte access by @joa in #106
• gguf: IQ3_S projection as .ph on sm_75 by @joa in #108
• gguf: qgemm_time times Q8_0 through qmma_t by @joa in #109
• lang: the int8 scaled dot issues m16n8k32 and whole m16n8k16 on Blackwell by @joa in #111
• gguf: a split .ph projection runs one CTA an SM on Blackwell by @joa in #112
• lang: Errors name source, not compiler ops, and label what they involve by @joa in #110
• lang: a register-resident tile is a DistType, its fragment accumulator one layout of it by @joa in #114
• lang: Regs and Gemm are DistType layouts, Blocked and Gemm by @joa in #115
• lang: Errors point at declarations and suggest var and barrier() by @joa in #113
• lang: generate(f), a tile computed where it is read by @joa in #116
• gguf: q50_qmma as .ph on sm_75, its weight generated per fragment word by @joa in #117
• lang: a let in the loop keeps the fragment accumulator by @joa in #118
• gguf: batch_check counts a near-tie the batch breaks the other way as agreement by @joa in #119
• lang: every operand of the int8 accumulator's updates may be generated by @joa in #120
• cluster: a generated tile is rejected in a cluster kernel by @joa in #121
• lang: Ops carry source locations, and race reports are diagnostics by @joa in #122
• website: A page on shared-memory races, and what an error looks like by @joa in #123
• website: generate in the language guide by @joa in #124
• website: A supported platforms page with the tested cards by @joa in #125
• gguf: Q8_0 experts stream on the device by @joa in #126
• gguf: The lookahead spends a block's one copy on its predicted miss by @joa in #127

Full Changelog: v0.1.2...v0.2.0