규칙 11.13

A fuse scope, fuse(body) from typeshade/fuse (change 0058), runs the kernel calls (Rule 8.21) its body makes as one plan, and must give every array they write, and every result, the bits the same calls give run apart (Rule 11.8). device.fuse(body) of a device owner (Rule 11.8, change 0071) runs the calls body makes as one plan on that device, under every rule of this section and with the tiers of that device; a call in the body that takes a Resident of another owner is a TypeError naming both owners, and a fuse call inside the body joins the device’s scope. Its recorder is available without importing typeshade/fuse (change 0085). The body runs synchronously exactly once. At its return the plan reserves its queue position before a cached import loads the package-owned executor. Later calls and reads cannot pass that reservation. A loading failure rejects every recorded call, remains on every written resident and frees scratch storage. Nested scopes, asynchronous-body refusals and synchronous body errors retain the rules below. device.fuse(body, { chunk: { key, units } }) (change 0074) runs one chunk of a long kernel, units of the work of the kernel loop key names, as a plan alone, at the place fuse was called (Rule 11.8): at that place the device waits for the work it submitted before, as device.wait() does, so the measured time is the chunk’s alone, starts a wall clock, runs the plan, fences it as device.wait() does, stops the clock and learns the rate of key (Rule 11.8), and the report gains chunk, { key, units, elapsedMs }; a call queued after the chunk is in neither its plan nor its time, and the promise resolves after the fence. A key that is not a non-empty string and units that is not a whole number of 1 or more are each a TypeError naming the field, and a chunked plan inside the body of another fuse is refused as a use of the kernel queue, since it runs alone. A plan with no chunk runs as this section gives and does not wait for the fence, and fuse(body) of typeshade/fuse takes no chunk. The runtime runs body once, synchronously. It records each kernel call whose written arrays are all Residents (Rule 11.8), a scratch array among them, and each entry call (Rule 8.24) whose storage arrays with no size are all Residents, and runs none. An entry call forms a group of its own in the plan: no kernel call joins it, and it joins none (fact F8 of the planner core/fuse-group.ts). Its dispatch is one dispatch in the plan’s order, and the report names it by the entry’s name. A kernel call that takes a texture (Rule 8.23, change 0077) also forms a group of its own, since the planner records no texel fact: on WebGPU its group’s why is b takes a texture (a group on the CPU tier carries none, surface §71), and it runs as it runs outside a scope, on the plan’s owner and the textures of that owner. A kernel call whose loop pushes onto an array (change 0078) forms a group of its own too, its why b pushes onto an array: its count, scan and store are one item of the plan. A kernel call that takes a Map or a Set (change 0080) forms a group of its own, its why b takes a Map or a Set, or, when it runs on the CPU tier, the reason it does. An entry call with an indirect count (Rule 8.24, change 0075) is recorded and grouped as any entry call is. The plan reads its counts array as a read of the call, so the call runs after the last call before it that wrote the array, and the report’s group gains indirect: true. On WebGPU dispatches counts its dispatch once, whatever the counts hold, since the plan does not read them. On the CPU tier and WebGL2 the call reads the counts from the host copy when it runs (Rule 11.8), and a read of the device that this needs is a submission. In the body, a draw, a kernel call that writes a host array, an entry call that takes a storage array with no size as a host value or writes a host value, read(), write() and destroy() of a Resident and every other use of the kernel queue are a TypeError that names the call, since each would run before the calls recorded earlier or would hold a host value the plan does not own; an async body is a TypeError; a body that throws leaves nothing to run; and a fuse call joins the scope that is open, so one plan runs. When body returns, the plan runs as one item of the kernel queue, at the place fuse was called (Rule 11.8): every call and every read() made after fuse was called sees what the plan wrote. Every recorded call is checked before the first runs, so a call that fails its check fails the plan before any call has run. The promise fuse returns resolves to the report when the plan has run, and rejects with the first error, which also stays on each Resident a recorded call writes, for read() to throw; the promise of each recorded call resolves to what the call resolves to outside a scope, or rejects with that error.

The report’s type, FuseReport, is exported from typeshade/fuse, from typeshade/runtime and from the package root, and mergeReports(reports), from typeshade/runtime and the root, is the report of several plans as one, as a plan of several groups reports its own: the dispatches and the submissions summed, the groups in the order given, and the tier the most capable report ran on, WebGPU before WebGL2 before the CPU tier; with no report it is the report of a plan that did nothing. It is arithmetic over the records and adds no planner dependency of its own (Rule 11.11; change 0071, Amendment 2). A merged report has no chunk row, since it is no one chunk; device.chunkRates holds the rates (change 0074). On WebGPU the plan joins consecutive recorded calls into groups, and a group of several calls is one dispatch of one compute entry, which the runtime builds from the loop IR in the faces of the calls (Rule 8.22); a reduction at the end of a group adds the dispatches that fold its partials. The entry is emitted under the flavor of the emulated double of the scope’s owner (Rule 11.8, change 0073), from the loop IR before the fp64 pass: the _fp64 guard is declared and bound under 'float' and absent under 'integer', and the two flavors are two entries, each built once for its shape. A call joins the group before it when each of these holds against every call already in the group:

• F1: the call is of a kernel function that lowers, was built with typeshade({ fuse: true }) and has one loop, and that loop does not scatter, has no continue of its own and runs at least one iteration.
• F2: the loops have one start, one trip count and one step.
• F3: an array that a call in the group wrote and the call reads is written and read at exactly the loop’s counter i.
• F4: no array the call writes is one a call in the group reads or writes; an array that calls only read is no conflict, and one array has one type in a group.
• F5: the group binds no more distinct arrays than the device allows storage buffers to a stage (maxStorageBuffersPerShaderStage in the device’s report from callDevice(), change 0051, 8 where the report has none), counting one partials binding for each variable a reduction folds and no scratch array that the entry holds as a local.
• F6: a loop that reduces (Rule 7.2) is the last call of a group and no call joins after it; its iteration runs the iterations of the calls before it first, as a prologue, and its folds of the partials and its tail are as the compiler lowered them.
• F7: the calls declare one workgroup size, or none. A call with @workgroup_size (Rule 8.22) declares its size, and every call of a group that declares one declares the same size. A call with none has no preference: it takes the size of its group, which is 64 when no call of the group declares one. A group that ends in a reduction runs at 256 (Rule 7.2), so a call that declares a size joins it only when the size is 256. A call with no attribute joins a reduction as it did before.

A call that joins no group starts one of its own. On WebGPU the report gives each group workgroupShapes, one triple [n, 1, 1] for each dispatch of the group in dispatch order: a fused group or a single-loop call has one, and a lone call of a function with more than one loop has one for each loop. The fold levels of a reduction are not entries; they run at the size of their reducing loop. A group that does not run on WebGPU has no workgroupShapes. Each group also gives its gpuTime, the GPU time of its compute passes in nanoseconds, when the device measured it with timestamp queries (change 0066). The points are these:

• The call layer’s device requests timestamp-query when the adapter reports it. The request is optional: an adapter without the feature gives a device without it, and no refusal applies.
• Each compute pass that a plan records on WebGPU writes two timestamps, at its beginning and at its end, through the timestampWrites of beginComputePass. Only the batch of a plan does this: a call outside a scope writes none. A group’s gpuTime is the sum, over its passes, of the end time minus the beginning time.
• The dispatch batch of the plan resolves the timestamps and copies them to a map-read buffer in the same command encoder as the passes. The plan makes no extra submission, so submissions is the same with the feature and without it.
• After the submission, the plan maps the copy and converts its 64-bit values to numbers of nanoseconds. It sets each group’s gpuTime before the promise of fuse() resolves.
• A group has no gpuTime when the device lacks the feature, when the group runs on WebGL2 or on the CPU tier, when creating the query set throws, when one submission of the plan holds more than 2,048 passes (a query set holds 4,096 queries), or when mapping the copy rejects. The last three cases leave the results of the plan and every other report field as they are.
• A group’s time is the GPU time its passes take. Gaps between passes are in no group’s time. A browser may round a timestamp to a coarser step, and the value is what the device reports.

Only WebGPU fuses: on WebGL2 the calls of a scope run apart, each as it runs outside a scope, and the CPU tier never fuses (Rule 11.7); a module built without the option carries no facts, so its calls run apart. The entry of a group is built once for each shape of group, which is the faces, the slot each array parameter takes and the slots that are locals, with no length and no value, and it is kept with the face of the first call of the group; a group whose entry cannot be built runs its calls apart. A group that ends in a reduction runs as the reduction runs outside a scope (reduceOnGpu): the loop’s dispatch and one dispatch for each level of partials join the plan’s submission. The levels share one device buffer of partials for each variable and each has a uniform of its own, so no level reads the partials back before the next. The same submission copies the folded value of each variable, and the plan makes the result of the call from that copy after its last submission. A call that runs on WebGL2 or on the CPU tier reads its arrays from the host, so the plan submits what it holds before it. The report is { tier, dispatches, submissions, groups }, with chunk for a chunked plan of a device (change 0074). tier is the best tier any call of the plan ran on: WebGPU, then WebGL2, then the CPU tier. dispatches counts the dispatches the plan made on WebGPU and the draws it made on WebGL2, a fused group once; submissions counts its queue.submit calls and the reads of a Resident from the device that a call on the host needed, and the copy of a reduction’s folded value is in a submission already counted and adds none. groups lists the groups in order, each with the names of its kernel functions (calls) and whether it is fused, and indirect: true for the group of an entry call with an indirect count (change 0075). A group that stands apart from the one before it also has why, in the author’s names: <a> and <b> run as two dispatches because <reason>, with an access that is not at i given as the access and its line in the words TS8070 uses (mul and blur run as two dispatches because blur reads `t[i - 1]` (line 9)); three reasons are whole sentences instead, WebGL2 runs the calls apart, <b> runs on the CPU, <the reason the face gives> and <b> runs apart from <a>, which runs on the CPU (or on WebGL2); a group with nothing to say has none. scope.scratch(length, element?) makes a scratch array, a Resident for an intermediate that the host never reads, of 'f32' (the default), 'i32', 'u32' or 'vec4f' elements; it holds zeros when the plan starts. A 'vec4f' scratch array holds length vectors, which are 4 × length numbers of its Resident<Float32Array>, and a kernel call takes it where it takes an array<vec4f>; a call that takes it where it takes an array of another element is a TypeError. The entry of a fused group holds a scratch array as a local, with no device buffer and no host array, when every call of the plan that names it is in that group, exactly one call of the group writes it, every access of it is at exactly i and no call reads its length (t.length); F5 counts it as a buffer until the last call of the plan that names it has joined. A 'vec4f' scratch array is held as a local of vec4<f32>. Otherwise the plan makes it a zeroed array for the scope, a device buffer when a call on WebGPU binds it and a host array when a call on the host needs it, and frees it when the plan has run, on success and on failure. A 'vec4f' device buffer is 16 bytes for each vector, so 16 × length bytes. Its read() rejects with a TypeError and its write() throws one, each naming it as scratch, and its destroy() does nothing; a kernel call that takes it outside the body of the scope that made it, or after its plan has run, is a TypeError; and scratch() throws a TypeError for a length that is not a whole number from 0 to 4294967295, for any element other than 'f32', 'i32', 'u32' and 'vec4f', and when the body has returned. A Resident from resident() is never a local: it keeps its buffer and its write, so the host reads it after the scope. The fused entry computes each operation in the order and with the rounding of the calls apart; what a driver may do with an entry, such as regroup its operations, is the room surface §38 lists, as it is for any entry. The fused entry is WGSL text that the GPU compiles, and no tier builds code with new Function (Rule 11.7). typeshade/fuse stands beside the load-time emitter (Rule 11.11).

Constants selected by an entry call (Rule 8.24, change 0084) are captured when the body records it. They preserve its own group, order, resident bindings and the owner’s fp64 flavor. The report carries its resolved workgroup shape. Selecting constants adds no submission, fence or readback to the plan.