Module Ir.Cuda_mma

The inline-PTX arms. Both accumulate in the architected m16n8 f32 per-lane register layout, which is what lets their persistent-fragment scope hold registers across an outer reduction (gh-ocannl-1063).

The nvcuda::wmma C++-API arms (tensorize-mma T3).

Every tensor-core rendering the CUDA hooks can emit. all_of_arm enumerates them, so a table over arms (the arch-target markers, the floors' minimum) covers a new arm automatically.

The arm the hooks render a storage triple with in an emission scope, under the current numerics policy; None: no tensor-core form, the scalar fallback renders. Only the precision- and policy-level choice: the hooks still decline on extents, strides, address spaces, layouts and arm_floor, and a decline is final -- no hook falls through to another arm. An update-only statement against a fragment array is resolved at Mma_fragment_scope, the scope that declared the array. The one scope-dependent triple is uniform f16 under Numerics.Fp16_wide: inline-PTX m16n8k16 per statement, Wmma_f16_wide in a fragment scope; a per-statement arm never has a converted destination.

The arm's arch floor, the minimum compute capability whose PTX contains its instructions: the only place CUDA's mma cutoffs are written.

The marker the arm's rendering comment carries in parentheses (e.g. wmma-tf32, mma-bf16), published as Backend_intf.mma_arm.arm_name; gpu_arch_options reads it back.

The intrinsic tile (m, n, k): 16x8x32 for fp8, 16x8x16 for m16n8k16, 16x16x8 for tf32 and 16x16x16 for the other wmma arms.

The precision the arm's accumulator resides at: f32 for both inline-PTX instructions whatever the storage, the fragment's for wmma.

Which operands the arm loads with the warp-cooperative ldmatrix, given whether each is eligible at all (shared memory in the swizzled 16-byte-unit layout) and stored transposed. fp8 registers are contiguous only in A's row-major and B's transposed orientations; the 16-bit arms use ldmatrix in both; wmma never. Shared by emission and its census, and the per-arm place an operand-load description for the cost model hangs (gh-ocannl-1254).

Whether an operand can be loaded with ldmatrix at all (gh-ocannl-481 item 3): the swizzled 16-byte-unit layout in shared memory, the only combination whose per-lane row addresses ldmatrix can both reach and de-conflict. The emission's eligibility test.

arm_ldmatrix_operands of the scope's arm (resolve_arm) over each operand's ldmatrix_eligible: which operands an accepted call loads with ldmatrix. The backend's C_syntax_config.mma_operand_loads, so the rendering census and Backend_intf.codegen_capabilities.mma_operand_loads — what the cost model's load tie-break reads (gh-ocannl-1254) — both answer with the emission's own rule. (false, false) for a triple with no arm.

Whether the arm's emission reloads each operand's intrinsic fragment for every tile of the other operand's free axis: the inline-PTX bodies load both fragments inside their m x n x k tile loops, wmma loads each fragment once per reduction step and reuses it.

The scope's arm (resolve_arm) as its emission loads operands: arm_tile — the arm's own tile, not the descriptor's (uniform f16 under Fp16_wide emits m16n8k16 per statement under an advertised 16x16x16) — and arm_reloads_operands; None with no arm. The backend's C_syntax_config.mma_operand_fragments, published as Backend_intf.codegen_capabilities.mma_operand_fragments for the cost model's operand-load count (gh-ocannl-1254).

A wmma arm's fragment types and stride constraints; its tile is arm_tile.

Whether a wmma arm renders a statement on a device of compute capability cc: every layout plain, arm_tile dividing the extents, the leading dimensions meeting wmma_info's stride multiples, every addressed operand in device or shared memory, and cc at the arm's arm_floor. d is `Resident for an update-only statement against the fragment array a scope declared (its accumulator is not addressed), and `Addressed (ldd, place) for a self-contained statement's destination or a fragment scope's backing target. The one predicate both of Cuda_backend's wmma sites decide by (gh-ocannl-1456); a self-contained statement additionally never has a converted destination (wc_d_cvt).

wmma_frag_typ kind role elt layout: the nvcuda::wmma::fragment type of role (matrix_a, matrix_b, accumulator) over the arm's arm_tile blocks of elt, with layout for an operand.

The reduction loop of an accepted m x n x k wmma statement, as both of Cuda_backend's hooks render it: load_matrix_syncs of the a and b fragments from __mma_ap / __mma_bp (col_major for transposed storage, ta / tb; tf32 then converts each element) and mma_sync into the mt x nt accumulator array acc. operand_fragments describes its loads, which test/operations/mma_format_arm_tables counts.

An inline-PTX arm's destination element type and its (widen, narrow) boundary conversions (None for fp8's f32 destination, moved directly).

Whether an inline-PTX arm renders a statement's operands on a device of compute capability cc: arm_tile dividing the extents, every operand plain or ldmatrix_eligible, in device or shared memory, and cc at the arm's arm_floor. The one operand predicate both of Cuda_backend's inline-PTX sites decide by (gh-ocannl-1511); the destination's conditions stay with them.

An accepted m x n x k inline-PTX statement as Cuda_backend's mma_syntax renders it (gh-ocannl-1511): its block-comment title and the lines after the __mma_dp / __mma_ap / __mma_bp pointer declarations -- the __mi x __ni x __ki loops, each step gathering (or ldmatrix-loading) an a and a b fragment and issuing one mma.sync. resident names the register fragment array an update-only call accumulates into (None: a self-contained statement that loads and stores d at its storage type, inline_ptx_d_storage). a_swz / b_swz: the operand is ldmatrix_eligible; the arm's arm_ldmatrix_operands decides which of them load through ldmatrix, and the others gather per lane through the swizzled address map. operand_fragments and operand_loads describe its loads, which test/operations/mma_format_arm_tables counts.

The d boundary of an inline-PTX register fragment scope (gh-ocannl-1063): frag is float frag[mt][nt][4], one m16n8 f32 accumulator tile per (__mi, __ni), filled from (`Load) or written back to (`Store) the storage at __mma_dp (leading dimension ldd), converted per inline_ptx_d_storage.

resolve_arm projected to its marker, accumulator and floor (gh-ocannl-1153): the backend's C_syntax.C_syntax_config.mma_arm.

The cp.async staging arm's arch floor (gh-ocannl-487 phase 2): the backend renders it, and capability proposes software-pipelined depths, exactly at and above it.

The tile-MMA descriptor of a device of compute capability cc; None below every arm's floor. Each entry is named by the arm it seeds and takes that arm's floor and tile.

The --gpu-architecture options one nvrtc compile of the given CUDA source gets, given the minimum compute capability across the attached devices: a pure function of the source's arch markers. A FLOOR target (forward-JIT-compiled on every later GPU) is the max of the triggered floors -- each arm's (<arm_marker>) at its arm_floor, nvcuda::wmma at the lowest wmma floor, cp.async. at async_copy_floor, half/bf16 arithmetic at theirs -- raised to compute_75 where every device can load that (CUDA 13 dropped offline compilation below it). The (mma-mxfp8) marker instead selects a compute_<cc>a FAMILY target, only on a device of that family (gh-ocannl-481 item 3, D4). A source with no marker gets no option: nvrtc's own default arch. ~target:Native (default Floor) gives every source the device's compute_<cc> instead, the family marker its compute_<cc>a.

One nvrtc program a backend compiled.

with_ptx_target target f runs f with every CUDA compile in its extent targeting target (current_ptx_target), restoring the previous target afterwards, and returns f's result with the programs the backend compiled in the extent, oldest first: empty on a backend that compiles no PTX, so a test gates its target claims on a compile having been recorded, never on a backend name. Process-wide, not domain-local: for a test comparing the two targets in one process. The CUDA kernel code cache keys on the complete option list, so a kernel compiled at one target is never reused at the other (a cache hit compiles nothing, so records nothing).

The cuda_ptx_target setting: floor (the default) or native (gh-ocannl-1705), which reruns every compile of a process -- a whole suite -- at the native target. Raises Invalid_argument on any other value.

The innermost with_ptx_target's target; outside every extent, configured_ptx_target.

The arch of a PTX module's .target directive, without its , debug suffix.

Called by the CUDA backend with every program's PTX text, as nvrtc produced it; kept only inside a with_ptx_target extent.