Module Ir.Schedule

Applied left to right by apply.

Builds a Split with fresh outer and inner index symbols (via Indexing.get_symbol) and returns them, so subsequent ops in a programmatically built schedule can reference the new loops.

Builds a Partition with one fresh index symbol per segment (via Indexing.get_symbol) and returns them in segment order, so subsequent ops in a programmatically built schedule can address individual segments.

The Partition a Pad of an axis of extent to a multiple of block prescribes, once that axis is Split by block into block_loop and an inner loop: block_loop partitioned at its last full block, so every guard the pad left on the inner loop (block*block_loop + inner < extent, e.g. a zero-fringe Stage's edge guard on each packed element) is interval-decided in the full blocks and folds there, and survives only in the edge block (gh-ocannl-1631). Empty when block divides extent (no pad) or the axis has no full block. The one spelling of this split for the packed GEBP builders: the autotuner's packed sketch (Sketch_families) and bin/bench_gebp.ml (gh-ocannl-1670).

Derives candidate Partition breakpoints of the axis loop from the guards already present in its body: statement If conditions and scalar Where conditions (e.g. the virtualizer's per-component range guards of an inlined concatenation, Split's remainder guard, or a clamped window's range guards, gh-ocannl-504) are scanned for comparisons whose two sides differ by k*axis + off. With everything besides axis constant, such a comparison flips truth value at one point of the axis range. Non-axis symbols bound by other loops (inside the axis loop — e.g. the window symbol of a clamped-window guard — or enclosing it) are bounded by their loop ranges, putting off in an interval: the comparison is then always-true / mixed / always-false over the axis range, and both transition points are recorded — the mixed (boundary) segments are exactly delimited while the decided segments fold their guards. Returns the collected flip points that fall strictly inside the loop range, sorted and deduplicated (possibly empty — e.g. when every guard is already interval-decided); partitioning at them makes every such guard interval-decided within each segment (for interval-ranged comparisons: outside the mixed segments), so apply's trailing simplify erases them. Loops are located the way apply rewrites them (gh-ocannl-668): inside a Local_scope body, where the accumulation mints of a materializing Unroll or a Partition put the segment and inner loops; and over EVERY loop binding axis, not the first — a materializing Unroll leaves one copy per unrolled step with a different constant substituted for the unrolled index, so one source guard's copies flip at different points and the result is their union (each copy filtered by its own range). Raises Invalid_argument when no loop binds axis.

fold_mma ~query ~width mints the lane and block symbols of a Fold_mma.

coalesce ~outer ~inner mints the merged loop's symbol of a Coalesce.

Builds an Expand_zero with one fresh symbol per axis of tn (forcing tn's dims) and returns the symbols for subsequent Split/Retype ops.

Builds a Privatize whose tile resides at accum_prec of target's storage precision (forcing it). Pass the rendering backend's resolution, (Context.codegen_capabilities ctx).accum_prec — the same function its code generation widens the serial rendering's accumulator by — so the privatized reduction's width is the policy's, not the schedule's (gh-ocannl-1116).

The detail of Privatize's v1 refusal of target: under the accumulation loop it is accessed at more than one index vector. Unlike Stage's (gh-ocannl-1147), this refusal is not relaxed: the target is written under the loop, so a second accessor at another vector would read a stale target between the tile's init-load and its store-back unless proved disjoint from every cell the loop writes. No tuned workload reaches it (gh-ocannl-1395's census, pinned by test/operations/privatize_census), which classifies Autotune.privatize_extension's drops by this text.

Builds a Split_reduce with fresh block and inner index symbols plus one fresh combine symbol per axis of target (forcing target's dims), and returns them (op, block, inner, combine) so subsequent ops can annotate the pass-1 block loop and the combine nest's loops.

Low_level.fused_accumulation: the fused form of a Tile_mma fallback's accumulation into d[di], None for an integer accumulation (gh-ocannl-824) or no product accumulated into d[di]. Tensorize no longer writes its fallback through it (gh-ocannl-1457): the renderer fuses the fallback unless the backend's compile pins contraction off. The sketch seeding asks it whether the register tile can render a site.

Builds a Tensorize with a fresh lane symbol (via Indexing.get_symbol) and returns it. ?tile is the register-tile geometry request (gh-ocannl-619), absent by default.

add_symbol idx s is idx plus s with coefficient 1: the per-tile coordinate Tensorize and Fuse_epilogue add to an accumulator's base index. None -- which each caller refuses in its own words -- for a Concat and for a Sub_axis: the latter has no per-axis affine form, and extended as the empty sum it would turn into an ordinary Iterator s coordinate where the flattened component after it carries the address (gh-ocannl-1165).

The loops that would have to enclose the reduction for this Split_reduce to be recognized: the accumulation cell's symbols that are currently bound inside the reduction loop (gh-ocannl-537). This is the one rejection cause a loop interchange can remove — OCANNL lowers conv bias/weight gradients with the accumulated channel loops innermost and the reduction loops (batch, y, x) outside them, the exact inverse of the static form's pinning discipline — so autotune seeding hoists these symbols with a Swap chain and re-probes, rather than dropping the site. The answer comes from the recognizer itself (a hermetic probe of its own apply), not from a re-implementation of the discipline or a parse of its message.

[] for every other outcome: a legal op, a non-Split_reduce op, and any rejection an interchange cannot remove (the cell mentioning the reduction loop itself, an enclosing loop that pins no component, an unrecognized accumulation form).

Why Fuse_epilogue on target would fail on this code — the recognizer's own refusal message, probed hermetically like split_reduce_hoist — or None when an eligible elementwise tail immediately follows the reduction over target and the fusion applies. The autotune family tree judges its fused-flavor branch with it (the check runs on the pre-schedule code, where the plain accumulation-nest site applies), so a site without a fusable tail carries the recognizer's reason as a construction-time refutation rather than silently minting no twins.

Option.is_none (fuse_epilogue_witness ~target opt): whether the fusion applies.

Whether the node is eligible as a hoisted Stage source: declared value-constant (Tnode.known_host_constant) with registered host-init data to pack from. Shared by the autotune sketch and by Schedule_cache.canonicalize, which renders it per tensor node so that same-shape programs differing in operand constancy do not share cached schedules.

The constancy half of a hoisted Stage's source preconditions: Tnode.known_host_constant, or a known_constant placement in the lineage.

gh-ocannl-1554: whether some schedule can take every read of the node off it — a hoisted Stage reads a host-packed tile instead and emits no load nest, so a candidate whose store is filtered to its code footprint holds the tile, sized by the staged accesses, and not the node. Exactly the hoisted-source preconditions apply checks on the node itself (registered host-init data and hoist_known_constant).

Whether the node's placement is settled On_device in this lineage: the one placement the shared epilogue placement of Fuse_epilogue leaves alone.

Applies the ops left to right to the optimized code, then elides the staging barriers a later Tensorize made redundant, then re-runs Low_level.simplify_llc (which folds remainder and edge guards the loop extents prove) and, when a materializing Unroll duplicated code, CSE + cross-statement hoisting. Stage registers its tile in the traced store (and workgroup_shared when shared); the optimization context and merge node are never changed. Raises Invalid_argument when an op references a loop that does not exist at its point in the schedule, or violates an op precondition (see optop) -- the rendering of the typed refusal apply_classified passes through. An empty schedule is the identity. Every op is applied with the program the schedule started from in hand, which Privatize judges its tile width on, and its wide tiles are re-checked on the finalized program (privatized_tile_widens).

Each op is refused, naming it, when the program it leaves breaks the scope grammar (gh-ocannl-1467): a Local_scope body is hoisted out of its position, so it may hold neither a hardware-annotated loop (Low_level.scope_hardware_loop) nor an effect beyond its own locals (Low_level.scope_purity_violation) — e.g. a hardware Retype of a loop inside the scope a materializing Unroll minted, or a Stage whose loads land there. Only a violation the op added is refused; an entry program already breaking the grammar (hand-built IR) is left to the compile gates. Likewise when a Privatize tile minted earlier in the same application gets a hardware-typed loop inside its per-thread lifetime (see Privatize in optop); a tile minted by an earlier apply is not seen.

The barrier elision (gh-ocannl-487, widened to every staged anchor by gh-ocannl-567) is a synchronization-only rewrite — bitwise identical output, on every backend path. It rests on the Tile_mma emission contract, whose two halves are NOT symmetric: every rendering form ENDS the intrinsic block with a workgroup barrier, so a staging barrier that follows one (inside a staged anchor's body, or right after the loop) is dropped at any pipeline depth; but only some forms OPEN one per iteration — the fragment-scope form opens it once, on the scope wrapping the anchor loop — so a barrier is never elided against a following Tile_mma. Hence the loop-carried arm (a rotor body's leading barrier, against the previous iteration's bracket) stays pipelined-only, and a depth-1 k-block keeps the barrier between its loads and its compute.

Internal candidate-facing variant of apply. An op precondition refusal (an op referencing a loop absent at its point in the schedule, or violating a precondition of optop) raises Schedule_outcome.Cause_at (Transform, Illegal_schedule { check = "Schedule.apply"; _ }), which a search counts as a declined candidate. Every other exception escaping the transforms or the finalization passes -- a bare Invalid_argument, an array bound, a List.zip_exn -- is a transform bug and propagates untyped, hence fatal under strict_failure_classification (gh-ocannl-1323).

The re-check of a Privatize tile minted at the accumulator residency (gh-ocannl-1151): every write of tile in the program is either a transfer (it does not read the tile) or a self-accumulating update whose base Low_level.accum_base_widens — a Tile_mma's scalar fallback judged like any statement. apply asks it of the finalized program and refuses a tile failing it. Exposed for tests.

Test seam (gh-ocannl-1323): with_transform_fault fault body runs body with fault called on each op just before apply / apply_classified applies it, so a test can make a transform raise what a transform bug raises — a genuine bug has no stable trigger. The calling domain's previous fault is restored on return or exception; the default is a no-op and no configuration selects it.

gh-ocannl-1231: how an IR breaks the staging scope's invariant for one of its pipelined tiles (Low_level.pipelined_tile) — the two positional facts the renderer's buffer rotation needs.

The first violation of the staging scope's invariant in opt.llc, over every tile of opt.pipelined: every loop binding a rotor is Serial, and every read of a pipelined tile sits inside a loop binding its rotor. None when opt has no pipelined tiles. These are the facts C_syntax checks while rendering (its typed Backend_codegen declines remain as the backstop); op_legality refuses the moves that would break them.

The decline feature naming the violation's kind — the key a search report aggregates Unsupported declines under.

The renderer's typed decline for a violation it meets while rendering (an IR that bypassed op_legality): Unsupported with pipeline_violation_feature and site ^ ": " followed by pipeline_violation_message. C_syntax builds both of its pipelined-tile declines through it, so the oracle and the renderer state each fact in one wording; pipelined_move_legality pins that, on the same IR, they reach the same violation.

How a Tile_mma into a marked fragment falls outside every scope the renderer recognizes — rendered, then, as the lane-0 scalar fallback.

One violation per Tile_mma into a fragment of opt.simdgroup_fragments that no recognized scope covers, in program order. A scope is Low_level.fragment_region over a marked fragment — the renderer's recognition — and coverage follows C_syntax.pp_ll's own route to it (each Seq tried whole, a recognized scope's trailing statements rendered one by one). [] when opt marks no fragment.

The nodes the code writes: Set-family and Zero_out statements, Tile_mma accumulators, and the statements inside Local_scope bodies and scan-loop inits.

gh-494 waypoint 3: the op-legality oracle. A schedule op is a thread-pairing transform over the routine's access relations; its obligation is decided by the Affine queries — pairing the op's loop symbol(s) as thread identity must leave every write-involving access pair of every written node Disjoint or Same_thread (with the reduction reassociation license for Vectorized retypes and Swaps of accumulations). Three-valued and sound in both proven directions: Op_legal proves the annotation race-free; Op_illegal proves a violation (e.g. a materialized node's unguarded write provably independent of the retyped axis); Op_unknown means "compile and see" — the op may be valid under semantics the queries do not model (per-thread scratch copies, renderer fallbacks) — and must never be treated as a rejection. Hardware annotations interleave sibling statements' threads with no grid-wide synchronization, so a hardware retype/split of a loop sharing an accessed node across its statement boundary (with a write on either side) reports Op_unknown: such pairs need the default annotator's aligned-mapping analysis, which the single-op oracle does not model. Vectorized retypes and Swaps only reorder within the nest and keep the intra-nest scope.

Tensorize's role assignment is decided: the micro-kernel recognition of apply (a pure function of the code) is probed, so an invalid role permutation — e.g. the reduction loop assigned to i/j, violating the discipline that the k role is the only rmw-carrying loop of the accumulation — is a proven Op_illegal; on structural success the affine queries decide i/j iteration independence under the reduction-reassociation license (the intrinsic reassociates k), with the hardware-lane cross-statement downgrade to Op_unknown as above. Stage is likewise probed hermetically (a raising apply is Op_illegal), then its implicit contract — the staged tile covers the reads it replaces within the staging scope — is the containment query Affine.read_covered_before over the fresh tile's accesses: a covered non-shared (packing) stage is Op_legal; shared staging (barrier placement and launch geometry are validated downstream) and hoisted staging (link-time host-side packing) report Op_unknown. Privatize/Expand_zero/Fuse_epilogue report Op_unknown; their own apply-time preconditions remain in force.

Before any of that, on code carrying pipelined tiles (gh-ocannl-1231): an op whose application leaves a pipeline_violation is a proven Op_illegal carrying pipeline_violation_message — the staging scope's invariant, stated once for every op shape, so that the beam's menu, seeding's schedule_legality and the renderer agree. Likewise, on code marking simdgroup fragments (gh-ocannl-1342): an op whose application leaves more fragment_scope_violations than it found is Op_illegal with fragment_scope_violation_message of the first it added. First of all, on code where some Local_scope body holds a loop: an op whose application adds a scope-grammar violation (gh-ocannl-1467, see apply) is Op_illegal with the refusal apply would raise — the race queries alone call a Grid retype of a scope-nested loop legal. The Privatize lifetime check needs the application's record of its tiles, so a single op judged here does not get it; schedule_legality does, as it applies the chain.

Per-op verdicts, each against the code with the preceding ops applied (on a hermetic copy — checking never mutates the argument). Stops after a proven-illegal op or a refused application; an op whose application is refused reports Op_illegal with the refusal. Any other exception propagates: it is a transform bug, not a verdict (gh-ocannl-1323).

The cross-nest analysis behind default_gpu (the CPU preset proves nests separately), as data: per top-level loop nest (paired with the nest statement itself), the outermost loops that may carry hardware geometry — already trimmed to the common aligned prefix each dependency component admits, so chain position k of two linked nests denotes the same hardware thread coordinate and the two extents at position k are equal. None when the analysis bails, i.e. when the kernel is not safely parallelizable at all.

max_chain caps each nest's chain length; the default 2 mirrors the presets' one-Grid + one-Workgroup shape. A caller supplying its own geometry per chain position passes the arity it needs — a batched matmul site's chain is its batch loops plus row plus column, and capping at 2 made every rank-3+ site's companion coverage decline (gh-ocannl-569). The alignment rule itself is arity-independent.

The presets consume this to emit their own Grid/Workgroup-per-nest geometry. A sketch pipeline (autotune) consumes it to cover the nests it does not build — a tensorized accumulation nest binds two Grid slots and a Workgroup lane, geometry no preset emits, and its companion nests must be annotated with the matching slot structure or Low_level.validate_parallel rejects their writes (gh-ocannl-521). Callers supplying their own geometry must keep the positional pairing: the alignment argument is exactly that thread (c0, c1, ...) covers the same index slice in every linked nest.

expanded_zeros names nodes whose whole-node Zero_out the caller's schedule will turn into a per-element nest (Expand_zero) carrying the caller's geometry. Such a statement is skipped for the query instead of bailing it as a bare materialized write — the write it stands for is not bare by the time the code is validated.

How the default GPU schedule's lane geometry treats a preamble reduction — a serial loop accumulating into the fused attention backward's dp scope local ahead of the lane loop, dp = Σ_e dO·v (gh-ocannl-1124); only that local qualifies (Online_softmax.reassociable_local), since the cooperative form reassociates under the fused backward's approximate-tier license. Preamble_refused: such a nest keeps its plain plan (the gh-ocannl-1003 stage-1 rule). Preamble_duplicated: the nest takes lanes and every lane recomputes the reduction serially, in its summation order. Preamble_cooperative: a reduction whose extent is the lane's whole one-loop workgroup and k of the device's simdgroup_width, k at most lane_all_reduce_simdgroups_for, is retyped Workgroup_reduce, sharing the lane's slot, and renders as an all-reduce leaving the sum in every lane (a butterfly within each simdgroup, then, past one, the partials through workgroup-shared memory between two barriers, gh-ocannl-1168); a kernel holding any other admitted reduction keeps its plain plan, never duplicated lanes (a renderer that still cannot shuffle runs the loop in every lane, the retype's serial meaning). Config gpu_lane_preamble_reduction, whose auto default resolves per device (lane_preamble_reduction_for).

Config gpu_lane_preamble_reduction (auto, the default, is None | refused | duplicated | cooperative).

The treatment default_gpu applies at limits absent an explicit ?preamble_reduction: the configured mode, or under auto the device's economics — Preamble_cooperative where Backend_intf.hardware_limits' lane_scalar_recompute_cheap holds (measured: Metal, CUDA), Preamble_refused otherwise (HIP, the C backends, anything unmeasured), since the lanes recompute each pair's scalar preamble once per lane (gh-ocannl-1124).

Config gpu_lane_all_reduce_simdgroups (auto, the default, is None | a positive count).

The most simdgroups a Preamble_cooperative reduction may span at limits absent an explicit ?all_reduce_simdgroups: the configured count, or under auto the device's measured Backend_intf.hardware_limits' lane_all_reduce_simdgroups (1 wherever unmeasured, so a head wider than one simdgroup keeps its plain plan; gh-ocannl-1168).

How default_gpu weighs its lane geometry (gh-ocannl-1167). Every lane recomputes its nest's per-pair scalar preamble, which the plain plan pays once per thread. Lanes_cut: taken wherever a kernel admits it, and fission_keep_mapping cuts a kernel rather than merge a lane nest out of its lanes. Lanes_admitted: taken wherever a kernel admits it, but merges between nests with a lane path are judged on the plain plans, so such a merge costing a nest nothing but its lanes is taken; a neighbour with no lane path is still cut off rather than cost a nest its lanes (gh-ocannl-1239, mapping_probe). Lanes_off: never taken, so the preamble-reduction treatment has nothing to apply to. Config gpu_serial_lanes, whose auto default resolves per device (serial_lanes_for).

Config gpu_serial_lanes (auto, the default, is None | cut | admitted | off).

The configured mode, or under auto the device's economics: Lanes_cut where Backend_intf.hardware_limits' lane_scalar_recompute_cheap holds (measured: Metal, CUDA), Lanes_admitted otherwise (HIP: the forward value pass wins on lanes in its own kernel, the fused backward's dV loses when cut from dK for them).

What the canonical value of the configuration key key (Utils.canonical_config_value) resolves to on the device limits describes, for Utils.config_class_fingerprint ~resolve (gh-ocannl-1272): an auto lane mode becomes the mode lane_preamble_reduction_for, lane_all_reduce_simdgroups_for or serial_lanes_for picks at limits, and an explicit one stays itself; a gpu_schedule_block_size becomes its clamp to the device's caps. Any other key, and a value its parser rejects, is returned unchanged.

The default GPU annotator preset (schedule-ir-optops §6): for each top-level loop nest whose parallelism is provable from the lowered code alone, annotate every loop of its proved parallel chain (gh-ocannl-1133) — the leading loops Grid (Grid slots >= 2 fold onto .z), the innermost the Workgroup lane, split Grid blocks x Workgroup by block_size when it is longer (default from config gpu_schedule_block_size = 256, clamped to limits' Backend_intf.hardware_limits.max_threads_per_workgroup and .x workgroup cap when given — the configured block size is a target, the device's workgroup capacity a hard cap). While the workgroup holds fewer than workgroup_fill threads (default from config gpu_schedule_workgroup_fill = 256), the loops just above the lane join it as further Workgroup dimensions (at most three, the product within the block size, each extent within its dimension's cap). The plan is taken only when every chain-carrying nest of the kernel fits one common hardware topology (a nest one Grid slot short may split its lane into a one-block Grid loop), the launch fits the device's per-dimension caps, and no nest does less useful parallel work than under the two-loop presets below, with more work somewhere; otherwise the presets apply. The presets annotate exactly one Grid and one Workgroup loop per nest (splitting a single parallel loop by block_size; a second loop longer than the block size keeps its outer part Serial). A loop is parallelizable when its index occurs as a plain Iterator component in every materialized write vector beneath it — the same coverage property Low_level.validate_parallel enforces, used generatively — and the kernel passes a conservative race analysis (all accesses to written nodes agree on parallel-index components, no Zero_out of materialized nodes, no barriers or opaque statements; reduction loops stay serial). Both plans map proved chains: the lane plans propose each nest's whole chain to that analysis, the presets a pair — when the leading parallel extent is below min_parallel, a later pair with no fewer grid groups and a larger grid-times-clamped-workgroup product. Skipped leading loops remain serial. This choice precedes the race analysis; if it fails, or alignment loses groups, active lanes, or the launch threshold, the original outermost pair is used instead. Expanded whole-node zeros (zero_expansion) use the same plans. Unless lanes is false (default: the serial_lanes_for mode at limits is not Lanes_off), a lane geometry takes priority over both (gh-ocannl-1003): when every nest carrying a chain has a parallel loop under a serial loop past loop-free lane-uniform scalar work (declarations and assignments of scope locals -- the online-softmax hoist's value pass, for (b, s, h) { for t { p := P[s, t]; for e { O[s, e] += p * V[t, e] } } }; a preamble holding an inlined reduction is excluded, since every lane would recompute it; a preamble reduction loop into a scope local is admitted per ?preamble_reduction, see lane_preamble_reduction, spanning at most ?all_reduce_simdgroups simdgroups when cooperative), the chain extends through that preamble uncapped: the loops above the serial loop become Grid loops (slots >= 2 fold onto .z) and the loop past it a Workgroup lane, Grid (b, s, h) -> Serial t -> Workgroup e — taken when the planner finds such nests one common topology, the launch fits the device's caps, and it has more threads than the plain plan's geometry. The race analysis is the same: thread identity is the tuple of chain symbols wherever they sit, the serial loop runs whole in every thread, and Low_level.validate_parallel already accepts a hardware loop under a serial one. Cross-nest producer/consumer (or WAW/WAR) pairs over a written node are allowed only when aligned: the linked nests' chains are trimmed to a common equal-extent prefix — identical annotation geometry, so each hardware thread covers the same index slice in every linked nest — and per axis position the paired accesses either both use plain Iterators of same-chain-position parallel symbols or neither mentions one; otherwise the analysis bails. For non-materialized (per-thread copy) scratch the edge additionally requires value thread-invariance at the chosen trim: a chain symbol feeding a scratch write's value without pinning the written cell would leave each consumer thread's copy holding its own chunk's last value where the serial reference holds the last chunk's, so the trim search serializes that loop (gh-494; direct syntactic dependence only). Returns the empty schedule when any check fails or when the largest parallelizable nest has fewer than min_parallel iterations (default from config gpu_schedule_min_parallel = 64: a kernel launches either way, so any real parallelism beats the serial 1x1 fallback — a single GPU thread is 1-2 orders of magnitude slower than a CPU core; the remaining small threshold keeps sub-simdgroup-scale programs fully serial so their segments coalesce and placements stay unchanged).

The default CPU annotator preset (gh-ocannl-1458): per top-level nest, at most one loop is retyped to Grid -- pool-backed Grid rendering in the C backend (docs/proposals/gh-ocannl-164.md) partitions it into contiguous chunks on a process-global native thread pool, and a Workgroup split would only add loop structure that runs serially inside a chunk. The candidate is the nest's first path loop whose index is a plain Iterator in every materialized write -- in every scratch write for a nest writing only scratch (Local, one function-scope array on the C backends; gh-ocannl-1343) -- with at least two iterations and the per-nest hazard analysis passing over that ONE symbol: the loop actually made parallel is the one proven independent.

No cross-nest alignment, unlike default_gpu. The C renderer runs each pool-parallel Grid loop as one synchronous dispatch (dispatch_apply / #pragma omp parallel for) that joins before the next statement, and everything outside those loops once, on the calling thread, in program order. So two accesses in different top-level statements are ordered by a join for every placement -- materialized buffers and shared function-scope locals alike; a Local is privatized per chunk only when all its accesses in the kernel sit inside that loop's body, so no privatized value crosses statements; within one dispatch the hazard proof, C_syntax.parallel_grid_safe (locals) and the renderer's store-separation check decide; a Grid loop under serial loops forks and joins per instance. Low_level.validate_parallel's ~fork_join mode is the validator's side of the same model: writes outside every Grid loop are not subject to coverage, so a serial nest may sit beside a parallel one. A statement the access scan cannot model, or a nest with no candidate or a failed proof, stays serial without affecting its siblings.

Profitability is per dispatch: a candidate is retyped only when one run of it does at least min_work scalar updates (Cost_model.fork_join_updates, the renderer's launch-floor units; enclosing serial iterations and sibling nests do not count). Default from config cpu_schedule_min_work = 16384; a non-positive value disables the gate (the tuner's ungated presets pass ~min_work:0).

Per-channel vector accumulators for direct conv nests (gh-ocannl-1630): per top-level statement that is a perfect Serial accumulation nest of write loops then reduction loops, whose destination's minor axis is the innermost write loop oc, with one operand reading oc (and otherwise only reduction loops, not written in the routine) and some operand reading a sliding window, at single or double storage throughout: Swaps sinking oc below every reduction loop (each cell keeps its terms in order), a packing Stage of the oc operand channel-minor at the routine root (skipped when it is channel-minor already), a Privatize of the destination over the outermost reduction loop at its storage precision, and oc retyped Vectorized. Bitwise identical to the scalar-chain rendering. A site whose ops do not apply after base (on a scratch copy of the store) is left out, and so is everything when the sites do not apply together.

The default CPU preset: default_cpu, followed, under config cpu_schedule_channel_vectors (default true), by channel_vectors over it. What maybe_default_schedule and maybe_default_schedules apply on CPU backends, and the tuner's segment presets.

Config schedule_log_launches: backend compile logs one stderr line per compiled segment with its launch grid/block dims and statement count — for diffing what two compiles of nominally identical code actually emit.

All tensor nodes the code references (scope ids' backing nodes included) — the set a fission segment's traced store is filtered to — and whether it reads the merge buffer.

gh-ocannl-1554: whether some schedule can keep this existing node off the device, the premise on which the device-memory floor (Backends.link_floor) leaves it out. The one schedule op that turns an existing node Local is Fuse_epilogue's shared placement, of a fusion target (written, read by the fused tail) not settled_on_device; it exists only where the backend states workgroup_shared_placement, and charges the node's whole size to the workgroup memory that check_hardware_limits_classified holds to max_workgroup_memory_bytes for every kernel a compile emits, so a node larger than that limit cannot leave the device. None as the limit bounds nothing.

Whether a search on a backend with facts can read the default pipeline's configuration key key (gh-ocannl-1331): the GPU-only keys (automatic_gpu_schedule, gpu_schedule_*, the lane modes, gpu_fission_keep_mapping) exactly where Backend_intf.binds_hardware_grid holds, the CPU-only ones (automatic_cpu_schedule, cpu_schedule_min_work) where Backend_intf.renders_on_host_pool does, and any other key everywhere. Autotune.search_shaping_key_read adds the seeded plan's keys.

Whether maybe_default_schedule (and maybe_default_schedules) would apply a default preset on this backend rather than the identity: the backend binds Grid loops, the corresponding config gate (automatic_gpu_schedule / automatic_cpu_schedule) is on, and runtime kernel logging (debug_log_from_routines) is off. Callers that substitute their own default-schedule choice (e.g. the cost-model default selection of gh-ocannl-491) must respect this gate.

The config schedule_fission gate consulted by maybe_default_schedules: whether the untuned default pipeline fissions into per-segment kernels (subject to automatic_schedule_active), or applies the whole-routine maybe_default_schedule instead.

A stable summary of the policy and configuration that shape the untuned default pipeline on this backend: the automatic_schedule_active gate ("inactive" when it is off — the untuned default is then the unscheduled serial form), the policy's version tags (bumped when a change redefines the untuned default), the schedule_fission gate, and the preset thresholds (gpu_schedule_block_size / gpu_schedule_min_parallel / cpu_schedule_min_work). Diagnostics recorded against one default pipeline (e.g. the autotuner's default_ms, gh-ocannl-552) compare fingerprints to detect that a config change redefined the default (Codex P2 on PR #279). Does not cover per-device hardware limits — cache consumers already key per backend.

The implicit transform applied by backend compile when the caller passes no ?lowered_transform: apply of default_gpu on GPU backends, of default_cpu on CPU backends, the identity otherwise. limits (default Backend_intf.no_hardware_limits) should be the compiling backend's Backend_intf.Backend_device_common.hardware_limits. Disabled by config automatic_gpu_schedule=false / automatic_cpu_schedule=false respectively, and skipped when runtime kernel logging (debug_log_from_routines) is active, to keep logs serial and deterministic.

The expand-and-annotate schedule maybe_default_schedules applies to a fission segment of materialized whole-node Zero_outs on GPU backends: expand_zero plus the same geometry policy as default_gpu — the lane plans over each node's non-singleton axes (its whole loop set is a proved chain: each cell is written once), under one common topology across the segment's nodes, else the two-loop presets. Below min_parallel (largest node) the zeros stay whole-node (a serial kernel renders them as memset). Exposed for callers (e.g. the autotuner) that replicate the default fission pipeline with custom per-segment schedules.

The hardware mapping schedule gives each top-level statement of the code, as (groups, active threads) of the statement's OWN loops, read off the ops: what schedule-aware fission (fission_keep_mapping) compares between a merged kernel and a statement alone. A Workgroup_reduce retype shares the slot of the statement's widest Workgroup loop (the lane geometry's cooperative preamble reduction, gh-ocannl-1124), adding threads only past its width.

The probes of schedule-aware fission's merge rule (gh-ocannl-1126): mp_kernel, the schedule a merged kernel will receive, against which a merge is refused when it gives some statement fewer groups or active threads of its own loops than it gets alone; and, where lanes are only Lanes_admitted (gh-ocannl-1167), mp_lane_trade, the same schedule without the lane geometry, on which a merge is judged instead -- unless the merged kernel would hold a loop nest with no lane path beside one with a lane path, which must pass on both (gh-ocannl-1239: such a kernel declines the lanes for a neighbour that could never take them).

mapping_probe ?lane_trade kernel: the probes, no lane trade unless given.

The keep_mapping probes maybe_default_schedules passes to fission_scheduled: default_gpu at limits on a GPU backend while config gpu_fission_keep_mapping is on (the default), None otherwise -- with its lane geometry unless serial_lanes_for is Lanes_off, and under Lanes_admitted with the lane-free default_gpu as the lane trade, so a merge costing a nest only its lanes is taken where every nest of the kernel has a lane path. A caller replicating the default segmentation (the autotuner's fissioned candidates) passes the same, so its segments are the untuned pipeline's.

The kernel-fission pipeline underlying maybe_default_schedules, with caller-supplied per-segment schedules and a per-segment result: the routine's top-level statements are partitioned at cross-workgroup dependency edges exactly as described there (edges the aligned cross-nest rule of default_gpu proves race-free without losing any nest's standalone parallelism do not cut), preset is called on each `Normal segment's (pre-schedule) optimized slice and zero_sched on each `Zeros segment's nodes, and each result tuple carries the segment kind, the pre-schedule segment, the schedule chosen for it, and the scheduled segment (apply of the schedule). `Solo segments (opaque to the analysis, or coalesced runs of unannotated segments) get the empty schedule.

fold_zeros (default false, gh-ocannl-1175): a `Zeros segment holding a single zero that initializes the reduction opening the next segment folds into that segment, expanded by zero_sched's Expand_zero, when the merge rules below re-admit the whole segment with the per-cell zero at its head; no other boundary moves, so the fold never merges two reductions the zero kept apart. Only the autotuner's and model selector's per-segment sketch candidates pass it, so the fold ships only where a timed (or modeled) sketch pipeline carries it: the untuned default keeps the separate zero kernel, because an untiled accumulation with its zero folded measured 7% slower on HIP (unified memory) although its emitted source only drops the opening read. A replayed segmentation refolds where its record says.

When fission does not apply (single segment, unfissionable crossings, or everything coalesces back) the result is a single `Normal tuple over the whole routine with preset's schedule. Callers compile each scheduled segment as its own kernel in order (the plural transform seam of backend compile); see maybe_default_schedules for the synchronization contract.

arity_cuts (default false, gh-ocannl-574): segment for the full-arity sketch pipelines instead of the default presets. The no-parallelism-loss guard normally compares chains under the presets' Grid+Workgroup cap, so a merge that trims a rank-3+ site's minor axis reads as lossless — correct for the default annotators, which never bind more than that prefix, but the GPU matmul sketches annotate the site's whole chain and cover every materialized-writing nest of the kernel with the same geometry (companion_geometry, gh-ocannl-521/569). Under arity_cuts chains are analyzed uncapped and a merge additionally requires every materialized-writing nest to keep an identical extent list, so a companion that cannot follow the site's arity — a reduction over the site's minor axis and the reduction target's initialization nest, the lm_head's max-logits row — is cut into its own downstream kernel (stream order supplies the synchronization) and the site's kernel seeds at full arity. The finer segmentation costs launches the default pipeline does not want to pay unconditionally, so this is a candidate-generation mode (the autotuner times it), never the default.

keep_mapping (gh-ocannl-1126): the default GPU schedule, as fission_keep_mapping probes it -- the schedule each candidate kernel would actually receive, except that where lanes are Lanes_admitted a merge between nests with a lane path is judged without the lane geometry, so no boundary is kept for lanes between them (mapping_probe, gh-ocannl-1239). A merge the rules above admit -- an aligned dependent merge, and a conflict-free one, which they admit unconditionally -- is still refused when some statement of the merged kernel gets fewer Grid groups or fewer active threads of its OWN loops under keep_mapping than it gets in a kernel of its own: the kernel boundary is kept rather than a nest's mapping lost. Only the refused merges add cuts, so a merge that keeps every mapping (an elementwise tail over the same chain) still saves its launch. The same rule holds a `Zeros segment together (gh-ocannl-1169): a zero joins one only when no node of the merged segment gets fewer groups or active threads under zero_sched than alone -- a union of zeros with no common lane-plan topology sends every node to the two-loop presets. Each maximal run of whole-node zeros is stably sorted by the zeroed node's rank and dimensions in every mode (the zeros commute), so the cuts fall at shape changes rather than at every alternation of a backward pass's gradient zeros. Ignored under arity_cuts. None (the default): the legality rules alone decide.

promote_locals (default false): promote statement-crossing Local scratch to On_device before segmentation. A nest whose only writes land in Local scratch gets no parallel chain (the annotator's coverage property quantifies over materialized writes), and Local producer/consumer edges do not cut — so such producers either drag their segment down to a serial 1x1 launch or are redundantly re-executed by every hardware thread; small reduction intermediates (layer-norm statistics, softmax max/denominator) land exactly here via the Local stack threshold. Promotion lets the ordinary materialized machinery apply; promotions fission does not end up needing (single-kernel fallbacks, or all accesses within one segment after coalescing) are restored. maybe_default_schedules passes true on GPU backends, where a serial nest costs orders of magnitude more than on CPU.

A recorded segment's kind: a derived segment's (`Normal, `Zeros, `Solo), or `Unfolded (gh-ocannl-1480) -- the unfissioned routine's one segment whose head is a whole-node zero left whole although its fold is provable. A `Normal segment opening with such a zero is a folded one and replays with the zero expanded; an `Unfolded one replays as a `Normal segment with the zero as it stands, so a replay rebuilds exactly the code its recording compile ran.

A fission segmentation as data (gh-ocannl-1164): the kind and the length, in units, of each segment in order — a unit being one top-level statement with the comments before it, with each maximal run of whole-node zeros stably sorted by the zeroed node's shape, so the units are a function of the code alone. One segment over every unit is the unfissioned routine. The schedule cache persists it with a fissioned winner, so a replay cuts where the winner was cut instead of re-deriving the cuts under the current policy.

The kind a recorded segment is replayed as: `Unfolded is a `Normal segment.

fission_scheduled, also returning the segmentation of its result (one element per tuple), or — given replay = (segmentation, schedule) — applying that recorded segmentation instead of deriving one. A recorded segmentation is neither re-grouped nor coalesced, so arity_cuts and keep_mapping do not reach it; the promotions and scope-local replicas are recomputed for its boundaries exactly as for derived ones, and the ith segment's (pre-schedule) slice gets schedule i whatever its kind — by position, because segments of one structure can differ in placements, hence in kind and schedule. Where schedule i is None the segment gets its kind's schedule as in a derived segmentation (preset, zero_sched, or none for a `Solo): a per-segment substitution over a recorded segmentation (the autotuner's sketch candidates, gh-ocannl-1243) names only the segments it schedules. A segmentation that does not fit the routine — its lengths do not sum to the routine's units, a segment is empty, or a boundary separates a scope-local definition from a use it cannot be replicated for — raises Schedule_outcome.Cause_at with an Illegal_schedule cause, a candidate's ordinary decline; it is never repaired into another segmentation. Whatever raises, the promotions made so far are restored first.

Like maybe_default_schedule, but with kernel fission: the routine's top-level statements are partitioned into segments at the cross-workgroup dependency edges the default annotator's interference analysis would otherwise reject wholesale — materialized producer/consumer (and WAW/WAR) pairs of sibling statements, bare materialized writes, materialized whole-node Zero_outs, statements opaque to the analysis — and each segment receives the default schedule independently, on its own launch geometry. The caller compiles each returned optimized as its own kernel and runs them in order on the routine's stream, chaining a device-side event at each boundary — that event supplies the grid-wide synchronization at each cut (queue FIFO alone does not order overlapping command buffers over Metal's untracked resources). Scalar scope-locals hoisted across statements are replicated into consuming segments when provably value-preserving (segments merge back and run serially otherwise); Local-placed scratch whose live range crosses a cut is promoted to On_device in the compile's placement fork. On GPU, segments consisting of materialized Zero_outs are expanded (expand_zero) and annotated like ordinary nests. Adjacent segments that end up unannotated are coalesced, so an all-serial routine stays a single kernel. Returns a singleton equal to maybe_default_schedule's result when fission does not apply (non-GPU/CPU backend, automatic scheduling disabled, config schedule_fission=false, kernel logging active, or nothing to split).

As much of a candidate's launch geometry as its caller knows (gh-ocannl-709): the two grid dimensions backends cap and the workgroup's three, None meaning "not predicted here" — which exempts that dimension rather than refusing on it. Five fields, not six: grid.(0) is 2^31-scale on every backend that binds hardware axes, so no backend has a cap to report for it. check_hardware_limits_classified fills all five from the lowered code (launch_geometry_of_dims); autotune's seeding pre-filter predicts what it can from the parameters it is about to commit to. A prediction is a lower bound — it describes the site's own nest, while Low_level.launch_dims maxes over the kernel's zeroing and companion nests too — so an under-prediction costs a compile the gate then declines, and only an over-prediction could withhold a legal candidate.

All-None: predicts nothing, refuses nothing. The base a family builds its prediction on.

The geometry a lowered kernel will actually launch with — every field Some.

The geometry a kernel launches with when its nests' hardware loops have these extents — per nest, its Grid and its Workgroup loop extents in nest order, outermost first (gh-ocannl-1133). The one encoding of the positional slot rule outside Low_level.launch_dims, which reads the same rule off lowered code: among loops of one kind the innermost binds .x, the next .y, the next .z; the maximum is taken per slot across nests, and Grid slots >= 2 fold the product of their per-slot maxima onto .z. Every field is Some; None when that product overflows. The default annotators' lane plans and the sketch families' predictions (Autotune.matmul_launch_geometry / Autotune.conv_launch_geometry) both read it.

The first dimension of a geometry the device refuses. lx_phrase is the verb phrase both callers render — e.g. "requests a .z workgroup extent of 128, exceeding the device limit of 64" — so a gate detail and a seeding refutation witness say the same thing about the same candidate.

The one static reading of a device's per-dimension launch caps (gh-ocannl-709), consulted both by check_hardware_limits_classified (on the lowered code's geometry) and by autotune's seeding pre-filter (on a candidate's predicted geometry), so the two cannot disagree about what a device permits and a cap is never encoded twice. None = every predicted dimension fits. Rows are tested in the order .x/.y/.z workgroup, .y grid, .z grid fold; a None cap (Backend_intf.hardware_limits.max_workgroup_dims on the C backends, Backend_intf.hardware_limits.max_grid_yz on Metal) exempts its dimensions. The workgroup thread PRODUCT is deliberately not a row: it is not a per-dimension geometry question, and only the gate checks it.

The workgroup-shared charge of a kernel declaring arrays (each one array's declared bytes) and tile-MMA emission scopes of the (d, a, b) storage triples mma_triples: every array rounded up to the backend's Backend_intf.codegen_capabilities.workgroup_allocation_alignment (gh-ocannl-1330), plus the backend's per-kernel Backend_intf.codegen_capabilities.mma_kernel_workgroup_bytes over the triples, once (gh-ocannl-1255). What workgroup_memory_bytes charges the code it reads, and what a seeding prediction of it charges the tiles and scopes every completion of a sketch declares (Sketch_families), so the two round and count alike.

The static workgroup-shared bytes the kernel will declare, charged by kernel_workgroup_bytes (Metal rounds every threadgroup array up to 4 bytes, gh-ocannl-1330): its workgroup_shared tiles (a pipelined tile, one array of rotating copies, counted once per copy) plus, once for the kernel, the rendering backend's Backend_intf.codegen_capabilities.mma_kernel_workgroup_bytes over the storage triples of its tile-MMA emission scopes — Metal declares one coordinate table per kernel when any of them converts its destination boundary (gh-ocannl-1205, gh-ocannl-1255). A scope is one Tile_mma statement, each copy listed where a loop transform duplicated it; a simdgroup-fragment scope is rendered only around exactly one Tile_mma, so it is the same list. Plus, per lane all-reduce site that stages per-simdgroup partials (Low_level.lane_all_reduce_sites' lar_cross_simdgroup, gh-ocannl-1168), the backend's Backend_intf.codegen_capabilities.lane_all_reduce_workgroup_bytes for its extent at the larger of its local's compute and accumulator precisions. Plus, per warp-shuffle reduction into a cell (Low_level.warp_reduce_cell_sites, gh-ocannl-1452), the backend's Backend_intf.codegen_capabilities.warp_reduce_workgroup_bytes for its extent at the accumulator residency of the cell's storage precision: the per-warp partials its multi-warp form declares. Pass the compiling context's Context.codegen_capabilities; the capability-free Backend_intf.no_codegen_capabilities counts the tiles alone. Forces no node's size_in_bytes (Tnode.size_in_bytes_peek).

The per-work-item scratch bytes the kernel will declare (gh-ocannl-1313): one function-scope array per node of its traced store that settles Local (Tnode.Placements.is_local_peek), sized at its storage precision, workgroup-shared tiles excepted. Settles no placement and forces no node's size_in_bytes (Tnode.size_in_bytes_peek). An upper estimate of what codegen declares — the CPU backends' per-chunk privatization and Metal's simdgroup fragments are rendering decisions that can drop arrays from the function scope — and of the device compiler's frame, which may drop an array it proves dead; checked by check_hardware_limits_classified against Backend_intf.hardware_limits' max_thread_private_bytes.

Validates the scheduled kernel name against the device limits, raising Utils.User_error on violation: the launch's workgroup size (the product of Low_level.launch_dims' block dimensions) against Backend_intf.hardware_limits.max_threads_per_workgroup, the workgroup_memory_bytes estimate against Backend_intf.hardware_limits.max_workgroup_memory_bytes, the thread_private_bytes estimate against Backend_intf.hardware_limits.max_thread_private_bytes, and then the launch geometry dimension by dimension through the shared launch_geometry_excess — the workgroup's .x/.y/.z extents against Backend_intf.hardware_limits.max_workgroup_dims (a separate hardware fact from the thread product: CUDA caps maxThreadsDim.z at 64 while the product cap is 1024 — gh-ocannl-679), and both 16-bit-capped grid dimensions, the .y extent (grid.(1), a blocktiled matmul's row-block count) and the folded .z extent (grid.(2), the product of the Grid slots >= 2 — gh-ocannl-643), against Backend_intf.hardware_limits.max_grid_yz. grid.(0) is deliberately ungated: it is 2^31-scale wherever hardware axes bind. Backend compile calls this after any ?lowered_transform (or the default annotator, which already respects the limits), turning driver-level launch failures into early, named errors. A no-op for all-None limits.

Test seam (gh-ocannl-1330): runs f with check_hardware_limits_classified judging by adjust capabilities instead of the capabilities it is handed, in the calling domain only, restoring the previous adjustment on return or exception. The default is the identity and no configuration selects another. What it is for: reaching a refusal downstream of the gate (the driver's own) with a kernel the exact gate would refuse first — test/operations/metal_threadgroup_overflow removes the per-allocation rounding. Not a production seam: a schedule the gate misjudges is a defect of the estimate.

Internal candidate-facing variant of check_hardware_limits; transports excess thread, workgroup-memory, per-work-item scratch (Schedule_outcome.resource.Thread_scratch at phase Hardware_limits, before any codegen — the post-link HIP budget check reports the same resource at Backend_link), per-workgroup-dimension, .y-grid or folded-.z-grid requests as typed Schedule_outcome.cause.Resource_exceeded causes, one variant per launch dimension (Schedule_outcome.resource.Workgroup_x_extent .. Schedule_outcome.resource.Grid_z_extent) so an autotune search's declines group by the dimension that asked too much.