GPU, and whether lex-gpu should be its own language

Status: measured, and the measurement decides it.

The question arrived as three: can cancho run on a GPU, how would a native one work, and should it be a separate language. The first two are design; the third looked like taste until the numbers came in.

§2 is the result and it has a surprise in it. A bounds check is free even in vectorised code — an overflow trap is not. So the memory half of this language’s no-undefined-behaviour claim survives on a GPU at no cost, and the arithmetic half cannot survive at all.

1. Today: nothing, and the backend is why

Cranelift has no GPU target — no PTX, no SPIR-V, no AMDGPU. Emitting for a GPU means LLVM, which ROADMAP.md currently places outside the next slices and purity.md §4.2 argued down rather than up.

Checked (#68): backend-limits.md §1.3. src/isa/ in cranelift-codegen 0.121.2 is aarch64, pulley_shared, riscv64, s390x and x64 — and pulley is Cranelift’s own portable interpreter rather than a device.

So this document is not about a feature that is nearly there. It is about what the design would have to give up, priced.

2. What each guarantee costs a vectoriser

The same reduction — sum a million-element buffer, 200 rounds — written four ways in C and twice in cancho, with the guards switched independently. benches/reduce.c and benches/reduce_{checked,wrapping}.cho, which the conformance suite holds to the same answer.

SIMD instructions counted in the emitted run, not inferred from the clock. clang 18, -O2; minimum of nine runs:

Three things fall out, and none of them was obvious beforehand.

2.1 A bounds check is free

1.01×, and the SIMD count does not move. The index condition is provably true inside the loop the compiler already proved bounded, so it is deleted outright. A bounds check never was a reassociation barrier — it is a branch on a constant, and a branch on a constant is not a branch.

overflow-cost.md §3.2 established that a check costs the vectoriser. It measured the overflow one and generalised. Half of that generalisation is wrong, and this is the half: memory safety is not what costs anything here.

Corrected (#64): that last sentence is the over-generalisation again, one document later. check-cost.md measured all eight loop-body checks, and memory safety is not uniformly free: s[lo..hi] costs 2.16×, the most expensive integer check in the language. It is the same kind of check as s[i] and it traps the same way.

What makes s[i] free is that i < n is the loop’s own condition, so the compiler discharges it. Move the same two comparisons onto bounds loaded from memory and the cost is 2.16×; move them back onto the induction variable and it is 0.99×. Provable is free, and loaded is not — and this row happened to measure the one check in the language whose condition the loop always proves.

2.2 An overflow trap costs 1.46×, and it costs it as SIMD

Ten SIMD instructions to zero. This reproduces §3.2’s finding on a new kernel and a new compiler version: an addition that may trap is an addition whose order is observable, so a reduction cannot be split across lanes.

2.3 Deleting the trap does not reach C

This is the one that decides the third question. cancho without traps is 2.27× off vectorised C, and still emits no SIMD at all — it is 1.55× slower than even scalar clang. The trap accounts for 1.83× of cancho’s own cost (252.1 → 137.6); everything remaining is Cranelift.

So removing the trap is necessary and not sufficient. A GPU-shaped cancho on the current backend would be scalar, which on a GPU is the same as not having one.

And it will stay scalar — backend-limits.md §1.4. Cranelift has SIMD: vector types, rewrite rules, per-target lowerings. What it has no pass for is producing them from scalar code, because its SIMD is there for WebAssembly’s v128 rather than for a vectoriser. So the 2.27× attributed to the backend here is a property of its design and not a version it has not reached.

Corrected (parallelism.md §3.1): the table above is Cranelift’s, and Cranelift was the default when it was written. With --backend llvm, the default now, the same two programs on the same machine: wrapping 64-67 ms against C’s 66-69 (SIMD present: 54 instructions), checked 93-95 ms (1.4x C, scalar). So “deleting the trap does not reach C” is a statement about Cranelift: on LLVM it does, and the trap costs 1.4x on this memory-bound kernel and 2.6x on a cache-resident one. The finding that survives is the one about the trap: a checked loop is still not reordered, on either backend.

3. How a native GPU cancho would work

Worth writing down, because three of the pieces already exist and cost nothing to reuse.

A kernel is a function whose row is empty. A kernel may not open a file, syscall or allocate; here that is not a convention but the absence of Io, Fs, Heap and Ffi from a signature, and purity.md measured 35% of functions already proving it. No kernel keyword is needed: launch names the roots the way main already names emission roots, and the check is on the row.

Regions are address spaces. region a { … } is a lexically scoped arena whose references cannot escape — 11 must-reject fixtures enforce exactly that — which is the lifetime CUDA __shared__ has and does not check. Same rule, same checker, nothing new.

The device is a capability. Gpu(device) beside Fs(prefix), refinable by the existing narrow, with gpu_launch(d) and gpu_copy(d) as labels. cancho authority then names the device, and a program never handed one provably never launched.

4. What it would cost the design

4.1 The trap model is the thesis, and the hardware does not have one

defined-behaviour.md §2.1: an operation with no right answer stops. A GPU cannot stop — no signals, no per-lane abort. Three answers:

1. Wrapping arithmetic in kernels — abandons the claim exactly where the compute is.

2. Poison, reduced at the launch boundary — each lane sets a flag, launch answers some lane trapped. The claim survives; where does not.

   Measured (#65): poison.md, and the awkward part is which half it saves. Element-wise arithmetic under poison is free on a wide vector ISA — a + b inside a loop costs 0.98× once the overflow test is spelled as sign logic rather than as an opaque builtin. A reduction is not saved and cannot be: four lanes compute four different partial sums, so “no partial sum overflowed” is not the same property after vectorising, and poison there is worse than trapping. A reduction is what a GPU is for.

3. Refuse trapping operations in kernels — a narrower language inside the braces.

Only (2) keeps the thesis, and §2.2 prices what (1) and (3) are buying: 1.46× on a mature backend, which is most of the reason to own a GPU.

But §2.1 halves the bill. Bounds checking is free, so a kernel dialect keeps array safety in full and gives up only overflow trapping. That is a much smaller concession than “no undefined behaviour does not fit a GPU” — it is arithmetic undefined behaviour, and it is recoverable by (2) at a cost this document has not measured because nothing here can run a kernel yet.

5. So: should lex-gpu be its own language?

Not yet, and the reason is §2.3 rather than taste.

The case for independence is real: a language whose failure model is trap and one whose failure model is poison are two languages, not one with a flag. That is the same argument that justified cancho existing beside lex-lang — different layer, different job.

But the measurement says the split buys nothing where it matters. Deleting the trap leaves cancho scalar; the thing that actually buys GPU speed is LLVM, and LLVM is equally required whether the kernels are a cancho dialect or a separate language. And the two genuinely novel parts — checked barriers and poison-as-an-effect — are new work in either design, because cancho does not have them to inherit.

What is worth saying about the shape: cancho is already six crates, so a sibling could depend on cancho-syntax and cancho-types and own only the dialect, the lane checker and the backend. Roughly 80% of the front end shared and none of the back end. That is a much better arrangement than a fork, and ROADMAP.md’s risk register names the alternative: “second-system trap — during a port every change lands twice.” A third repository makes it three.

And the plainest argument: cancho’s own README says not a usable language yet. A third sibling before the second is usable is a recognisable way for a project to end.

5.1 What would change this

6. Open