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.refThe 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_layoutmodule type C_syntax_config = sig ... endval 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 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.
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.
module Pure_C_config (Input : sig ... end) : sig ... endmodule C_syntax (B : C_syntax_config) : sig ... end