Ir.Low_levelmodule Scope_id : sig ... endval sexp_of_scope_id : scope_id -> Sexplib0.Sexp.tval hash_fold_scope_id :
Ppx_hash_lib.Std.Hash.state ->
scope_id ->
Ppx_hash_lib.Std.Hash.stateval hash_scope_id : scope_id -> Ppx_hash_lib.Std.Hash.hash_valueHow 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.
val sexp_of_axis_type : axis_type -> Sexplib0.Sexp.tval axis_type_of_sexp : Sexplib0.Sexp.t -> axis_typeval axis_type_label : axis_type -> Base.stringLoop keyword used by the human-readable printers: plain "for" for Serial, "for@<axis>" otherwise.
type t = | Noop| Comment of Base.string| Staged_compilation of Base.unit -> PPrint.document| Seq of t * t| For_loop of {index : Indexing.symbol;from_ : Base.int;to_ : Base.int;body : t;trace_it : Base.bool;axis : axis_type;}| Zero_out of Tnode.t| Set of {tn : Tnode.t;idcs : Indexing.axis_index Base.array;llsc : scalar_t;mutable debug : Base.string;}| Set_dynamic of {tn : Tnode.t;idcs : Indexing.axis_index Base.array;Static everywhere except dyn_axis (a Fixed_idx 0 placeholder there).
dyn_axis : Base.int;Which idcs slot is replaced by dyn_value at codegen time.
dyn_value : scalar_arg;Integer-valued index spliced into the row-major offset at dyn_axis.
llsc : scalar_t;mutable debug : Base.string;}A scatter: like Set but the write lands at a runtime row of axis dyn_axis — the write counterpart of scalar_t.Get_dynamic. gh-466: produced only by rewrite_one_hot_reductions (transposed one-hot pattern, the embedding-table gradient); never constructed by Assignments lowering. Schedule analyses must treat this write as statically unknown: loops whose index reaches dyn_value carry a cross-iteration write dependency and must stay serial (the deterministic no-atomics invariant).
| Set_from_vec of {tn : Tnode.t;idcs : Indexing.axis_index Base.array;length : Base.int;vec_unop : Ops.vec_unop;arg : scalar_arg;mutable debug : Base.string;}| Set_local of scope_id * scalar_t| Declare_local of {id : scope_id;needs_init : Base.bool;}| Workgroup_barrierWorkgroup-scoped synchronization (__syncthreads() / threadgroup_barrier). An opaque effectful statement: no CSE, hoisting, or code motion across it. Grid-scoped synchronization is deliberately not representable.
| If of {cond : scalar_arg;body : t;}Guarded statement: body executes iff cond is nonzero (renders as if (cond != 0) { body }). Introduced by launch-extent guards on hardware-annotated loops (docs/proposals/axis-types-for-loops.md §2); simplify_llc erases a guard whose condition an interval proves. A conditional write is never a definite write; virtualization treats guarded computations as non-inlineable in v1.
| Tile_mma of {d : Tnode.t * Indexing.axis_index Base.array;Accumulator block base.
*)a : Tnode.t * Indexing.axis_index Base.array;b : Tnode.t * Indexing.axis_index Base.array;m : Base.int;n : Base.int;k : Base.int;Covered block extents (multiples of the backend's intrinsic tile).
*)lane : Indexing.symbol;The cooperating Workgroup axis (extent = SIMD width).
fallback : t;Semantically equivalent scalar micro-kernel over fresh serial symbols.
*)}Cooperative tile multiply-accumulate (docs/proposals/tensorize-mma.md): d[i,j] += Σ_{l<k} a[i,l] * b[l,j] for i < m, j < n, relative to the operands' base index vectors, executed jointly by the threads of the lane axis (tensor cores / simdgroup_matrix). Each operand's tile spans its tnode's last two axes (row-major); the base indices must not mention lane. Constructed by schedule transforms only (Schedule.optop.Tensorize), after the optimization pipeline. Backends without an MMA hook render fallback under an if (lane == 0) guard. Validates like Workgroup_barrier plus a write of d for the coverage rule; see validate_parallel.
Cases: t -- code, scalar_t -- single number at some precision.
and scalar_t = | Local_scope of {id : scope_id;body : t;orig_indices : Indexing.axis_index Base.array;}| Get_local of scope_id| Get of Tnode.t * Indexing.axis_index Base.array| Get_dynamic of {tn : Tnode.t;The gathered table; treated as a read of tn, like Get.
idcs : Indexing.axis_index Base.array;Static everywhere except dyn_axis.
dyn_axis : Base.int;Which idcs slot is replaced by dyn_value at codegen time.
dyn_value : scalar_arg;Integer-valued index spliced into the row-major offset at dyn_axis. gh-343: produced only by rewrite_one_hot_reductions; never escapes low-level / backend codegen.
}| Get_merge_buffer of Tnode.t * Indexing.axis_index Base.array| Ternop of Ops.ternop * scalar_arg * scalar_arg * scalar_arg| Binop of Ops.binop * scalar_arg * scalar_arg| Unop of Ops.unop * scalar_arg| Constant of Base.float| Constant_bits of Base.int64Direct bit representation, primarily for uint4x32
*)| Embed_index of Indexing.axis_indexThe argument precision is preserved in heterogeneous precision operation arguments, and is ignored (overridden) in homogeneous precision operations.
val sexp_of_t : t -> Sexplib0.Sexp.tval sexp_of_scalar_t : scalar_t -> Sexplib0.Sexp.tval sexp_of_scalar_arg : scalar_arg -> Sexplib0.Sexp.tval equal_scalar_arg : scalar_arg -> scalar_arg -> Base.boolval compare_scalar_arg : scalar_arg -> scalar_arg -> Base.intval loop_over_dims :
Base.int Base.array ->
body:(Indexing.axis_index Base.array -> t) ->
tval unroll_dims :
Base.int Base.array ->
body:(Indexing.axis_index Base.array -> offset:Base.int -> t) ->
tval loop_over_padding_region :
dims:Base.int Base.array ->
padding:Ops.axis_padding Base.array ->
body:(Indexing.axis_index Base.array -> t) ->
tGenerate 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. The middle region continues recursing to find padding in other dimensions.
val has_accumulation : t -> Base.boolWhether 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).
Phase B of docs/proposals/axis-types-for-loops.md. Hardware slot assignment is positional, not stored in the IR: among a kernel's annotated loops of one kind, the innermost binds .x (slot 0), the next .y, then .z; Workgroup and Workgroup_reduce share the block/threadgroup slot space.
Arrays of length 3 (.x, .y, .z); all-1s for all-Serial code.
val sexp_of_launch_dims : launch_dims -> Sexplib0.Sexp.tval equal_launch_dims : launch_dims -> launch_dims -> Base.booltype hardware_axis_info = {ha_index : Indexing.symbol;ha_kind : [ `Grid | `Workgroup ];ha_slot : Base.int;Positional: the innermost same-kind loop binds .x = slot 0.
ha_from_ : Base.int;ha_extent : Base.int;to_ - from_ + 1.
}val hardware_axes : t -> hardware_axis_info Base.listAll hardware-annotated loops in pre-order, with their positional slots.
val launch_dims : t -> launch_dimsPer-slot maximum extents over the kernel's annotated loops.
val validate_parallel : Tnode.Placements.t -> t -> Base.unitBackend-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). Cannot prove iteration independence — that is the annotating pass's obligation.
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.
type virtualize_settings = {mutable enable_device_only : Base.bool;mutable max_visits : Base.int;mutable max_inline_reduction : Base.int;Recompute-cost cap for inlining: a node whose setters have enclosing reduction loops (loops not appearing in the setter's indices) with a trip-count product exceeding this value is never virtualized. Negative values disable the cap.
*)mutable max_tracing_dim : Base.int;mutable inline_scalar_constexprs : Base.bool;mutable inline_simple_computations : Base.bool;mutable inline_complex_computations : Base.bool;}val virtualize_settings : virtualize_settingsval sexp_of_visits : visits -> Sexplib0.Sexp.tval visits_of_sexp : Sexplib0.Sexp.t -> visitsval visits : Base.int -> visitsval recurrent : visitsval is_visits : visits -> Base.boolval is_recurrent : visits -> Base.boolval visits_val : visits -> Base.int Base.optionval recurrent_val : visits -> Base.unit Base.optionmodule Variants_of_visits : sig ... endtype traced_array = {tn : Tnode.t;assignments : Base.int Base.array Base.Hash_set.t;accesses : (Base.int Base.array, visits) Base.Hashtbl.t;mutable zero_initialized_by_code : Base.bool;mutable zeroed_out : Base.bool;mutable read_before_write : Base.bool;The node is read before it is written (i.e. it is recurrent).
*)mutable read_only : Base.bool;Surprisingly, the notions of read-only and of constant memory mode come apart: small hosted constants are not read-only because they are initialized on devices by being assigned to; and a volatile memory mode is read-only from the devices' perspective.
*)mutable is_scalar_constexpr : Base.bool;True only if the tensor node has all axes of dimension 1, is either zeroed-out or assigned before accessed, is assigned at most once, and from an expression involving only constants or tensor nodes that were at the time is_scalar_constexpr.
*)mutable is_accessing : Base.bool;False only if the tensor node is built from index embeddings and scalar constant expressions.
*)mutable is_complex : Base.bool;True only if the tensor node is built from a genuinely complex scalar computation (one that accesses other non-constexpr computations). Sharing a loop symbol with another tensor does not, by itself, make a node complex (see #134).
*)mutable prefers_virtual_one_hot : Base.bool;True when at least one setter for this tensor is a one-hot selector assignment, i.e. a Cmpeq between the embedded range iterator and a loop-variable-free expression. When has_non_one_hot_setter is false this tensor is exempt from the visit-count Never_virtual rule (task-73617488).
mutable has_non_one_hot_setter : Base.bool;True when at least one setter is NOT a one-hot selector (including Set_from_vec). A tensor with prefers_virtual_one_hot && not has_non_one_hot_setter is the candidate for the one-hot virtualizer exemption.
mutable is_range_producer : Base.bool;True when at least one Set assigns this tensor from a bare Embed_index scalar, i.e. the tensor is a Range_over_offsets producer. Used by the indirect arm of is_one_hot_selector_assignment to prove that a Get(rtn, [k]) will inline to Embed_index k rather than arbitrary values (task-73617488).
mutable inline_reduction_extent : Base.int;The largest product of trip counts of loops that enclose one of the node's setters without appearing in its indices (i.e. reduction loops). Inlining the computation replays these loops at every read site; compared against virtualize_settings.max_inline_reduction.
mutable read_by_other : Base.bool;True when some statement other than the node's own setters reads the node. Unlike accesses, same-cell reads count, while a setter's own read-modify-write does not. Gates the recompute-cost guard: a node never read in the routine has no inlining cost, so it must stay eligible for virtual dead-code elimination.
}val sexp_of_traced_array : traced_array -> Sexplib0.Sexp.tval get_node :
(Tnode.t, traced_array) Base.Hashtbl.t ->
Tnode.t ->
traced_arraytype traced_store = (Tnode.t, traced_array) Base.Hashtbl.tval sexp_of_traced_store : traced_store -> Sexplib0.Sexp.ttype optimize_ctx = {computations : (Tnode.t,
(Indexing.axis_index Base.array Base.option * t) Base.list)
Base.Hashtbl.t;The computations (of the tensor node) are retrieved for optimization just as they are populated, so that the inlined code corresponds precisely to the changes to the arrays that would happen up till that point. Within the code blocks paired with an index tuple, all assignments and accesses must happen via the index tuple; if this is not the case for some assignment, the node cannot be virtual. Currently, we only allow for-loop symbols in assignment indices of virtual nodes.
*)placements : Tnode.Placements.t;Per-compilation-lineage memory-mode resolution (docs/proposals/context-scoped-memory-modes.md): the pipeline's placement decisions (Virtual / Local / On_device) land here, seeded by and never written back to the tnodes' declared intent (Tnode.t.memory_mode).
}val sexp_of_optimize_ctx : optimize_ctx -> Sexplib0.Sexp.tval empty_optimize_ctx : Base.unit -> optimize_ctxval copy_optimize_ctx : optimize_ctx -> optimize_ctxA 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.
type optimized = {traced_store : traced_store;optimize_ctx : optimize_ctx;llc : t;merge_node : Tnode.t Base.option;}val sexp_of_optimized : optimized -> Sexplib0.Sexp.tval optimize :
optimize_ctx ->
unoptim_ll_source:(PPrint.document -> Base.unit) Base.option ->
ll_source:(PPrint.document -> Base.unit) Base.option ->
name:Base.string ->
Indexing.static_symbol Base.list ->
t ->
optimizedreads_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.
val simplify_llc : Indexing.static_symbol Base.list -> t -> tTop-down algebraic simplification with interval-driven comparison folding (in particular, it erases If guards whose conditions the loop extents prove). 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.
val rewrite_one_hot_reductions :
?static_indices:Indexing.static_symbol Base.list ->
t ->
tgh-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.
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.
val input_and_output_nodes :
optimized ->
(Base.Set.M(Ir.Tnode).t * Base.Set.M(Ir.Tnode).t) * Tnode.t Base.optionInputs 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.
val function_header_doc :
?name:Base.string ->
?static_indices:Indexing.static_symbol Base.list ->
Base.unit ->
PPrint.documentval to_doc_cstyle :
?name:Base.string ->
?static_indices:Indexing.static_symbol Base.list ->
Base.unit ->
t ->
PPrint.documentAdheres more to the C syntax, outputs implicit type casts.
val to_doc :
?name:Base.string ->
?static_indices:Indexing.static_symbol Base.list ->
Base.unit ->
t ->
PPrint.documentAdheres to the %cd syntax.