Module Ir.Low_level

How a loop's iterations map to hardware; see docs/proposals/axis-types-for-loops.md. Serial is an ordinary for-loop. Grid / Workgroup bind the loop index to a GPU grid / workgroup (block, threadgroup) hardware index instead of looping; Workgroup_reduce is a Workgroup axis participating in a workgroup-cooperative reduction (see the contract below). Unrolled is emitted as the repeated body with substituted constants. Vectorized renders eligible bodies as explicit SIMD code — elementwise statements via vector extensions / packed loads (gh-ocannl-164 / gh-ocannl-463), a single recognized accumulation as independent accumulator chains with a horizontal reduce at exit on CPU backends (gh-ocannl-468) — and everything else as a serial loop annotated with the backend's vectorization pragmas when it provides them (a plain, un-annotated serial loop for accumulating bodies, whose loop-carried dependency the pragmas would deny) — like the hardware kinds, the annotating pass asserts iteration independence or, for a recognized accumulation, licenses reassociating it. Hardware slots are positional: among a kernel's loops of one kind, the innermost binds .x, then .y, .z. Annotated loops must have from_ = 0 and iterations with no cross-iteration dependencies (Vectorized accumulations again excepted). Workgroup_reduce is the labelled exception; its body must either stage its communication explicitly through workgroup-shared nodes and barriers (rendered by binding the index like Workgroup), or be a single accumulation statement acc = op(acc, contrib) over an associative-commutative op with the accumulator's indices free of the loop index — the renderer then owns the communication: warp/simdgroup shuffles on GPU backends (gh-ocannl-462), the plain serial loop on CPU backends. Like Vectorized, the annotation licenses reassociating the (floating-point) reduction.

Loop keyword used by the human-readable printers: plain "for" for Serial, "for@<axis>" otherwise.

Which pass minted a Local_scope (gh-ocannl-687). The construct has two producers, and a consumer that walks the IR looking for schedulable structure means only one of them:

• Inlined_computation -- virtualization's inline of a virtual node's computation at a read site, and the rewrites that carry those scopes along. The loops inside such a body are the inlined node's own iteration space, re-instantiated per use site; no Schedule op has ever targeted them.
• Schedule_minted -- the accumulator localization built by Schedule's materializing Unroll and by Partition (gh-ocannl-639), and by C_syntax.try_localize_serial_reduce: a running value for a MATERIALIZED cell, whose body holds the very per-step / per-segment loops Schedule.rewrite_loop retargets.

This is the fact Autotune.collect_loops needs: it enumerates loops inside Schedule_minted scopes only, so the action menu does not spend its per-unit budget proposing splits, swaps, unrolls and vectorize retypes for inlined interpolation and reduction loops that no schedule op has ever been able to reach.

Deliberately NOT the mechanism behind the scope-target contract input_scope_ids serves: that one asks whether a scope was in the program a given optimize call was HANDED, which is a per-call fact -- a virtualizer-minted scope handed back into a second optimize still carries Inlined_computation and still is not that call's to retract -- and hand-built IR has no honest way to spell "not mine".

The iteration order of a t.Scan_loop (gh-ocannl-696): Forward runs the index from from_ up to to_, Backward from to_ down to from_. Part of the construct from day one because every adjoint of a forward scan is a backward scan over the same range.

Cases: t -- code, scalar_t -- single number at some precision.

The argument precision is preserved in heterogeneous precision operation arguments, and is ignored (overridden) in homogeneous precision operations.

One loop-carried scalar of a t.Scan_loop: read as prev, written as next, rotated prev := next after each iteration, prev := init before the first. Both ids name the same virtual node, whose precision is the state's precision.

Ordered read-only Low_level traversal. Every policy decision is mandatory. Statements are visited before children and after their reads; Seq is left-to-right. Tile_mma descends through its equivalent fallback. A Prune skips children AND the after hook. Hooks may prune Tile_mma or Scan_loop when the consumer treats them as opaque.

The tensor node a statement writes itself (not through a sub-statement): Set, Set_dynamic, Set_from_vec, Zero_out, and a Tile_mma's accumulator. An exhaustive match: with Access_fold.fold's descent, the one place a new statement constructor must be classified for the write-sensitive walks (gh-ocannl-1451).

The one statement-position walk (gh-ocannl-1451): loops, scans, guards and sequences, in source order; never an expression (so never a Local_scope body or a scan's synthetic local assignments) and never a Tile_mma's fallback (the statement counts atomically as its write). Scan bodies are entered (gh-ocannl-1134); a scan binder is reported with axis = None. Dead loops are entered: the walk is structural.

The one read collector (gh-ocannl-1451): the tensor nodes one evaluation of llsc, the right-hand side of a setter of self, loads -- the transitive inline-fanin guard's contributors (gh-573), shared by trace_node_facts' per-setter setter_reads and the guard's replay of stored computations. self and merge-buffer reads are excluded, a gather counts its table, projection-discarded operands are not evaluated, a Where charges its wider arm (the first on a tie) -- or, with blends (a backend whose blends_value_selects fact holds, gh-ocannl-1775), the union of its arms when speculable_select admits it: the blend evaluates both. Admission is judged at the arms' own annotations and ahead of inlining, an over-approximation of the renderer's (which also needs a float compute precision, and on HIP no NaN-literal arm) that errs toward the union only. And Local_scope bodies -- hoisted, unconditional -- charge the statement whichever arm they sit in. Dead loops replay nothing; a Tile_mma is opaque.

evaluation_reads of a statement: one replay of a stored computation of self.

gh-ocannl-777: what one replay of a stored computation of self reads, for Context.observability. Only the statements the inliner's instantiation keeps are visited — loops (a dead one replays nothing), self's own setters, scope-local assignments — so the sibling setters a shared-loop body carries cannot contribute a read or a symbol.

gh-563: the one canonical rendering of lowered code, shared by both digest consumers — analysis_cache_stats' cache (keyed inside optimize) and Schedule_cache.canonicalize (schedule replay across sessions).

Recognizes And or eager Mul of two signed-index comparisons (or folded 0/1 constants). This identifies the Boolean product used by range guards for narrowing and partitioning; arbitrary numeric multiplication is opaque. It does not prove that index arithmetic is total: the simplifier separately checks the active domains before replacing And with Mul.

Generate loops that iterate only over the padding margins of a tensor. For dimensions with padding, generates separate loops for left margin, middle (recursing), and right margin, each binding its own fresh symbol (gh-ocannl-1384). The middle region continues recursing to find padding in other dimensions.

Whether the tree carries a read-modify-write accumulation: some Set (resp. Set_local) reads its own target — a loop-carried dependency through memory when the written cell does not vary with an enclosing loop. Conservative: Local_scope contents count as reading anything, and Tile_mma and (gh-466) Set_dynamic accumulate by construction. Used by the autotune menu and by codegen fallbacks that must not assert iteration independence (e.g. vectorization pragmas) over an accumulating body (gh-ocannl-468).

Whether the scalar reads the node anywhere (gh-ocannl-639) — certifies that an accumulation's contribution is free of the accumulator's node, which licenses holding the accumulator out of memory across the reduction.

The statement-body counterpart of scalar_touches_tn: any read or write of the node. Blind to a Staged_compilation callback's accesses: pair it with has_staged_compilation wherever a "no" has to mean "never accessed".

The nodes code_touches_tn reports, each once, in the order the code first touches them (gh-ocannl-1473). A function of the code alone, independent of node uids and of any hash-table traversal, so two fresh builds of one program list corresponding nodes at the same positions -- the order kernel parameters and local declarations are emitted in. Equally blind to Staged_compilation.

Whether the code holds a Staged_compilation anywhere, scalar positions included: a callback whose emitted code no analysis can inspect, so it may access any node.

The accumulation-update statement shape tn[idcs] = op(tn[idcs], contrib) (or its FMA form) over an associative-commutative op, with contrib free of tn; returns (op, contrib). The single source of truth shared by C_syntax's widened renderings and Schedule.Unroll ~materialize:true's scope-form unrolling (gh-ocannl-639), so "what counts as an accumulation" cannot drift between the schedule transform and the emission.

Retarget an accum_update_parts-shaped update's read of tn[idcs] to the scope local id. Raises on any other shape.

Whether the scalar reads the scope local id, descending into nested scope bodies — the scope-local counterpart of scalar_touches_tn. Distinguishes a scope-opening init (a Set_local whose value is free of the local) from a self-referential update.

The reduction operator of a scope-local update in either spelling: the plain accum_local_update_parts form, or virtualization's guarded-read form Where (index-only cond, update, Get_local id) (possibly nested per guard condition). The residency classifier's recognizer (gh-ocannl-663): the guarded form is a reduction for accumulator-width purposes but does NOT decompose into an unguarded (op, contrib), so it is deliberately not part of accum_local_update_parts or of scope_updates_reduce_op's hoist license.

Whether a value select — a Where, or a gated operator (Relu_gate, Satur01_gate) — may evaluate every one of its arms whatever cond says, so that a backend can render it as a branchless blend (gh-ocannl-1622; C_syntax's select_blend_syntax hook). Pass the arms at the select's result precision.

The invariant it rests on: lowering emits a read that can be out of range only under a guard of its own, and every such guard either compares embedded indices — gh-504's clamp conditions, virtualization's range and unit-solve conditions, Schedule's tile, fragment and pad masks — or protects a Get_dynamic (gh-343's gather guard, whose condition is a data value). So:

• cond must not mention an Embed_index (read syntactically: a guard is spliced in place, never routed through a Local_scope); a select over an index comparison stays short-circuiting even when it is a user's mask, which is safe and only forgoes the blend;
• no arm may contain a Get_dynamic, nor a Div / Mod that renders at an integer precision (which can trap) — judged at the precision C_syntax renders each operand at, not at the operand's own annotation, which a homogeneous operator overrides;
• a nested conditional operator is walked whole, arms included: evaluated eagerly, its own condition can hold where the enclosing one did not, so everything under it is reachable. (An index guard nested inside stays ?: and still protects its read: its condition says exactly whether the read is in range, whatever the enclosing select says.)

Local_scope bodies are hoisted ahead of the statement and run unconditionally already, so a scope reference is a plain local read. Admitting an instance evaluates an arm the classifiers (Ops.ternop_conditionality, Ops.binop_conditionality) call conditional: their floors stay floors, Affine's a_gated stays a conservative over-approximation, and Cost_model's upper walk, which charges both arms, stays an upper bound (flagged approximate: it reads no backend, and CUDA and Metal still short-circuit). The inline fan-in guard's read collector (evaluation_reads) charges a single arm only where the backend short-circuits: on a backend whose blends_value_selects fact holds it charges an admitted select's arms by their union (gh-ocannl-1775).

gh-ocannl-1785: whether the expression, evaluated at an INTEGER compute precision, is float-valued -- the one reading of Ops.unop_integer_valuation and its siblings over an expression tree. A real-valued operator is; an operator that follows its operands is when one of its value operands is; a projection is its selected operand; everything else (a leaf, a Local_scope read, a truth value, Mod, the heterogeneous and refused operators) is not. C_syntax renders by it -- computing a float-valued subexpression in float and converting it once, at the store or scope holding it or at a Mod operand -- and the simplifier keeps the scope boundaries and folds that conversion depends on. Meaningless at a float precision, where every value is a float.

Whether an RNG lane conversion occurs anywhere in the expression (not inside a Local_scope). Such a conversion picks its result type and the random bits it consumes from the precision it renders at, so the expression renders at its target's storage precision and an accumulation carrying it keeps a storage-width accumulator.

Whether code generation holds this update's accumulator at the backend's accumulator residency: an accum_update_parts accumulation, free of RNG conversions (mentions_rng_conversion), with routine logging off. Any other update — a non-reduction recurrence, a mixed-operator update, an RNG-bearing one — narrows to storage at every step in the serial rendering, and a Schedule.Privatize tile taking it over stays at storage precision too (gh-ocannl-1116).

For a base peel_accum_nest reached: whether code generation holds it at the accumulator residency — accum_update_widens for a raw update; for a scope-form base, routine logging off and no RNG conversion in its scope local's assignments.

The reduce-shaped update of a scope LOCAL, local = op(local, contrib) (or its FMA form) with contrib free of the local — subst_accum_read's output shape; returns (op, contrib). The SIMD reduction rendering uses it to fold vector chains into a widened accumulator's scope local (gh-ocannl-639), and peel_accum_nest's scope-form validation is built on it.

How a node's storage is shared across the threads a hardware binding creates: device-resident (every thread of the launch), workgroup-shared (the threads of one block), or thread-private (a per-thread local array, which no binding can race). The tags are spelled as the tile-MMA address space C_syntax.mma_space spells them, so the fragment-operand sites answer the same question by coercing this classifier rather than repeating it.

One coordinate of thread identity: the hardware dimension a bound loop occupies (hardware_axis_info's kind and slot).

The legality of a hardware binding, as a query on the written cell (gh-ocannl-959). active is the launch's thread identity, one coordinate per slot some register-bound loop of extent above one occupies; thread names every bound loop and its slot — the register-bound ones and cc's pool-parallel outermost Grid loops, a per-loop binding; deferred the bound loops whose separation another call judges (a reduce lane left to its own rendering: it still covers its slot). The first store in the kernel to storage the threads share that two threads may own, with the reason: an active slot with no bound loop enclosing the store (every coordinate of the slot executes it), or a cell whose index map does not Affine.separates the enclosing thread symbols over every loop symbol in scope — mentioning a bound axis is not separating it (acc[i + j] under two bound axes, acc[i + k] under a bound i and a serial k). Guards narrow the environment as gh-ocannl-566's simplifier narrows it: If (i < 16) shrinks the range the radix argument sees and If (i == 0) pins i to one thread along that axis alone. A Grid axis is not a thread of workgroup-shared storage (one block, one copy). What is judged is the store's own threads: whether other statements' threads touch the cell across a barrier is dependence analysis, the obligation of whoever mints the pin or the binding, as it was for every store before this rule. Reads play no part: a store every thread performs to one cell is a write-write race whatever the values. A dead level and a false guard execute nothing; a Set_dynamic is judged by its static slots, a Set_from_vec by its aligned run blocks, a Zero_out as a store of every cell, a Tile_mma through its fallback with its own lane excused (gh-ocannl-960). The whole kernel is walked. The renderer asks this at every binding it emits: for the kernel's Grid/Workgroup axes before rendering, and for a Workgroup_reduce lane the warp shuffle cannot own, before falling through to the plain binding.

Whether the tree holds a SELF-RECURRENCE: some Set whose value reads the very cell it writes (gh-ocannl-733). This is what makes the localizing peel a live question at a level, and it is the predicate the peel census gates on. Narrower than has_accumulation in two ways, both deliberate: the recurrence is on the CELL rather than on the node, and a Local_scope counts only when its body actually reads that cell — has_accumulation must count every scope conservatively, being the predicate that decides whether iteration independence may be asserted, and censusing on it recorded every non-reduction virtualized scope as a declined reduction site. A dynamic gather may recur unless unequal fixed indices outside its runtime axis prove the cells disjoint; symbolic slots remain potentially aliasing. The selector's reads count too.

What peel_accum_nest decided about one If it peeled through (gh-ocannl-733). Two shapes that render identically can earn different verdicts here, which is the whole point of reporting them: a test pinning the rendered FORM cannot tell "the guard was confined to the peeled levels" from "the guard mentioned an enclosing lane and the cell separated it".

Why peel_accum_nest stopped, when it did (gh-ocannl-733).

What one peel_accum_nest call decided, beyond whether it succeeded (gh-ocannl-733). Reported to the optional ~report callback on every call, success or refusal; C_syntax turns it into the per-routine peel census, so that a test can pin which DECISION produced a kernel and not only which form was rendered.

Peel a single-statement reduction nest down to its accumulation base (gh-ocannl-639). Levels are Serial/Unrolled/Vectorized loops, loops extra_level vouches for (codegen passes a predicate accepting hardware-annotated reduction loops its backend serializes — the schedule mints pass nothing, since wrapping a hardware-annotated loop in a scope at transform time would break the schedule on backends that do bind the hardware dimension), and pure-index-guarded Ifs (the gh-490 If (i < s) shape and its constant-bound sibling — data-dependent guards stay opaque), each containing nothing else; the base is a raw accum_update_parts-shaped update (`Update) or the scope form a previous rewrite minted (`Scope: the scope id and the update statements after the opening init, validated to carry ONE reduction operator — mixed-operator sequences are not reductions and keep their per-iteration narrowing), with the accumulated cell invariant across the peeled levels (free_of seeds the invariance check with the caller's own loop). rebuild re-wraps a replacement base statement in the peeled levels. The ONE definition shared by C_syntax's localizing serial fallback and the schedule mints, so transform and emission cannot drift.

A DEAD level (to_ < from_) is never peeled, and that refusal is load-bearing rather than tidiness: the body of a dead loop performs no accesses at all — the routine-interface walk propagates liveness as live && to_ >= from_, so a node reached only under one is absent from the parameters and need not be allocated — while every form this peel licenses reads and writes the accumulated cell OUTSIDE the levels, unconditionally. Peeling a dead level would invent accesses the program does not make, possibly on an identifier the interface never declared.

Whether a GUARD may join the peeled levels is Affine.peel_guard's and Affine.separates's answer, not this function's (gh-ocannl-722) — rebuild keeps the guard around the accumulating update only, so the localized form loads and stores outside it, and the engine is where the two hazards that creates are stated: a guard mentioning no peeled symbol is fixed for the whole nest, so the hoist invents both accesses, and one mentioning an enclosing loop symbol selects among that level's iterations, so lanes the guard excludes would write their unchanged local back over the accumulating lane's result. The second is admitted when the accumulated cell SEPARATES those enclosing symbols — each instance then owning a distinct cell makes the invented load/store pair private and idempotent (gh-ocannl-721) — which is why the cell reaches the decision and the peel defers it to the base.

loop_ranges is loop_ranges of the enclosing program, and supplies both the classification and the ranges: a guard symbol in it that is not peeled is an enclosing level's index, while one outside it is a static index parameter or a runtime extent and is harmless — which is what keeps gh-490's runtime-extent guard (Assignments.extent_guard's i < s, whose bound is a static symbol rather than a constant) peelable. Required rather than defaulted, and derived from the program rather than certified by the caller, so that no call site can forget it: the mints of Schedule need it as much as codegen does, since a refused mint there turns segment seams into narrowing points instead of merely declining an optimization.

report, when given, is called EXACTLY ONCE per call with the peel_report of what was decided — how many levels were peeled, which verdict each peeled guard earned, and the refusal where there was one (gh-ocannl-733). The result alone cannot answer that: two nests differing only in whether the accumulated cell mentions the enclosing index peel a different number of levels under a different guard verdict and render the same localized kernel. Codegen passes a reporter that accumulates the per-routine peel census; the schedule mints pass none.

Why decide_accum_width declined a base peel_accum_nest reached.

How the accumulator of one reduction level is held: decide_accum_width's answer.

The accumulator-width decision of the reduction level binding index over from_..to_ with body t — peel_accum_nest over the body, then the declines the serial rendering applies to the base it reaches (accum_skip). The ONE owner of that decision (gh-ocannl-1151): code generation asks it at every level that can localize (C_syntax.decide_accum_width, whose warp-shuffle, SIMD-grid and localizing-serial renderings all consume the answer, gh-ocannl-754), and Schedule.Privatize asks it of the program it was handed to choose its tile's width. extra_level is the serialized-slot predicate: a hardware-annotated loop the backend renders serially (no hardware index for its slot) is a peelable level like any other — code generation passes its own (C_syntax.serialized_hardware); loop_ranges is loop_ranges of the enclosing program, as for peel_accum_nest. Pure but for the routine-logging flag it reads.

Whether the decision holds the accumulator at the backend's residency: a base the level localizes (pinned = None) that accum_base_widens. A scope-form base whose scope local an RNG conversion reaches is localized at storage width — codegen's census types that scope at storage — so it localizes without widening.

The accumulating writes of tn in program order — a Set reading its own cell (raw or scope-form, has_accumulating_cell), a Set_dynamic reading its node — through the statement positions serial_accum_widths walks (loops, guards, sequences; never a Tile_mma or a scope body), so the two lists align entry by entry.

For each of accumulating_writes of tn, in the same order, what the program's SERIAL rendering does with its accumulator — decide_accum_width asked at exactly the levels C_syntax.pp_ll asks it at, top-down, with serialized as the serialized-slot predicate: Serial, Unrolled and Vectorized levels always, Workgroup / Workgroup_reduce levels where serialized says the backend renders them serially, Grid levels never. The first level whose decision localizes a base of tn answers for that base; a raw write no level localizes narrows per step, a scope-form one is typed at the residency by codegen's census all the same. Schedule.Privatize reads its tile width off this, asked of the program its schedule was applied to (gh-ocannl-1151).

gh-ocannl-998: the accumulation into d[di] in the fused form a tile multiply-accumulate computes. An FMA accumulating into d[di] is returned as is and a floating plain add of a product is returned as the FMA; None for an integer accumulation (gh-ocannl-824) or no product accumulated into d[di]. The sketch seeding asks it whether the register tile can render a site (both forms can be tiled), and fuse_tile_fallback applies it.

A Tile_mma fallback with every accumulation into d in its fused_accumulation form: what C_syntax renders a fallback as unless the backend's compile pins contraction off (C_syntax_config.compiler_contraction_off, gh-ocannl-1457). Schedule.Tensorize leaves the fallback in the form the contraction licence left it, so under a contraction-off compile the tile and its fallback compute the serial loop's two roundings.

A whole-node zero, or its unconditional covering positive-zero Expand_zero loop nest.

The zero-forwarding proof shared by code generation and Schedule.Privatize: one unconditional same-cell RMW pair covering the whole node. Returns its write access and every loop that repeats a cell; the accepting accumulator scope must contain all those loops. allow_workgroup_barriers is for Privatize, which retains the barriers while replacing the target accesses with private-tile accesses. Opaque effects and Tile_mma are always refused.

zero_initializer_target, also admitting the launch-extent guards of a hardware-mapped zero nest: a guard bounding an enclosing loop's own symbol trims that loop's range before the covering query (gh-ocannl-1257).

seed_first_write ~tn stmt: stmt, the first statement to touch tn after a covering zero of it, rewritten so that its accumulation into tn opens from zero itself (gh-ocannl-1257) — value-identical, given that zero, down to the accumulator's residency. A constant zero where no loop repeats a cell; the localized scope form where code generation would localize, opening at zero or at the cell guarded by the first iteration of the repeating loops outside the scope; otherwise the opening read guarded by the first iteration of the repeating (Serial/Unrolled) loops. A guard is placed only where the loops it ranges over iterate no more times than the node has cells: it costs an evaluation per iteration per cell to save one pass over the cells. None when zero_seed_candidate refuses stmt, when where code generation localizes depends on which hardware levels it serializes, when a needed guard does not pay, or when no form applies.

elide_covered_zeros segments, over a routine's kernel bodies in launch order: drops every covering zero (zero_fill_target) whose node's next access is a statement seed_first_write accepts, rewriting that statement (gh-ocannl-1257). Nothing opaque may lie between the two. None marks a segment left with no statement. Post-schedule: the rewritten accumulations are no longer in the shapes schedule transforms and sketch detection expect.

Arrays of length 3 (.x, .y, .z); all-1s for all-Serial code.

The launch slot space a loop annotation binds: Grid the blocks, Workgroup and Workgroup_reduce the threads of a block, the rest none (iterated inside one work item).

All hardware-annotated loops in pre-order, with their positional slots.

The first hardware-annotated (Grid, Workgroup, Workgroup_reduce) loop inside a Local_scope body anywhere in llc — a scope body is hoisted out of its position and is in no kernel's hardware-axis table — with the scope whose body holds it. The one predicate behind validate_parallel's rejection and Schedule's refusal of the op that puts such a loop there (gh-ocannl-1467).

The lane all-reduce sites of a kernel (gh-ocannl-1168), in pre-order: Workgroup_reduce loops whose body (past their own launch guard) is ONE reduce-shaped update of a scope local reassociable admits (Online_softmax.reassociable_local). C_syntax.try_lane_all_reduce renders a site across simdgroups only where lar_cross_simdgroup holds, and Schedule.workgroup_memory_bytes counts the partials of exactly those sites, so the estimate and the kernel read one predicate.

The two _partials forms; the other two insert warp shuffles only.

gh-ocannl-1238 / gh-ocannl-1334: the statements of a kernel whose rendering emits a lane collective the IR does not contain -- a workgroup barrier, or warp shuffles -- on a backend whose warp shuffles span warp_size lanes ([] when warp_size <= 0: the C backends render neither). The ONE notion of "this statement renders a lane collective" shared by validate_parallel, which refuses a non-uniform entry before anything renders, and by the renderers: C_syntax.try_warp_reduce renders either form only for a listed, uniform entry, and C_syntax.try_lane_all_reduce renders its butterfly only for a listed entry and its partials by lane_all_reduce_sites, derived from the same walk.

Lane-form entries are listed where the admission holds: the butterfly at exactly one warp under warp-uniform control, the partials where lar_cross_simdgroup holds past one simdgroup. Both admissions ask for uniformity themselves, so the form declines to the serial loop under divergent control and the validator refuses none of its entries. Cell-form entries are every Workgroup_reduce loop whose body (past its own launch guard) is ONE accumulation statement into a cell invariant in the loop -- C_syntax's shared accumulator decision at a single statement, asked of the same peel_accum_nest -- uniform or not: that form has neither a barrier-free nor a shuffle-free rendering. Past one warp they are Warp_reduce_partials, within one Warp_reduce_shuffle. An entry the renderer refuses for other reasons (partial warps, a narrow residency, ...) is listed too, so the list over-approximates only by statements that never render.

gh-ocannl-1452: the cell-form statements of inserted_collectives, in pre-order and at every extent, with the node each accumulates into: the Workgroup_reduce loops C_syntax.try_warp_reduce renders as a warp-shuffle tree (a duplicated copy is its own site). Past one warp each declares its per-warp partials in workgroup-shared memory, at the accumulator residency of the cell's storage precision; Schedule.workgroup_memory_bytes charges each through the backend's warp_reduce_workgroup_bytes, which is 0 within one warp and wherever the renderer refuses the site.

Per-slot maximum extents over the kernel's annotated loops; grid.(2) is the product of the per-slot maxima of grid slots >= 2 (the .z fold, see the section comment).

The binding arithmetic of a Grid loop at slot >= 2 under the .z fold: (stride, cap) such that the loop binds (z / stride) % cap — stride the product of the per-slot maxima of grid slots in [2, slot), cap = None (no modulo) when no higher grid slot exists in the kernel. (1, None) — the bare .z register — for the common single-slot-2 case.

gh-ocannl-584: the predicate form of validate_scope_bodies — None when every Local_scope body in llc is pure, else a description of the first violation. hoist_cross_statement_cse uses it to guard its own precondition (it is the one pass that can move an effect out of a scope, so an impure body reaching it would be laundered past the codegen gate); tests use it to assert that a body was left where it was.

gh-ocannl-584: enforces the scope-purity contract stated at scalar_t.Local_scope — a scope body's only effect is on the locals it owns (its own scope id, plus ids Declare_locald lexically within it). Raises Invalid_argument on anything else inside a Local_scope body at any nesting depth: a tensor-node write (Set, Set_from_vec, Set_dynamic, Zero_out, Tile_mma), a Set_local of a sibling or enclosing scope's local, a Workgroup_barrier, or a Staged_compilation. Applied at both ends of the pipeline — optimize on the way in (before a pass can launder a violation out of any body) and C_syntax.compile_proc on the way out. The raw analysis entry points analyze_proc / specialize_proc deliberately do not validate, being the probes that must stay conservative on IR they may not trust. The optimization pipeline satisfies the contract by construction.

gh-ocannl-696: the well-formedness contract of t.Scan_loop -- one state node DECLARED virtual per carried pair, ids pairwise distinct across the list, inits free of carried state and of the scan's own index, each next written exactly once as a top-level body statement and read only by later statements, no write of a prev, no binder or reference of a carried id beyond its scan, no opaque callback inside, and a non-empty range. Raises Invalid_argument naming the scan and the clause. Like validate_scope_bodies it runs at both ends of the pipeline: optimize on the way in, backend codegen on the way out.

gh-ocannl-1402: the loop-binder lint. Each loop symbol is bound by at most one loop (For_loop or Scan_loop) in the whole routine -- sibling loops included, not only nested ones -- and never by a loop when it is one of the routine's static indices; every symbol an index mentions (Iterator, Affine, Concat, Embed_index, a Local_scope's call-site indices, a Tile_mma's lane) is bound by an enclosing loop or is a static index. Raises Invalid_argument naming the symbol and the loops involved. Every flat symbol-keyed consumer -- extent tables, Schedule.statement_mappings, the canonical render behind both digest caches -- relies on this, and misreads a violation silently: concat segments (gh-ocannl-765) and padding-margin strips (gh-ocannl-1384) each broke it once. Applied by optimize on the way in, ahead of the analysis cache, so lowering's generators are held to it; hand-built IR must mint a fresh symbol per loop too. Entry-only by design, unlike validate_scope_bodies: past optimize, schedule rewrites emit sibling nests that share symbols (the staging transfer nests), and the renderer and the post-optimize analyses accept that shape.

gh-ocannl-1485: the predicate form of validate_no_slice_aliases — None when no storage access in the code names a slice-alias view (Tnode.is_alias), else a description of the first one, naming the construct, the view and its parent.

gh-ocannl-1485: a slice-alias view owns no buffer, so every access of one must be rewritten to its parent (Assignments.resolve_alias); a path that forgets the rewrite reads or writes nothing, silently (gh-ocannl-1475). Raises Invalid_argument, prefixed with site, when any construct that addresses a node's buffer names a view: Zero_out, Set, Set_dynamic, Set_from_vec, Get, Get_dynamic, Get_merge_buffer (a view has no buffer to transfer) or any of a Tile_mma's three nodes, at any depth. Scope ids name registers and are not checked. Applied at three points: per assignment in Assignments.to_low_level (naming the assignment whose lowering produced the access), by optimize on the way in (covering the rewrite tier and hand-built IR), and in C_syntax.compile_proc on the way out (covering schedule transforms and a ?prelowered routine that skipped optimize).

gh-ocannl-1709: the predicate form of validate_operand_precisions — None when every homogeneous operand slot's annotation is descriptive, else a description of the first one that is not.

gh-ocannl-1709: a homogeneous operator (every binop but Uint4x32_to_prec_uniform_lane, every unop but Uint4x32_to_prec_uniform1, FMA, Mul3, and Where's two arms) evaluates its operands at its result precision — the destination node's for a Set, the local's for a Set_local, a Where condition's or a guard's own annotation inside one — so the precision a scalar_arg carries there is never semantics. It may still describe: the evaluation precision itself, or the operand's own precision (scalar_precision), which is what apply_op and the assignments lowering write (an f32 sum of bf16 loads annotates its operands bf16). Any other annotation reads as a request to compute at that precision which the renderer would silently drop, so this raises Invalid_argument, prefixed with site, naming the destination node, the annotation, the evaluation and the operand's own precision, and the remedy (to compute at the annotated precision, store to a node or local of that precision). Heterogeneous slots (a Where condition, a guard, a Get_dynamic index, a Set_from_vec argument, the heterogeneous operators' operands) keep their annotation and are not checked against anything. Applied at two points: by optimize on the way in (covering the assignments lowering, the rewrite tier and hand-built IR routed through optimize), and in C_syntax.compile_proc on the way out (gh-ocannl-1739: covering the optimizer's own output, schedule transforms and a ?prelowered routine that skipped optimize).

gh-ocannl-1739: rewrites each homogeneous-slot annotation validate_operand_precisions would refuse to the slot's evaluation precision, operands before the operators over them (an operator's own precision is read off its first operand's annotation); every other annotation is left as it is. The renderer never reads such an annotation, so this changes no emitted value. Sound only over code whose builder's annotations already passed the check, so that nothing it rewrites is a request: the optimizer applies it to its output (its input passed the entry gate), and Schedule.apply to a schedule's result when the program the schedule started from passes the check -- a ?prelowered routine carrying a request is left for the exit gate to refuse.

Backend-independent well-formedness of hardware annotations (axis-types proposal §2); a no-op for all-Serial code. Raises Invalid_argument on structural violations: nonzero from_, more than 3 slots per kind, annotated loops inside Local_scope bodies, barriers under divergent extents or If guards, writes to materialized nodes not nested under annotated loops covering every active (non-unit) hardware dimension — launch dimensions are global to the kernel, so an uncovered dimension executes the write once per hardware index — and whole-node Zero_out of materialized nodes in multi-threaded kernels (nesting never distributes it). With warp_size (and the lane all-reduce's reassociable, Online_softmax.reassociable_local) it also refuses a lane collective the renderer inserts (inserted_collectives) under control flow that is not uniform across the lanes it spans (gh-ocannl-1238 for barriers, gh-ocannl-1334 for warp shuffles); the default warp_size of 0 is a backend rendering no warp shuffles. fork_join (default false, hardware-thread semantics) states the C backends' execution model (gh-ocannl-1458): each pool-parallel Grid loop is one synchronous dispatch joined before the next statement, and code outside every Grid loop runs once, on the calling thread — so a materialized write or whole-node Zero_out nested under no Grid loop is exempt from the coverage rules, while writes under a Grid loop keep them. Does not prove iteration independence: an annotating pass proves it for the annotations it mints, and the renderer asks unseparated_thread_write of every binding it actually emits.

Internal backend-facing variant of validate_parallel; transports a validation Invalid_argument as a typed Schedule_outcome.cause.Illegal_schedule. warp_size and reassociable are required here: the callers are the ones about to render, or to rank candidates for a backend, and know its shuffle width.

Wraps bodies of annotated loops whose extent is below their slot's launch dimension in If (index < extent) guards, for the kinds the backend binds in hardware.

gh-ocannl-616: the namespace of the scratch nodes footprint-scoped materialization mints — how a test tells them from the program's own nodes in an optimized record's traced store.

One setter's fan-in contributors under both renderings of a value select (gh-ocannl-1775): the analysis is cached across backends, so it keeps both and the guard picks by the backend's blends_value_selects fact.

A private copy of a traced store: the table AND its traced_array records (whose fields are mutable), so writes through the copy never reach the original or a sibling copy. Hashtbl.copy alone shares the records.

A shallow-copy fork of the lineage state (computations and placements tables): the copy sees everything decided so far; its later mutations are invisible to the original and to sibling copies. Backend compile forks the incoming context's optimize_ctx through this, so sibling candidate compiles from one frontier are hermetic.

Whether the tag is one of decide_placements' three heuristic caps: a flippable policy prior rather than a legality or observability verdict.

The configuration key whose cap forced the decision, when a cap did — the setting to raise to try to keep the node virtual instead of materializing it.

Records the Inline preference (gh-555) for the nodes, withdrawing any Footprint preference they carried: the two are one node's mutually exclusive readings, so the later request wins — a search trying a node's sibling flips one after the other must not accumulate them (gh-ocannl-616).

The Footprint counterpart of prefer_inline.

Granularity of the XOR remap applied to a swizzled node's minor axis (gh-ocannl-481 item 3, D1). Both flavors are per-row bijections of the minor axis, so the IR-level semantics are identical; they differ in the unit the XOR permutes and therefore in which access pattern they de-conflict.

gh-ocannl-616/1017: the three placements of a node in the inlining lattice — materialized (a full-domain buffer), inlined (recomputed at every read) and footprint-scoped (a routine-private scratch over the readers' footprint, filled right after the producer's last write). A flip moves a node from its default reading to another one.

The reading's lowercase name — "materialize", "inline" or "footprint" — for logs, arm labels and test output.

Splits decided readings into the (materialized, inline, footprint) lists that Context.lowered_for_decisions takes, keeping the input order within each list.

One reading a search can flip a flip_candidate's node to.

gh-555, one record per node since gh-ocannl-1017: one searchable inlining decision dimension of a compile — a node whose placement the default policy decided, with the reading it decided and the readings a search can flip it to. The alternatives are the node's mutually exclusive readings, so a consumer decides a candidate as a whole — keep the default, or take exactly one alternative — and a top-N cut counts candidates. The default is a field rather than an inference from which alternatives exist: an inherited footprint-scoped node and a cap-materialized node with no smaller footprint both offer `Inline alone, from opposite defaults.

Applying a flip: `Materialize via Context.decide_materialized, `Inline via Context.decide_inline, `Footprint via Context.decide_footprint. An `Inline or `Footprint flip's legality is settled only when the virtualizer replays: a rejected flip reproduces the materialized placement.

gh-ocannl-637, gh-ocannl-1011: the seam through which the cost model prices flip_candidates' alternatives — given the routine's static indices and, per node, the WORLD a read of it is instantiated in (placements and the node's computations), the exact op count of ONE read of a candidate (None when the model's count is only a bound, or the inliner cannot serve the read). specialize_proc supplies the world: a node's stored computations under the placements as the routine's walk left them (before cleanup commits its own reading), or, for a node a heuristic cap materialized before the walk, the computations the walk itself stores when it is not materialized — the walk re-run over the node's raw setter statements in a scratch copy of the lineage — with the placements that re-run left. Every reading prices that one instantiation, replayed by the inliner at a synthetic read (instantiate_at_synthetic_read), so the guards it emits, its binding choices and the pipeline's rewrites are in the priced code. Exact for the most general in-bounds reader; what a particular reader folds away (a constant substituted for an index, a sub-image collapsing a loop) or shares (sibling readers whose instantiations hoist_cross_statement_cse merges) only lowers what executes, so the product with the per-cell read multiplicity is a bound in the same sense the traced proxy's is. Cost_model registers it at module initialization; the default prices nothing, so every candidate carries the traced proxy. A pricer must be pure and consume only the candidate node's world. specialize_proc resolves and validates each node's private world once per compile, also deciding fa_refused. Its scalar price is memoized across specializations sharing an analysis, keyed by the resolved computations, effective placements and node facts, and current static domains. Replacing the registered function invalidates its prices. The memo retains no world tables.

gh-ocannl-1011: the passes virtualized code receives in specialize_proc, in order — cleanup_virtual_llc (placement commitments; a scope over a node committed non-virtual after the scope was minted retracts to a read), simplify_llc under the static indices' interval environment, the hosted-constant link-time conversion, rewrite_one_hot_reductions, eliminate_common_subexpressions and hoist_cross_statement_cse. One function, so the routine and the flip-candidate pricer's synthetic instantiation cannot receive different pipelines. Commits placements in the given table.

gh-ocannl-1011: the code ONE read of the node executes — the given computations (a stored template list, newest first) instantiated by inline_computation's own core at a synthetic read site, wrapped as the scope a reader's statement carries (Set_local of it), and put through post_virtualization_pipeline, all on a copy of placements in which the node is virtual. The read site is the most general in-bounds reader: a fresh symbol per axis, ranging over the axis, except where every value-carrying computation writes the same position free of non-static symbols (a fixed slice, a static index — the only position such a component can serve). Side-effect free: nothing is committed, and the symbols and scope ids it mints are negative, never from the counters generated code is numbered by. Error carries the inliner's rejection code (or the pipeline's message) where the read cannot be served.

gh-487: a software-pipelined (double-buffered) staged tile — codegen allocates pt_depth rotating copies of the tile and renders every access with a buffer-selection term rotated by the pt_rotor loop counter: reads select copy rotor mod depth, writes copy (rotor + 1) mod depth (the schedule emits the loads one iteration ahead), and writes outside the rotor loop (the prologue load) select copy 0. The IR keeps the tile's single-copy dims and indices — the rotation is a physical-layout choice like swizzle_kind, invisible to IR-level semantics — so the pipelined rendering is bitwise identical to the unpipelined one.

gh-ocannl-498: the backend-agnostic byte footprint implied by an optimized routine's placement vector. The allocator-side scorer lives in Backends.score_footprint.

blends_value_selects is the rendering backend's fact of that name (Backend_intf.backend_facts; Backends.lower_assignments passes it): where true, the inline fan-in guard charges a select speculable_select admits by the union of its arms, which the backend's blend evaluates (gh-ocannl-1775). Default false, the short-circuiting reading: hand-built IR optimized without a backend in mind.

Decision-independent analysis of a lowered routine (gh-555 step 1): the structural per-node facts and the lazily-materialized affine access metrics — everything the optimization pipeline consumes that does not depend on the lineage's placement decisions, plus a bounded memo of scalar recompute prices keyed by validated placement worlds (gh-ocannl-1113).

Compute the analysis once for a routine. optimize is analyze_proc followed by specialize_proc (plus the pretty-printing callbacks).

The decision-dependent tail of the pipeline: placement decisions (decide_placements under the given lineage's placements and inline preferences, and the backend's select rendering as in optimize: the analysis keeps both charges, so one analysis serves either), virtualization, cleanup, simplification and CSE. Cheap to replay per candidate over one shared analysis (gh-555): sibling calls with hermetic optimize_ctx forks (see copy_optimize_ctx) produce hermetic optimized results — the traced store is record-copied per call.

The gh-ocannl-805 seam assertion run after virtualization and before cleanup: every tensor-buffer read (Get / Get_dynamic) in a statement cleanup will keep, including reads inside a kept inlined Local_scope body, must name a node with a placement decision. Reads in a virtual candidate's setter are excluded because cleanup drops the whole setter without visiting its right-hand side. Exposed so hand-built-IR tests can inject the otherwise-unreachable negative control directly at this boundary.

A structural digest of llc that two lowerings of one routine agree on (gh-ocannl-1526): tensor nodes by identity (uid), or by tn_token when given (to rename the nodes a compile mints); loop binders, local-scope ids and free symbols -- each lowering mints its own -- alpha-renamed by first occurrence, syms pre-registered in order (to name the free index symbols of a stored computation's index tuple). Comments are skipped.

gh-560: (hits, misses) of the process-global analysis cache consulted by optimize: sibling candidate compiles of one routine share one analyze_proc result — keyed by a canonical digest of the raw lowered code and the static indices (tensor nodes and static symbols by identity, loop binders and local-scope ids alpha-renamed) — and replay only specialize_proc. Cumulative counters, for tests and diagnostics.

Drops the analysis cache's entries (the stats persist). Entries retain their routines' lowered code, hence their tensor nodes; the cache clears itself before accessibility snapshots (Tnode.print_accessible_headers) and callers that tear down a session (Tensor.unsafe_reinitialize) clear it to release the nodes promptly. Never needed for correctness: entries keyed by stale nodes cannot alias fresh ones (Tnode.uid is never reused).

reads_scope_before_set id body returns true if id is read (via Get_local) before the first definitely-executed Set_local id in body. Use this at code-generation time to decide whether a Local_scope or Declare_local declaration needs a zero initializer.

The licences of the rounding-changing algebra over the reals (gh-ocannl-998), one field per licence, each true when that family of rewrites may run:

One step of the licensed algebra at the root of the expression: the first licensed rewrite that matches (first-match order matters), not yet re-simplified, or None. A ToPowOf is expected with its operands already simplified.

Top-down algebraic simplification with interval-driven comparison folding (in particular, it erases If guards whose conditions the loop extents prove). The interval environment is narrowed by every enclosing If condition that is a conjunction of integer-affine index comparisons (gh-ocannl-566), so a guard the statement guard proves folds too — what is simplified under a condition is valid only where that condition holds. Called internally by optimize; exposed for Schedule.apply, whose transforms construct guards after the pipeline's simplify already ran (docs/proposals/schedule-ir-optops.md §2), and for testing. Pure and idempotent.

Exact: every rewrite preserves IEEE results (the x * 1 and x - (+0) identities, constant folding, guard folds); integer arithmetic is reassociated as fp_algebra_step does with no licence. A single-assignment Local_scope is collapsed into its reader only where its rounding to its own precision is not part of the value (gh-ocannl-1788): the slot evaluates at the scope's precision, or the value is a load whose storage precision the scope's widens exactly (Ops.widens_exactly). A leaf it does not rewrite keeps its annotation, so a ?prelowered routine's request on a load still reaches validate_operand_precisions at codegen. The licensed floating-point algebra runs once, on the raw lowering, in the rewrite tier member Fp_algebra (gh-ocannl-998), so neither this pass nor the schedule transforms that rerun it can contract or reassociate across an inlining boundary.

Runs the thunk with interval folding (simplify_llc and the transforms' trailing simplification) blind to tensor nodes' bounds candidates, on the calling domain: a fold then rests on precisions, loop ranges and static indices alone, so the result is a function of the code (gh-ocannl-1480). A bounds candidate is process state that grows as routines link, and the autotuner's replay of a saved schedule must rebuild the very code it was saved against. Less folding is always sound: a guard left in place is only checked at run time.

gh-343: rewrites the narrow one-hot embedding pattern -- an Add reduction over a loop variable k whose body selects an embedding-table row via k == index_expr (a logical one-hot) -- into a guarded dynamic gather (Get_dynamic) that reads the table row at index_expr directly, with an in-range guard returning 0 out of [0, vocab_size) to preserve the one-hot semantics. The guard is constructed generically and interval analysis (docs/proposals/interval-analysis-scalar-t.md) erases the conjuncts it can prove -- from the index precision's machine range, loop extents seeded from static_indices, and settled per-tensor bounds (Tnode.bounds_state).

gh-466: also rewrites the transposed one-hot pattern -- the embedding-table gradient for k in [0, V): tn[.., k, ..] += (k == index_expr) * g where the loop variable indexes the written tensor itself -- into a guarded dynamic scatter-accumulate (Set_dynamic): if in_range(index_expr): tn[.., index_expr, ..] += g, dropping the O(V) per-position work (llm.c's deterministic encoder backward, docs/research/llmc-lessons.md B5). The enclosing position loops keep their original serial order and the schedule analyses never parallelize over a dynamically-written node, preserving determinism without atomics.

Unmatched or unsupported reductions are left unchanged. Called internally by optimize between simplify_llc and eliminate_common_subexpressions; exposed for testing.

Whether a scalar reads the node, descending into scope bodies and dynamic index values.

Eliminates common subexpressions within each statement's scalar expression tree. Replaces duplicate Local_scope nodes (structurally identical modulo scope_id) with Get_local references to the first occurrence. Called internally by optimize; exposed for testing.

Hoists shared Local_scope computations from sibling statements to the enclosing scope. When two or more sibling statements share an alpha-equivalent Local_scope node, the computation is extracted as a Declare_local + body preceding the first user, and all occurrences are replaced with Get_local.

Inputs are the materialized read-only and read-before-write (within the code) non-constant non-merge nodes. They are inputs in a broad sense, as they could be recurrent nodes or parameters. Outputs are all the materialized nodes written-to by the code. The last returned component is the input merge node, if used in the code. Reads entering the code only through a cross-routine inlined computation count: the traced store is completed from the final optimized code (gh-610).

Every For_loop and Scan_loop binding within the code, ONE inclusive range per symbol: the hull of all its bindings. Loop symbols are not unique within a routine: validate_loop_binders holds only at optimize's entry, and schedule rewrites emit sibling nests sharing a symbol by design (staging transfer nests, a materializing unroll's copies), with possibly different bounds. A per-binding list would let a first-binding lookup drop the siblings' values, the hole gh-ocannl-1282 went through (gh-ocannl-1305). The box environment for Affine queries, and what peel_accum_nest classifies its guard symbols against: one absent from it is bound outside every loop (a static index parameter, a runtime extent) and cannot select among any level's iterations.

The value-dependence symbols of statement-level scalar scope-locals, whole-code: per scope id, the union of the symbols its statement-level assignments depend on, transitively through Get_local references (such a local may be assigned in one statement and read in another). Assignments inside Local_scope bodies are not recorded — a scope id is re-instantiated at every use site with per-site loop symbols, and scope-internal flow is covered lexically by the value scans. Consumed by the setter value scans so a value routed through a scope-local is not laundered of its symbols (gh-494 per-thread value-variance).

Loop symbols a scalar expression's value depends on syntactically — index symbols of reads, embedded indices, dynamic-index sub-expressions — resolving scope-locals through locals (from scope_value_syms).

gh-494 waypoint 1: the routine's tensor-node accesses as explicit affine relations (Affine.access), extracted from (typically optimized) code, in program order (a statement's right-hand-side reads precede its write; Local_scope bodies are descended into at their use site; Tile_mma is traversed through its scalar fallback); read accesses carry their scalar gatedness (a_gated). Beside them, from the same walk (gh-ocannl-1016), the Affine.statement_effect rows: everything the code does that no tensor-node access carries — scope-local writes, reads and declarations (a local named by its node and scope id, so a hazard query can match a read to a write), Local_scope bodies, barriers, opaque Staged_compilation (whose accesses are not enumerated: callers needing exhaustiveness check for its row), Tile_mma as a construct, and merge-buffer reads. A "between two points, what else runs?" question is a query over the pair — never a fresh walk over the raw code.

The access half of affine_relations.

The effect half of affine_relations.

The read-before-write fact behind the placement decisions and the routine interface: whether every read of a node in accesses (typically affine_accesses) is covered by prior writes (`Covered), covered only thanks to the read-modify-write exemption — a read at its enclosing statement's own write position (`Covered_rmw_exempt) — or not covered (`Unknown, with a witness). ?write_eligible restricts which writes may cover which reads. Exposed so hand-built IR tests can assert the verdict itself where the pipeline's later passes would mask a wrong one (gh-ocannl-1244).

gh-ocannl-489 liveness-based buffer aliasing: per-tnode access span over the final (post-schedule, post-fission) code of a routine, as a closed interval of positions; the input is the routine's kernels in execution order (a singleton when not fissioned). With stmt_serial:true every top-level statement gets its own position — sound only for backends where consecutive top-level statements of one compiled procedure are fully synchronized (the C backends); with stmt_serial:false all statements of a segment share one position, since GPU kernels have no grid-wide synchronization between top-level statements. Returns None when the code contains Staged_compilation (opaque accesses: no aliasing plan can be trusted).

gh-ocannl-489 follow-up: sinks each top-level Zero_out to just before the first later top-level statement accessing the zeroed node (Train.grad_update's up-front zero_grads block otherwise starts every gradient's live span at that block, nesting the backprop chain's intervals and defeating the arena planner). Sound: a Zero_out commutes with statements not accessing the node; it never crosses such an access, a Staged_compilation, or a Workgroup_barrier. Apply to whole-routine code BEFORE scheduling/fission.

The code-name minter for the nodes llcs mention. A node whose label ident is in blacklist, starts with one of reserved_prefixes, or satisfies reserved (a namespace given by rule, such as the implementation's __x / _X, gh-ocannl-1338), is treated as a repeating ident: it takes the disambiguated n<id>_<label> form rather than the bare label, so it cannot equal a name the emitter declares. A prefix reserves a namespace whose members are minted on the fly -- the C-family emitters' ocannl_vec<lanes><suffix> typedefs, one per lane count (gh-ocannl-1102) -- which no finite blacklist can list.

Adheres more to the C syntax, outputs implicit type casts.

Adheres to the %cd syntax.

One of the accumulator transfers Schedule.contract_tensorized_accumulator synthesizes around a persistent fragment's reduction — lane { if lane == 0 { serial copy loops { fragment <- target } } } for into_fragment, the store back otherwise — as (lane, lane width, fragment, target, target indices, copy symbols innermost-first). The renderer's fragment-scope recognition and the cost model's operand-load count (gh-ocannl-1254) both read regions through it.

A pad-masked contraction's guarded transfers (the init's Where, the store-back's Ifs) are not transfers, by design: the intrinsic loads and stores move whole tiles, which a padded target does not have. A staged GPU pipeline's masked loop-site contraction instead stages them through a padded scratch tile and wraps the reduction in a plain region whose target is that tile (gh-ocannl-1736, Schedule.contract_tensorized_accumulator); elsewhere the masked tile is never marked as a simdgroup fragment. A marked fragment's region that this function does not recognize is refused by Schedule.fragment_scope_violations.

The nodes llc writes under a lane-selecting guard (if lane == 0 of an enclosing Workgroup loop, the form of fragment_transfer's transfers): one hardware thread writing on behalf of its workgroup, whose cells other threads of the workgroup read after a barrier. Such a node carries values between threads within one kernel, which a no-alias kernel parameter rules out (gh-ocannl-1598): C_syntax leaves it unqualified. In first-write order, deduplicated.

The transfer target's leading-dimension stride in elements: the stride of the axis carrying the row copy symbol (the outermost of fragment_transfer's copy symbols), the minor dim without one.

The persistent-fragment region Schedule.contract_tensorized_accumulator synthesizes, recognized at the head of a statement sequence: init transfer, reduction, store transfer agreeing on fragment, target, lane, width and indices, with exactly one Tile_mma into the fragment. The renderer's recognition (C_syntax, which also requires the fragment to be marked in optimized.simdgroup_fragments) and the cost model's operand-load count (gh-ocannl-1254) share it.

A tile-MMA operand's layout as the emission hooks receive it: `Plain unswizzled, `Swizzled_b128 for a b128-swizzled rank-2 tile addressed from its origin with its minor dim as leading dimension; `Decline (with the reason) for any other swizzle, which renders the scalar fallback without reaching a hook. Shared by the renderer and the cost model (gh-ocannl-1254).