· 5 min read

A compiler made of resumable lanes

Gpufuck turns functional-language semantics into bounded GPU state machines, batches independent programs across lanes, and reads resolved Core back to emit ordinary WebAssembly on the host.

An IBM 082 punch-card sorter stands beneath a row of output bins holding colored cards.
waelder, CC BY-SA 3.0

Gpufuck's semantic shader opens with @workgroup_size(1), which on a GPU is a small heresy. One invocation owns one compilation lane, loads one durable state record, and advances it until a bounded quantum expires — no workgroup cooperating on a program in sight.

That odd choice makes the repository understandable. Gpufuck flattens functional programs into fixed-width tables, runs name resolution and type inference as resumable machines, batches independent programs across lanes, then gives resolved Core to TypeScript for WebAssembly emission. The arrangement gives regular semantic work to the GPU while the host keeps the irregular edges.

The GPU sees tables#

A frontend keeps its parser, module discovery, ownership rules, and diagnostic wording, then lowers an accepted program into the Functional Surface, where the surface builder interns names and writes fixed-width u32 records. Whatever your language argues about — Rust moves, OCaml's sequential scope — it argues about before this point. The GPU never learns what your keywords were.

The generated Haskell tree trace shows the handoff. mapTree enters the packed ABI without an annotation, while GPU resolution and inference replace source names with numeric evidence and leave the entry concrete.

Surface evidenceResolved CoreWhat changed
$argument0Local depth=0A lexical name becomes a position relative to its binder.
mapTreeGlobal definition=d1A recursive global becomes an index into the linked definition table.
BranchConstructor c1:BranchA nominal constructor keeps its identity through a numeric constructor index.
gpuMain: <inferred>entry=d3; type=i32Inference leaves the public entry with a concrete first-order boundary.
Four rows from tree.trace.md — each one is a name's last appearance before it becomes an index.

A dispatch is a time slice#

WGSL offers no recursion and no comfortable heap of pointer-linked compiler objects. Gpufuck stores the active phase, cursors, source spans, fuel, and unfinished work in buffers. The relevant loop in compiler_shader.ts is almost aggressively plain:

compiler_shader.ts — compile_lane, abridgedwgsl
fn compile_lane() {  if state.status != STATUS_PENDING { return; }  initialize_compilation();  let dispatch_start_steps = state.total_steps;  loop {    if state.status != STATUS_PENDING { return; }    if state.total_steps >= state.maximum_steps {      state.status = STATUS_STEP_LIMIT;      return;    }    if state.total_steps - dispatch_start_steps >=      state.maximum_steps_per_dispatch { return; }    state.total_steps += 1u;    advance_compilation();  }}
  1. 1A completed or failed lane costs no more semantic work.
  2. 2The dispatch records its own starting fuel before entering the machine.
  3. 3The total budget bounds the complete compilation across resumed dispatches.
  4. 4The quantum returns control to the host while the program remains pending.
  5. 5One transition advances one explicit validation, resolution, or inference phase.
One lane advances a durable compiler state until completion, failure, total fuel exhaustion, or the dispatch quantum.

Type inference uses the same discipline for union-find traversal, occurs checks, generalization, indexed refinements, and coverage. Input-sized work becomes durable frames that survive dispatches. There is an honesty to this loop: a charged step performs bounded semantic work instead of hiding a million-record scan behind the word “step.”

Cancellation remains cooperative because WebGPU cannot interrupt a command after submission. Smaller quanta shorten the interval before the host can observe an abort; a quantum of one makes every transition a submission and readback boundary. Whether that trade stings in practice I can't say — the latency benchmarks exist, but the repository avoids pretending one adapter's numbers apply to another.

The host keeps the machine alive#

Inference keeps logical arenas for types, environments, frames, refinements, scratch data, and output. When one fills, the shader stops and names the arena that ran out — more courtesy than most out-of-memory paths ever manage. The host runner enlarges that region, copies live records, patches the bindings, and resumes the same machine state.

  1. Dispatch

    A lane advances by at most the current semantic quantum.

  2. Observe

    The host reads status, transition counts, output size, and any exhausted arena.

  3. Grow

    Only the exhausted logical arena is enlarged, with live records copied to new bases.

  4. Rebind

    The replacement buffers resume with the same phase, results, and semantic fuel.

After "Rebind": a pending lane enters the scheduler again — back to "Dispatch".

The host/device loop used while inference remains pending or one logical workspace arena must grow.

The workspace tests begin with deliberately tiny capacities and force every arena to exhaust. They verify that only the named capacity grows and that the resumed transition count stays unchanged. Memory repair happens between semantic steps, so growth does not spend compiler fuel or restart inference.

Width arrives between programs#

One lane advances one branch-heavy program. Width comes from putting independent programs beside it. compileBatch() packs modules into lanes, while the dispatch scheduler lets sibling promises meet and coalesces their command encoders into one submission.

One test keeps 512 copies of the repository's Brainfuck compiler fixture in a single GPU pack — 512 compilers running at once is the whole pitch condensed into one fixture. Another test makes one lane exhaust its output arena while a sibling finishes normally; the exceptional lane takes the growth path and the completed result remains valid. Tiny one-off programs still pay submission and readback overhead, which is why the benchmarks include batches instead of offering one grand speed number.

The GPU stops at Core#

Semantic success leaves a GpuFunctionalModule owning resolved buffers. Wasm emission reads those Core nodes back and runs capture analysis, storage planning, reachability, and specialization on the host. The artifact boundary calls the TypeScript emitter directly, with no WAT round trip and no GPU dependency in the resulting module.

The current sharp edge follows from that split. Compilation has no CPU semantic fallback; the TypeScript inferencer is a differential oracle for tests, and the development guide explicitly rejects turning it into an implicit production path. A machine without WebGPU can run the emitted artifact, but it cannot build one through this package.

The Haskell trace finishes with the names gone: mapTree is d1, Branch is c1, and gpuMain has settled at i32. Nothing in the emitted module remembers WebGPU either. The artifact runs on any Wasm host; only the build machine ever needed a GPU.