Supported platforms
Phobos runs on NVIDIA GPUs of compute capability 7.5 or newer, Turing and every architecture after it, on Windows x64 and Linux x64. Compiling a kernel to PTX needs no GPU at all.
Tested cards
Section titled “Tested cards”Phobos is developed on an RTX 2080 SUPER, and every number in Benchmarks was measured on it. It has also been run and benchmarked on more cards, rented through vast.ai under Linux:
| Card | Architecture | Chip | Memory |
|---|---|---|---|
| RTX 2080 SUPER | Turing | sm_75 |
8 GB |
| RTX 3090 | Ampere | sm_86 |
24 GB |
| RTX 4060 | Ada | sm_89 |
8 GB |
| RTX 4090 | Ada | sm_89 |
24 GB |
| RTX 5090 | Blackwell | sm_120 |
32 GB |
The tuning is done on the 2080 SUPER; a default that matters for speed is also measured on a 3090 and a 4090 before it ships. The only 4060 for rent ran the kernel benchmark and not the model one: a desktop under WSL2 that kept the card with 1.2 GB in use and up to 4% utilization when idle, so its rates carry the desktop’s share of the card.
Measured on each card
Section titled “Measured on each card”Every card but the 4060 ran the same two benchmarks. Qwen3.5-4B Q4_K_M ran
against llama.cpp b11681 through scripts/bench.py: five rounds, the engines
interleaved, three repetitions each, the sizes of the
Benchmarks run. Then phobos-kbench ran twice,
each run autotuning its kernels afresh, against cuBLAS, compiling them for the
card’s own chip, which its JSON report records, with 64-bit indices from
sm_80 on. On a rented card it ran on a
second rental, built from the commit the card names. The 2080 SUPER’s model figures
are the ones Benchmarks publishes.
A ratio compares two engines on one card in one session. Rates on different
cards are not comparable to each other: the hosts differ, and so do the power
limits a host sets. The 4090’s model run was at 350 W, below the card’s stock
450 W; the 3090’s kbench runs were at 300 W of its 350 W. A
rented card was checked for another tenant before the run (idle power,
utilization, and how long the driver has spent capping its power), and
bench.py checks it again between rounds. On the 3090 that check reads the
card’s own power state: it holds about 150 W for two seconds after a process
exits, at 0% utilization, and is back at 21 W within ten
(trace).
Its run therefore raised the check’s power threshold from 120 W to 200 W; the
utilization check was unchanged.
Qwen3.5-4B-Q4_K_M, Phobos against llama.cpp on the same card, tokens per second
| card | tg128 | pp128 | pp512 |
|---|---|---|---|
| RTX 2080 SUPER | 1.16x Phobos 122 +/- 0.3 llama.cpp 105 +/- 0.1 | 1.39x Phobos 3,300 +/- 17.0 llama.cpp 2,381 +/- 4.0 | 1.17x Phobos 3,676 +/- 3.2 llama.cpp 3,137 +/- 1.0 |
| RTX 3090 | 1.12x Phobos 199 +/- 0.1 llama.cpp 177 +/- 0.8 | 1.31x Phobos 5,949 +/- 56.7 llama.cpp 4,538 +/- 18.1 | 1.01x Phobos 7,121 +/- 10.2 llama.cpp 7,078 +/- 47.9 |
| RTX 4060 | not measured | not measured | not measured |
| RTX 4090 | 1.05x Phobos 234 +/- 0.3 llama.cpp 223 +/- 0.1 | 1.21x Phobos 9,901 +/- 54.7 llama.cpp 8,168 +/- 109 | 0.98x Phobos 15,149 +/- 133 llama.cpp 15,453 +/- 151 |
| RTX 5090 | 1.01x Phobos 318 +/- 0.4 llama.cpp 315 +/- 0.5 | 1.50x Phobos 13,653 +/- 94.8 llama.cpp 9,117 +/- 22.8 | 0.98x Phobos 17,674 +/- 89.4 llama.cpp 18,053 +/- 49.9 |
Each cell is Phobos's rate over llama.cpp's, then both rates, the mean over rounds with its standard error.
Every test, every card
| card | test | Phobos | llama.cpp | ratio |
|---|---|---|---|---|
| RTX 2080 SUPER | pp128 | 3,300 +/- 17.0 | 2,381 +/- 4.0 | 1.39 |
| pp512 | 3,676 +/- 3.2 | 3,137 +/- 1.0 | 1.17 | |
| tg32 | 122 +/- 0.4 | 102 +/- 0.1 | 1.20 | |
| tg128 | 122 +/- 0.3 | 105 +/- 0.1 | 1.16 | |
| tg512 | 122 +/- 0.3 | 105 +/- 0.1 | 1.16 | |
| tg1024 | 121 +/- 0.3 | 105 +/- 0.1 | 1.16 | |
| tg2048 | 121 +/- 0.2 | 104 +/- 0.1 | 1.16 | |
| RTX 3090 | pp128 | 5,949 +/- 56.7 | 4,538 +/- 18.1 | 1.31 |
| pp512 | 7,121 +/- 10.2 | 7,078 +/- 47.9 | 1.01 | |
| tg32 | 199 +/- 0.2 | 172 +/- 0.7 | 1.16 | |
| tg128 | 199 +/- 0.1 | 177 +/- 0.8 | 1.12 | |
| tg512 | 199 +/- 0.1 | 177 +/- 0.7 | 1.12 | |
| tg1024 | 198 +/- 0.1 | 177 +/- 0.6 | 1.12 | |
| tg2048 | 197 +/- 0.0 | 176 +/- 0.3 | 1.12 | |
| RTX 4090 | pp128 | 9,901 +/- 54.7 | 8,168 +/- 109 | 1.21 |
| pp512 | 15,149 +/- 133 | 15,453 +/- 151 | 0.98 | |
| tg32 | 232 +/- 1.5 | 216 +/- 0.5 | 1.07 | |
| tg128 | 234 +/- 0.3 | 223 +/- 0.1 | 1.05 | |
| tg512 | 234 +/- 0.0 | 224 +/- 0.1 | 1.05 | |
| tg1024 | 234 +/- 0.0 | 224 +/- 0.0 | 1.05 | |
| tg2048 | 233 +/- 0.0 | 223 +/- 0.0 | 1.05 | |
| RTX 5090 | pp128 | 13,653 +/- 94.8 | 9,117 +/- 22.8 | 1.50 |
| pp512 | 17,674 +/- 89.4 | 18,053 +/- 49.9 | 0.98 | |
| tg32 | 318 +/- 0.7 | 315 +/- 0.6 | 1.01 | |
| tg128 | 318 +/- 0.4 | 315 +/- 0.5 | 1.01 | |
| tg512 | 319 +/- 0.4 | 327 +/- 0.4 | 0.97 | |
| tg1024 | 319 +/- 0.4 | 327 +/- 0.2 | 0.98 | |
| tg2048 | 317 +/- 0.3 | 327 +/- 0.3 | 0.97 |
phobos-kbench, Phobos's share of cuBLAS on the same card, then Phobos's own throughput
| kernel | RTX 2080 SUPER | RTX 3090 | RTX 4060 | RTX 4090 | RTX 5090 |
|---|---|---|---|---|---|
| saxpy | 100% +/- 0.0 71.6 GFLOP/s | 91% +/- 0.2 127.0 GFLOP/s | 95% +/- 0.0 41.3 GFLOP/s | 99% +/- 0.0 153.4 GFLOP/s | 94% +/- 0.1 245.3 GFLOP/s |
| gemm, f32 | 87% +/- 0.3 9.6 TFLOP/s | 87% +/- 0.3 17.2 TFLOP/s | 100% +/- 0.4 7.9 TFLOP/s | 92% +/- 0.4 40.3 TFLOP/s | 87% +/- 0.3 51.4 TFLOP/s |
| gemm, f32, split-K | 96% +/- 3.6 7.8 TFLOP/s | 84% +/- 0.3 12.6 TFLOP/s | 110% +/- 0.2 7.3 TFLOP/s | 227% +/- 1.1 36.5 TFLOP/s | 235% +/- 4.2 37.9 TFLOP/s |
| gemm, f32 data on tensor cores | 56% +/- 0.0 21.7 TFLOP/s | 70% +/- 0.0 36.9 TFLOP/s | 49% +/- 0.1 13.2 TFLOP/s | 72% +/- 0.5 87.2 TFLOP/s | 57% +/- 0.2 96.6 TFLOP/s |
| gemm, f16 data, f32 result | 83% +/- 0.0 36.7 TFLOP/s | 99% +/- 0.1 58.1 TFLOP/s | 90% +/- 0.1 28.0 TFLOP/s | 92% +/- 0.7 139.9 TFLOP/s | 92% +/- 0.1 186.2 TFLOP/s |
| gemm, f16 | 60% +/- 0.9 40.9 TFLOP/s | 63% +/- 0.2 64.0 TFLOP/s | 52% +/- 0.0 28.6 TFLOP/s | 70% +/- 1.4 175.3 TFLOP/s | 64% +/- 0.2 181.1 TFLOP/s |
| flash attention, f32 | no cuBLAS 7.8 TFLOP/s | no cuBLAS 12.8 TFLOP/s | no cuBLAS 12.3 TFLOP/s | no cuBLAS 19.5 TFLOP/s | no cuBLAS 17.4 TFLOP/s |
| flash attention, f16 | no cuBLAS 10.4 TFLOP/s | no cuBLAS 17.2 TFLOP/s | no cuBLAS 14.6 TFLOP/s | no cuBLAS 28.0 TFLOP/s | no cuBLAS 26.0 TFLOP/s |
A share is the mean over phobos-kbench runs of each run's Phobos over cuBLAS, with its standard error; one run has none. Each run times kernels alone between CUDA events and takes the median of rounds in which Phobos and cuBLAS alternate. Each gemm is held against the cuBLAS call with its own storage, multiply and accumulate types; for the tensor-core f32 gemm, which reads f32 and rounds while staging, that iscublasGemmEx on f16 copies of A and B, with the conversion in its time. Flash attention has no cuBLAS counterpart, and saxpy is memory-bound.
Where each card's numbers come from
| card | chip | memory | power limit | host | driver | Phobos | runs |
|---|---|---|---|---|---|---|---|
| RTX 2080 SUPER | sm_75 | 8 GB | 250 W | the development workstation, Windows 11 | 617.14 | bbc8898 (kbench ca08e29), 2026-10-16 | bench.py, 5 rounds x 3 reps; kbench, kbench |
| RTX 3090 | sm_86 | 24 GB | 350 W (kbench 300 W) | vast.ai, California (kbench vast.ai, Utah) | 580.178.04 (kbench 595.71.05) | ffcc8c1 (kbench 79131b9), 2026-10-16 | bench.py, 5 rounds x 3 reps; kbench, kbench |
| RTX 4060 | sm_89 | 8 GB | 115 W | vast.ai, Turkey, WSL2 on a Windows desktop | 610.88 | 79131b9, 2026-10-18 | no bench.py run; kbench, kbench |
| RTX 4090 | sm_89 | 24 GB | 350 W (kbench 450 W) | vast.ai, Spain (kbench vast.ai, Poland) | 595.99.02 (kbench 580.178.04) | ffcc8c1 (kbench 79131b9), 2026-10-16 | bench.py, 5 rounds x 3 reps; kbench, kbench |
| RTX 5090 | sm_120 | 32 GB | 575 W (kbench 560 W) | vast.ai, North Carolina (kbench vast.ai, Italy) | 615.71.09 (kbench 595.84) | ffcc8c1 (kbench 79131b9), 2026-10-16 | bench.py, 5 rounds x 3 reps; kbench, kbench |
A release also carries compiled kernels for sm_80 and sm_90, the data
center chips of the Ampere and Hopper generations. They have not been tested
or tuned on a card. Any other card of compute capability 7.5 or newer runs the
PTX for its chip, compiled on first use if the release’s cache does not have
it (see Kernel cache).
Operating systems and drivers
Section titled “Operating systems and drivers”- Windows x64, or Linux x64 with glibc 2.28 or newer.
- A driver from the CUDA 13 line, R580 or newer. Phobos emits PTX ISA 9.0, which older drivers refuse to load.
Install lists what a build from source needs on each.
What differs between chips
Section titled “What differs between chips”The language is the same on every chip, and a kernel compiles for any of them;
phobos-compile --chip picks one. The compiler chooses the instructions per
chip, and the kernel cache keeps one folder of compiled kernels per chip. A few
differences show through:
- Asynchronous copies. From
sm_80on, a pipelined loop prefetches its next stage withcp.async. Onsm_75each thread holds its share of the next stage in registers instead. See@pipeline. bf16converts natively fromsm_80on and is emulated onsm_75.- Shared memory per block under
@dynshared: 64 KB onsm_75, 163 KB onsm_80, 99 KB onsm_86,sm_89andsm_120. - Index width. Some paths, such as tensor-core loads with
ldmatrixandmma.synconsm_80and newer, need 64-bit indices, and a module that uses them is widened. See Compile and benchmark.
A performance figure holds for the card it was measured on. One measured on the 2080 SUPER says little about the speed on another card, so the published benchmarks name theirs.
Other backends
Section titled “Other backends”Phobos runs only on NVIDIA GPUs today. The compiler lowers through MLIR and LLVM to PTX, and the runtime drives the card through CUDA.
- ROCm and SPIR-V. Both are LLVM targets, so an AMD backend, or a portable one through SPIR-V, is possible in principle. Neither is planned for now: the author has no hardware to test them on, and a backend that cannot be checked on a real card is not one to ship.
- Metal. Under evaluation. Apple publishes no official documentation of its GPU instruction set, so LLVM and MLIR have no Metal target. A backend would have to emit Metal source and hand it to Apple’s compiler, which is not how Phobos generates code elsewhere and would not produce code as efficient as the NVIDIA path.
- WebGPU. Of interest for teaching rather than speed. As with Metal, a backend would emit source, WGSL, and leave the compiling to the browser, so it would not match the NVIDIA path either. What it would give is a kernel that runs in a browser: trying one with no CUDA install, on any GPU the browser can use, and a kernel visualizer that steps through a kernel as it actually runs rather than through a recorded trace.