Phobos

Tile-based GPU kernel language & LLM inference engine

Write tiles like Triton. Set up warps and lanes like Gluon. One language.

Get startedInferenceGitHub

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<unroll U> 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 visualizerctx in the docsTiles 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.

The compiler's stations for gemm_fp32.phThe compiler's stations for the gemm kernel: the source lines, the AST, the graph IR as a graph of ops and values with the k loop as a region and its carried accumulator, the shared-memory passes, the MLIR region tree, the PTX instruction mix and the launch grid.FIG / GEMM_FP32 / SOURCEexamples/gemm_fp32.phthe k loop13var acc: tile<f32>[TILE_M, TILE_N] = 0.015for<unroll U> kt in range(0, K, TILE_K) {16var a = A[pm * TILE_M :+ TILE_M, kt :+ TILE_K]17var b = B[kt :+ TILE_K, pn * TILE_N :+ TILE_N]18acc += dot(a, b)19}22C[pm * TILE_M :+ TILE_M, pn * TILE_N :+ TILE_N] = alpha * acc + beta * c_oldFIG / GEMM_FP32 / ASTkernel gemmattributes+4params+5let pm+2let pn+2var acc+1for ktUunroll0startKendTILE_Kstepvar aindexA:+*startpmTILE_MTILE_Mlen:+ktstartTILE_Klenvar bindexB:+ktstartTILE_Klen:+*startpnTILE_NTILE_Nlen+=accdot( )ablet c_old+12store C+1940 nodes. Statements outside the for are folded; +n is what each one hidesFIG / GEMM_FP32 / GRAPH IRtensor, globalaccumulator, registersindexdivisibility trapfor kt in 0..K step 416 tripscarried acc, back to the next tripABCalphabetaassume_align AM = dim 0 AK = dim 1 Aassume_align BN = dim 1 Bassume_align CM % 32 != 0N % 32 != 0K % 4 != 0pm = program_id 0pn = program_id 1gemm_initfor ktkt%30 carriedpm * 32slice A[32, 4]pn * 32slice B[4, 32]gemm_dotyieldfor resultpm * 32pn * 32slice C[32, 32]gemm_storeFIG / GEMM_FP32 WITHOUT @ALIGNED / SHAREDwithout @aligned: 6 buffers, 8192 B peak, 9 barriersas written: 4 buffers, 2048 B, 1 barrierprogram orderfor kt, 16 trips: body and remainder5 barriers kept4 elidedshared memory, by byte offset04096460851208192 Bacc 32x32 f32, 4096 Babc_old 32x32 f32, 4096 Bc_old takes the bytes a and b helda buffer drops in at its first op, fades after its last; a point is about 23 bytesFIG / GEMM_FP32 / MLIR REGIONSgpu.func @gemm 34621scf.for loop 1812916scf.if branch 89scf.for contract 29scf.if branch 16scf.if 32arith 198vector 74memref 37gpu 15scf 19llvm 3one point per opFIG / GEMM_FP32 / PTX MIX.target sm_75.maxntid 256424 instructionsone point per instructionfma128int116mul.f32mov30bra26ld.shared16add.f16st.shared15ld.param11setp11ld.global10st.global4bar.sync3trap3cvt2ret1static counts over the PTX text, not executed instructionsFIG / GEMM_FP32 / LAUNCHCTA (0, 0)CTA (1, 0)CTA (0, 1)CTA (1, 1)grid 2 x 2 x 1256 threads a CTA8 warps of 32 lanes2048 B shared a CTAM = N = K = 64sm_75the host's choice for these dimsThe compiler's stations for gemm_fp32.phThe compiler's stations for the gemm kernel: the source lines, the AST, the graph IR as a graph of ops and values with the k loop as a region and its carried accumulator, the shared-memory passes, the MLIR region tree, the PTX instruction mix and the launch grid.SOURCEexamples/gemm_fp32.phthe k loopvar acc: tile<f32>[TILE_M, TILE_N] = 0.0for<unroll U> 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)}C[pm * TILE_M :+ TILE_M, pn * TILE_N :+ TILE_N] =alpha * acc + beta * c_oldASTkernel gemmattributes+4params+5let pm+2let pn+2var acc+1for ktUunroll0startKendTILE_Kstepvar aindexA:+*startpmTILE_MTILE_Mlen:+ktstartTILE_Klenvar bindexB:+ktstartTILE_Klen:+*startpnTILE_NTILE_Nlen+=accdot( )ablet c_old+12store C+19GRAPH IRfor kt16 tripsABCalphabetafor ktslice AdotaccstoreSHAREDwithout @aligned6 buffers, 8192 B, 9 barriersas written: 4 buffers, 2048 B, 1 barrierprogram orderfor kt, 16 trips5 barriers kept4 elidedshared memory040968192 Baccabc_oldc_old reuses a, bMLIR REGIONSgpu.func @gemm 346scf.for copy 21scf.for loop 1812916scf.if branch 89scf.for 29scf.if branch 16scf.if branch 32arith 198vector 74memref 37gpu 15scf 19llvm 3PTX MIX.target sm_75 .maxntid 256424 instructions, a point eachfma 128int 116mul.f 32mov 30bra 26ld.shared 16add.f 16st.shared 15ld.param 11setp 11ld.global 10st.global 4bar.sync 3trap 3cvt 2ret 1static counts, not executedLAUNCH(0, 0)(1, 0)(0, 1)(1, 1)grid 2 x 2 x 1256 threads a CTA8 warps of 32 lanes2048 B shared a CTAM = N = K = 64sm_75

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

ASTParsed and checked, 40 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, 346 ops in the kernel's regions, lowered through NVVM to LLVM IR.

PTX424 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.

  1. .ph source: The kernel as written: an accumulator, a k loop over two tiles, and the store.
  2. AST: Parsed and checked, 40 nodes: the for kept whole, every other statement folded to one.
  3. graph IR: Ops and the values between them, before any memory is placed: 31 nodes, 36 edges, one value carried round the loop.
  4. passes: Staging, pipelining, shared-memory placement and barrier elision, shown for the build without @aligned.
  5. MLIR: The gpu, scf, vector, memref and arith dialects, 346 ops in the kernel's regions, lowered through NVVM to LLVM IR.
  6. PTX: 424 instructions from LLVM's NVPTX backend for sm_75, by family.
  7. 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 at88% 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-17

  • Phobos
  • cuBLAS (cublasSaxpy_v2, cublasSgemm_v2, cublasGemmEx on f16 copies of A and B, cublasGemmEx)
saxpyn = 2^25, memory-bound, in GB/s
gemm, f32M = N = K = 4096, f32, peak 11.4 TFLOP/s
gemm, f32 data on tensor coresM = N = K = 4096, f32 A and B, rounded to f16 in the kernel, f16 tensor cores, f32 accumulation, peak 45.7 TFLOP/s
gemm, f16 data, f32 resultM = N = K = 4096, f16 A and B, f16 tensor cores, f32 accumulation, peak 45.7 TFLOP/s
gemm, f16M = N = K = 4096, f16 A, B and C, 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

    20 kept, 9 elided

    A barrier follows a tile store only when other threads next touch those bytes. Of flash attention's 29 candidate sites, 9 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.29x the f32 gemm

    The same gemm body with @tensorcore: 22.0 TFLOP/s against 9.6 TFLOP/s, operands rounded to f16. mma.sync and ldmatrix with swizzled staging, or WMMA.

    @tensorcore

  • A search, not a guess

    0 gemm configurations

    @autotune names the tile sizes, the unroll and the CTA size. phobos-kbench compiles and checks every configuration, ranks them on warm clocks, then races the fastest over 30 interleaved rounds.

    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 b60d28c, 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.

The decode token loopSchematic. A GGUF file supplies the weights to the layers. The tokenizer turns the prompt into tokens for the layers, the layers produce logits for the sampler, and each sampled token streams out and goes back into the layers as the next input.FIG / DECODE / TOKEN LOOPschematicGGUF filetokenizerlayersfused decode kernelsampleroutTiles move through shared memory.

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 33.2 tokens/s while llama.cpp drops to 5.6. Once the memory is free, the cache grows back.

An elastic expert cacheMeasured on Qwen3.6-35B-A3B, replaying a coding agent's session. The card's 8 GiB hold the rest of the model, its context and an expert cache of 41 experts a block. Another process takes 2 GiB; the cache shrinks in steps to 8 experts a block and the CPU threads compute more of the misses. When the process exits the cache grows back to 39. Decode before, during and after: Phobos 44.2, 33.2, 43.1 tokens/s, llama.cpp 29.3, 5.6, 19.3.FIG / MOE / ELASTIC CACHEschematiccard memory, 8 GiBmodel, contextcache 41, 27, 13, 8, 39 a blockanother process, 2 GiBCPU threadscompute misseswall time, one replayPhobos 81.4 sllama.cpp 372.7 sdecode, tokens/sPhobosllama.cppbefore2 GiB takenafter44.233.243.129.35.619.3

The process exits and the cache grows back to 39 experts a block. Phobos decodes at 43.1 tokens/s, llama.cpp at 19.3. The replay took 81.4 s on Phobos and 372.7 s on llama.cpp.

  • the card
  • CPU, misses
  • another process
Qwen3.6-35B-A3B on an RTX 2080 SUPER, results committed 2026-10-16, median wall time of the replay, decode rates the mean over 3 and 5 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 bbc8898, measured 2026-10-16

  • 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
MiniCPM-V-4_7-35B-A3B-Q4_K_M
MiniCPM-V-4_7-35B-A3B-Q8_0
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

  1. Download a release

    Prebuilt phobos-cli, phobos-compile, phobos-bench and phobos-cache for Windows and Linux, x64.

    curl -fsSLO https://github.com/joa/phobos-lang/releases/download/v0.2.0/phobos-v0.2.0-linux-x64.tar.gz && tar -xzf phobos-v0.2.0-linux-x64.tar.gz && rm phobos-v0.2.0-linux-x64.tar.gz
    Latest release

  2. Compile a kernel

    Compile a kernel source, such as gemm_fp32.ph from the repository's examples, to PTX for sm_75 by default. No GPU needed.

    phobos-compile gemm_fp32.ph

  3. Serve a model

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

    phobos-cli --listen 127.0.0.1:8080 --gguf minicpm5-1b-Q8_0.gguf --temp 0.9 --top-p 0.95 -n 32768
Read the getting started guide