규칙 8.22
이 페이지에서

규칙 8.22

8장, 함수와 진입점 테스트

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.

규칙 본문과 그 아래 항목은 설계 문서에 적힌 컴파일러 원문이라 한국어 페이지에서도 영어로 둡니다.

근거

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.

출처

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

검증 방식

테스트로 확인합니다. 테스트나 게이트 스크립트, CI 워크플로가 이 규칙을 가리키고, 아래 파일 가운데 하나라도 규칙을 더는 가리키지 않으면 추적 검사가 실패합니다.

컴파일러가 이 규칙을 적용하는 곳을 규칙이 직접 적은 내용입니다.

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.

커밋 c66579bf에서 이 규칙을 확인하는 파일입니다. 파일마다 규칙을 처음 가리키는 줄로 연결했습니다.

설명하는 절

커밋 c66579bf의 표면 문서에서 이 규칙을 설명하는 절입니다.

적용하는 오류 코드

규칙이 Enforced by 항목에서 가리키는 진단 코드와, 레지스트리 설명에서 이 규칙을 가리키는 진단 코드입니다.

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.

함께 보기

소스

커밋 c66579bf의 규칙 원문:

이 페이지 편집 문제 보고