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_at prec cond body guards body on cond being nonzero (Ir.Low_level.t.If), the condition read at prec — the caller's claim, unchecked here. Prefer if_, which reads the condition's own precision, or if_idx.

if_ cond body is if_at with the condition read at its OWN precision (Ir.Low_level.scalar_precision, the function the operand-precision check reads an operand's own precision with): a value read (the usual flag tensor or local) at its node's precision, and a comparison cmp built at the index precision its operands carry. A fixed annotation could not serve both: under a single guard an index comparison evaluates at single, and its Constant operands' index-precision annotation is neither that nor their own single, which optimize refuses (gh-ocannl-1709). if_idx pins index precision whatever the condition is.

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.

declare_local id declares the scope-local scalar id (Ir.Low_level.t.Declare_local) — the statement a virtualizer or Privatize emits ahead of the Set_locals of a lifted local. ~needs_init (default false, the computation writes the local before reading it) asks codegen to zero the declaration, which an accumulating local read on its first update needs.

set_from_vec ~length ~vec_unop tn idcs arg stores length consecutive lanes of the vector operation vec_unop applied to arg into tn, starting at the cell idcs (Ir.Low_level.t.Set_from_vec) — the packed form the lowering mints for e.g. the uint4x32-counter uniform. length is explicit: it may be fewer than the operation's Ir.Ops.vec_unop_lanes (a peeled tail store), and arg carries its OWN precision, typically uint4x32 for a counter block.

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.

define_local id v declares the scope local id and assigns it v: two statements, the Declare_local + Set_local pair a rewrite such as online softmax emits. Mint id once per local (Ir.Low_level.get_scope returns a fresh id on every call) and read it with Get_local id.

scope tn v is a Local_scope over tn whose body assigns v once: the inlined computation of tn (read at orig_indices, default none), in the single-assignment form simplify collapses into its reader where the collapse is exact (gh-ocannl-1788). Mints a fresh scope id. Over a materialized tn the form exists only past the optimizer: build its program with Ll_test.optimize_scoped.

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).

binop_own op a b is binop with each operand annotated at its OWN precision (Ir.Low_level.scalar_precision), built by Ir.Low_level.apply_op, the constructor the assignments lowering uses: what a store whose precision is not single takes. A binary operator computes at its destination's precision whatever its operands say, and optimize refuses an annotation that is neither that precision nor the operand's own (gh-ocannl-1709), so binop's single under a half or bfloat16 store is refused while an own-precision annotation is accepted under every store.

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

ic n is the integer constant n as a scalar.

cmp_at prec op a b applies a comparison (Cmplt, Cmple, Cmpeq, Cmpne) or a connective (And, Or) computed at prec (gh-ocannl-1741): the precision of the store it feeds, as a homogeneous binop's operands are, so a flag kept in a single local is built with cmp_at single where cmp's index-precision annotation would be refused.

cmp op a b is cmp_at at index precision: a guard's comparison or connective.

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. Under a store that is not single, use where_own.

where_own cond then_ else_ is where_ with every operand at its OWN precision, as binop_own builds a binop: the arms are accepted under any store, and the condition is read as if_ reads one (a cmp-built comparison at index precision, a value read at its node's precision).

if_idx cond body is if_ with the condition read at INDEX precision whatever it is: the standing of a launch-extent or fringe guard. For a cmp-built condition it builds what if_ builds; it differs on a value-read condition, which it converts to an index.