From PatLang source to real GPU kernels: a Triton lowering pipeline, honestly scoped

For new readers

This page covers PatLang's GPU story: a real lowering pass that compiles a bounded subset of ordinary PatLang function bodies into Triton's own kernel language and runs them on an actual CUDA GPU, verified against plain-PatLang computation at every step. It's honest about what kind of achievement this is — targeting an existing, mature compiler (Triton) rather than building a GPU backend from first principles — and about the real bugs found chasing it, including one that silently reported success while computing nothing at all.

The starting point was an aside: a question about porting an unrelated Rust deep-learning framework (Ferrotorch) into PatLang. That wasn't realistic — GPU backends, a JIT fusion compiler, and 28 crates of scale aren't a weekend port — but it surfaced a much more tractable, already-half-started idea sitting in this codebase: “we were already considering the GPU side of PatLang, initially e.g. a Triton equivalent”. What follows is that idea, taken seriously.

Why Triton, not raw PTX

Triton is NVIDIA/OpenAI's own kernel-authoring language: a Python-embedded DSL that compiles down to real GPU machine code via its own mature, actively-developed compiler. The design choice made early was to target Triton's Python DSL as a compilation target — PatLang source gets translated into real @triton.jit Python text, which Triton's own compiler then turns into an actual running kernel — rather than PatLang emitting PTX or machine code directly. Direct-to-PTX emission was explicitly discussed and deliberately deferred: “we can consider a more direct route later”. This is the honest “cheat” at the centre of the whole effort: the hard, decades-of-engineering part of GPU code generation — register allocation, memory coalescing, warp scheduling — is Triton's problem, not PatLang's. What PatLang had to build was the lowering pass in between: walking real PatLang source and emitting the Python text Triton expects.

Slice 1: proving the plumbing, before any real lowering existed

The very first slice deliberately proved the architecture with a fixed, hand-written Triton kernel — the canonical elementwise vector-add — embedded as a literal Python string. PatLang wrote that string to disk, ran it via a subprocess (exec_capture, the same pattern already used elsewhere in this codebase for shelling out to rustc), and parsed the JSON result back into ordinary PatLang values. Nothing about a PatLang-authored kernel body was translated yet — this slice proved only that PatLang could drive a real GPU computation end to end and get a numerically correct answer back, verified against a plain PatLang loop computing the same sum, on an actual RTX 5070 Ti.

Slices 2–3: a real lowering pass, then multi-dimensional grids and matmul

The next slice replaced the fixed kernel text with a genuine AST-walking lowerer (self_hosting/lib/lower_triton.patlang): an ordinary PatLang function — parsed with the exact same lexer and parser every other PatLang program uses, no new grammar invented — gets walked statement by statement and translated into real Triton Python. Deliberately no new keyword: “a new keyword would ripple into lexer/parser/lower/codegen AND their Rust mirror ... AND the language-spec BDD gate ... far more than this slice needs”. A kernel is just a function whose body stays inside a checked, bounded subset:

make a function called add_kernel takes x, y, out
  store(out, offs, x[offs] + y[offs])
end

offs and mask are implicit block-scope variables — every kernel body is wrapped in the same program-id/offset/mask preamble real Triton kernels use. The covered subset grew to arithmetic, comparisons, counted loops, block reductions (sum, max, argmax...), atomics, and a small math-function allowlist — anything outside it is a clear rejection at lowering time, not a silent miscompile.

The next slice added real multi-dimensional grids and matmul via block pointers — tl.make_block_ptr, tl.advance, tl.dot, a tiled accumulator loop over the K dimension — first as a fixed, hand-generated canonical shape (the same “prove the architecture with a fixed kernel first” discipline Slice 1 used), then, several slices later, genuinely AST-lowered from real PatLang source using a small vocabulary of block-pointer primitives (make_block_ptr, load_block, advance_block, dot, zeros_2d, store_block). A 2×3 @ 3×2 matmul was checked three ways at once — the fixed-generated kernel, the AST-lowered kernel, and plain PatLang — all agreeing, then stress-tested against a deliberately non-square, non-block-aligned 37×29 @ 29×41 case to genuinely exercise the boundary-padding path. Both passed on the first real run.

One source file, two execution paths

A plain comment pragma lets a single function serve as both ordinary, directly-callable CPU code and a GPU kernel, from one definition:

# @gpu_kernel
make a function called add_kernel takes x, y, out
  store(out, offs, x[offs] + y[offs])
end

The marked function is fully ordinary PatLang — the compiler has no special handling for it at all — and the marker only matters to a small text-level scanner that extracts marked function bodies and builds a name->source registry, matching this codebase's own line-based include-expansion style rather than inventing a new AST-level concept. This let a real demo call the same function both ways and check the results directly against each other, instead of hand-duplicating an expected-value computation the way every earlier demo had to.

Where it went wrong first: a Mandelbrot demo that lied

Asked for a demo that would compute a Mandelbrot set both on CPU and GPU and time the two paths, the first version reported 100% exact pixel match. That result was wrong, and not caught until a later session segment. Two real bugs, stacked:

A Triton compile-time rejection, silently swallowed. The kernel seeded its per-pixel accumulator as a bare scalar (let zr = 0.0), then reassigned it every loop iteration via a per-lane vector operation. Triton requires a loop-carried variable's type to stay consistent across iterations — scalar-then-vector is a genuine compile error (“Loop-carried variable zr has initial type fp32 but is re-assigned to <[BLOCK], fp32> in loop!”). It passed the lowerer's own subset-checker and produced syntactically valid Python; Triton's own compiler rejected it; the subprocess-capture plumbing surfaced that only as an empty result, never a visible error. The kernel was computing nothing at all.
A CPU reference bug, once the first one was fixed and the kernel actually ran. The plain-PatLang comparison function returned 0 for every single pixel — it referenced a top-level let from inside a named function, and named functions in PatLang don't close over top-level lets at all, only anonymous closures do. A previously-documented gotcha from Slice 1, re-hit by not applying the lesson to new code.

Both bugs were invisible to the test suite's own pass/fail summary until individual GPU-versus-CPU values were printed side by side for real coordinates. Fixed, the demo now genuinely agrees pixel-for-pixel, and the incident directly motivated closing the underlying safety gap: the lowerer now does a conservative shape-inference pass, rejecting any if/elif whose condition depends on a per-lane value before emitting any Python for it, rather than letting Triton discover the problem later. select(cond, a, b) (→ tl.where) is the correct primitive for per-lane branching; the checker now enforces that a bare if can't be used for it by mistake.

Measuring what actually costs time

Once the demo was genuinely correct, the honest timing question followed: a single one-shot kernel call takes roughly 2–3 seconds end to end — almost entirely subprocess spawn and torch/triton import, not kernel compute. Running the same kernel 100 times inside one spawned process (paying that import cost once, not per call) brought the real per-call cost down to roughly 0.01–0.1 milliseconds — a difference of three to four orders of magnitude between “the number a naive benchmark reports” and “what the GPU is actually doing.”

That same instinct led to a genuinely persistent GPU session (self_hosting/lib/triton_session.patlang): one Python/torch/triton worker process stays alive across many kernel launches, with tensors resident on the GPU between them, modeled directly on this codebase's own existing signals infrastructure (spawn a process, connect over a local TCP port, one-line-JSON request → one-line-JSON reply) rather than its durable, SQLite-backed message queue — a synchronous “launch this kernel, get this result” call wants a live round trip, not a polled, persisted row. A two-kernel chain (add, then double) now runs with the intermediate tensor never leaving the GPU between launches. Five one-shot calls: roughly ten seconds. Five launches inside one session, after a one-time setup cost: a few milliseconds.

A detour into the runtime itself: a genuinely quadratic list

Scaling the Mandelbrot demo up to see where GPU compute actually starts winning surfaced something with nothing to do with Triton at all: building a plain PatLang list via the ordinary let xs = list_push(xs, v) loop idiom was genuinely O(n²) — 8,192 elements in about half a second, 32,768 elements in over seventeen. The root cause, traced precisely rather than guessed at: reading a local variable for a function-call argument always cloned it, leaving the variable's own reference alive until the later reassignment overwrote it — which only happens after the call returns. By the time list_push ran, at least two live references to the same list always existed, so the underlying copy-on-write mechanism was forced to deep-clone the entire backing array on every single push, no matter how uniquely the caller actually held its own copy.

The fix touched the bytecode interpreter shared by every compiled PatLang program and every interpreted run: a conservative, whole-function analysis marking a variable read as safe to move rather than clone when it's provably the last read before an immediate reassignment to the same variable, plus a second fix so list_push itself stopped creating an extra reference of its own regardless. Verified directly: the same 65,536-element loop that used to hang past any reasonable timeout now completes in 12 milliseconds; a million elements, 187. Filed and fixed as GitHub issue #61.

Rounding it out: autotuning, and matmul from real source

Three further pieces closed out the arc: enforcing the vector-condition safety gap the Mandelbrot bug had exposed; @triton.autotune support, letting a kernel expose several candidate block sizes and have Triton itself benchmark and pick the fastest at first launch (rejected outright for reduction kernels, where varying the block size would silently change which elements a given launch actually reduces over); and generalizing matmul from the one fixed, hand-generated shape into something a kernel author can actually write in PatLang source, using the new block-pointer primitive vocabulary. One genuine gap surfaced and fixed along the way: the counted-loop lowering had always assumed a step of one, which cannot express matmul's own K-dimension loop (while k0 < K do ... let k0 = k0 + BLOCK_K end) — fixed with a small step-detection pass that stays fully backward-compatible with every existing step-one kernel.

A smaller, unrelated gap turned up while building the persistent session: PatLang's tcp_connect is fatal on failure, with no non-fatal equivalent to the existing tcp_try_listen — meaning a straightforward “poll until a listener is ready” retry loop wasn't actually expressible. Filed and fixed as GitHub issue #62, mirroring tcp_try_listen's own contract exactly on the other side of the same problem.

What this is, and isn't

Every demo in this arc is checked against an independently-computed plain-PatLang answer, not just “did it run without erroring” — a discipline that directly caught the Mandelbrot silent-failure bug once, and is applied everywhere else specifically because of that. But the honest framing matters as much as the result: this is a lowering pass onto Triton's own compiler and PyTorch's own runtime, not a GPU backend built from first principles. No PTX emission, no register allocation, no memory model implemented from scratch. The interesting, hard part of turning source code into fast GPU machine code was solved by someone else's compiler, reused rather than reinvented — a legitimate, common way real production DSLs reach a GPU, and one worth naming plainly rather than overselling.