Module C_syntax.Pure_C_config

The low-level procedures this functor application will render: one per kernel of the compilation unit (a singleton for a plain compile, one entry per routine for a batch). The functor reads them for whole-unit analyses -- the identifier census, the scope-local verdicts -- while each kernel's own rendering goes through compile_proc, which takes the full Low_level.optimized record.

gh-ocannl-1603: Some w when a kernel whose top-level statements weigh more than 16 * w (Kernel_outline.groups) is rendered OUTLINED -- its statements packed in order into parts of about w each, every part a noinline function the kernel calls in turn -- because a C compiler's cost grows faster than linearly in a function's size. None renders every kernel whole. Only a host C backend can outline: a part is an ordinary function call.

Whether typ_of_prec accepts f64 tensor storage. Explicit so capability queries need not deliberately raise through a rendering function.

File-scope names the backend's own vendor headers declare (or its compiler predefines) that a kernel DEFINITION cannot take, beyond the C floor c_stdlib_idents already covers: global type names, macros, dialect keywords, and C-linkage functions (gh-ocannl-707). Constrains routine names only (C_syntax.kernel_ident), for the reason c_stdlib_idents does: a parameter or a local may shadow these, a file-scope definition may not. Names in the implementation's reserved namespace (__x, _X) need no entry: implementation_reserved closes that class by rule. [] for a dialect whose headers stop at the C floor.

Prepended to every identifier a tensor node takes in an emitted kernel -- a parameter, a pool-pointer local, a kernel-scope or shared array -- and to nothing else: not to the node's code name (Tnode.t.code_name, what debug output and error messages show), not to the text of a routine log line, not to scope locals (v<scope>_<name> already carries a prefix) (gh-ocannl-1338).

A GPU dialect's vendor headers declare global types (float4, uint4, half2) and macros (M_PI_F, MAXFLOAT, CUDARTAPI) outside the reserved namespace, and the kernel body spells the types: a parameter named uint4 shadows the type for every later device uint4* of the kernel, and a node named after a macro is macro-expanded before it is declared. A fixed prefix closes that class by construction, where a blacklist would have to track every header the toolchain's runtime compiler pulls in, including built-in ones no file holds. "" for cc, whose only headers are the C library's: their macros and the types the renderer spells are already in ident_blacklist, and the compiler's predefined macros are closed by rule (implementation_reserved, system_macro_idents).

How materialized in-context tensor nodes are passed to the kernel. `Per_param (the default, used by C and CUDA) emits one typed pointer parameter per node, whose host side binds pool_base + offset -- byte-identical to the pre-pooling codegen. `Pooled n (Metal) emits a fixed n byte-pointer pool parameters plus one (pool_index, byte_offset) slot table, and a kernel prologue that forms each node's typed pointer by casting (pools at slot.pool) + offset. This collapses O(num_nodes) buffer bindings to n + 1, which Metal needs to stay under its ~31 argument-buffer binding limit.

Format specifier for printing floating point numbers in debug logs.

Function to convert potentially floating-point numeric values for logging.

Lane i (0 to 3) of a uint4x32 value, as a %u routine-log argument. A uint4x32 value is a struct (uint4x32_t, C/CUDA/HIP) or a vector (uint4, Metal), never a number, so neither float_log_style nor styled_log_arg applies to it: a log statement prints its four lanes as unsigned integers instead. Coupled to typ_of_prec's spelling of Uint4x32_prec.

How a value of loop_index_type is passed to a log statement: the printf conversion specification to splice into the format string, and the argument document, cast where the backend's variadic logging call needs a different type than the loop index's own.

A loop index is int32_t/int normally but 64-bit under large_models, and the conversion has to track that width: passing a 64-bit argument to %d is a variadic type mismatch -- undefined behaviour on the printf paths, a compile error on Metal's os_log, and wrong digits for values past 32 bits either way.

Coupled to loop_index_type and to pp_log_statement: a backend that spells the loop index as its own type, or logs through a call whose conversions are not C's, overrides this alongside.

The dialect of the builtin parameter guards a kernel's prelude carries under the test seam Param_guards.with_guards (gh-ocannl-1768); None (the Pure_C_config default) guards nothing.

A branchless rendering of a value select at this compute precision (gh-ocannl-1622), in one of the select_blend forms: it must evaluate both arms unconditionally — each rendered before any ?:, && or || (operand_conditionality_violations checks it) — and have exactly the C type and value of the short-circuiting spelling it replaces, bit for bit. pp_scalar_op uses it for a Where, Relu_gate or Satur01_gate that Low_level.speculable_select admits, keeping the condition exactly as the short-circuiting spelling tests it.

Why: C forbids evaluating the unselected arm of ?:, so where an arm holds a load the compiler cannot prove dereferenceable it emits a conditional branch, and a step's time follows the data (cifar_resnet's inlined ReLU and max-pool backward selects on cc, gh-ocannl-1557). With both arms already values, the compiler if-converts. AMDGPU branches too (gh-ocannl-1695), so HIP blends over auto locals. None (the Pure_C_config default) keeps the short-circuiting ternop_syntax / binop_syntax spelling: CUDA's, which ptxas if-converts itself, and Metal's, whose compiler no tool here can read.

The precision the arithmetic over a node runs at, given the precision the node is stored at (gh-ocannl-517). Identity by default: a backend with native arithmetic at every storage width computes where it stores, which is what the GPU backends do (__nv_bfloat16, MSL's bfloat/half, and the 16-bit tensor-core shapes that consume them).

The CPU backends have no 16-bit arithmetic, so they map the narrow floats to f32 (subject to Ir.Numerics.t.narrow_compute_f32). Reads then widen once at the load and the result narrows once at the store, instead of every operator round-tripping through f32 and rounding its result to the narrow format — the "16-bit storage, f32 compute" of gh-ocannl-517.

Only the register precision of an assignment's intermediates is at stake: this function is never consulted for a declaration, a kernel parameter, or a buffer's element type, which always take the storage precision. It must be a function of the storage precision alone — identical across sibling autotune candidates — or schedule transforms would stop being numerics-preserving.

The precision a reduction accumulator resides at across its whole nest, given the precision the accumulated node is stored at (gh-ocannl-663). Where it differs from the storage precision, every serial-rendered form of a recognized accumulation — the plain serial fallback, unrolled and partitioned nests, and reduction-shaped scope locals — holds the accumulator here and narrows once at the store, so a reduction's effective accumulation width is a per-backend policy, never a property of which schedule or rendering ran (gh-ocannl-639's guarantee, extended beyond compute_prec).

On the CPU backends this is compute_prec: the accumulator is an assignment intermediate like any other. The GPU backends compute where they store (compute_prec is the identity on their native narrow arithmetic) but accumulate per their tensor-unit format triples, and this hook is where a backend mirrors those: CUDA's bf16 mma legs hold f32 per-lane registers across the whole k extent (the hardware has no bf16 accumulate), so its serial legs must widen bf16 the same way, while Metal's uniform-bf16 tile accumulates in bf16 fragments, so its serial legs keep bf16 residency — as HIP's do only under Numerics.Bf16_narrow: under the default Bf16_auto HIP swaps its uniform-bf16 arm to an f32 accumulator fragment and widens its serial legs to match (gh-ocannl-1051). fp8 has an accumulator format on no backend and follows the CPU policy (f32) everywhere; f16 accumulates natively at f16 in every seeded GPU triple and stays put.

The recognized accumulation update renders wholly at this precision, contribution included — gh-ocannl-639's rendering shape, kept on GPU deliberately: operand widenings are exact, a narrow-by-narrow product is exact at the wider precision (the same full-precision-product semantics the tensor units apply per element), and the FMA form stays a single fused operation. Statements outside recognized accumulations keep compute_prec.

Must resolve at least as wide as compute_prec (asserted at codegen setup: narrowing an intermediate below its own arithmetic precision would round-trip every update), and like it must be a function of the storage precision alone. When overriding compute_prec, override this together with it — the two are bound at include time, so a stale pairing does not track the override.

Whether the explicit vector-arithmetic renderings (`Vec_extensions Vectorized loops and reductions) can operate at this compute precision. f32 and f64 everywhere; fp16 additionally on CPU targets with native 16-bit arithmetic (gh-ocannl-516), which is exactly where compute_prec leaves Half_prec alone. A storage precision this rejects can still be vectorized when compute_prec maps it to one this accepts -- that is gh-ocannl-517's convert-on-load/store. The `Packed_struct rendering does no vector arithmetic and does not consult this: its own gate admits every float precision with a full-width vec_typ_of_prec pack type that is also the storage precision of the nodes it packs (gh-ocannl-1295).

The hardware register expression an annotated loop's index binds to (e.g. "blockIdx.x", "gid.y"), or None when the backend cannot bind this axis in hardware — the loop then renders as a serial for (a legal implementation absent barriers; see docs/proposals/axis-types-for-loops.md §2/§5). Slots are positional: 0 = .x, 1 = .y, 2 = .z.

Workgroup barrier statement (__syncthreads(); / threadgroup_barrier(...);); None makes Workgroup_barrier a compile-time error (serialization cannot implement a barrier).

gh-ocannl-487 phase 2: asynchronous global→workgroup-shared copies (CUDA cp.async) for the staging loads of software-pipelined tiles (Low_level.optimized.pipelined). When provided, an eligible staging Set (a raw same-precision copy of a materialized global into an async-eligible pipelined tile) renders as ac_copy instead of a load+store through registers, so the prefetch issued for iteration k+1 genuinely overlaps the compute of k. Completion is uniform, not per-group: the rotor loop's body is prefixed with ac_wait_all followed by a workgroup barrier (re-inserting, for the async arm, exactly the phase opener that Schedule.elide_staged_barriers elides for synchronous stores — those are published by the previous iteration's trailing bracket, an async copy needs its wait BEFORE the publishing barrier). Per-statement eligibility is opportunistic: an ineligible staging statement (precision conversion, a surviving fringe ternary, a non-global source) keeps the plain store, which the same barrier publishes — correctness never depends on which statements the arm accepted. None (the default, and every backend but CUDA today) keeps the portable synchronous rendering everywhere.

Pool-backed Grid rendering (docs/proposals/gh-ocannl-164.md): how to render an eligible outermost Grid loop when hardware_index does not bind it. `Dispatch emits libdispatch's dispatch_apply over contiguous chunks (macOS; blocks extension), `Openmp a #pragma omp parallel for over the chunk loop; both runtimes own a single process-global thread pool, so no pool state lives in the compiled kernel. `None keeps the serial fallback. Eligibility is decided per loop by compile_proc (see parallel_grid_safe); Workgroup loops always stay serial inside a chunk.

Target chunk count for parallel_grid_syntax (e.g. a small multiple of the core count); the actual count is capped by the loop extent. Values <= 1 disable parallel rendering.

Declaration prefix for workgroup-shared placements (__shared__ / threadgroup ); None makes a non-empty workgroup_shared set a compile-time error.

Whether this backend's kernel compile is pinned so its compiler never contracts a a * b + c into a fused multiply-add on its own (gh-ocannl-1457), derived from the same option builder the compile uses: cc's cc_backend_fp_contract=off (-ffp-contract=off, placed after any -ffast-math), CUDA's --fmad=false (cuda_fast_math=false), HIP's -ffp-contract=off (hip_fast_math=false). Only an explicit pin counts, never a toolchain default; Metal exposes no contraction switch, so it is false there (the Pure_C_config default). Read per call: the settings behind it are mutable within a process.

When true, a Tile_mma fallback renders in the form the contraction licence left it, and the register-tiled rendering spells a plain-add fallback as an unfused vector multiply then add -- both the serial loop's two roundings, which the compiler is forbidden to fuse. When false, the fallback renders fused (Low_level.fuse_tile_fallback), as every tile multiply-accumulate computes.

Workaround for a Metal shader-compiler miscompilation of serial accumulations (observed on macOS 15/Metal 3.1-3.2 through macOS 26/Metal 4), in both spellings a reduction can take. Named for what it covers rather than for either spelling: it was volatile_scalar_rmw while only the read-modify-write arm existed (gh-ocannl-782 renamed it).

The original manifestation: a serial loop accumulating into a loop-invariant address of a kernel-parameter-derived pointer — acc[k] = acc[k] + f(i) with k free of i — can execute as if the load were hoisted above the loop and the store sunk below it without carrying the accumulation, leaving only the last iteration's contribution (scalar losses collapsed to the last sample's CE; w.grad accumulated only the last batch element). Standalone repro: benchmarks/runners/ocannl/bench_metal_bug.ml.

The localized manifestation (gh-ocannl-731): after the serial-reduction localizer the accumulator is a scope local and the node is stored once, and the same pass corrupts THAT form instead — by a data-independent additive constant. Standalone repro and one-factor-at-a-time matrix: benchmarks/runners/ocannl/bench_metal_bug_local.ml, whose findings are why the rule keys on statement shape rather than on anything finer. Neither the pooled slot table nor __restrict nor the placement of the preceding device store is the trigger: dropping every dynamic load (pointers built from literal offsets straight off a kernel parameter) miscompiles identically, and so does moving the preceding store to an unrelated cell. What does remove it, besides the qualifier, is having the accumulating loop read no device pointer at all.

When true: device reads in reduction-shaped scope-local updates and their controlling guards, and in Set statements that read the written node at an index invariant across at least one enclosing serial for loop, use expression-level volatile pointer casts. The cast exists only while rendering the accumulating expression: accumulator declarations, opening reads, vectorized/packed paths and MMA reads stay plain. Pointwise updates and non-accumulator locals stay untouched. Both decisions are reported per routine by the volatility census (volatility_summary).

The form matters: qualifying the accumulator itself cost 1.06x on a memory-bound per-thread reduction, 2.15x on an accumulator-bound dependency chain and 4.1x on a long single-threaded scalar-loss reduction on an M4 Max. The confined volatile-read form measured 1.03x on that loss shape while keeping the standalone reproducer matrix correct row-for-row (bench_metal_bug_local, gh-ocannl-820).

No-alias qualifier for kernel pointer parameters and, in the pooled style, for the derived per-node pointers (restrict / __restrict__ / __restrict); None emits no qualifier. Sound because kernel parameters are buffer-owning roots addressing disjoint (sub-)ranges: alias views are rewritten to parent accesses at assignments lowering and never reach compile_proc's parameter list (asserted there; gh-ocannl-164). The merge buffer stays unqualified — a streaming merge mode could point it at a live same-device buffer.

Lines emitted verbatim before a Vectorized-typed loop's for statement (gh-ocannl-164), e.g. guarded #pragma clang loop vectorize(enable) / #pragma GCC ivdep. An empty list renders the loop as a plain serial for — the legal fallback, mirroring hardware_index = None. Used when explicit vector emission (vector_file.bytes) is disabled or the loop body is ineligible for it.

The vector register file of the explicit-SIMD renderings, its width and its register count as one value (gh-ocannl-1399), so that a backend overriding the width cannot inherit a count of 0 and silently lose the register tile.

The width (vector_file.bytes): vector register width in bytes for explicit SIMD rendering of Vectorized loops via GCC/Clang vector extensions (the Vectorized codegen follow-up of gh-ocannl-164 / docs/proposals/watch-ocannl-README-md-347818d3.md): eligible loop bodies emit vector-typed loads, arithmetic and stores in lanes = bytes / element size chunks plus a serial remainder loop, instead of relying on the compiler's auto-vectorizer (which e.g. cannot reassociate strict-FP reductions — the Vectorized retype carries that permission, like Swap). A recognized accumulation body renders as independent accumulator chains with a horizontal reduce at loop exit (gh-ocannl-468; `Vec_extensions only). 0 disables explicit emission (vectorize_pragma fallback).

The count (vector_file.registers, gh-ocannl-1285): what the register-tiled Tile_mma's column cap Register_tile.rn_cap fits to, the same file the backend reports as hardware_limits.simd_vector so the seeding and this renderer agree. 0 where the width is 0, and on the `Packed_struct backends, which render no register tile.

How eligible Vectorized loops emit explicit vector code when vector_file.bytes > 0. `Vec_extensions (CPU): GCC/Clang vector_size types, unaligned __builtin_memcpy loads/stores, vector-infix arithmetic. `Packed_struct (GPU, gh-ocannl-463; llm.c's Packed128, llmc/cuda_utils.cuh): the backend's vec_typ_of_prec aggregate is loaded and stored through reinterpret_cast at guaranteed-aligned offsets — the 128-bit LDG/STS transactions that bandwidth-bound kernels need — while the arithmetic stays scalar in a per-lane loop over the pack's .v payload (on GPU the payoff is memory transactions, not SIMD ALUs; per-lane fmaf/fma also matches the serial path's rounding exactly). `Packed_struct eligibility additionally requires every vector-accessed node to be materialized (device buffers and pool offsets are Ops.buffer_alignment-aligned, stack and workgroup-shared arrays only element-aligned) and every access's non-loop offset contribution to be a lane multiple.

Declaration suffix aligning stack-allocated local arrays for SIMD access, e.g. __attribute__((aligned(32))) (gh-ocannl-164). Applies to the plain stack-array branch only, never to workgroup-shared placements.

Whether a Local-placement array past the stack budget is declared on the heap (gh-ocannl-1669): the routine's function-scope arrays beyond routine_stack_locals_bytes, a pool-parallel Grid chunk's privatized tiles beyond chunk_stack_locals_bytes, smallest first on the stack. A heap array is malloced (calloced where its declaration zero-initializes) Ops.buffer_alignment-aligned where the stack array would be declared, and freed at the end of the same scope. false (the Pure_C_config default) for the GPU backends, whose thread-private arrays have no heap to go to.

SIMD-group (warp) width for the warp-shuffle rendering of Workgroup_reduce accumulation loops (gh-ocannl-462; llm.c's warpReduceSum/blockReduce idiom). Backends setting this to a nonzero power of two must define ocannl_shfl_xor(value, lane_mask) overloads in their builtins for the supported accumulator precisions (single, and double where it exists), bind workgroup slot 0 in hardware_index, and provide barrier_syntax plus shared_decl_prefix (needed by the two-phase multi-warp form). Those two overloads are all the rendering asks for at any storage width: the shuffled value resides at accum_prec of the storage precision (gh-ocannl-682), and a storage precision whose residency is neither f32 nor f64 is refused — as is an RNG-bearing contribution wherever that residency is wider than storage, since the serial rendering of one keeps a narrow accumulator (Low_level.accum_skip). 0 disables the rendering: Workgroup_reduce loops render like Workgroup — hardware binding, or the serial fallback (which is the correct meaning of a recognized accumulation body on CPU backends).

Cooperative tile-MMA emission for Low_level.Tile_mma (docs/proposals/tensorize-mma.md §4): given the per-operand precisions (the backend decides which combinations its units support — Metal advertises uniform storage triples but can select an f32 accumulator internally for its wide-f16 arm; CUDA wmma's flagship combination is mixed f16×f16→f32), the transposed-storage flags ta/tb (the operand's stored layout is the transpose of its role — load tiles with the hardware transpose flag and swapped offset arithmetic), the covered block extents m/n/k, and per operand its leading-dimension stride in elements, its address space and its physical layout (mma_layout) — plus, for d, a pointer expression to the tile base (already offset) — emit the intrinsic sequence (fragment declarations / loads / mma steps / stores) executed by every lane of the enclosing lane loop. Return None to decline a particular call (unsupported precision combination, extents not multiples of the intrinsic tile, thread-space operand, a swizzled layout the arm has no load form for) — the caller then renders the scalar fallback under an if (lane == 0) guard, which is also the path when the whole hook is None (cc, and any backend until wired).

Acceptance is thus decided without the a/b tile addresses: an accepting arm returns an mma_emission that the caller applies to them once it stands where they are renderable. Callers that only need to know whether a call is supported — the fragment scope deciding whether to alias its accumulator back to the backing target — test the outer option and never apply the emission.

Accepting a `Swizzled_b128 operand is a promise that it was consumed through a swizzle-aware load. mma_operand_loads distinguishes that instruction from swizzle-aware gathers in the rendering census. An arm without a swizzle-aware load must decline the call.

Which of an accepted mma_syntax call's a and b operands load through a warp-cooperative matrix load (ldmatrix) rather than per-lane gathers. Shared-swizzled operands can also use byte gathers, so layout alone cannot establish ldmatrix emission: derive it from the rule the emission itself applies. Read twice, so the two cannot drift: by the rendering census (Mma_intrinsics_ldmatrix when either is true) and, published as Backend_intf.codegen_capabilities.mma_operand_loads, by the cost model's load tie-break (gh-ocannl-1254). Backends without that instruction inherit (false, false).

How an accepted call for these storage precisions, in this emission scope, loads its operand fragments (Backend_intf.mma_fragments); published as Backend_intf.codegen_capabilities.mma_operand_fragments (gh-ocannl-1254). State it beside the loops the hook emits. Default None: undescribed, so the cost model counts no loads.

The arm mma_syntax (Mma_per_statement) or mma_fragment_syntax (Mma_fragment_scope) selects for these storage precisions under the current numerics policy (gh-ocannl-1153), published as Backend_intf.codegen_capabilities.mma_arm and so into the schedule cache's identity. Derive it from the very combination table and floors those hooks dispatch on — a second copy would let a code change to the hooks' arm choice leave the identity alone. Where several arms compete for a triple (CUDA), resolve them once and have the hooks match on the result (Cuda_mma.resolve_arm, gh-ocannl-1232) rather than restate a dispatch order here. None for a combination neither hook has an arm for, and everywhere on a backend without them.

Rendering of a marked cross-reduction accumulator lifetime. The callback receives the backing target and one representative Tile_mma's operands/shape, so it can decline before forcing body. When accepted, forcing body renders its Tile_mma with d identified as `Fragment fragment, allowing the backend to emit update-only MMA steps between one outer fragment load and store.

Only target carries a pointer: it is the one operand this emission addresses (the fragment load and store bracketing the reduction). a and b arrive as mma_source — extents, space and layout to decide acceptance by, no address — because they are addressed by the nested Tile_mmas, each at its own position inside the reduction loop.

Kernel-scope lines — declarations and their initialization — that the tile-MMA emissions of one kernel share, given the (d, a, b) storage triples of the emissions the hooks accepted in it (gh-ocannl-1255). compile_proc places them once, after the local declarations and before the main logic; every emission scope reading them opens with a barrier, which orders the initialization before its reads. Metal's converted destination boundary is the one non-empty case: its coordinate table, the same 0..63 for every scope, lives in threadgroup memory, so one copy per kernel serves all of them. Default: none.

The workgroup-shared bytes mma_kernel_prologue declares for these triples, derived from the same predicate, so Schedule.check_hardware_limits_classified's estimate cannot drift from the kernel. The estimate passes the triples of EVERY Tile_mma statement of the kernel — including those a hook then declines (extents, spaces, layouts), which the renderer does not pass — so it is an upper bound. Default 0.

The granule, in bytes, the device's allocator rounds each shared_decl_prefix array up to (gh-ocannl-1330), published as Backend_intf.codegen_capabilities.workgroup_allocation_alignment for the schedule-level estimate to round each allocation by. Metal: 4 (measured). Default 1: the declared sizes add up exactly.

Kernel parameter for logging, if any. E.g., (Some ("int", "log_id")) or (Some ("const char*", "log_file_name")).

Whether the logging setup involves opening/closing a FILE* (e.g., for fprintf).

Generates a C log statement.

• log_param_c_expr_doc: Document for the C expression of the log parameter (e.g., string "log_id" or string "log_file_name"), if kernel_log_param is Some).
• base_message_literal: The raw, unescaped, unquoted base printf-style format string (e.g., "index %s = %d\n").
• args_docs: Documents for the C expressions of the arguments to the format string. The implementation should handle quoting base_message_literal, choosing the log function (printf, fprintf, os_log), and prepending any necessary prefixes (like a log_id or captured_log_prefix) to the format string and arguments.