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.

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.

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.

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.

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.

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

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.

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.