Module Ir.Backend_intf

Where the tenant's buffer starts in the arena's pool.

The arena's bump pointer just after the tenant's allocation.

The residency hint the arena's slab was allocated with; only nodes with the same hint share it.

Newest first. Retirement trims the tail until it reaches a tenant still held by a surviving pool owner.

Current maps of upload-lineage leaves owning each pool. References let a linear upload update its owner without adding one; forks retain independently. Lifecycle descendants have a separate table and retain the existing leaf-release precondition.

Threads cooperating in one tile-MMA instruction (CUDA warp / Metal simdgroup width).

The canonical intrinsic tile shape (m, n, k) (8×8×8 for Metal simdgroup_matrix, 16×16×16 for CUDA wmma), used where schedule construction has no typed operand site; a Low_level.t.Tile_mma's block extents must be multiples of the tile of the format actually emitted. Typed matmul/conv sketch seeds use mma_format_tiles below.

Per (a-operand, b-operand, accumulator) format intrinsic tile shapes, for formats whose tile diverges from mma_tile as well as the ones matching it (e.g. CUDA fp8 16×8×32, tf32 16×16×8). Typed autotune seeds use the matching entry for divisibility; whether a given call ultimately emits is still decided by the backend's mma_syntax hook plus the Numerics policy.

The accumulator format is part of the key because it is NOT free to choose: the operand pair that a backend supports against an f32 accumulator is generally not the pair it supports against a narrow one. CUDA is the case that made this explicit (gh-ocannl-545): nvcuda::wmma pairs bf16 operands with a float accumulator only, so keying on the operands alone made the autotuner seed — and time, and rank — 36 candidates per arm on a uniformly-bf16 network that every one of them rendered as the lane-0 scalar fallback.

The emission scopes in which the backend's uniform-f16 arm — an f16-storage destination with f16 operands — holds the accumulator in f32 and converts once at that scope's d boundary, as Numerics.fp16_mode.Fp16_wide requires (gh-ocannl-680, gh-ocannl-836).

CUDA sm_80+ advertises both: Mma_per_statement through its inline-PTX mma.sync.m16n8k16 arm, and Mma_fragment_scope since gh-ocannl-925 through nvcuda::wmma float accumulator fragments whose f16 d crosses a converted boundary addressed by a coordinate table (wmma's element order is unspecified, so no fragment-to- fragment copy). HIP advertises both scopes since gh-ocannl-789 (rocWMMA's (f16, f16, f32) fragments and converted boundary), and Metal advertises both since gh-ocannl-837 (mixed-type simdgroup_multiply_accumulate plus a thread_elements() boundary copy). Withholding only the unsupported scope preserves legal tensorized schedules without letting an outer k split introduce extra f16 narrowing boundaries.

The bf16 twin of mma_f16_wide_acc_scopes: the emission scopes in which the backend's uniform-bf16 arm holds the accumulator in f32 and converts once at that scope's d boundary, as Numerics.bf16_mode.Bf16_wide requires (gh-ocannl-838). HIP advertises both scopes (rocWMMA's (bf16, bf16, f32) fragments behind the same converted boundary as the wide f16 arm); CUDA sm_80+ both (its inline-PTX m16n8k16 arm, whose accumulate is f32 in hardware, per statement, and since gh-ocannl-1063 that arm's per-lane registers held across the outer reduction); Metal both (a simdgroup_float8x8 accumulator over bfloat operand fragments behind the wide-f16 arm's thread_elements() boundary, gh-ocannl-923).

Format triples whose cooperatively staged operand tiles the backend can read in a non-row-major layout, and which layout (gh-ocannl-481 item 3, D3). Autotune's staged mma sketches seed a swizzled twin per staged seed exactly for the advertised triples — the tuner, not a heuristic, then decides whether the bank-conflict fix beats the plain tile.

Keyed by format triple for the same reason as mma_format_tiles, and pre-filtered for the same reason (gh-ocannl-479): eligibility is per operand AND per orientation, and the orientation the staged sketches mint is each role's own. CUDA's fp8 arm, for instance, feeds A from ldmatrix in that orientation; B's four bytes per register are strided, so its row-major staged tile uses swizzle-aware byte gathers instead (gh-ocannl-1073). Both operands tensorize, allowing the advertised swizzled twin. Empty everywhere the question does not arise.

Software-pipelining depths beyond the unpipelined 1 that autotune's staged mma/conv sketches propose as twins of each staged seed (Schedule.Stage ~pipeline_depth, gh-ocannl-487) — a list, not a flag, so the search has a dimension. The portable double-buffered rendering is backend-generic, but a depth is advertised only where the arm has been validated on hardware: Metal ([2], phase 1, portable form) and CUDA ([2] on sm_80+, phase 2, the cp.async arm); HIP stays empty until its LDS async-copy arm lands. Empty on CPU backends (cooperative staging is not renderable there).

Maximum number of distinct device-memory pools one routine may bind. None when the backend imposes no pool binding budget. Counts pools, not tensor nodes: host uploads and device copies into one context share arenas. Metal enforces this at link time.

Upper bound on the number of threads in one workgroup (CUDA thread block / Metal threadgroup); None when the backend imposes no limit (the C backends render annotated loops serially).

Capacity in bytes of the workgroup-shared memory (CUDA __shared__ / Metal threadgroup); None when the backend imposes no limit.

The per-work-item budget for one kernel's declared scratch arrays: the largest private (stack) frame the backend's compiler accepts, less a reserve for what else lands on the frame (gh-ocannl-1313). The kernel's routine-scoped scratch, the function-scope arrays of its Local-placed nodes (Schedule.thread_private_bytes), is checked against it pre-driver by Schedule.check_hardware_limits_classified, as Schedule_outcome.Thread_scratch. Without the check, a kernel over the limit reaches the device compiler, which refuses it untyped after compiling the whole kernel: HIP's hiprtc took 52 s to report stack frame size (360544) exceeds limit (262136) on the serial base form of a GPT-2 training step, whose 89 Local arrays declare 360452 bytes, and under strict failure classification that ended the search (outcome pre-search-failure). HIP: LLVM AMDGPU's frame limit of 262136 bytes, on every AMD part seen (gfx1102, gfx1151), less a 1 KiB reserve (that kernel's frame held 92 bytes beyond its arrays; spills add more), independent of the post-link scratch-budget check that HIP applies to what does compile. None where unmeasured, which checks nothing, and on the C backends. Omitted from the record's sexp when None, so adding it moved no other backend's schedule-cache key.

Per-dimension upper bounds on the launch's workgroup shape — the caps on .x, .y and .z of Low_level.launch_dims' block, in that order. Beside, not instead of, max_threads_per_workgroup: that one caps the product, and the two are not the same fact. CUDA's maxThreadsDim is (1024, 1024, 64) — the .z component is 16x smaller than the product cap — so a workgroup of 2 x 2 x 128 has a perfectly legal 512-thread product and is still an invalid launch configuration (gh-ocannl-679). Workgroup slots are capped at 3 and the innermost binds .x, so the outermost annotated loop's extent lands on .z directly; no fold is involved.

Unlike max_grid_yz this carries all three bounds rather than one shared one, because here the dimensions genuinely differ: on CUDA .z is the odd one out, while HIP (queried max_threads_dim) and Metal (maxThreadsPerThreadgroup's three components) report three equal values. None on the C backends, which render annotated loops serially.

A tuple rather than an int array, on two independent grounds. It is the shape mma_tile already uses for a 3-D quantity here; and it is immutable, which this record needs from every field it has. The GPU backends memoize their hardware_limits behind a lazy and Context.hardware_limits returns that record itself, so one mutable cell anywhere in it would let a caller deriving tighter limits for a custom schedule write through into the process-wide singleton — after which compiles reject legal kernels or admit illegal ones, with nothing to point at. A tuple also makes the arity exactly the three Workgroup slots, so no reader has to bounds-check a length the type does not promise.

Enforced pre-driver by Schedule.check_hardware_limits_classified (as Schedule_outcome.Workgroup_x_extent / _y_ / _z_extent); Schedule.default_gpu and Schedule.zero_expansion clamp their block size against the .x entry too, and check the .y/.z entries a widened workgroup reaches, so the gate is a backstop rather than the first line of defence.

Upper bound on each of the launch's .y and .z grid dimensions: the row-block count (grid.(1)) and the folded batch-extent product (grid.(2), the dimension Grid slots >= 2 fold onto — gh-ocannl-643, Low_level's hardware-axis section comment). One field rather than two because it is one hardware fact: CUDA and HIP cap gridDim.y and gridDim.z at the same 65535 while gridDim.x is 2^31-scale (on CUDA an architectural constant, unlike the queried per-device limits above; HIP queries it, conservatively as the smaller of the two components). What a caller does about an excess differs per dimension, and that distinction lives in the typed cause instead (Schedule_outcome.Grid_y_extent vs. Grid_z_extent). None when the backend imposes no such limit (Metal's threadgroups-per-grid dimensions are not 16-bit, and the C backends render annotated loops serially). Both dimensions are checked pre-driver by Schedule.check_hardware_limits_classified; the autotune batch-grid twins also consult the .z reading at seeding so an over-cap candidate is never proposed.

Tile-MMA units (simdgroup_matrix / tensor cores); None when the backend has none wired — Tile_mma statements then render their scalar fallback.

The vector register file of the C backends' explicit vector-extension renderings (Vectorized loops, the register-tiled Tile_mma micro-kernel): its width in bytes and its register count, one value so that neither can be stated without the other (gh-ocannl-1399). Vector_file.none when the backend does no such rendering (GPU backends bind hardware axes instead). Carried here so schedule construction (autotune's seeding pre-filter, gh-ocannl-479) can statically rule out candidates the renderer must decline, e.g. a micro-kernel column extent below one vector's lane count.

The count is 32 on AArch64 and on AVX-512 targets (at every width, since AVX-512VL gives the narrower forms the same 32), 16 on other x86 targets (SSE, AVX2) -- the target the kernels are compiled for, probed under their flags. It is not a function of the width: a 16-byte file is 32 registers on NEON and 16 on SSE, and Register_tile.rn_cap fits the register tile's columns to the count (gh-ocannl-1285), so the seeding and the renderer must read the same file.

Advisory peak arithmetic throughput in FLOP/s (single-precision, FMA counted as two), the hardware envelope of the analytic cost model (gh-ocannl-491): rough documented constants or cheap device queries — the model ranks candidate schedules, it does not predict runtimes. Never gates compilation and never overrides a measured timing; None when the backend offers no estimate. Known single-precision bias (gh-ocannl-575): a pure-fp16 kernel on a native_fp16_arithmetic target has twice this ceiling, so its roofline flops leg over-estimates — harmless for ranking because a site's candidates all share one policy-resolved compute precision (footprint widths, by contrast, are exact: they come off each node's own storage precision).

Advisory peak main-memory bandwidth in bytes/s, the other leg of the roofline envelope (gh-ocannl-491). Same contract as peak_flops — advisory, rough, never load-bearing for correctness; None when the backend offers no estimate — with one bias requirement (gh-ocannl-578): the value must be a class ceiling, at least what any machine of the backend's class can sustain. Streaming kernels with exact byte counts are routine now (the calibration pass, packed initializations), and each one achieving more than the advisory trips the gh-514 agreement warning — while under autotune_bound_pruning an understated leg over-prunes. Overstatement only loosens an advisory bound; calibrated model_peak_* overrides beat these wherever fidelity matters.

Untimed (see limits_key): the cache capacity of the cost model's ranking score (gh-ocannl-1253, Ir.Cost_model.ranking_seconds) — once a kernel's working set, per-block global read bytes times resident blocks, exceeds it, global reads are charged per block rather than as distinct cells. A FITTED constant of that proxy, not the device's L2 size: the working set sums per-block footprints that resident blocks share, so the value that separates cache-resident from DRAM-bound kernels sits well above any physical cache. Per-class advisories on CUDA and HIP, and a fitted one on the C backends (gh-ocannl-1687, Schedulers.cpu_mma_limits), overridable by config model_cache_capacity; None elsewhere: no DRAM tier (see cache_bandwidth). Read only by the untimed choosers, never by schedule construction, codegen or timing — so it is not in the schedule-cache key (limits_key).

Untimed: how many threads the whole device keeps resident at once (multiprocessors times threads per multiprocessor), which divided by a kernel's workgroup size is the ranking score's resident-block count. None where unknown: every block then counts as resident. Same readers and same cache identity as cache_capacity_bytes.

Untimed: the rate, in bytes/s, at which the cache serves re-reads within capacity — the ranking score's second tier (gh-ocannl-1253): cross-block re-reads of a kernel whose working set fits cache_capacity_bytes are charged at this rate rather than for free; it is the requested-rate plateau of the capacity sweep below the capacity (tools/calibrate_cache_capacity.exe). Overridable by config model_cache_bandwidth; None where unmeasured, which leaves re-reads within capacity uncharged. Same readers and same cache identity as cache_capacity_bytes.

Untimed: the effective scalar operation rate, in operations/s (summary.flops' units, an FMA counting two), of a register-blocktiled kernel, its operand staging included — the ranking score's compute leg for scalar work (gh-ocannl-1442, Ir.Cost_model.ranking_seconds), where peak_flops stays the bound's ceiling; accepted Tile_mma work is priced at mma_rate. The slope of tools/calibrate_issue_rate.exe's sweep of the seeded weight-gradient sketch — or, where fill_threads is set, the full-fill rate its launch sweep fits jointly with it (gh-ocannl-1479): the two are one calibration, and a pre-fill rate under a fill model double-counts the calibration kernel's partial fill. Read as at most peak_flops. Overridable by config model_issue_rate; None where unmeasured, which prices scalar work at peak_flops as before. Same readers and same cache identity as cache_capacity_bytes.

Untimed: the DRAM rate kernels achieve, in bytes/s — the ranking score's DRAM tier and distinct-byte charge (gh-ocannl-1442), where peak_memory_bandwidth stays the bound's class ceiling. The STREAM fit of tools/calibrate_bandwidth.exe. Read as at most peak_memory_bandwidth. Overridable by config model_dram_bandwidth; None where unmeasured, which prices DRAM traffic at peak_memory_bandwidth as before. Same readers and same cache identity as cache_capacity_bytes.

Untimed: the effective rate, in operations/s (an FMA counting two), of Tile_mma work the backend's mma_accepts hook accepts — the ranking score's compute leg for work on the matrix units (gh-ocannl-1495, Ir.Cost_model.ranking_seconds). A FITTED constant of the emitted tensorized kernels, operand staging and shared-memory loads included, not the units' datasheet peak: the geometric mean of the rates the staged tensorized leaves of a site achieve (benchmarks/gh1442_agreement.py --fit-mma). Read as at most peak_flops. Overridable by config model_mma_rate; None where unmeasured, which prices tile-MMA work at peak_flops as before. Same readers and same cache identity as cache_capacity_bytes.

Untimed: how many per-lane operand gather instructions a lane group issues per second from DEVICE memory — the rate of the ranking score's device-gather term (gh-ocannl-1495, Ir.Cost_model.device_operand_gathers): a Tile_mma operand the arm gathers element by element straight from global memory (an unstaged inline-PTX operand) costs its count over this rate, at the Tile_mma contexts' fill. Its unit is the operand-load count's (executions per lane group times instructions per lane). Shared-memory gathers, ldmatrix and whole-fragment loads are not counted. A FITTED constant: the rate that makes a site's unstaged tensorized leaves' excess over mma_rate come out as measured (benchmarks/gh1442_agreement.py --fit-mma). Overridable by config model_device_gather_rate; None where unmeasured, which adds no term. Same readers and same cache identity as cache_capacity_bytes.

Untimed: the device's multiprocessor (SM / CU / GPU core) count — the P of the ranking score's launch fill (gh-ocannl-1479, Ir.Cost_model.context_fill). Queried on CUDA and HIP; on Metal, which publishes none, read from the IORegistry's gpu-core-count (gh-ocannl-1638; None on a non-AGX GPU); None elsewhere. Overridable by config model_multiprocessors. The launch fill applies only when both this and fill_threads resolve. Same readers and same cache identity as cache_capacity_bytes.

Untimed: the threads per multiprocessor at which a register-blocktiled kernel's scalar issue saturates — the t of the launch fill (gh-ocannl-1479). A FITTED constant of the launch sweep of tools/calibrate_issue_rate.exe, jointly with issue_rate, which it normalizes (the rate is then a full-fill rate); not the physical residency. Per-class advisories where measured, overridable by config model_fill_threads; None elsewhere, which leaves the compute leg as before. Same readers and same cache identity as cache_capacity_bytes.

Untimed: how many excess 32-byte memory sectors a device serves per second -- the rate of the ranking score's memory-transaction term (gh-ocannl-1600, Ir.Cost_model.memory_transactions): per launch block (its lane groups of simdgroup_width lanes sharing what L1 serves), the sectors a global access touches beyond the compact count for the same distinct bytes, at the access's own context fill, added to the compute leg. What orders a layout whose lanes read contiguous cells against a twin whose lanes stride. A FITTED constant (benchmarks/gh1600_txn_fit.py), per device class. Overridable by config model_transaction_rate; None where unmeasured, which adds no term and keeps the coalesced layouts out of model_default's candidate set. Same readers and same cache identity as cache_capacity_bytes.

Untimed: the workgroup-shared memory one multiprocessor (SM / WGP) holds for all its resident blocks -- the capacity of the ranking score's occupancy factor (gh-ocannl-1600, Ir.Cost_model.occupancy): a kernel declaring W workgroup bytes in blocks of T threads keeps min (capacity / W) (threads per multiprocessor / T) blocks resident, and its compute leg is divided by the fraction of the multiprocessor's thread slots they fill (the threads per multiprocessor are max_resident_threads over multiprocessors). Also the second half of model_default's gate on the coalesced layouts, beside transaction_rate: their deep-K tiles are told apart only by it. The physical value where advisory (HIP's RDNA wave32 parts: 128 KiB of LDS per WGP, the unit HIP counts as a multiprocessor); None where the occupancy does not explain the measured order (CUDA, gh-ocannl-1600) or is unmeasured. Overridable by config model_workgroup_bytes_per_multiprocessor. Same readers and same cache identity as cache_capacity_bytes.

Whether 16-bit float arithmetic executes natively at twice f32's lane count (gh-ocannl-516: ARMv8.2-FP16, AVX512-FP16). false covers both "no _Float16 on this target" and the middle case that matters for ranking: the type exists and computes correctly, but the compiler implements it by promoting to float, so the lane count does not double and candidates must not be seeded as if it did. Whether the type exists at all is a separate, purely textual question the emitted C answers for itself (HAS_NATIVE_FLOAT16); this field is about throughput.

Always false on the GPU backends, whose 16-bit story is their native types and tensor-core shapes rather than a CPU vector width.

The SIMD-group (warp) width the backend's warp-shuffle renderings assume -- its C_syntax configuration's warp_size: 32 on Metal, CUDA and HIP (whose shuffles pass an explicit width of 32). None where kernels render no shuffles (the C backends). The default GPU schedule retypes a cooperative lane reduction only at a multiple of this width, at most lane_all_reduce_simdgroups of them (gh-ocannl-1124, gh-ocannl-1168). Where a transaction_rate is set, also the execution width the memory-transaction count groups lanes by (gh-ocannl-1600): a backend sets that rate only on parts that execute this width (HIP: the RDNA wave32 parts).

Device economics (gh-ocannl-1124): whether scalar work every lane of a lane geometry recomputes redundantly -- the per-pair preamble a lane nest repeats in each of its lanes (the fused attention backward's p, dp, ds) -- is cheap relative to a thread's serial walk over the channel loop that the plain plan pays it once for. Measured: true on Metal (M4 Max) and CUDA (RTX, rog-nv), whose dK and dQ nests on lanes cut the D1 training step by 2-9%; false on HIP (gfx1151), where the same lanes cost 7-20% against the plain dK plan. false -- today's behaviour, the plain plan -- wherever unmeasured, and on the C backends, which run no lane geometry. Read by the default GPU schedule's gpu_lane_preamble_reduction = auto and its gpu_serial_lanes = auto (gh-ocannl-1167): where false, the stage-1 lanes of gh-ocannl-1003 are still taken in a kernel that admits them (the forward value pass won on gfx1151), but are no reason to cut a kernel (the fused backward's dV cut from dK for them lost to the merged plain kernel).

Device economics (gh-ocannl-1168): the most simdgroups a cooperative lane all-reduce is worth spanning. At one simdgroup the all-reduce is a register butterfly; across k > 1 it stages each simdgroup's partial in workgroup-shared memory between two barriers per (query, key) pair, which may eat the lanes' win -- so a head wider than simdgroup_width (GPT-2's 64) gets lanes only where that was measured to pay. 1 (one simdgroup, the gh-ocannl-1124 rule) wherever unmeasured, and on the C backends, which run no lane geometry. Measured: 4 on Metal (M4 Max; the D1 training step at head width 64 0.97x / 0.92x / 0.82x at seq 128 / 512 / 1024, at width 128 0.69x at seq 512, against the plain plan: benchmarks/gh1133_cells.sh d1ar4 vs d1ar1); 4 on CUDA (RTX, rog-nv; at head width 64 0.84x / 0.52x / 0.41x at seq 128 / 512 / 1024, at width 128 0.37x at seq 512, the same A/B, gh-ocannl-1256). Read by the default GPU schedule's gpu_lane_all_reduce_simdgroups = auto.

Automatic key-block size for online_softmax_block=auto (gh-ocannl-1171); 0 selects the two-pass rewrite. Metal and CPU keep 16 from the gh-ocannl-1003 block sweep (benchmarks/report-gh1003-block-fold.md). CUDA's tuned f32 transformer measured 4.49 ms at 16 versus 2.44 ms at 0 with tf32 retained (gh-ocannl-1194); CUDA and unmeasured devices conservatively select 0. HIP gfx1151 unified measured 7.22 ms at 16 versus 4.72 ms at 0; other HIP and unmeasured targets conservatively select 0. Explicit numeric sizes bypass this policy. Consumed before lowering analyses, so the resolved choice is in the code digest.

Compact signature of the worker pool timings execute on (w8P, w24, ...), filled by the CPU backends from the pool-uniformity policy (gh-ocannl-530). Enters the autotune disk-cache key the way the numerics tag does: schedules crowned on one pool do not transfer to another, so a policy flip or a different external pinning must re-tune rather than replay. None (GPU backends) leaves the cache key unchanged.

Compact signature of this backend's codegen configuration: the settings the backend consults when rendering and compiling a kernel, which are therefore invisible to the canonical digest of the lowered code (gh-ocannl-572). Same contract as worker_pool_tag — it enters the autotune disk-cache key, so a knob flip re-tunes instead of replaying a winner crowned in another codegen regime, which is the hazard gh-ocannl-568 measured at 5.9x. Fill it from resolved values, not raw settings: what "auto" resolves to is a per-machine fact, and crowns do not transfer across machines either. None where the backend has no such knobs.

The arm's own marker, as its rendering comment carries it (e.g. wmma-tf32, mma-bf16): which arm, not only what it accumulates at.

The precision the arm's accumulator resides at.

The device-architecture floor the arm itself checks before emitting, in the backend's own numbering (CUDA compute capability). None where the backend gates tensor units at the device level only (hardware_limits.mma).

The intrinsic tile (tm, tn, tk) the arm emits: an a fragment is tm x tk, a b fragment tk x tn.

true: both fragments are loaded inside the m/tm x n/tn x k/tk tile loops, so each is reloaded for every tile of the other operand's free axis (CUDA's inline-PTX arms); false: each fragment loads once per reduction step and is reused across that axis (CUDA's wmma arms).

true: each fragment not loaded through ldmatrix comes in whole, by one cooperative matrix load (CUDA's wmma load_matrix_sync); false: by per-lane element gathers (the inline-PTX arms).

Whether the backend dialect can represent f64 tensor storage. This is explicit rather than an exception probe against typ_of_prec.

The resolved compute precision for a storage precision, from the same C_syntax_config.compute_prec function code generation uses.

The resolved accumulator precision for a storage precision, from the same C_syntax_config.accum_prec function code generation uses.

Whether eligible pipelined staging copies use a dialect-specific asynchronous copy arm. A portable synchronous depth-2 pipeline is not this capability.

Whether the backend's kernel compile pins contraction off, from the same C_syntax_config.compiler_contraction_off code generation reads (gh-ocannl-1457): then a Tile_mma fallback, and the register tile, keep the serial loop's two roundings. Read per call, since the settings behind it are mutable within a process.

The arm the backend's mma_syntax (for Mma_per_statement) or mma_fragment_syntax (for Mma_fragment_scope) selects for these storage precisions, derived from the same combination table those hooks consult; None where the table has no arm for them. Only the precision- and policy-level choice: extents, strides, address spaces, layouts and the device can still make a call decline.

Which of the a and b operands the backend's mma arm for these storage precisions, in this emission scope, loads with a warp-cooperative matrix load (ldmatrix on CUDA) rather than per-lane gathers, given each operand's place and whether it is stored transposed (gh-ocannl-1254). The very rule the emission and its rendering census apply (CUDA: Cuda_mma.operand_loads over Cuda_mma.arm_ldmatrix_operands), so a cost model reading it reads the decision the kernel makes. Like mma_arm, only the arm- and operand-level choice: a call the hook declines on extents or the device renders the scalar fallback whatever this says. (false, false) on a backend without such a load.

Whether the backend's mma hook accepts a statement with these storage precisions, orientations, extents and, per operand, leading dimension and place — asked of the hook itself (a C backend's C_syntax_config.mma_syntax, which decides acceptance without the operand addresses), so it holds every condition the emission checks: layouts, spaces, strides, extents, the device's floor (gh-ocannl-1254). With fragment_target — the backing target of a persistent accumulator, its leading dimension and place — the question is the fragment-scope rendering, in the renderer's order: the enclosing mma_fragment_syntax hook accepting that target and the operands, then the update-only statement (d's own leading dimension) against the fragment array. false on a backend without the hook.

How the backend's mma arm for these storage precisions, in this emission scope, brings its operands into registers: its intrinsic tile and whether it reloads each operand fragment for every tile of the other's free axis (gh-ocannl-1254). The arm's own description, read by the cost model's operand-load count, which has no count where this is None: no arm, or a backend that does not describe its emission (the default).

Static workgroup-shared bytes a kernel declares beyond the staged tiles for its tile-MMA emission scopes, given their (d, a, b) storage triples: the backend's C_syntax_config.mma_kernel_workgroup_bytes, which it derives from the same predicate as the kernel-scope declaration (C_syntax_config.mma_kernel_prologue). A per-KERNEL term, not a per-scope one: non-zero only for Metal's converted destination boundary, whose coordinate table is declared once per kernel however many scopes read it (gh-ocannl-1205, gh-ocannl-1255). Schedule.check_hardware_limits_classified adds it once to the staged tiles.

The granule the device's allocator rounds each workgroup-shared array up to, in bytes: the backend's C_syntax_config.workgroup_allocation_alignment. Metal rounds every threadgroup array to 4 bytes (measured on an M4 Max: two odd-length half tiles whose exact bytes total the limit are refused, the driver requesting 4 bytes more, gh-ocannl-1226), so Schedule.workgroup_memory_bytes rounds each allocation it counts up to this before summing (gh-ocannl-1330); 1 where the sum of exact sizes is what the device allocates.

Static workgroup-shared bytes one lane all-reduce site declares for its per-simdgroup partials (gh-ocannl-1168): a reduction of extent lanes into a scope local residing at prec, rendered across more than one simdgroup. The renderer (C_syntax.try_lane_all_reduce) decides by this same field, so the two cannot drift; 0 where the site renders within one simdgroup, serially, or on a backend without shuffles, barriers or shared declarations. Only the backend-level half: whether a site may stage partials at all is Low_level.lane_all_reduce_sites' structural predicate, which Schedule.workgroup_memory_bytes reads beside it.

Static workgroup-shared bytes one warp-shuffle reduction into a cell declares for its per-warp partials (gh-ocannl-1452): a Workgroup_reduce accumulation of extent lanes whose shuffles stage at the accumulator residency prec, rendered across more than one warp (C_syntax.try_warp_reduce's wred_partials_ array, one slot per warp). The renderer checks its declaration against this same field, so the two cannot drift; 0 where the reduction fits one warp, or where the multi-warp form cannot render (no shuffles, barriers or shared declarations, an extent not a power-of-two count of at most warp-size whole warps, a narrow residency) — the renderer refuses those loudly, so no kernel declares them. Only the backend-level half: which statements are such reductions is Low_level.warp_reduce_cell_sites' structural predicate, which Schedule.workgroup_memory_bytes reads beside it.

A node can be placed in workgroup-shared memory (the shared epilogue placement of Schedule's Fuse_epilogue): the premise of Schedule.may_keep_off_device. CPU codegen rejects a non-empty workgroup_shared.

The untuned compile seeds its plan (Autotune.seeded_default, gh-ocannl-1183): the backends that plan's before/after was measured on. Also scopes the seeded plan's numerics gate in Schedule_cache.codegen_tag (gh-ocannl-1405).

A queued timing window refused as contended is retried once: the device scheduler's transient bursts (Metal).

A dispersed timing window is attributed through the host-load signal rather than refused outright (gh-ocannl-1558): the device's co-tenants are host processes, which host CPU accounting sees.

The dialect's tensor-node identifier prefix (C_syntax_config.node_ident_prefix, gh-ocannl-1338): "t_" where the vendor headers declare global names outside the reserved namespace.

The extension, dot included, of the routine source a backend writes as its build_files/ artifact: what each backend's artifact writer appends, and what test/config/config_reads.ml's code_extension -- linked by tests that link no backend -- is held equal to by test/operations/marker_backend_vocabulary.

The dialect renders a value select that Low_level.speculable_select admits as a branchless blend, evaluating every arm (C_syntax_config.select_blend_syntax, gh-ocannl-1622, gh-ocannl-1695): cc and HIP, at the float compute precisions their blends cover. CUDA's ptxas if-converts the short-circuiting spelling itself and Metal keeps ?:. The renderers derive their blend from this fact, and the inline fan-in guard reads it before codegen (gh-ocannl-1775): on a blending backend a speculable Where is charged the union of its arms' loads, elsewhere its wider arm.

field-hardware_limits.issue_rate.

field-hardware_limits.dram_bandwidth.

field-hardware_limits.mma_rate.

field-hardware_limits.device_gather_rate.

field-hardware_limits.multiprocessors.

field-hardware_limits.fill_threads.

field-hardware_limits.transaction_rate.

field-hardware_limits.workgroup_bytes_per_multiprocessor.

The group atom naming the dump, e.g. "cuda_devices".

One (key, value) assoc per device, in ordinal order.

The materialized read-only and read-before-write (within the routine) non-constant nodes. They are inputs in a broad sense, as they could be recurrent nodes or parameters.

Similar to inputs, for the merge buffer. The execution-dependency ledger consumes this as a read edge on the transfer that filled the transient slab, but it is not an ordinary context input requiring initialization.

All the materialized nodes written-to by the routine.

The number of the represented backend's device, in the range from 0 to the number of the backend's devices - 1.

A unique identifier among all device instances of all backends. Note that multiple device_id (distinct device instances) might refer to the same physical device.

The merge buffer's reserved single-tenant pool location, or None if not yet allocated. The slab can be reused (grown in place) for nodes that fit.

Byte capacity of the reserved merge-buffer pool; drives the grow decision.

The completion event for the most recent updating (writing to) a node via this device, while it is not done: each new write drops the done entries first (gh-ocannl-1747), so the table does not keep the nodes of graphs nobody holds.

The tensor node that was most recently scheduled to be in the device's merge buffer. See also updating_for.

Per-device cache for read-only/constant buffer allocations: device-wide for host literals, per root otherwise (Constant_key). Either kind's pool is freed, and its entries dropped, when no live release unit binds it (constant_owners); what is left at device teardown is freed there.

For each constant pool (one holds either Rooted entries of one root or Shared ones, never both, and nothing else), how many live release units bind it (field-context.held_constant_pools), across roots for a Shared pool; the release taking it to 0 frees the pool and drops its cache entries (gh-ocannl-1654 for Rooted, gh-ocannl-1762 for Shared: dropping a literal's strong Tnode.t key is what lets its node, and through Host_inits its host array, be collected). Guarded by constant_mutex.

The lock of constant_owners and of the cache entries a release drops.

Deterministic per-device pool-id counter, advanced by the shared allocator seam in tnode iteration order. Pool id 0 is reserved for the merge buffer; tnode pools start at 1.

This map contains the deterministic buffer locations used in this context or an ancestor context.

Pools this context has already released. Retained across a failed-finalize retry so a cleanup that freed some pools before raising never calls the backend free twice.

Latest map in this linear upload chain. Uploading from an older map forks its ownership into an independent leaf; linear uploads supersede their argument without an extra owner.

This upload lineage's arenas and pool reference counts, fresh in Device.make_child.

What a release does besides freeing pools: drop device resources a routine linked into this context retains across runs (the CUDA/HIP graph execs of a graph-captured segment batch, gh-ocannl-1445). Run by every finalize of a value that shares the list -- a linear upload chain and its upload forks -- so each hook must be idempotent and must leave its routine runnable (rebuilding what it dropped), since a fork can outlive the value it came from. Fresh in Device.make_child.

The optimization context threaded through compilation: all OCANNL backends compile through the Low_level IR, so this is concretely Low_level.optimize_ctx (the abstraction for hypothetical assignments-level backends was retired; the Assignments.comp -> code seam can be reintroduced if such a backend ever materializes).

The constant pools this release unit binds -- its root's Rooted ones and the Shared host literals it resolved -- each counted once in field-device.constant_owners when the unit was created and released when it is finalized. Shared by every value of the unit (record copies keep it).

The independently created context this one derives from, as far as constants go: minted by Device.make_context, inherited by every derived context. Keys the non-literal entries of field-device.constant_buffer_cache (Constant_key).

The tensor node that a Backend.device_to_device transfer with into_merge_buffer:Copy placed (or will place) into this context's device's merge buffer. It is a static, immutably-chained fact carried producer -> consumer: linking a consumer whose code expects a merge-buffer node verifies it against this field at link time. A transfer with into_merge_buffer:No does not touch the merge buffer and inherits the parent's value.

Bytes of the context-owned working pools.

Bytes of the new per-device constant pools.

How many pools, working and constant.

Whether the liveness planner (buffer_aliasing, gh-ocannl-489) found nodes it may overlap. The delta then depends on the kernel sequence, not only on the node set: a segment is the GPU planner's position, so one kernel overlaps nothing that a fissioned schedule of the same lowering can.

The floor's working bytes.

The floor's new-constant bytes.

How many pools the probe requests: each group's bytes cut at the pool cap.

Working bytes left out because some schedule can keep the node off the device (Schedule.may_keep_off_device: the shared epilogue placement).

New-constant bytes left out because a hoisted Stage can replace the node by a packed tile (Schedule.may_replace_by_packed_tile).

Under buffer_aliasing: bytes of overlap-eligible nodes past the largest of each precision class, which some kernel sequence may overlap.