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
whileand a multiplicative step are refused. - R2: no
returninside it, and nobreakthat belongs to it (one inside a nested loop orswitchbelongs to that). - R3: every write (an assignment, an
inoutargument, an atomic, atextureStore) lands on a name declared in the body; on an outer array ata*i + c(aa nonzero constant, the same at every write to that array, with eachcin[0, |a|)or onec), at one index whose coefficient ofiis loop-invariant, or ati*W + xover a nested loop ofxfrom 0 to< W; on a texture atvec2(i % W, i / W); on a reduction variable written only ass op= e,s = s op e,s = min(s, e)ors = max(s, e)with oneopof+ * & | ^ min maxandenot readings; or on an integer array written only asa[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
consolecall 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
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:
-
examples/loop-examples.test.ts:2 -
src/compiler/ts/host-kernel.test.ts:2 -
src/compiler/ts/kernel-corpus.test.ts:1 -
src/compiler/ts/kernel-loops.test.ts:1 -
src/compiler/ts/kernel-loops.ts:1 -
src/core/host-kernel.ts:212 -
src/core/kernel-tree.test.ts:1 -
src/core/kernel-tree.ts:1 -
src/core/passes/kernel-lower.ts:1 -
src/core/passes/parallel-loop.test.ts:1 -
src/core/passes/parallel-loop.ts:1
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:
-
TS8070KERNEL_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:
docs/language-design.md:859(the design document)reqs/rules/RULE-0822.md(its traceability item)