Parameter Compile_driver.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.

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

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.

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 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 HIP's and Metal's uniform-bf16 tiles accumulate in bf16 fragments, so their serial legs keep bf16 residency. 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 renderings (Vectorized loops) 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 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.

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_bytes) is disabled or the loop body is ineligible for it.

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

How eligible Vectorized loops emit explicit vector code when vector_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.

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 (see accum_pinned_to_storage_prec). 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: the caller records the call as Mma_intrinsics_ldmatrix on that basis. An arm without such a load form must decline the call.

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