Module Ir.Backend_intf

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.

Whether the backend's uniform-f16 arm — an f16-storage destination with f16 operands — can hold the accumulator in f32 and convert once at the d boundary, which is what the Numerics.fp16_mode.Fp16_wide policy requires of every rendering (gh-ocannl-680). CUDA's inline-PTX mma.sync.m16n8k16 arm can (sm_80+, the architecturally-defined fragment layouts are shared by .f16 and .bf16); HIP can since gh-ocannl-789 (rocWMMA's (f16, f16, f32) fragments, the boundary converted elementwise through a destination-typed accumulator fragment); Metal cannot (simdgroup_matrix is uniform-precision only, structurally). Where false, the seeding gate in Sketch_families.mma_tile_for_precisions withholds uniform-f16 seeds under the wide policy — the serial legs then carry the f32 residency via accum_prec, keeping the width schedule-uniform per the gh-ocannl-545/663 discipline.

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, can feed A from ldmatrix in that orientation but not B — 4 fp8 bytes of a B register are strided there — so a swizzled fp8 twin would be timed and ranked as a tensorized candidate while rendering the scalar fallback. 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).

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.

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

Vector register width in bytes used by the C backends' explicit vector-extension renderings (Vectorized loops, the register-tiled Tile_mma micro-kernel); 0 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.

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.

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.

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

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.

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.

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