cancho-gpu
Alpha · runs a 27B model on a Metal GPU

GPU kernels that say what they can do.

Write a tile kernel once. The compiler checks it, refuses what is wrong and names the rule, and writes CUDA C and Metal. The same kernels run Qwen3.8‑27B on a Mac's GPU.

Try itSee the evidence
A real session, recorded: the commands run for real, only the typing is simulated. Kernel compiled and a wrong one refused, then Qwen3.8-27B (NVFP4, weights from Ollama's store) generating 64 tokens, greedy, no chat template, on an Apple M4 Max.
Transcript
# cancho-gpu: typed tile kernels, one source for CUDA C and Metal

$ cat swiglu.chx
// a = silu(g) * u over a row of n: the gated activation of a SwiGLU MLP.
algo swiglu(n)
  in  g: f32[1, n]
  in  u: f32[1, n]
  out y: f32[1, n]
{
  let gt = load g
  let ut = load u
  let sg = sigmoid &gt
  let a  = gt * sg
  store a * ut -> y
}

schedule swiglu for apple-m-series { threads 256 }
schedule swiglu for nvidia-ada     { threads 256 }

$ cancho-gpu emit swiglu.chx out n=17408
swiglu_17408.cu: 256 threads, grid 1x1, 0 B of tiles
swiglu_17408.metal: 256 threads, grid 1x1, 0 B of tiles

# what it writes for Apple GPUs (the CUDA C is written next to it)
$ sed -n '/kernel void/,/^}/p' out/swiglu_17408.metal | head -22
kernel void swiglu_17408(
    device const float* p0_g [[buffer(0)]],
    device const float* p1_u [[buffer(1)]],
    device float* p2_y [[buffer(2)]],
    uint tid [[thread_index_in_threadgroup]],
    uint3 tgpos [[threadgroup_position_in_grid]])
{
    const uint gid = tgpos.x;
    float v2[68];
    for (uint k = 0; k < 68u; ++k) {
        const uint e = k * 256u + tid;
        if (e < 17408u) {
            v2[k] = float(1.0f / (1.0f + precise::exp(-float(float(p0_g[(uint)(0) * 17408u + (uint)(0) * 1u + ((e))])))));
        }
    }
    float v3[68];
    for (uint k = 0; k < 68u; ++k) {
        const uint e = k * 256u + tid;
        if (e < 17408u) {
            v3[k] = float(float(float(p0_g[(uint)(0) * 17408u + (uint)(0) * 1u + ((e))])) * float(v2[k]));
        }
    }

# a wrong program is refused, with the rule it broke
$ cat leak.chx
algo k(n)
  in  x: f32[1, n]
  in  w: f32[1, n]
  out y: f32[1, n]
{
  let a = load x
  let b = load w
  store sigmoid a -> y
}
schedule k for nvidia-ada { threads 32 }

$ cancho-gpu check leak.chx n=8
error[leak]: `w.1` is never consumed; drop it, yield it or move it into an op

# and the kernels run a 27B model on this Mac's GPU
$ cancho-gpu-run35 ~/.ollama/models kernels "Q: Why is the sky blue?\nA:" 64 --stream --timing

The sky appears blue due to a phenomenon called **Rayleigh scattering**.

When sunlight enters the Earth's atmosphere, it collides with gas molecules (primarily nitrogen and oxygen). Sunlight is composed of all colors of the visible spectrum, which have different wavelengths. Shorter wavelengths (like blue and violet) scatter
timing: load_ms=4580 prefill_tokens=11 prefill_ms=311 decode_tokens=63 decode_ms=2808

Refuses wrong programs

Every tile is consumed exactly once; bounds, shapes and the threadgroup budget are checked, and a refusal names the rule it broke.

One source, two targets

The same .chx kernel becomes CUDA C and Metal, byte-for-byte what the independent Rust implementation writes: 182 emitting cases, 26 refused by both.

Held to an oracle

Goldens come from the Rust, never from this compiler. A CPU interpreter says what a kernel computes; on a Mac the GPU agrees with it to about 1e-6.

Runs a real model

Qwen3.8-27B decodes at 25 tokens a second on a Metal GPU (33 with speculative decoding), with the same tokens and top five as a reference's log-probabilities.

Honest about what is missing

CUDA has run only against a mock driver, speed on the 27B is about half of Ollama's, and attention stops at 1364 positions. Each is in the status section, with its issue.

An authority you can read

The report lists exactly what the compiler may do, and the gate fails if it widens. The one program that calls foreign code says so.

How it works.

A kernel is checked before it is lowered; a refusal never reaches a GPU. The interpreter is the reference for what a kernel computes.

.chx kerneltyped tiles Checkbounds · shapes · budget Lowerper target CUDA C Metal Refusedthe rule it broke, by name Interpreter (CPU)the GPU is compared with it Device pathcompiles and runs it Weights and textGGUF · safetensors · BPE · sampler Model runnerevery layer is a compiled kernel

Look at one kernel.

The real source of swiglu.chx, what the compiler writes for each target, and what it says to a program that breaks a rule.

// a = silu(g) * u over a row of n: the gated activation of a SwiGLU MLP.
algo swiglu(n)
  in  g: f32[1, n]
  in  u: f32[1, n]
  out y: f32[1, n]
{
  let gt = load g
  let ut = load u
  let sg = sigmoid &gt
  let a  = gt * sg
  store a * ut -> y
}

schedule swiglu for apple-m-series { threads 256 }
schedule swiglu for nvidia-ada     { threads 256 }
Fifteen lines. Every tile is a value that must be consumed once; the schedule says how many threads a threadgroup has.

The Metal and CUDA text is the compiler's actual output for n=17408 (cancho-gpu emit tests/chx/swiglu.chx out n=17408), the kernel function only; the refusal is tests/reject/leak.chx.

A real model on the same kernels.

Qwen3.8-27B (NVFP4, 48 gated-delta layers) reads Ollama's store and runs on a Metal GPU. It is not the fastest way to run this model, and the page says so.

Decode speed, tokens a second, greedy, 128 tokens of a prose continuation, three runs each, one machine (Apple M4 Max, 64 GB), one invocation, with other jobs running on the machine (load average 14 at the start; the runs themselves varied by under 2%) (docs/benchmark.md). Ollama decodes this model with speculative decoding from the model's own draft head; cancho-gpu does too in the speculative row (a chain of two drafts, three rows checked a pass, the same tokens as plain decoding: #26), and is not there yet. MLX-LM uses a different 4-bit quantisation. Prompt processing takes eight tokens at a time: a 49-token prompt in 0.65 s, against 2.2 s a token at a time (one run each).

Correct, not just fast

Three prompts, eight steps each: the same token and top five as the reference's golden log-probabilities, logit gaps within 1e-4.

Small first

Qwen2.5-0.5B gives the same greedy tokens as Hugging Face transformers in float32, every top-5 logit within 1e-4, at about 6 ms a token.

Measured, with the misses

An experiment that looked like a win and was a bug, and one that did not move the pipeline at all, are in #24 and #26.

A kernel compiler should be small enough to audit.

QuestionWhat cancho-gpu does
Can the compiler call into C or touch the network?No. Its report lists args, err_write, fs_read, fs_write, heap and io_write, and says it never touches the network or foreign code. tests/authority.sh fails the gate if that changes.
How do I know the output is right?The Rust emitter is the oracle: goldens are recorded from it, never from this compiler, and a live differential runs the same cases against a build of it. A CPU interpreter checks what a kernel computes, and a device program compares the GPU against it.
Does it run on a GPU?Metal: yes, on Apple Silicon, with the result inside the 2e-2 tolerance of the interpreter. CUDA: only against a mock driver so far, because no machine here has an NVIDIA GPU. The page and the README say which is which.
Is it an inference runtime?A first one. cancho-gpu-run and cancho-gpu-run35 run Qwen2.5-0.5B and Qwen3.8-27B: every layer is a compiled kernel on a Metal GPU. The 0.5B model gives the same greedy tokens as Hugging Face transformers; the 27B (NVFP4 weights from Ollama's store, 48 gated-delta layers) matches a reference implementation's golden log-probabilities to 1e-4 and decodes at about 25 tokens a second (33 with speculative decoding), against about 43 for Ollama in the same benchmark run. It also serves an OpenAI-compatible chat endpoint with streaming (one request at a time).

Every claim has a measurement, or says it has none.

The oracle is external, so "it works" is a comparison, not an opinion. Where a claim turned out false, the design document was corrected in place.

208cases byte-identical with the reference implementation: 182 emit, 26 refused by both
24interpreter runs with the same summary as the Rust's, f32 order included
~5,000hostile files and texts across the GGUF, safetensors and tokenizer readers: none traps or hangs; every refusal names its rule
0differences from the oracles on 5 real GGUF models, 3,379 Qwen3.8 tensors, and 3,811 records through 4 real tokenizers; the sampler's picks match on 60,000+ rows

The design document holds how it is built and what it cost; the device document what is and is not proved on a GPU; the runtime document where each reader differs from its oracle, and the kernels document the context limit.

Where it fits.

Use it for

  • Compiling .chx tile kernels to CUDA C and Metal with a small, auditable toolchain
  • Checking a kernel on the CPU before a GPU is involved
  • Trying a cancho program that drives a GPU through a thin C shim

Not yet

  • CUDA on real hardware, Hopper, AMD
  • Speed on the 27B model, long contexts (it needs a change to the language), concurrent requests
  • Attention over more than 1364 positions (it needs a change to the language)
  • Speed: reach and agreement are measured, timing is not

Try it.

You need the cancho compiler at the commit pinned in cancho.toml, and Rust to build it.

git clone https://github.com/alpibrusl/cancho
git clone https://github.com/alpibrusl/cancho-gpu && cd cancho-gpu
REV=$(sed -n 's/^cancho = "\([0-9a-f]*\)".*/\1/p' cancho.toml)
(cd ../cancho && git checkout "$REV" && cargo build --release -p cancho)
export CANCHO=$PWD/../cancho/target/release/cancho
$CANCHO build
mkdir -p out
build/cancho-gpu emit tests/chx/softmax.chx out r=16 n=256

Other commands: cancho-gpu check <file.chx> [name=value ...] [--run] and cancho-gpu gguf <file.gguf> [--sum]. scripts/gate.sh is everything CI runs.

Status: alpha