v0.3.0

Highlights

Phobos in Python. import phobos.lang as ph spells a kernel in Python: a @ph.kernel is a kernel, a @ph.fn a function, and the front end reads the module's source without running it. Every location the compiler reports, errors, the race check, traces and MLIR locations, is a Python line. A launch binds its arguments to the contract phobos-compile now writes beside the PTX, through ctypes on the driver, and takes any __cuda_array_interface__ array. @ph.kernel(key=...) autotunes over the kernel's own @autotune space, checking every variant against the caller's reference before timing it, and keeps the winner in the kernel cache. Triton's tutorials 01, 02, 03, 05 and 06 are ported. ph2py prints any .ph in the Python spelling, and each archive ships it with python/phobos. See Python.

A standard library. Inlined .ph functions every kernel can call: activations, norms, row and column reductions, attention, the losses of torch.nn.functional, counter-based random numbers and dropout, batch norm, quantization and matmul helpers, with extent(x, d) for a tensor's size. lang_check runs them on the device against a host reference. --no-std leaves it out. See Standard library.

Loops pipeline by default. A GEMM's operands, a dynamic extent split into whole tiles and a loop carrying a register accumulator double-buffer without an attribute wherever the stages fit; @pipeline asserts it and @pipeline(off) vetoes it. The register-staged tensor-core GEMMs prefetch two tiles ahead, the f32 GEMM stages through registers on every chip, a GEMM's k-loop takes for<unroll n>, @launch takes @autotune constants, and @index_bits lets a tuner search the index width. SGEMM reaches 88% of cuBLAS on the RTX 2080 SUPER at M=N=K=4096.

The compiler places a persistent kernel's grid barriers. In a @persistent kernel, var of a tensor is scratch the launch supplies; the compiler works out which blocks touch which memory, puts a grid barrier only where one block reads what another wrote, and holds scratch every block writes alike in shared memory. The fuse pass now writes its decode chains from a library of stages and leaves barriers and scratch placement to the compiler.

A sturdier server. --ctx sets the context limit, which otherwise defaults to what fits the card, and a prompt past it is a 400. A full queue answers 429, a disconnected client's request is dropped, a failed generation is an error rather than a stop, and --api-key requires a bearer token. A malformed GGUF file is an error, not a crash, and NaN or infinite logits fail the draw. The --listen dashboard shows the card's power draw and temperature.

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.

No head-to-head benchmark was run for these notes. The published figures are in Benchmarks, measured on the RTX 2080 SUPER against llama.cpp b11681 on the tree before this tag.

What's Changed

• gguf: The kernel handle cache keys on the module's driver handle by @joa in #137
• gguf: IQ3_XXS and IQ3_S take their .ph projection on sm_86 and sm_89 too by @joa in #131
• lang: Passes over the built IR run in one slot between build and verify by @joa in #128
• lang: The fragment accumulator is a pass over the built IR, not an AST match by @joa in #129
• lang: A dot's f16 product stays in mma.sync fragments when every use can read it there by @joa in #130
• lang: grid_barrier() takes the barrier word a persistent launch appends by @joa in #132
• lang: Kernel scratch, block dependence, and the fused chains from a stage library by @joa in #133
• lang: A persistent kernel's grid barriers and scratch placement are the compiler's by @joa in #134
• gguf: The fuse pass writes its chains and the compiler places their barriers by @joa in #135
• gguf: The fuse pass walks its units in grid-stride loops by @joa in #136
• ci: Run the tests that need no GPU by @joa in #138
• ci: Move GitHub actions to their Node 24 majors by @joa in #139
• kbench: Kernels compile for the card's chip, at 64-bit indices from sm_80 by @joa in #168
• server: Check an optional API key as a bearer token by @joa in #169
• lang: Loops pipeline by default; @pipeline asserts and @pipeline(off) vetoes by @joa in #172
• server: refuse a prompt longer than the context, add --ctx by @joa in #171
• server: full queue answers 429, disconnected clients are dropped by @joa in #170
• server: a failed generation is an error, not a 200 stop by @joa in #173
• lang: deeply nested source is a diagnostic, not a stack overflow by @joa in #174
• kbench: A @dynshared kernel launches with its dynamic shared memory by @joa in #175
• lang: bound the size of an inlined call by @joa in #176
• bench: cold expert cache and distinct tokens per repetition by @joa in #177
• gguf: a malformed file is an error, not a crash by @joa in #178
• inference: NaN or infinite logits fail the request instead of sampling by @joa in #179
• gguf: qwen35 refuses rope scaling and unsupported rotary sections by @joa in #180
• kernels: refuse tensors past a 32-bit index; cluster leaves compile for the node's chip by @joa in #181
• website: the front page counts 512 gemm configurations, not 0 by @joa in #182
• bench: -b and -ub as llama-bench has them, and --n-depth by @joa in #185
• Python Frontend by @joa in #186
• release: 0.3.0 by @joa in #187

Full Changelog: v0.2.0...v0.3.0