Rule 11.8

A kernel function’s call (Rule 8.21) runs on the first tier of configure({ prefer }) that can run it, WebGPU, WebGL2 and the CPU tier being the default order: WebGPU when the function lowers (Rule 8.22) and there is a device, WebGL2 when every loop writes one array of 4-byte elements (f32, i32, u32) at exactly i and every scalar parameter is a number or a vector, each loop one fragment program into an R32UI target that starts out holding the array, so an iteration that writes nothing leaves its element, and the CPU tier always; a list of one tier makes that tier required, and a call no tier on the list can run throws an Error naming each tier’s reason. Every tier computes a reduction in the tree order (Rule 7.2). Kernel calls run one after another in the order they were made, whether or not the caller awaits them. resident(value) wraps a host value (Rule 8.21) once in a Resident, which holds its own copy: a call on WebGPU uploads it the first time and binds the same buffer at every call after, reading nothing back, and await r.read() waits for the calls made before it and returns a new copy of what it holds; r.write(value) replaces what it holds and r.destroy() releases its buffer, each after the calls made before it, or at once when none waits, a use after destroy() being a TypeError; a call that fails while nobody awaits it leaves its error on each Resident it writes, which read() throws. One Resident must not be two parameters of one call. The same Resident is a buffer of the program runtime too, bound by any draw or dispatch of a runtime on the same device, and configure({ runtime }) puts every call on that runtime’s device (Rule 11.11); a program runtime’s Texture passed where a call takes an image is bound as it is, and a draw that falls to WebGL2 refuses it with a TypeError. A @compute entry’s call (Rule 8.24) runs in the same order as kernel calls, and on the first tier of configure({ prefer }) that can run it: WebGPU when there is a device, and the CPU tier when the entry reaches no barrier and no texture; WebGL2 has no compute stage and never runs one, and a call no tier on the list can run throws, a TypeError naming what the entry reaches when the CPU tier refused it, an Error naming each tier’s reason otherwise. A storage array binding with no size may be a Resident, bound on WebGPU as its device buffer and read back by nobody, and a call whose written bindings are all resident has a second signature, returning void; a Resident for any other binding, or as two bindings of one call, is a TypeError. A draw’s tiers are fixed: WebGPU, then WebGL2, then the CPU tier, where a canvas keeps the tier of its first draw; a draw may read a Resident for a storage array with no size, and then runs in the order of kernel and entry calls, seeing what the calls before it wrote, with the handle bound as its device buffer on WebGPU and its copy brought up to date on the CPU tier; a Resident for any other binding is a TypeError.