Phobos: tile-based GPU kernel language & LLM inference engine

Tiles and lanes

Whole tiles, as in Triton, and warps and lanes, as in Gluon, in one language and in one kernel. Every picture is sized from the compiler's own traces.

Tiles, like Triton

for kt in range(0, K, TILE_K) {  var a = A[pm * TILE_M :+ TILE_M, kt :+ TILE_K]  var b = B[kt :+ TILE_K, pn * TILE_N :+ TILE_N]  acc += dot(a, b)}

TILE_M 32, TILE_N 32, TILE_K 4, 256 threads, schematic

Lanes, like Gluon

for k in range(0, K, ctx.lanes * 4) {  let at = k + ctx.lane * 4  acc += sum(W[row, at :+ 4] * x[0, at :+ 4])}let total = sum<over ctx.lane>(acc)

32 lanes, 4 elements each per step, schematic

Mixed, like Phobos

var xs = x[0 :+ 1, :]if row < N {  for k in range(0, K, ctx.lanes * 4) {    let at = k + ctx.lane * 4    acc += sum(W[row, at :+ 4] * xs[0, at :+ 4])  }}let total = sum<over ctx.lane>(acc)

4 warps share one tile of x, schematic

Step through it in the visualizer ctx in the docs Tiles and lanes in one kernel

One compiler, source to PTX

Seven stations between a kernel file and the card, for gemm_fp32.ph, each drawn from what the compiler holds there. The same points travel the whole way.

.ph sourceThe kernel as written: an accumulator, a k loop over two tiles, and the store.

ASTParsed and checked, 39 nodes: the for kept whole, every other statement folded to one.

graph IROps and the values between them, before any memory is placed: 31 nodes, 36 edges, one value carried round the loop.

passesStaging, pipelining, shared-memory placement and barrier elision, shown for the build without @aligned.

MLIRThe gpu, scf, vector, memref and arith dialects, 335 ops in the kernel's regions, lowered through NVVM to LLVM IR.

PTX447 instructions from LLVM's NVPTX backend for sm_75, by family.

GPULoaded and launched through CUDA: 4 CTAs of 256 threads, the grid the host picks for M = N = K = 64.

.ph source: The kernel as written: an accumulator, a k loop over two tiles, and the store.
AST: Parsed and checked, 39 nodes: the for kept whole, every other statement folded to one.
graph IR: Ops and the values between them, before any memory is placed: 31 nodes, 36 edges, one value carried round the loop.
passes: Staging, pipelining, shared-memory placement and barrier elision, shown for the build without @aligned.
MLIR: The gpu, scf, vector, memref and arith dialects, 335 ops in the kernel's regions, lowered through NVVM to LLVM IR.
PTX: 447 instructions from LLVM's NVPTX backend for sm_75, by family.
GPU: Loaded and launched through CUDA: 4 CTAs of 256 threads, the grid the host picks for M = N = K = 64.

Measured against cuBLAS

Each bundled kernel is autotuned and timed beside cuBLAS on the same card, in the same run. The f32 gemm runs at76% of cublasSgemm_v2.

Kernel throughput against cuBLAS, higher is better

RTX 2080 SUPER, one phobos-kbench run, each kernel autotuned, results/results.csv committed 2026-10-08

Phobos
cuBLAS (cublasSaxpy_v2, cublasSgemm_v2, cublasHgemm)

saxpyn = 2^25, memory-bound, in GB/s

gemm, f32M = N = K = 4096, f32, peak 11.4 TFLOP/s

gemm, f32 on tensor coresM = N = K = 4096, f16 tensor cores, f32 accumulation, peak 45.7 TFLOP/s

gemm, f16M = N = K = 4096, f16 tensor cores, peak 91.4 TFLOP/s

flash attention, f32Nq = Nk = 4096, D = 64, f32, peak 11.4 TFLOP/s

flash attention, f16Nq = Nk = 4096, D = 64, f16 tensor cores, f32 accumulation, peak 45.7 TFLOP/s

A track is the card's peak for the kernel's precision as phobos-kbench derives it from SM count and clock, so a bar's length is its share of that peak; saxpy's track is the faster of its two rates. The flash attention kernels carry @tensorcore and run their products on tensor cores where the tile shapes qualify; phobos-kbench rates the f32 one against the f32 peak. Measured once per kernel, so there are no error bars.

What the compiler does on its own

Shared memory by liveness
6 tiles, 10 KiB in 8 KiB
A tile lives from its declaration to its last use, and tiles whose lifetimes do not overlap share bytes. Shown for gemm_fp32 without @aligned.
Step through it
Barriers where threads meet
21 kept, 8 elided
A barrier follows a tile store only when other threads next touch those bytes. Of flash attention's 29 candidate sites, 8 need none.
Tiles and shared memory
Double-buffered loops
2 to 4 staging tiles
The next k-slice loads while this one computes, on by default. On sm_80 the copies become cp.async: 32 in the gemm_fp16 trace, 0 on sm_75.
@pipeline
Tensor cores from one attribute
2.58x the f32 gemm
The same gemm body with @tensorcore: 21.4 TFLOP/s against 8.3 TFLOP/s, operands rounded to f16. mma.sync and ldmatrix with swizzled staging, or WMMA.
@tensorcore
A search, not a guess
64 gemm configurations
@autotune names the tile sizes. phobos-kbench compiles and checks every configuration, then races the four fastest over 30 interleaved rounds. for<unroll n> takes a tuned n too.
Autotuning
One instruction seam
6 chips, sm_75 to sm_120
Only the target layer names an instruction. Kernels are compiled and cached for every chip ahead of time, with no GPU in the machine.
The kernel cache

Tile, barrier and copy counts are from the committed compiler traces (commit 5a313cd, small problem sizes); the speed-up is from the benchmark above.

The same language runs LLMs

Quantized GGUF models run end to end on kernels written in Phobos, served through an OpenAI-compatible API. Mixture-of-experts models that do not fit in VRAM keep an expert cache on the card and compute misses on the CPU.

One token at a time

The layers turn the last token into logits, the sampler picks the next one, and it goes back in. Decode fuses a chain of layer stages into one persistent kernel.

Each sampled token streams out and goes back into the layers as the next input.

An elastic expert cache

When another program takes 2 GiB of the card, the cache gives the memory back and the CPU threads take more of the misses, so nothing is paged: Phobos decodes at 30.8 tokens/s while llama.cpp drops to 9.0. Once the memory is free, the cache grows back.

The process exits and the cache grows back to 39 experts a block. Phobos decodes at 44.6 tokens/s, llama.cpp at 31.2. The replay took 52.4 s on Phobos and 131.7 s on llama.cpp.

the card
CPU, misses
another process

Qwen3.6-35B-A3B on an RTX 2080 SUPER, results committed 2026-10-06, median wall time of the replay, decode rates the mean over 2 and 3 replays. Cache sizes are the ones the run logged; the flow of misses is drawn, not measured.

Tokens per second against llama.cpp

NVIDIA GeForce RTX 2080 SUPER, driver 617.14, Phobos 3eec883, results committed 2026-10-08

Phobos
llama.cpp, or the fork its row names

Qwen3.5-0.8B-Q8_0

MiniCPM5-1B-Q8_0

Qwen3.5-4B-Q4_K_M

GLM-4.6V-Flash-Q4_K_M

Ternary-Bonsai-2-27B-PTQ1_0against PrismML's llama.cpp fork

Qwen3.8-27B-UD-IQ1_M

Qwen3.6-35B-A3B-UD-Q4_K_M

Decode over 128 tokens and a 128-token prompt. Each pair of bars has its own scale. The ratio is Phobos's rate over llama.cpp's.

Models, formats and benchmarks

Get started

Download a release
Prebuilt phobos-cli, phobos-compile, phobos-bench and phobos-cache for Windows and Linux, x64.
Latest release
Compile a kernel
From a checkout: compile a kernel source to PTX, for sm_75 by default. No GPU needed.
```
cargo run -r -p phobos-compile -- examples/gemm_fp32.ph
```

Serve a model

Download MiniCPM5-1B-Q8_0, then start the OpenAI-compatible server with the dashboard on the terminal. Needs an NVIDIA GPU.

cargo run --features cuda -r -p phobos-cli -- --listen 127.0.0.1:8080 --gguf ~\models\minicpm5-1b-Q8_0.gguf --temp 0.9 --top-p 0.95 -n 32768

Read the getting started guide

Tile-based GPU kernel language & LLM inference engine