Ir.C_syntaxmodule Lazy = Utils.Lazymodule Tn = Tnodetype mma_rendering = | Mma_intrinsics| Mma_intrinsics_ldmatrixThe intrinsic arms fed by warp-cooperative ldmatrix loads over a Swizzle_b128 staged tile (gh-ocannl-481 item 3) instead of per-lane gathers. Recorded distinctly because the gh-476 sweep must be able to pin which of the two load paths a timing measured.
| Mma_register_tiled| Mma_scalar_fallbackval mma_rendering_of_sexp : Sexplib0.Sexp.t -> mma_renderingval sexp_of_mma_rendering : mma_rendering -> Sexplib0.Sexp.tval compare_mma_rendering : mma_rendering -> mma_rendering -> Base.intval equal_mma_rendering : mma_rendering -> mma_rendering -> Base.boolval mma_census : (Base.string * mma_rendering) Base.list Base.refval is_tensorized_rendering : mma_rendering -> booltype tensorization = | Not_requestedCodegen emitted no Tile_mma statement at all: nothing about this routine claims tensor cores, and a timing of it is not a tensorized timing.
| TensorizedAt least one Tile_mma statement rendered to a real tensor-core or SIMD-register-tile emission (Mma_intrinsics, Mma_intrinsics_ldmatrix, Mma_register_tiled). The C backends' register-tiled rendering counts: it is the SIMD-tile emission those backends have, and the point of the label is fallback-vs-not, not which ISA.
| Scalar_fallbackTile_mma statements were emitted and every one of them declined to the lane-0 scalar fallback. A timing labeled tensorized in this state measured scalar code.
Whether a compiled routine's tensorization request was honoured (gh-ocannl-626). Three states, not a boolean, because "this routine has no tensor-core work in it" and "this routine asked for tensor-core work and got the scalar fallback" are opposite readings of a timing and the boolean that collapses them is the defect: an Tile_mma-free routine is not a failure, a wholly declined one is. Mixed renderings — some statements honoured, some declined — read as Tensorized, and the counts on mma_summary say by how much; the label answers "did any tensor-core / SIMD-tile emission happen", the counts answer "how much of what was asked for".
val tensorization_of_sexp : Sexplib0.Sexp.t -> tensorizationval sexp_of_tensorization : tensorization -> Sexplib0.Sexp.tval compare_tensorization : tensorization -> tensorization -> Base.intval equal_tensorization : tensorization -> tensorization -> Base.boolval tensorization_name : tensorization -> stringtype mma_summary = {renderings : (Base.string * mma_rendering) Base.list;The census entries of one compile, in emission order (kernel name, rendering). Fissioned segments of one routine contribute their kernels to the same summary.
*)tensorization : tensorization;statements : Base.int;Tile_mma statements emitted: List.length renderings.
scalar_fallbacks : Base.int;Of those, how many rendered as the lane-0 scalar fallback.
*)}What a compile's mma_census says about the routine it produced (gh-ocannl-626). Derived once, where the routine is compiled, so that "did this tensorize" is a property of the compiled routine rather than of whichever call site remembered to bracket the global.
val sexp_of_mma_summary : mma_summary -> Sexplib0.Sexp.tval summarize_census : (Base.string * mma_rendering) Base.List.t -> mma_summaryval empty_mma_summary : mma_summaryval merge_mma_summaries : mma_summary Base.List.t -> mma_summaryval mma_summary_string : mma_summary -> stringThe one-line census a timing report prints beside a measurement (gh-ocannl-626): the tensorization label first — a reader scanning a table for a mismatch reads that word, not a histogram — then the per-rendering counts. Shared so that every harness that names a variant after a rendering says "did this tensorize" the same way; the two bench copies had disagreed about the predicate (one counted the fallback, the other tested equality with Mma_register_tiled, which false-warns on any GPU backend).
val with_census : (unit -> 'a) -> 'a * mma_summaryRun f with the mma_census collecting, and return its result alongside the summary of what rendered during it (gh-ocannl-626). Both census refs are saved and restored, so calls nest and an inner compile does not disturb an outer collection; the summary covers only f's own entries. This replaces the hand-rolled Exn.protect bracket that had been copied to six call sites — and, applied inside Context.compile, makes collecting the census the default rather than something a timing harness must remember to do.
The census is a process global, as it always was: compiles are sequential on the main domain, and nothing here makes concurrent compiles from several domains attribute their renderings correctly.
type peel_skip = | Skip_debug_loggingdebug_log_from_routines: the per-step Set form is kept so the trace survives.
| Skip_dead_levelThe level being rendered is dead (to_ < from_).
| Skip_accum_pinnedThe peel reached a base whose update renders at storage precision (the rng carve-out, gh-ocannl-517), so localizing it would change the draw rather than move it.
*)val peel_skip_of_sexp : Sexplib0.Sexp.t -> peel_skipval sexp_of_peel_skip : peel_skip -> Sexplib0.Sexp.ttype peel_verdict = {levels : Base.int;Loop levels the localized scope spans: the level being rendered, plus those Low_level.peel_accum_nest took below it. One for a reduction localized at its own level; two where the accumulated cell let the peel swallow the enclosing level too.
guards : Low_level.peel_guard_verdict Base.list;The verdict of each If among them, outermost first.
}What a localizing site DECIDED (gh-ocannl-733). A named record rather than an inline one so a consumer can hold one: comparing the verdicts of two compiles is the whole use.
val peel_verdict_of_sexp : Sexplib0.Sexp.t -> peel_verdictval sexp_of_peel_verdict : peel_verdict -> Sexplib0.Sexp.tval equal_peel_verdict : peel_verdict -> peel_verdict -> Base.boolval compare_peel_verdict : peel_verdict -> peel_verdict -> Base.inttype peel_site = | Peel_localized of peel_verdictThe site localized, deciding this.
*)| Peel_refused of Low_level.peel_refusalLow_level.peel_accum_nest refused, with why.
| Peel_not_attempted of peel_skipCodegen declined without asking the peel, or after it.
*)val peel_site_of_sexp : Sexplib0.Sexp.t -> peel_siteval sexp_of_peel_site : peel_site -> Sexplib0.Sexp.tval peel_census : (Base.string * peel_site) Base.list Base.refval is_localized_peel : peel_site -> booltype peel_summary = {sites : (Base.string * peel_site) Base.list;The census entries of one compile, in emission order (kernel name, site). Fissioned segments of one routine contribute their kernels to the same summary.
*)localized : Base.int;declined : Base.int;Sites that did not localize, refused and not-attempted together.
*)}What a compile's peel_census says about the routine it produced (gh-ocannl-733).
val sexp_of_peel_summary : peel_summary -> Sexplib0.Sexp.tval summarize_peel_census :
(Base.string * peel_site) Base.List.t ->
peel_summaryval empty_peel_summary : peel_summaryval peel_summary_string : peel_summary -> stringThe one-line peel census a report prints beside a routine (gh-ocannl-733): how many reduction sites localized, and the distinct verdicts they earned.
val with_peel_census : (unit -> 'a) -> 'a * peel_summaryRun f with the peel_census collecting, and return its result alongside the summary of what the reduction peel decided during it (gh-ocannl-733). Nests additively and restores both refs, exactly as with_census does for the Tile_mma census — and for the same reason: an enclosing collection must still see an inner compile's sites.
type volatility_site = | Volatile_accumulation_reads of Base.stringA reduction-shaped scope local for which the backend selected expression-level volatile pointer casts on any device reads in its accumulating loop, by the local's emitted identifier: the localized accumulation form (gh-ocannl-731, gh-ocannl-820).
| Plain_accumulator of Base.stringA reduction-shaped scope local whose reads stay plain — either this backend does not request the workaround, or no materialized device read was emitted in its accumulating loop. The accumulator itself stays register-resident in both cases.
*)| Volatile_rmw_reads of Base.stringA device-memory read-modify-write at an address invariant across an enclosing serial loop, whose device reads use expression-level volatile pointer casts.
val volatility_site_of_sexp : Sexplib0.Sexp.t -> volatility_siteval sexp_of_volatility_site : volatility_site -> Sexplib0.Sexp.tval equal_volatility_site : volatility_site -> volatility_site -> Base.boolval compare_volatility_site : volatility_site -> volatility_site -> Base.inttype volatility_summary = {entries : (Base.string * volatility_site) Base.list;The census entries of one compile, in emission order (kernel name, site). Fissioned segments of one routine contribute their kernels to the same summary.
*)requested : Base.bool;Whether the backend this compile ran on asks for the workaround (C_syntax_config.volatile_serial_accumulation). false makes every localized site a Plain_accumulator and suppresses the device-memory RMW arm entirely.
volatile_accumulations : Base.int;plain_accumulations : Base.int;volatile_rmw_reads : Base.int;}What a compile's volatility census says about the routine it produced (gh-ocannl-782).
val sexp_of_volatility_summary : volatility_summary -> Sexplib0.Sexp.tval volatility_census : (Base.string * volatility_site) Base.list Base.refval summarize_volatility_census :
requested:Base.bool ->
(Base.string * volatility_site) Base.List.t ->
volatility_summaryval empty_volatility_summary : volatility_summaryval volatility_summary_string : volatility_summary -> stringThe one-line volatility census a report prints beside a routine (gh-ocannl-782/820): how many serial accumulations this routine carries and how many selected the volatile-read workaround.
val with_volatility_census : (unit -> 'a) -> 'a * volatility_summaryRun f with the volatility census collecting, and return its result alongside the summary of what the accumulation workaround decided during it (gh-ocannl-782). Nests additively and restores the refs, exactly as with_census and with_peel_census do.
The address space of a tile-MMA operand as the emission hooks see it.
The physical layout of a tile-MMA operand's storage (gh-ocannl-481 item 3, D2). This is the whole Stage -> emission contract: the emission never re-derives the layout, it trusts the component it is handed.
`Swizzled_elem operands never reach a hook — the Tile_mma rendering declines them centrally, since no intrinsic load form matches an element-granularity permutation. `Swizzled_b128 is only passed when the access is reconstructible from (ptr, ld) alone: the pointer is the tile origin of a rank-2 node whose minor dim is ld, so the element at (row, col) sits at row*ld + (((col/u) lxor (row land (ld/u - 1))) * u) with u = 16 / prec_in_bytes — everything else declines.
type mma_operand = PPrint.document * Base.int * mma_space * mma_layouttype mma_source = Base.int * mma_space * mma_layoutA tile-MMA input operand described but not addressed: leading-dimension stride in elements, address space, physical layout — no pointer. This is how the a and b operands reach both emission hooks, because a caller does not always have an address to give: a software-pipelined tile's live copy rotates per k-block (gh-ocannl-487), so its pointer only exists inside the rotor loop, and C_syntax_config.mma_fragment_syntax runs outside it.
Everything an arm decides acceptance by is here — the pointer never was a criterion.
The emitting half of an accepted C_syntax_config.mma_syntax call: the arm has committed, and only the a/b tile addresses are still outstanding. Splitting the hook this way makes Option.is_some (mma_syntax ...) a support predicate usable wherever a call is known but not yet addressed — with no separate acceptance implementation that could drift from the emitting one.
type async_copy_syntax = {ac_copy : dst:PPrint.document ->
src:PPrint.document ->
bytes:Base.int ->
PPrint.document;One element-sized asynchronous global→workgroup-shared copy statement; dst and src are element addresses (&ident[offset]) and bytes the element size — 4 or 8 today. The hardware also copies 16, but a 16-byte copy requires a 16-byte-aligned destination, and plain workgroup-shared declarations align to the element type only (4 for uint4x32_t) — so 16 stays out of the per-element eligibility until a rendering guarantees the alignment (Codex P2 on PR #317). The copy is byte-for-byte: eligibility (same storage precision on both sides, no value transformation) is the caller's check.
ac_wait_all : Base.string;Statement completing every asynchronous copy issued so far by the calling thread, committed or not (CUDA cp.async.wait_all). Cross-thread visibility still needs the workgroup barrier the caller emits right after.
}gh-ocannl-487 phase 2: asynchronous staging copies for software-pipelined tiles (see C_syntax_config.async_copy).
module type C_syntax_config = sig ... endWhether c lies exactly halfway between two adjacent f32 values, so that narrowing it to f32 is a tie that IEEE-754 breaks to even -- and any decimal near but not equal to c would instead break by whichever side it fell on.
Computed against the two neighbours rather than by masking mantissa bits, so that the f32 subnormal range (where the retained-bit count shrinks, and a bit mask would need a second case) is covered by the same three lines. Values that overflow f32 have no neighbours to sit between and are not ties.
A C source literal for the floating-point constant c, as a double-typed floating literal that parses back to exactly c on every C-family backend.
Three properties, each of which %.16g alone gets wrong (gh-ocannl-623). Two of them are not C's alone — a debug dump wants a floating literal that round-trips for the same reasons a kernel does — so they live in Utils.decimal_float_literal, which the IR printers share (gh-ocannl-713):
2. came out as the integer literal 2, value-preserving in the cast contexts the constants happen to sit in today — but -0. came out as "-0", the integer zero, hence +0.0 once cast, so a -0.0 in host data silently reached kernels as +0.0 wherever the constant fill inlines as scalar stores (gh-ocannl-615 fixed that one value). The rest of the class is latent rather than inert: an integer literal divides as an integer against another one, promotes as an integer, and would overflow its type once the digits outgrow long long.0.1 +. 0.2 at 16 digits is a different double. Values that already round-trip at 16 digits keep their exact previous spelling, which is why no codegen golden moves except by the appended .0.%.16g's words (inf, nan), which no C dialect parses. INFINITY and NAN are C99 math.h macros that MSL also provides; the CUDA and HIP preludes define them under #ifndef since nvrtc/hiprtc supply no standard headers. (-INFINITY) is parenthesized so it cannot glue into --INFINITY beside a subtraction. A NaN's payload and sign do not survive this (no dialect spells a payload portably); nothing in OCANNL depends on them, and a bit-exact constant has Low_level.scalar_t.Constant_bits available.Deliberately not carried here: a precision suffix. The literal is always double-typed and the narrowing is B.convert_precision's cast, which is a single rounding of the exact host double — the same conversion the host performs when it stores the value. An f-suffixed decimal would instead round the decimal straight to float. The cast is present for every non-double target: Ops.c_convert_precision and each backend's override return ("", "") only when from and to_ are the same precision.
That last argument has a hole, and is_f32_tie is what fills it: the cast only is a narrowing where the dialect has a double. MSL does not, so Metal rounds the decimal itself, at parse time, to f32. A decimal that round-trips as a double is not thereby exact, so the two readings — round c to f32, versus round a decimal near c to f32 — can differ, and they differ exactly when c sits on an f32 tie: the host takes ties-to-even while the dialect follows whichever side the decimal happens to land on. Which digit count is emitted then decides the value, so this is not a hazard the retry above introduced but one it moves. A hexadecimal literal takes it away instead of moving it: it is exact by construction, so there is no parse-time rounding for either reading to disagree about, and both round c itself. It is spelled only for ties — a handful of values that were otherwise a coin toss — so the ordinary constant keeps its readable decimal, and C99, CUDA, HIP and MSL all accept the form.
val default_binop_syntax :
Ops.prec ->
Ops.binop ->
PPrint.document ->
PPrint.document ->
PPrint.documentThe C-family rendering of a binary operation, from Ops.binop_c_syntax: prefix, first operand, infix, second operand (breaking after the operator), suffix.
Outside Pure_C_config because the GPU backends, which shadow binop_syntax wholesale (most ops need target-specific intrinsics or precision bridging), delegate here for the ops that are spelled the same in C, CUDA, HIP and MSL -- the comparisons and the logical connectives. Those then have one spelling (Ops.binop_c_syntax) and one layout (here) across all backends, instead of a copy per backend.
The RNG binops -- the two Threefry variants and the per-lane uniform conversion -- rendered as a call to the builtin of that name. Every C-family backend provides the same three builtins (builtins.c, Builtins_cuda, Builtins_metal) under the same precision contract: the Threefry ops produce a uint4x32 block, and the lane conversion consumes one to produce the target precision, so it is the one binop that rejects uint4x32 (its builtin already yields the target precision, bypassing the generic bfloat16/fp8 compute-in-single wrapping).
Backends pass their own two-argument call renderer -- the layouts differ in where a line break may fall -- and the backend name for the diagnostics. op must be one of the three ops above.
All maximal identifier-like substrings of s -- a run of alphanumerics and underscores starting at a letter or underscore. This decomposes a composite rendering like "(fabsf(floorf(" into ["fabsf"; "floorf"] rather than the concatenation "fabsffloorf".
A run whose first character is preceded by a digit or a . is the tail of a numeric literal (f in "0.0f", h in "1.0h"), not a name, and is skipped -- otherwise every single-letter literal suffix a backend emits would be reserved.
val all_precs : Ops.prec listThe precisions a backend's renderers are exercised over -- by op_syntax_idents and by operand_conditionality_violations. The operator sweeps use Ops' derived enumerations (Ops.all_of_binop and friends), so a newly added operator joins both checks automatically; Ops.prec cannot be derived the same way -- its constructors carry the phantom-typed precision witness -- so the exhaustive match below stands in for it: adding a precision is a build error here, and the fix is to extend this list (or to leave the precision out deliberately, as Void_prec is, since every renderer rejects it).
val _all_precs_is_complete : Ops.prec -> Base.unitval op_syntax_idents :
ternop_syntax:
(Ops.prec ->
Ops.ternop ->
PPrint.document ->
PPrint.document ->
PPrint.document ->
PPrint.ToBuffer.document) ->
binop_syntax:
(Ops.prec ->
Ops.binop ->
PPrint.document ->
PPrint.document ->
PPrint.ToBuffer.document) ->
unop_syntax:
(Ops.prec -> Ops.unop -> PPrint.document -> PPrint.ToBuffer.document) ->
vec_unop_syntax:
(Ops.prec -> Ops.vec_unop -> PPrint.document -> PPrint.ToBuffer.document) ->
convert_precision:
(from:Ops.prec -> to_:Ops.prec -> Base.String.t * Base.String.t) ->
Base.String.t listEvery function and type name a backend's operator rendering can emit, obtained by rendering each (precision, operator) pair over a placeholder operand and harvesting the identifiers.
Rendering the syntax functions is what makes this correct per backend: reading the names off Ops.unop_c_syntax instead would describe *C*, and the GPU backends shadow those functions wholesale. MSL spells Tanh_approx as tanh where C spells it tanhf, so the C-derived list left "tanh" free for a tensor node -- and Tensor.unop's ~op_label makes that an ordinary name, minted by every Operation.tanh. The Metal kernel then declared device float *__restrict tanh and the call on the next line resolved to that pointer (gh-ocannl-553). The same holds for exp, log, sqrt, sin, cos, trunc: unsuffixed in MSL, suffixed in C.
Backends reject some (precision, operator) pairs by raising, either while selecting the renderer or while applying it; those pairs contribute no names. Any escaping exception is swallowed rather than failing the compilation this list only guards.
val operand_conditionality_violations :
ternop_syntax:
(Ops.prec ->
Ops.ternop ->
PPrint.document ->
PPrint.document ->
PPrint.document ->
PPrint.ToBuffer.document) ->
binop_syntax:
(Ops.prec ->
Ops.binop ->
PPrint.document ->
PPrint.document ->
PPrint.ToBuffer.document) ->
string Base.List.tWhat a backend's operator renderings actually emit, checked against the operand-evaluation contract they are supposed to implement (Ops.binop_conditionality / Ops.ternop_conditionality) -- returns one message per disagreement, [] when they agree (gh-ocannl-582).
The classifier is what Low_level.affine_accesses and both Cost_model walks consume, and the renderers are the ground truth it describes; checking rather than consulting is deliberate, since a renderer cannot ask "am I allowed to short-circuit here" -- it either emits a conditional or it does not. Each operator is rendered over distinguishable placeholder operands at every precision, and the emitted text is read for the only C-family constructs that can skip an operand: ?:, && and ||. So:
Only_first / Only_second) must not mention its discarded operand -- every backend rejects one outright, which satisfies this trivially, and C_syntax.pp_scalar emits the selected operand alone before binop_syntax is ever reached;Gated_second must place one of those constructs before its second operand;Both_operands / All_three must emit none of them anywhere -- an Ops.Max spelled (a > b ? a : b) would evaluate a twice and b conditionally, and the cost model's floor would keep charging both;Cond_and_one_arm must put its arms on the two sides of a ?: alternative, which is what rules out MSL's select (a call: it evaluates both arms, so a range guard lowered to Where would still perform its out-of-range read).The operator sweep is Ops' derived enumeration, so an operator added to Ops.binop or Ops.ternop is checked here without anyone remembering to list it. A (precision, operator) pair the backend rejects by raising contributes nothing, as in op_syntax_idents.
The names defined by a backend's builtins table (the keys of its (name, definition, dependencies) entries), for its ident_blacklist. A node taking one of these names both shadows the definition and, since filter_and_prepend_builtins selects entries by searching the rendered kernel for their key, drags the definition into a kernel that never calls it.
The words the C language itself reserves, plus the scaffolding names this module's rendering emits unconditionally. Shared by every C-family backend through Pure_C_config, and by C_syntax.kernel_ident: a routine and a tensor node are both plain identifiers in the emitted source, so they are unsafe on exactly the same words.
asm is in the list even though C89/C99 reserve it only as a common extension: every compiler OCANNL emits for (gcc, clang, nvrtc, hiprtc, the Metal front end) treats it as a keyword, which is what made Block_comment "asm" emit void asm( and fail as an "internal" codegen bug (gh-ocannl-686).
Macro names the C-family preludes' unconditional includes define (<stdio.h>, <math.h>, <string.h>, <stdlib.h>, <stdint.h>). Macros are the half of the standard library that NOTHING can shadow -- a parameter or a local named NAN expands mid-declaration and the kernel fails to parse -- so these constrain every identifier, tensor-node names included, and belong in the shared ident_blacklist (gh-ocannl-686).
Function and type names the same unconditional includes declare. Unlike the macros above these are ordinary file-scope declarations, and C scoping makes the two cases genuinely different:
void printf(...) after <stdio.h> is a conflicting declaration and void exp(...) after <math.h> likewise -- an error regardless of whether the kernel ever calls either;exp is well-formed C and always has been.So this table constrains routine names only (C_syntax.kernel_ident) and is deliberately kept out of ident_blacklist: adding it there would rename tensor nodes -- Tensor.unop's ~op_labels are exactly these words (exp, log, sqrt, tanh) -- churning goldens to prevent a collision that cannot happen. The names a rendering actually CALLS are a separate concern already covered by op_syntax_idents, which is what stops a node from shadowing a callee the same kernel invokes.
The words C++ reserves on top of c_keywords. CUDA, HIP and MSL are all C++ dialects -- their kernels are parsed by a C++ front end even where the emitted body is plain C -- so their backends extend the blacklist with this table. Several entries are entirely plausible tensor labels or routine names (class, new, operator, bool, this, template), and a collision there is the same misattributed "bug in OCANNL" failure as gh-ocannl-686's asm.
C_syntax_config.full_printf_support for a dialect whose printf does support %g — every backend except Metal, whose MSL printf does not. Such a backend still gives the scaled-integer rendering up when asked for backend uniformity, so that its routine logs read the same as Metal's: the setting is what decides, and it is read here once rather than restated per backend.
module Pure_C_config (Input : sig ... end) : sig ... endmodule C_syntax (B : C_syntax_config) : sig ... end