Rule 8.22
On this page

Rule 8.22

Chapter 8, Functions and entry points Test

A kernel function is an exported function that is not an entry point and takes at least one array with no size (array<T>); its candidate loops are the for and for…of statements at the top level of its body, and no other loop is ever a candidate. Its body is scalar statements, then its loops, then an optional return of scalars; a statement outside its loops that writes an array, or that is neither a loop nor a scalar statement, runs the whole function on the CPU, and so does a loop that reads a variable an earlier loop reduces (“split the function”). A candidate loop runs as a GPU kernel only when this proof accepts it, on the IR, with no solver and no alias analysis:

  • R1: the loop is counted (Rule 7.5) with an additive constant step; a while and a multiplicative step are refused.
  • R2: no return inside it, and no break that belongs to it (one inside a nested loop or switch belongs to that).
  • R3: every write (an assignment, an inout argument, an atomic, a textureStore) lands on a name declared in the body; on an outer array at a*i + c (a a nonzero constant, the same at every write to that array, with each c in [0, |a|) or one c), at one index whose coefficient of i is loop-invariant, or at i*W + x over a nested loop of x from 0 to < W; on a texture at vec2(i % W, i / W); on a reduction variable written only as s op= e, s = s op e, s = min(s, e) or s = max(s, e) with one op of + * & | ^ min max and e not reading s; or on an integer array written only as a[k] op= e, a scatter reduction.
  • R4: an array the loop writes is read in the body only at an index it writes, a reduction variable is not read, and the value an atomic returns is not read.
  • R5: a function the body calls, transitively, writes no module variable and no binding.
  • R6: no barrier, no workgroup memory and no console call in the body or in what it calls.

A loop the proof refuses runs on the CPU, and the program is correct either way: the refusal is a warning, TS8070, on the loop, whose first sentence names the line and the author’s own names and whose second gives the remedy, one per loop for the first reason (Rule 12.4). A variable an accepted loop reduces is combined in the tree order on every tier (Rule 7.2): each iteration combines into it from the operator’s identity, the iterations’ values are folded 256 at a time by the workgroup tree (at stride 128, then 64, down to 1, slot t becomes slot[t] op slot[t + s], the last block padded with the identity), the partials the same way until one is left, and the variable is combined with that last; a loop that runs no iteration leaves it as it was.

The rule text and the parts under it are the compiler's own English, as the design document writes them.

Rationale

a loop is how a TypeScript developer writes work over an array, and docs/dx.md names dispatches and workgroups as the model the developer should not have to learn first; the proof says which loops are safe to run in any order, predictably, from rules a reader can apply by hand, rather than trusting the author’s assertion (Numba’s prange). Only a function that takes an array with no size is a kernel, so no loop that is per-invocation code today becomes a candidate: #252 measured the 48 loops of the corpus, and none is one. The proof reads an inout argument as a write at its call, which #252’s prototype first missed six times: a method that changes an outer object carries it to the next iteration.

Derives from

change 0013 in changes/ (roadmap item 15); #252; Rule 7.5; docs/dx.md.

How it is verified

Checked by a test. A test, a gate script or a CI workflow names this rule, and the traceability check fails when a file listed below stops naming it.

Where the rule says the compiler enforces it:

proveKernels in src/core/passes/parallel-loop.ts, whose facts src/compiler/ts/kernel-loops.ts words as TS8070 KERNEL_LOOP_ON_CPU in the author’s names; pinned by src/core/passes/parallel-loop.test.ts, every accepted form and refusal #252 measured, and by src/compiler/ts/kernel-loops.test.ts, each refusal’s text in the compiler and in the language service on the same source. The call’s dispatch is lowerKernel in src/core/passes/kernel-lower.ts, which lowers a function whose every loop the proof accepts to one @compute entry per loop, a loop that reduces to a workgroup tree of 256 and an entry that folds its partials, one dispatch per level, with the function’s result from its tail on the CPU tier, and a scatter into an integer array to an atomic on it; callKernel in src/core/host-kernel.ts runs it on WebGPU, and runs the function on the CPU tier where there is none or the function does not lower (a scatter with *, which no atomic does, a refused loop, a bool, a module binding; a function that writes a texture is not callable until #204 gives a storage texture its host value), an emulated f64 dispatched as two f32s with the module’s _fp64 guard bound and its reductions folded by the same tree. The tree order is kernelTree in src/core/kernel-tree.ts on the CPU backends. Pinned by src/compiler/ts/host-kernel.test.ts on the CPU tier against the oracle, by src/core/kernel-tree.test.ts for the tree order, and by the import journey, which calls six kernel functions on WebGPU in Chromium, two of them reductions held to the tree bit for bit and one a histogram by atomics.

The files that verify it at commit 26de7be8, each at the first line that names the rule:

Explained in

The sections of the surface document that explain this rule, at commit 26de7be8:

Error codes that enforce it

The diagnostic codes the rule names under Enforced by, or whose registry text names the rule:

TS8070 KERNEL_LOOP_ON_CPU
A warning on a loop of a kernel function (Rule 8.22, surface §65, proposal 0013) that the independence proof refused: the loop runs on the CPU, and the message names the line and the author's names that stop it, and the remedy.

See also

Source

The rule at commit 26de7be8:

Edit this page Report a problem