From canvas to WebGPU: the IR ladder

In the Studio you draw a workflow — nodes, wires, parameters. What runs, though, is not the drawing: it is GPU code, generated for your machine and executed right there in the browser tab. Between the two sits a ladder of intermediate representations, and the descent down that ladder is where the guarantees are made — feasibility settled before anything runs, determinism declared rather than assumed, every kernel earning its trust against a reference. The engine that owns this ladder is the microkernel from the previous article; here we watch it work.

Down the ladder, rung by rung#

Lowering proceeds in named, inspectable stages. The workflow document — the graph you drew — compiles into a typed program whose unit is the block (BlockIR): structure becomes explicit, which blocks run, in what order, under which control flow. Numeric work lowers further into a tensor-level operation graph, the rung where automatic differentiation happens (WebIR), then into backend-specific kernels — WGSL for WebGPU alongside the CPU reference (KernelIR) — and finally into an execution plan of concrete dispatches, buffers, and a memory plan (ExecPlan). When a run distributes across devices, one more rung — the WorkPlan — partitions it into verifiable shards.

Every rung is a first-class artifact you can open and read. There is no black box between the picture and the result; when you want to know what the machine is about to do, you inspect the rung that answers the question.

graph (mid-ladder, simplified)
  %x    = input "dataset"             shape [N, D]  dtype f32
  %w    = param "weights"             shape [D, K]  dtype f32
  %mean = reduce_mean %x axis=0       shape [D]
  %xc   = sub %x, broadcast %mean     shape [N, D]
  %y    = matmul %xc, %w              shape [N, K]

plan (computed before any dispatch)
  determinism_tier : prod             declared, never defaulted
  tolerance_class  : default          resolves to explicit thresholds per tier
  peak_memory      : parameters + activations + gradients
                     + optimizer_state + working
  verdict          : green            fits the declared budgets

Memory is settled before the run#

The plan exists so that feasibility is decided before the first dispatch. Memory is planned as an explicit breakdown by allocation class — parameters, the activations kept for the backward pass, gradients, optimizer state, and working memory for kernel temporaries — so “mystery GPU memory” has nowhere to hide. Mixed precision adds its f32 master weights to the picture. Large inputs are planned as declared streaming classes, with chunk sizes computed against the budget rather than discovered by crashing, and tiling handles the 128 MB WebGPU binding cap automatically. Those five classes, and what to change when they do not fit, are worked through on training in the browser.

Planning ends in one of three verdicts. Green runs within budget. Yellow runs with warnings — streaming more slowly than an in-memory pass, say. Red refuses, and the refusal is actionable: what failed, why, and what to change — a smaller batch, streaming enabled, a different precision. A mid-run out-of-memory surprise is treated as an architecture defect, not bad luck.

Three determinism tiers — never silently defaulted#

Every run declares how deterministic it must be. There are exactly three tiers:

• debug — bitwise-identical results across runs with the same seed on the same device. On a GPU this forces deterministic kernels: a fixed reduction order, no atomic float accumulation — it costs speed and buys certainty.
• prod — deterministic within a declared numeric tolerance class; the everyday posture for real work.
• fast — throughput first; nondeterminism permitted only where it is declared safe.

The tier travels with the work: declared in specs, selected at execution, recorded with the run. It is never silently defaulted — a workflow with no effective tier fails to compile with a named diagnostic instead of quietly assuming one. When composed parts declare different tiers, the strictest wins unless you explicitly override, and the override announces itself with a diagnostic whenever it actually drops a stricter tier.

What “the same result” means across GPUs#

Stated plainly: results computed on different GPUs are not bit-identical. Floating-point addition is not associative, and different devices order work differently. What holds across devices is agreement within the declared tolerance class — one of exact, tight, default, loose, each resolving to explicit thresholds per tier. So reproducibility is honestly two-tier: replayable on the oracle path (bit-exact) and convergent on the GPU path (tolerance-bounded). Bit-exact replay exists in exactly two places — the debug tier on the same device, and the CPU reference path, which is deterministic everywhere. A plan that cannot honour what you asked for — cross-device bit-exactness on GPUs, say — is refused with a diagnostic, never downgraded to a weaker guarantee that looks the same on screen.

Trust is earned against a reference#

Every built-in GPU operation is validated against a CPU reference implementation — the oracle — within its declared tolerance class. The oracle is the correctness anchor, not a runtime fallback: for a built-in op, a missing GPU kernel is a bug to fix, never a silent detour through slower code. At the bottom of the ladder the artifact under test is ordinary WGSL — readable and checkable:

@group(0) @binding(0) var<storage, read>       x    : array<f32>;
@group(0) @binding(1) var<storage, read>       mean : array<f32>;
@group(0) @binding(2) var<storage, read_write> out  : array<f32>;
@group(0) @binding(3) var<uniform>             dims : vec2<u32>;  // rows, cols

@compute @workgroup_size(64)
fn center_columns(@builtin(global_invocation_id) gid: vec3<u32>) {
  let i = gid.x;
  if (i >= dims.x * dims.y) { return; }
  let col = i % dims.y;
  // one element per invocation; conformance compares this kernel
  // to its CPU reference within the declared tolerance class
  out[i] = x[i] - mean[col];
}

User compute earns it too#

User-authored compute is held to the same discipline, because it is never self-attested. If you author a kernel, you also supply a reference; conformance generates cases from the declared signature and tolerance and proves the kernel matches. A custom gradient is grad-checked against finite differences of your own forward. Passing mints a ConformanceReceipt — profile-scoped, so a device with a different GPU envelope re-runs the check locally instead of inheriting a claim no GPU can make. The receipt gates trust, not execution: you can always iterate on an unconformed kernel on your own machine. This is the whole correctness regime, and at dispatch time the engine does not distinguish a community kernel from a built-in one.

No GPU? A loud CPU path, never a blank canvas#

Not every device has WebGPU — WebGPU rendering in particular is a capability flag, not a guaranteed floor, with Canvas2D and WebGL2 the guaranteed contexts. Scellis probes capability at startup and gives every outcome a defined, visible state. No GPU but a working CPU engine: whole workflows run on the CPU reference backend behind a loud, persistent banner — running on CPU, slower, heavy training limited. Neither: the Studio still opens for viewing and authoring — compose, inspect, share — and tells you exactly why nothing will execute. The one thing that never happens is a blank canvas or a silent non-answer.

That is the whole path of a run: a drawing lowers rung by rung into WGSL; memory and feasibility are settled before the first dispatch; determinism is a declared contract with honest cross-device semantics; and every kernel earns visible trust against a reference you can inspect. The same content-addressed artifacts that make a run trustworthy also make it reproducible by URL.