Placement is an effect: kernels, buffers, and a handler that decides where they run - #95
Merged
Merged
Conversation
Groundwork for treating placement as an effect. Two things the compiler
already knew but threw away are now carried through to the backends, and
the encoding every backend depends on is written down once.
`eir/layout.rs` holds the NaN-boxing constants that `value64.rs`,
`wasm.rs`, and `native.rs` each used to redeclare under a comment asking
the next person to keep them in sync. It also defines `DType` and
`BufferHeader` for the dense buffers kernels will exchange. A new
`abi_conformance` test compiles the same literals on all three backends
and compares raw bits.
That test immediately found three real divergences:
- `loon run --native` silently returned `()` for any program with a
`main` function. The entry point reaches `main` through Close +
Invoke, both unimplemented, and both emitted a unit placeholder
instead of failing. Operations that *use* an unrepresentable value
now error, as `End::TailInvoke` already did; ones that merely
construct one stay inert, since lowering emits a dead Close for
every named function.
- The WASM backend heap-allocated the nullary `None` instead of using
the immediate singleton. `value64.rs` documents that bit equality is
`None` equality and truthiness is a pure bit test, so a heap `None`
would have tested truthy there and falsy on the VM.
- The native backend encoded `None` as `()`, making the two
indistinguishable.
Ownership modes (Borrow/MutBorrow/Move) were inferred and then dropped
with the checker. They now ride on `Checker::fn_param_modes` and land on
`Func::param_modes` as In/InOut/Owned — the same directional information
Rust makes you write as `&T`/`&mut T`, inferred instead of annotated.
Mode inference also iterates to a fixed point: a single source-order
pass guessed `Move` for callees it had not reached yet, so the answer
depended on the order two functions happened to be written in.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
A kernel is an ordinary function that promised to stay small:
[kernel saxpy [i a x y out]
[put out i [+ [* a [at x i]] [at y i]]]]
`[kernel ...]` desugars to `fn` before anything else looks at it, so kernels
infer, lower, and run through exactly the same path as any other function.
What the keyword buys is a promise, checked in check/kernel.rs: no closures,
no allocation, no strings, no effects. That restriction is the safety
argument. The Rust offload work needs an `unsafe trait` for its partitioning
strategies because a kernel there can index a slice however it likes; here
the unsafe program cannot be written down.
Buffers (Obj::Buffer) are dense, unboxed, fixed-length arrays — the
representation that can leave the process, which persistent trees of
NaN-boxed words cannot. `at` reads, `put` writes in place.
`Place.run`, `Place.read`, `Place.pin`, `Place.stats` are effect operations.
Unhandled, they run serially right here, so a program that never mentions
placement still works. Handled, everything changes without the program
moving: os/place.oo has tracing, dry-run, and residency handlers in a few
lines each, and os/demo-place.oo runs one function under four of them to the
same answer. Because `Place.read` is an operation rather than an accessor, a
handler sees every synchronization point — which is what `Preload`/
`PreloadMut` exist to reconstruct in a language where it cannot.
Direction is inferred, not annotated: `at` is a borrow and `put` a mutable
borrow, so saxpy's modes come out In/In/In/In/InOut with no `&mut` anywhere.
Two bugs found on the way, both pre-existing and both silent:
- Closures called from higher-order builtins clobbered the caller's
register 0. `[let n 8] [map [fn [v] [* 2 v]] [range 0 n]]` left `n` as
14, the last value map computed. The VM re-enters itself for these calls
and named register 0 as the destination for a result that belongs to
Rust; there is now a RET_DISCARD sentinel and a parity regression test.
- `[len buf]` returned 0 through a catch-all rather than the length.
Also: effect operations still move their arguments, so "responding twice
with the same value is a compile error" survives. Place is the exception,
because placement reads its inputs and writes through its outputs — it never
takes ownership, and `[Place.read out]` afterwards is the entire point.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
`--place device` gives buffers a memory of their own. The arithmetic still
happens on the host — this is not a GPU yet — but the bookkeeping is real: a
buffer must be uploaded before a kernel can use it, results must come back
before the host can read them, and nothing survives a launch unless somebody
says it should. That last part is what makes the default honest. A device
that has not been told a buffer will be wanted again does not keep it.
Saying so is `Place.pin`, and that is the entire vocabulary a residency
policy needs:
[fn place/resident [thunk]
[handle [thunk]
[Place.run k n args] [do [Place.pin args] [resume [Place.run k n args]]]
[Place.read b] [do [let v [Place.read b]] [Place.unpin b] [resume v]]]]
Eight launches over one buffer, measured by os/demo-residency.oo:
no policy: uploads 8, resident hits 0, bytes in 128
place/resident: uploads 1, resident hits 7, bytes in 16
Same program, same answer. The Rust offload paper measures this same gap at
up to 400x and closes it with `Preload`/`PreloadMut` types at every call site
plus a transfer-hoisting pass inside LLVM that is still a prototype. The
information that pass tries to recover is already in the program here:
`Place.run` says what a launch touches and `Place.read` says when the host
wants an answer, so a handler can act on it.
Pinning deliberately does not mark a buffer resident. It says "keep this once
it is here", not "it is here already", so the first launch still pays for the
upload — otherwise a policy would look free in the accounting by declaring
itself so, which flatters a design instead of testing it.
Also: --place-stats prints the transfer table, and place/resident-only takes
the buffers to keep as an argument, for the case the paper flags as needing a
heuristic it did not ship.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
An EIR kernel becomes a @compute entry point: buffers are storage bindings, scalars are uniform fields, and the work index is global_invocation_id.x. saxpy comes out as @group(0) @binding(1) var<storage, read> b1: array<f32>; @group(0) @binding(3) var<storage, read_write> b3: array<f32>; ... r5 = b1[u32(r0)]; r6 = (params.s0 * r5); b3[u32(r0)] = r8; with `read` versus `read_write` decided by whether the body indexes the argument with `at` or writes it with `put` — the same fact that decides which way bytes move at runtime. `infer_arg_kinds` reads a launch signature off the body, so nothing has to be declared: an argument that is indexed is a buffer, one that is computed with is a scalar. Two decisions worth naming. Kernels are specialized per launch rather than emitted once for all uses, which is how a language with inferred types still produces a shader whose every binding has a definite type. And there is no relooper: WGSL has no goto, so the emitter wraps a switch on a block index in a loop, exactly as the WASM backend already does. A single-block kernel skips the dispatcher entirely and reads like the source did. Every kernel in the repo is parsed and type-checked by naga in CI, on a machine with no GPU. That is the automated cross-target validation the Rust offload paper reports as still missing — they found their host/device slice divergence by hand. It earned its keep immediately. Emitting the repo's own saxpy produced a shader referencing an undeclared identifier, because a buffer argument used as a number has no WGSL expression to become. That is now refused by argument number instead of emitted. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
wgpu turns the emitted WGSL into a compute pipeline and dispatches it. On this machine that is Metal on an Apple M4 Max; the same code path reaches Vulkan on Linux, DX12 on Windows, and WebGPU in a browser tab. That last one is the entry a compiler emitting PTX and AMDGCN cannot reach at all, and it comes from having picked a portable shading language rather than a vendor's. The end-to-end test is the one worth having: the same Loon kernel run through the interpreter and run on the GPU, compared. Not "close to a formula written separately in Rust" — the same program, both ways. A GPU that computes something different quickly is not an optimization. saxpy over 256 elements, f32 within 1e-4 a branching kernel, exercising the block dispatcher in the shader 10,000 elements across many workgroups, every one written `run` returns the bytes read back from each writable buffer rather than mutating host buffers in place, which keeps it free of any opinion about where the host's copy lives and makes the transfer something a test can look at directly. Only buffers declared writable come back — the same distinction the ownership pass draws, arriving at the hardware. Behind the `gpu` feature, off by default: it is a large dependency and every other placement mode works without it. Absence of a GPU is not a failure — tests print SKIPPED and pass, so nothing in the suite depends on hardware being present. It only gets checked on hardware when hardware is there. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
`Place.run` now dispatches to the GPU when the mode says so. The sample that has been running on the CPU since M1 runs unchanged: loon run samples/place/saxpy.oo --place cpu → #[0 5 10 15 20 25 30 35] loon run samples/place/saxpy.oo --place device → #[0 5 10 15 20 25 30 35] loon run samples/place/saxpy.oo --place gpu → #[0 5 10 15 20 25 30 35] Not one character of the program differs between those, and the residency handler written in M2 against no hardware at all now decides what an Apple M4 Max actually copies: eight launches over one buffer cost eight uploads with no policy and one upload under place/resident, on the GPU, same answer. Residency accounting is shared between the simulated device and the real one, because "how many times did these bytes cross the boundary" is a property of placement rather than of any particular device. That is why the policy did not need rewriting when real hardware arrived. The refusal test found a bug in this integration worth keeping: a buffer passed to a parameter the kernel multiplies by was silently becoming 0.0, because scalars and buffers were drawn from separate lists without checking that each argument matched the shape the kernel body implies. It now says which argument disagrees. A confident wrong answer is the exact failure this design exists to prevent, and it very nearly shipped inside it. Without the `gpu` feature, asking for `--place gpu` says the build has no GPU support rather than quietly running on the CPU. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
loon run samples/place/saxpy.oo --place gpu --record trace.oo loon replay trace.oo samples/place/saxpy.oo # no GPU feature compiled in Both print #[0 5 10 15 20 25 30 35]. Placement joins the recorded effects, so from the program's point of view a kernel launch is an operation that returned nothing and a read is one that produced some numbers. Recording those reproduces the run exactly on a machine with no GPU, in CI, or after the hardware it was written for stopped existing. Nothing about the recording is GPU-specific — a trace made against the simulated device replays the same way. `Place.stats` is deliberately excluded. It reports on the run currently happening, and a replayed run really did move no bytes; feeding back the original transfer counts would be a recording that lies about the execution it is part of. The replayed run reports zero launches, because it performed zero launches. Buffers now print as `#buf<f32 x 3>` rather than `<obj:3>`. A heap slot in output is unhelpful, and it would have made every recorded trace differ from the last for no reason. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
The residency accounting was describing a model, not the hardware. Every
launch was creating fresh GPU buffers, uploading, dispatching, and copying
results back — so `uploads 1` in the stats sat next to four real uploads on
the device. Fixed by giving the GPU a residency map keyed by the host's heap
slot: `Place.pin` keeps an allocation alive, a launch uploads only what is
not already there, and `Place.read` is the only thing that brings bytes home.
An upload counted is now an upload that happened, which makes the benchmark
worth reading. A chain of launches over one 4096-element buffer, Apple M4 Max
via Metal, identical program in both columns:
launches no policy place/resident speedup
8 28.9 ms 9.3 ms 3.1x
32 93.1 ms 15.4 ms 6.0x
128 356.7 ms 18.3 ms 19.5x
Two bugs fixed on the way, both of the kind that produce wrong answers rather
than slow ones:
- Evicting a dirty buffer discarded what the kernel had written. A cache
that drops dirty data is not a cache. Eviction now writes back first,
which is exactly the cost a policy of keeping nothing resident is
choosing to pay.
- `Place.read` collected values from the host copy *before* downloading, so
it returned stale data and reported a transfer that changed nothing. The
transfer now happens first, which is the whole content of the word
"synchronization".
Compiled shaders are cached, so a loop launching the same kernel does not
recompile it every iteration.
BENCHMARKS.md records all of this with the framing it needs: the CPU column
is Loon's own interpreter, the slowest reasonable baseline, and the ratio says
what there is to gain by leaving it — not how a generated shader compares to
a hand-written kernel. That comparison is not attempted and not claimed.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
docs/blog/2026-08-18-placement-is-an-effect.md and the changelog entry. The post's argument, in short: the two facts a GPU offload compiler works hardest to recover — what does this launch touch, and when does the host actually want to look — are both *events*, and we already have a way to spell those. Making Place.run and Place.read effect operations turns the transfer-hoisting pass into a nine-line handler, and turns the "heuristic we did not ship" for the Energy benchmark into an argument the user passes. It is explicit about what is not claimed. No comparison against hand-written CUDA is attempted. The CPU column in the benchmarks is Loon's own interpreter, so the GPU ratio says what there is to gain by leaving the interpreter, not anything about generated code quality. Reductions and atomics are outside the kernel subset, and f64 narrows to f32 on a device that has no 64-bit scalar. And the honest summary of the comparison: they built a compiler and this moved a seam. Those are different kinds of work. The seam is worth writing about because it makes residency, prefetch, eviction, tracing, simulation, and replay stop being compiler features somebody has to ship for you and start being code you can write on a Tuesday. Third in the arc: v0.7 effects end to end, v0.8 the kernel is a handler, v0.9 the GPU is a handler in the middle. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Kernels were running through the general interpreter: every value NaN-boxed,
every buffer access through the heap table, every arithmetic operation
checking what it was handed. Right for Loon, wrong for a loop doing the same
three floating-point operations a million times.
eir/kernel_exec.rs runs the numeric subset against raw slices instead. A
kernel outside the subset falls back to the VM, so nothing a kernel was
allowed to be has been narrowed — `supported()` is a question asked once
before a launch, not a requirement.
That makes the CPU baseline fair, which matters because "the GPU beats our
interpreter" was never an interesting claim. And because the executor works
over an index range on plain slices, `--place par` follows: split the range,
give each thread a disjoint piece of the output via split_at_mut, and let the
borrow checker be the thing that guarantees they do not overlap. The Rust
offload paper needs an `unsafe impl` of a partitioning strategy to promise
the same property.
A kernel that writes outside its own range fails with a message saying so,
rather than racing another thread.
The same kernel across four placements (Apple M4 Max):
elements cpu par gpu
1024 447 µs 566 µs 10.3 ms
262144 8.9 ms 3.3 ms 12.4 ms
1048576 36.4 ms 11.4 ms 19.2 ms
The last row is the interesting one: every core beats the GPU here. A machine
with this many fast cores, and a launch that pays submission and transfer
before computing anything, land that way. That is not a disappointing result
— you found it out by changing one word, because the program does not know
where it runs, and where the crossover sits is a property of the machine
rather than of the code.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Two corrections and a finding. The finding: a GPU dispatch submits to the queue and returns. Nothing waits until the host asks for data, and `Place.read` is the only way to ask — so a handler that defers reads defers every synchronization point in the program. 64 launches take 19.9 ms against 6.9 ms for one; if each blocked on the previous, that would be nearer 440 ms. The offload paper prototypes async transfers as a separate LLVM optimization. Here it is what happens when the synchronization point is something the program says out loud. The correction that matters: the post implied you can run one of these programs in a browser today. You cannot. WGSL is WebGPU's shading language and naga validates every kernel we emit, so the shader half really is portable — but Loon's wasm build still embeds the old tree-walking interpreter rather than the EIR VM, so there is nowhere for the program to run. That gap predates this work and has nothing to do with placement. The post now says the architecture points at a target the alternatives cannot reach, and does not claim to have arrived. Also corrected: Vulkan and DX12 are the same code path through wgpu, but I have not run either, so they are "should" rather than "does". Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
docs/plans/2026-08-18-placement-remaining.md, so the next person does not have to infer it from the code: what shipped, and the four things that did not — the browser (blocked on the wasm build embedding the old interpreter, not on anything about placement), reductions, 64-bit on a GPU, and compile-time rejection of scatter kernels. Also records that the placement stack itself compiles for wasm32-unknown-unknown today, which is the checkable half of the portability claim. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
`DType::gpu_ok` existed, was tested, and was never called. So `--place gpu`
silently narrowed an f64 buffer to f32 while the docs said the narrowing was
"reported rather than hidden". It is now refused:
error: argument 1 is a f64 buffer, and a GPU has no 64-bit number; run this
with --place cpu or --place par, or build the buffer with `buf` instead of
`buf-f64`
Handing back numbers of a precision the program never asked for, and has no
way to notice, is the confident-wrong-answer failure this design exists to
prevent. The same program runs on the CPU, which does have 64-bit numbers.
And the disjointness rule is now enforced where it belongs. A kernel may write
at its own index and nowhere else — that is what lets the parallel executor
hand each thread a slice and what lets a GPU run every work item at once. The
parallel executor caught violations at runtime; the checker now rejects them
at compile time (E0602), naming the offending index:
error: [E0602] kernel 'reverse' writes at an index other than its own
why: every work item runs at once, so each may only write element 'i';
writing elsewhere means two of them can reach the same element and
the result would depend on which got there first
fix: write at 'i', and read whatever else this element needs with `at`
Reading anywhere is still fine — gather is safe, scatter is not. The index
parameter can be called whatever the kernel calls it.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
A kernel may write at its own index and nowhere else, which sounds like it
rules out summing an array. It doesn't: a reduction is a map followed by a much
smaller reduction. Give each work item its own chunk of the input and its own
slot in a partials buffer and every item is still writing exactly one element —
its own. Combine the sixty-four partials on the host, where sending them to a
device would cost more than adding them.
So there is no `Place.reduce`, no workgroup-shared memory, and no exception
carved out of the disjointness rule. samples/place/reduce.oo sums 0..4096 to
8386560 on cpu, par, device, and gpu alike.
Getting there needed `loop` added to the kernel forms — `recur` was already
allowed, which is meaningless without the form that introduces it — and turned
up four bugs.
Two in `sum`, both pre-existing. The registry has always declared it
`Vec Num → Num` and the interpreter has always implemented that, but the
checker said `Vec Int → Int` with the comment "(approximate)", so summing
floats did not type check; and the EIR VM filtered the vector to integers and
returned an i64, so a vector of floats summed to 0. That is the worst shape of
wrong answer, because it looks like an empty sum rather than a failure.
Three in the WGSL emitter, each of which produced a shader that failed
validation — the good outcome, but only because the emitter is checked:
- The type-inference fixpoint stopped when no *new* register had a type, so
a loop-carried value kept whichever type its first predecessor gave it.
Types now widen monotonically and iterate until nothing changes.
- Branch targets were resolved across every function in the module. `BlockId`
is per-function, so a jump found whichever other function happened to have
a block with the same number and copied arguments into its parameters.
- A unit value reaching a float slot had no conversion, so the emitter wrote
a bare i32 register where WGSL wanted an f32.
The reduction kernel is now in the naga validation corpus, so all three stay
fixed.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
`crates/loon-wasm` gains an `eval_placed` export that runs on the EIR VM, and
`web/public/place.html` exercises it: pick a placement in a tab, see what
crossed the boundary.
placed on device: 4 launches over 32 work items;
9 uploads (288 B), 2 downloads (64 B), 3 resident hits
That is the residency handler — written in M2 against no hardware at all —
deciding what a browser copies. Verified in a real tab, not inferred from the
fact that it compiles.
Kernels and buffers exist only on the EIR VM, which is why a new export was
needed rather than a migration. The DOM-driving exports stay on the legacy
interpreter: its bridge is written against `Value`/`InterpError` rather than
the EIR's `Val`/`VmResult`, and the guide's examples use builtins such as
`push!` that the EIR VM does not implement. Adding an entry point beats
breaking documented pages.
WebGPU remains the missing half, and now for a precise reason rather than a
vague one: wgpu's device setup returns futures a browser resolves on its event
loop, and `eir::gpu` blocks on them with pollster, which cannot block on wasm.
`--place gpu` in a tab refuses and says why. The shaders are validated by naga
in CI, so both ends are known-good; the wire between them is what is absent.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
`run_parallel` handled exactly one output buffer and quietly ran everything
else sequentially. Asking for `--place par` and getting `--place cpu` was
indistinguishable from getting what you asked for, which is the kind of silence
this codebase spends most of its error messages avoiding.
Every output is now carved into per-thread slices. `iter_mut` yields disjoint
`&mut`s, so pieces taken from different buffers coexist without anyone
promising they don't overlap — `split_at_mut` returns that, and the borrow
checker enforces it.
A kernel splitting one input into two outputs now really does run across cores:
[kernel split [i src lo hi]
[let v [at src i]]
[put lo i [if [< v 0.0] v 0.0]]
[put hi i [if [< v 0.0] 0.0 v]]]
cpu, par, device, and gpu agree on the answer at 1000 elements.
Arguments are now passed to the executor already separated into scalars,
inputs, and outputs rather than pre-assembled, which is what lets the parallel
driver carve the outputs while sharing the inputs whole. `run_sequential`
takes the same shape so a caller does not assemble arguments two ways.
An output shorter than the launch is reported rather than silently covering
fewer elements.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
`place.html` was only reachable because I was serving `web/public` directly. The build copies a fixed list of static assets into `dist`, and the page was not on it — so on the deployed site it would have 404'd. Also removed the stray `loon_wasm*` copies I had left in `public/` root. The build says plainly not to put them there (`public/pkg/` is the single source of truth, and its contents are copied to the dist root), and they are gitignored, so they were local clutter that happened to make my test pass for the wrong reason. Verified against the built site rather than the source tree: device placement reports 9 uploads and 3 resident hits, and `--place gpu` refuses. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
The remaining-work doc had drifted: reductions, the browser, scatter rejection, and the 64-bit refusal were all filed under "not done" after being done. Restructured into three sections rather than two, because the middle one was the part being mis-shelved. "Decided, and deliberately restrictive" holds the rules the design rests on — write at your own index, no 64-bit on a GPU, no allocation or effects in a kernel. Those are not gaps waiting to be closed; they are the reason the parallel executor can hand out slices and the reason the unsafe program is unwriteable rather than merely rejected. What is left is genuinely left: WebGPU behind the browser needs either an asynchronous effect path in the VM or a worker blocking on Atomics.wait, and neither is really about placement; atomics would need a second kind of kernel with its own safety argument; and there is still no comparison against hand-written CUDA, which we have not measured and do not claim. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
`eir::device::Device` is the six things a placement backend has to do: name itself, make a buffer resident, say whether one is, dispatch, download, evict. wgpu implements it on a desktop. A browser will implement it by proxying to JavaScript, because WebGPU is reachable only through promises and this VM is synchronous. The VM now holds an `Rc<dyn Device>` and does not know which it has. A host can install one; failing that, the `gpu` feature opens wgpu; failing that, asking for `--place gpu` says there is nowhere to run rather than quietly running somewhere else. The refactor immediately turned up a real waste: every argument was marked dirty after a launch, so saxpy — two buffers read, one written — copied all three home and reported 96 bytes when 32 had changed. Only written buffers are dirty now. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
A Loon kernel now runs on the GPU in a browser tab:
placed on gpu: 4 launches over 32 work items;
9 uploads (288 B), 3 downloads (96 B), 3 resident hits
Same WGSL a desktop build hands to wgpu, same answer as every other placement,
with the residency handler deciding what gets copied.
The obstacle was that WebGPU is asynchronous — a device comes from a promise,
a readback goes through mapAsync — and Loon's VM is synchronous all the way
down: `Place.read` is an effect operation that returns a value, not a future.
Both facts cannot hold on one thread.
So they hold on two. The VM runs in a Web Worker; every device call posts a
request to the main thread and blocks on `Atomics.wait` until the reply lands
in a SharedArrayBuffer. The main thread, where the promises live, does the
WebGPU work and wakes the worker. The blocking is the trick, not a workaround:
it lets an asynchronous API sit under a synchronous language without either
pretending to be the other.
None of that reached the VM, because the device went behind `eir::device::
Device` first — name, ensure-resident, is-resident, dispatch, download, evict.
wgpu implements it on a laptop, a JS bridge implements it in a tab, and the VM
does not know which it has. That is the move placement makes at the language
level, one floor down.
crates/loon-wasm/src/gpu_bridge.rs the Rust side of the bridge
web/public/place-gpu.js WebGPU on the main thread
web/public/place-worker.js the VM, blocking
vercel.json COOP/COEP so SharedArrayBuffer exists
Verified in a real tab against the built site, not the source tree. Without
cross-origin isolation the page says so and the other placements still work.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Looking into what it would take to move the DOM exports onto the EIR VM turned
up something bigger than the DOM.
`set!` and `push!` are documented in four places — the collections guide, the
ownership guide, the builtin reference, and DESIGN.md — and neither exists on
the EIR VM, which is the default backend. `loon run` on a program from the
guide says "unbound symbol 'set!'".
On the interpreter they do exist, and `push!` does not mutate:
[let mut v #[1 2 3]]
[push! v 4]
[println v] ; #[1 2 3] — unchanged
[println [push! v 9]] ; #[1 2 3 9]
The guide's own example claims the opposite, and prints `#[]` on the
interpreter while failing to type check on the EIR VM. The documented
behaviour is currently true of no backend at all.
Not fixed here, and deliberately. Porting the interpreter's behaviour verbatim
would give a second backend a builtin whose name promises mutation and whose
implementation returns a copy; making it mutate is a language decision with a
real question attached (what a closure that captured the old value sees).
Written up in the remaining-work doc with both options rather than guessed at.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
I claimed that letting a browser answer `Place.read` asynchronously — and so
dropping the cross-origin isolation requirement — would need an asynchronous
effect path in the VM. Then I tried it. The VM already has one.
A handler clause does not have to call `resume`. If it hands `resume`
somewhere else and returns, the handled computation unwinds without ending,
and the continuation stays live in whoever caught it. Calling it later picks
the program up mid-expression:
work: starting
host: computation parked; the rest of it is mine now
host: ...doing something slow...
work: continued with 21
host: finished with 42
os/demo-park.oo, with tests pinning the ordering — "work: continued" appears
after the host has printed, which is only possible if the computation really
unwound and was restarted from outside. No VM support was added for any of
this. It is what reified escaping continuations already are.
Two things needed correcting as a result. The remaining-work doc said this was
"a real project rather than a small change"; the semantics were never the
problem. What is actually left is plumbing — a Vm that outlives one call,
since the continuation lives in its heap, and an export for the page to resume
through — and the post now says so.
I had assumed the hard part was the language and the easy part was the wiring.
It was the other way round, which is what comes of writing down what I thought
was true instead of trying it.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
The plumbing the last commit said was left. `eir::vm::Session` prepares a
program, runs it until it finishes or parks, and can be resumed later with the
value the host went off to fetch:
let mut s = Session::new(src, dir, Mode::Cpu)?;
let Step::Parked { .. } = s.start()? else { ... };
let k = s.pending().unwrap();
s.resume(k, s.vec_of_floats(&bytes_from_the_gpu))?;
`Host.park` is the effect a suspending handler performs; the VM keeps the
continuation and the request, and the handled computation unwinds. `Vm::
call_value` invokes a continuation or closure from Rust. The VM stays alive
across all of it, which is the one real requirement: a parked continuation is
a heap object, so dropping the VM would drop the rest of the program.
Writing the tests turned up the one genuinely surprising thing about parking,
now pinned by a test named after it: unwinding means the code *after* the
handle runs immediately, with whatever the clause returned, and only the
suspended part waits. So a deferring handler belongs outermost with nothing
meaningful after it, and the computation's value arrives through `resume`
rather than from the call that parked. That is written down in os/demo-park.oo
where somebody about to write such a handler will read it.
Also fixed: `Session::take_output` returned nothing, because `Vm::run` hands
its output out in the result and the session was discarding it between steps.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
`place_start` and `place_resume` expose the session to a page. A program that
parks returns `{done: false}`; the page does whatever it could not do
synchronously and calls back with the numbers.
Confirmed in a browser with crossOriginIsolated === false — no worker, no
SharedArrayBuffer, no COOP/COEP:
start: done=false ← parked at Place.read
resume: done=true out="read #[2 4 6]" value=12 ← finished with host data
So the browser GPU has two possible paths now. The worker one, which is what
the demo page uses and which is verified on real hardware, needs isolation
headers. This one needs nothing, and only the readback ever has to wait —
`writeBuffer` and `submit` are already synchronous.
Pointing the page at it is what remains, and it is a page change rather than a
capability one: the sample has to be written with the deferring handler
outermost, because parking unwinds to the `handle` and anything after it runs
immediately with the placeholder.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
|
The latest updates on your projects. Learn more about Vercel for GitHub.
|
main split the OS layer out to ecto/fxos, which deletes `os/` from this repo. The placement handlers and demos were living there and are not part of the OS, so they move to where the other placement samples already are: os/place.oo -> samples/place/lib.oo os/demo-place.oo -> samples/place/demo-handlers.oo os/demo-residency.oo -> samples/place/demo-residency.oo os/demo-park.oo -> samples/place/demo-park.oo `crates/loon-lang/tests/loon_os.rs` went with the split. The six placement tests it had gained are extracted to `tests/place_samples.rs`, resolved against the new directory; all six still pass. Two content conflicts, both the same shape — upstream and this branch each added a parameter to `run_file`, and each added modules to `eir`. Merged rather than picked, so `--unchecked` and `--place` compose: loon run samples/place/saxpy.oo --unchecked --place device 33 suites green, formatting clean, and every relocated sample runs. The residency demo still reports 8 uploads with no policy and 1 under place/resident. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Clippy's approx_constant is deny-by-default and caught two values in the golden f32 buffer written as 3.14159265 and -2.71828. They were never meant to be pi and e — the list is deliberately awkward bit patterns (signed zeros, a denormal, the extremes, a NaN with a payload) and those two were there as ordinary values with a full mantissa, to catch a backend that round-trips the special cases and mangles the mundane ones. Written as truncated constants they suggested they meant something. Replaced with values that are plainly arbitrary, and the comment now says what they are for. My miss: I ran fmt and the tests locally but never clippy, which is what CI runs. Checked now across the workspace and with the gpu feature; both clean. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Reading of GPU Offload in Rust: Portable, Safe, and Fast (Drehwald et al., arXiv 2608.13759), and what it suggested for Loon.
The paper reads data direction off
&T/&mut T, gets safe kernels via anunsafe traitfor partitioning strategies, and marks sync points withPreload/PreloadMut. What it's candid about not solving: its convenient interface is up to 400x slower than explicit transfers, closing that needs a transfer-hoisting pass inside LLVM that's still a prototype, and the "Energy" benchmark (6 kernels sharing ~15 arrays) needs a heuristic they didn't ship.Those two facts a compiler works hardest to recover — what does this launch touch and when does the host actually want to look — are both events. Loon has a way to spell that.
What this adds
Place.runandPlace.readare effect operations, so a handler sees every launch and every synchronization point. Residency is then library code, not a compiler pass —os/place.oo:[fn place/resident [thunk] [handle [thunk] [Place.run k n args] [do [Place.pin args] [resume [Place.run k n args]]] [Place.read b] [do [let v [Place.read b]] [Place.unpin b] [resume v]]]]On an Apple M4 Max via Metal, identical program in both columns:
Direction is inferred, not annotated: Loon already computed per-parameter ownership modes and threw them away.
atis a read,puta write-through, so saxpy's shader binds inputsreadand the outputread_writewith no&mutwritten anywhere.Kernels run on CPU, across cores, on a GPU via wgpu (Metal/Vulkan/DX12), and in a browser tab — all agreeing on the answer. They're recordable: a run on the GPU replays on a build with no GPU compiled in.
Bugs found, several pre-existing on main
The tests kept earning their keep. All of these produced wrong answers, not crashes:
loon run --nativesilently returned()for any program with amainfunction.map/filterclobbered the caller's register 0.[let n 8] [map [fn [v] [* 2 v]] [range 0 n]]leftnas 14.None, so it would have tested truthy there and falsy on the VM.sumfiltered vectors to integers — a vector of floats summed to 0, which looks like an empty sum rather than a failure. The checker also typed itVec Int → Intwith the comment "(approximate)", contradicting the registry's documentedVec Num → Num.[len buf]returned 0 through a catch-all.BlockIdis per-function), and a unit value reaching a float slot with no conversion.Deliberate restrictions
E0602). Reading anywhere is fine. This is what lets the parallel executor hand each thread asplit_at_mutslice — disjointness the borrow checker enforces rather than anunsafe implpromises.f64, so such a launch is refused by name rather than silently narrowed.What is not claimed
For the reviewer
docs/blog/2026-08-18-placement-is-an-effect.mdis the argument;docs/plans/2026-08-18-placement-remaining.mdis the honest boundary.set!andpush!are documented in four places and exist on neither the default backend; and on the interpreterpush!does not mutate despite its name, so the guide's own example is true of no backend. Both options are written up rather than guessed at.Atomics.wait(needs COOP/COEP). A second path needing no headers is built and verified — a handler that parks the continuation — but the demo page still uses the worker.gpufeature, 32 suites without;cargo fmt --checkclean; the whole stack compiles forwasm32-unknown-unknown.Try it:
🤖 Generated with Claude Code