Module Ll_builders

node_factory ~first_id ~dims () returns a maker of fresh single-precision tensor nodes with consecutive ids above first_id and default dimensions dims (overridable per node). Each test executable picks an id range of its own, so nodes stay distinguishable in debug output.

Declares tn materialized and observable: the executed legs seed and read back exactly these nodes, and observability is also what forbids the buffer-aliasing planner from handing their bytes to another node. Both are declared intent, settled before optimization, so neither perturbs a structural pin — see virtualize for what "declared intent" reaches.

Declares tn virtual — the standing of the scope-local scalars a virtualizer-emitted Local_scope owns.

This and materialize write the tnode's DECLARED INTENT (Tn.update_memory_mode, the memory_mode_intent field), not a lineage decision. Placement decisions live on the optimize_ctx's placements table, and Tn.Placements.get falls back to the declared intent for a node the lineage has not decided — which is the whole reason a test can hand optimize a node that is ALREADY virtual (or already materialized) before the analyses run, and the reason the passes read it back as such.

aff terms offset is the affine index sum (coeff * symbol) + offset.

if_ cond body guards body on cond being nonzero (Ir.Low_level.t.If). The condition is read at index precision only when it is an index expression; a value read (the usual flag tensor) keeps the node's precision, which is what single is here.

loop ~upto s body iterates s over 0 .. upto INCLUSIVE, mirroring Ir.Low_level.t.For_loop's own bounds; upto < 0 is a dead loop, which is a case worth building. ~axis declares the loop's hardware axis (Ir.Low_level.axis_type.Serial by default): the tests that judge a binding — a Grid block loop, a Workgroup lane, a Workgroup_reduce accumulation — name it here rather than spelling the record.

loop_n s n body iterates s over a range of WIDTH n, i.e. 0 .. n-1.

set_at tn idx llsc writes the single-axis cell idx — set over a one-element index array, which is the shape of every hand-built one-dimensional case.

carry ~init tn is one carried scalar over the state node tn, starting at init (a scalar that may read tensor nodes but no carried state).

prev cr reads the carried scalar's value from the previous iteration (its init on the first).

next cr reads the value the CURRENT iteration already assigned with set_next: the rotation is phi-style, so old and new values coexist inside one body.

set_next cr v assigns the carried scalar's next value -- exactly once per carried scalar, as a top-level statement of the body, which is the contract Ir.Low_level.validate_scan_loops enforces.

scan ~upto s ~carried body iterates s over from_ .. upto INCLUSIVE like loop, carrying carried across iterations; ~direction:Backward counts down instead.

gather ~tn ~idcs ~dyn_axis ~dyn_value reads tn at idcs with the dyn_axis row taken from the runtime value dyn_value (an index-valued scalar paired with the precision it is read at — iprec for an index computation, the node's own precision for a row number stored in a tensor). Counts as a read of tn, like get.

scatter ~tn ~idcs ~dyn_axis ~dyn_value llsc writes llsc into that same cell: set with the dyn_axis row supplied at runtime. Loops whose index reaches dyn_value carry a cross-iteration write dependency, so schedule analyses must treat the write as statically unknown — which is much of what makes this shape worth building by hand.

scatter_add ~tn ~idcs ~dyn_axis ~dyn_value addend is the accumulating form tn[.., dyn_value, ..] += addend — the shape rewrite_one_hot_reductions actually mints for the embedding-table gradient. The read-back is an explicit gather of the written cell at the node's storage precision, which is what makes the accumulation visible to read-tracking and to has_accumulation; addend carries its OWN precision, as the matched gradient argument does there — a mixed-precision accumulation (an f32 gradient into a bf16 table) is a shape worth building, and relabelling the addend with the target's precision would build different IR from the one the pipeline mints.

Cooperative tile multiply-accumulate

Ir.Low_level.t.Tile_mma: d[i,j] += Σ_{l<k} a[i,l] * b[l,j] over a block of the declared extents, executed jointly by the threads of a Workgroup lane axis (tensor cores / simdgroup_matrix / the register-tiled CPU GEBP kernel), carrying a scalar micro-kernel fallback the renderer falls back to when it declines the block. Hand-built IR is the only way to put one in front of a pass: schedule transforms mint Tile_mma AFTER the optimization pipeline, and Ir.Low_level.optimize rejects one outright — so a test that wants a tile in a routine builds the scalar twin, optimizes THAT, and substitutes the tile into the result.

ldd/lda/ldb default to the declared extents read as a contiguous row-major block — ldd = n, lda = if ta then m else k, ldb = if tb then k else n — a purely syntactic default off the tile's own geometry, NOT a read of the operands' dimensions: an operand whose tile is a window into a wider array, or whose tile major axis sits outside its minor two (a batched site, gh-ocannl-528), passes its stride explicitly. lane defaults to a fresh symbol and tile to None, the renderer's own choice of C-tile geometry (gh-ocannl-619).

embed_idx idx embeds an arbitrary index expression, where embed takes a symbol.

ic n is the integer constant n as a scalar.

cmp op a b applies an index-precision binary operator — a comparison (Cmplt, Cmple, Cmpeq, Cmpne) or a connective (And, Or).

where_ cond then_ else_ is the Where ternop, its condition read at index precision and its arms at single — the shape a zero-fringe guard renders as.

if_idx cond body is if_ with the condition read at INDEX precision: the standing of a launch-extent or fringe guard, whose condition is an index expression rather than a value read.