Wrappers

For ops whose operand semantics break the chain default, PTX.jl ships hand-written wrapper methods on the matching Operation{...} singleton. The user-facing call site is identical to the chain default — the wrapper just provides a typed method that takes priority over the @generated chain dispatch when its argument types match.

There are four patterns that force a wrapper:

1. Mixed address-space pointer constraints. Shared-AS pointers need the r (32-bit) constraint; the chain emits l (64-bit) for any LLVMPtr. cp.async, ldmatrix, stmatrix, mbarrier use this override to satisfy ptxas.
2. Multi-output return. The chain returns one value; ldmatrix x2/x4, mma fragments, and shfl-with-pred-output all return tuples.
3. Special operand layout. Braced register-vector operands ({$N, $N+1, ...}), tensor-coord forms ([ptr, {c0, c1, c2}]), tied accumulators for wgmma — none have a chain-default rendering.
4. Compile-time-constant operand types. TMA's <N>d rank, mma shape/dtype, fragment counts — these need to be pinned at the registration boundary so each variant gets its own typed method.

Each wrapper file follows the same shape: a small declarative table maps (shape, dtype, …) to register counts and constraints, a _register function builds the family-local spec (asm template, llvmcall body, or intrinsic name) and emits a typed method per valid combination from a for loop. Bookkeeping is shared, not per-file: every generator records what it defined in the wrapper registry (src/wrappers/registry.jl) — one WrapperRecord per (family, op, mods, tier, intrinsic) key behind a single idempotency set, so re-invoking a registration helper cannot inflate the inventories the test count pins protect. Tier-2 generators obtain their IntrinsicCall through wrapper_intrinsic_call, which validates the generated name against the pinned NVVM registry and records it in one step; asm/core-IR generators call register_wrapper! directly. The test-side oracles read the inventories back through wrapper_records / wrapper_intrinsic_names / wrapper_asm_forms / wrapper_missing_intrinsics.

Literal spellings are declared with the @optype_str definition macro (@inline optype"tcgen05.alloc.cta_group::1.sync.aligned.b32"(...) = ...); only generator loops that build the mods tuple programmatically define methods on the Operation{op, mods} singleton directly.

Adding a new dtype/shape combination is one entry in the table plus one line in the loop.

Family overview

Schema-driven structured results

setp, lop3, match.sync, and elect.sync no longer rely on a small set of hand-written dual-result methods. src/ledgers/structured_results.jl expands the complete PTX 9.3 grammar into 1,134 reviewed schemas and the generic chain emitter consumes those schemas directly. This keeps scalar and grouped siblings in the same closed boundary: modifier legality, source carriers, result tuple shape, sink positions, and target metadata cannot drift between a wrapper and the fallback.

General setp uses the Julia-only leading .dual selector for the optional compare/complement result. Packed .f16x2 and .bf16x2 always return the two lane predicates. BoolOp lop3 returns (UInt32, Bool) and requires a Val{N} LUT with N in 0:255; match.all.sync.*.pred selects the optional predicate, and elect.sync always returns (UInt32, Bool). The two warp-collective families use the same convergent nomerge LLVM call-site path as the existing MMA and mbarrier wrappers. The single-predicate queries testp (source-type tail) and isspacep (unbracketed .u32/.u64 generic-address value) live in the same ledger because the generic trailing-dtype rule would misread their destination ABI.

For the mbarrier full-report form, reportValue is a PTX .b8 destination. This agrees with the CUDA Runtime API's description of cudaFabricOpStatusSourceMbarrierV1 as 1-byte aligned and 1-byte wide, and is enforced by ptxas evidence at sm_90 and sm_121. The Julia result is a zero-extended UInt16 carrier because LLVM's NVPTX inline-asm interface has no i8 register constraint; only its low byte is the opaque status value.

Arrival calls that spell destination _ use the synthetic .sink selector, for example ptx"mbarrier.arrive.sink.release.cluster.b64"(remote_addr). This distinction cannot be inferred from argument types: the otherwise identical local/generic form returns a UInt64 state token. It is especially important for a generic address produced for a remote cluster mbarrier, where the ISA requires the sink destination. Explicit shared::cluster forms are always sinks and need no selector. Base-plus-constant-offset addresses remain valid; tensor-coordinate address lists are rejected rather than truncated.

The simple WGMMA synchronization ops (wgmma.fence, wgmma.commit_group, wgmma.wait_group) still flow through the reviewed generic chain contract. Every tcgen05 form is typed-wrapper-only, including the ptxas-covered pointer forms of alloc and non-multicast commit and the before/after thread-sync fences. An exact method is authoritative and a dispatch miss fails before generic asm rendering. This is necessary because the shared opcode spans address destinations, register vectors, sinks, descriptors, fences, and MMA operands with incompatible schemas.

MX target evidence is deliberately split by PTX target class. The explicit .scale_vec::* spellings are compiled offline through ptxas at sm_100a and sm_110a; the equivalent .block16/.block32 aliases are compiled at sm_100f and sm_110f. Those checks validate source schema, inline-asm constraints, target gating, and assembler acceptance without launching a kernel. CUDACore 6.2.1 cannot emit an sm_110* module, so those cases emit the identical body at the matching sm_100a/sm_100f feature level, assert that only the .target directive changes, and invoke CUDA 13.3 ptxas directly on the sm_110a/sm_110f text. No tcgen05 runtime claim is made for the available CC 12.1 GB10, which does not implement tcgen05.

Ordered sparse MMA metadata

mma.sp::ordered_metadata is not a looser spelling of ordinary sparse MMA. For 2:4 f16/bf16/FP8 storage, each metadata nibble must list its two retained element indices in increasing order; other orderings have undefined behavior. (tf32 uses one retained index per 1:2 chunk.) PTX.jl cannot infer or repair that property from an opaque UInt32 metadata word, so callers must establish it while pruning and packing. The runtime evidence in test/gpu/ampere/gemm_sparse.jl does this by sorting each retained pair before encoding the nibble.

The typed surface deliberately mirrors only the 12 sparse MMA ABIs already supported by mma.sp. The ordered qualifier was introduced in PTX 8.5 and the classic f16/bf16/tf32 forms target sm_80+; the FP8 forms require sm_89+. Selector legality is enforced by exact Val dispatch: k16 16-bit and k8 tf32 accept Val(0) through Val(3), k32 16-bit and k16 tf32 accept Val(0) or Val(1), and k64 FP8 accepts only Val(0). A wrong selector, fragment width, or unreviewed shape/type combination fails before the generic scalar emitter. Like all mma.sync operations, these calls are warp-collective and carry convergent nomerge; they remain memory-free arithmetic operations.

The classic integer sparse surface is separate from tcgen05.mma.sp. It uses register fragments rather than tensor-memory descriptors and contains exactly 64 forms: u8/s8 at m16n8k32 and m16n8k64, u4/s4 at m16n8k64 and m16n8k128, all four A/B signedness pairs, with and without .satfinite, and both ordinary and ordered metadata. A/B register tuples contain 2/2, 4/4, 2/2, and 4/4 packed UInt32 values respectively; C/D are always four Int32 values, and metadata is one UInt32. The k32 u8/s8 and k64 u4/s4 products use selector Val(0) or Val(1); k64 u8/s8 and k128 u4/s4 require Val(0). For ordered u4/s4 storage, each 4:8 group is encoded as two ordered 2-element retained subgroups; the wrapper cannot validate those opaque metadata nibbles. Ordinary integer sparse MMA requires PTX 7.1, ordered metadata requires PTX 8.5, and all forms require sm_80+.

Extended-precision arithmetic

The extended-precision wrapper is a correctness boundary, not only an ergonomic overload. PTX CC.CF is implicit state that LLVM cannot name as an input/output constraint. Scalar typed wrappers reify it as Bool; aggregate helpers keep the entire chain in one asm unit. Generic, raw, and transpiled instruction-at-a-time fallbacks are rejected instead of emitting optimizer- unsafe calls.

The public helpers are PTX.add_with_carry, PTX.sub_with_borrow, and PTX.mul_wide; their full API documentation is in the reference.

The blocks are per-thread (convergent=false), carry sideeffect and ~{cc} but no false memory clobber, and use early-clobber word outputs so a low result cannot alias a high input that a later instruction still needs. Subtraction seeds and materializes the flag through sub.cc/subc, directly mirroring PTX borrow semantics instead of relying on a cross-family flag representation.

Host-side descriptor builders

Two opcodes consume packed 64-bit shared-memory descriptors plus (Blackwell only) a 32-bit instruction descriptor. PTX.jl ships pure bit-packing helpers for both:

These are not exported but are part of the documented API. Access them as PTX.wgmma_descriptor, PTX.tcgen05_descriptor, etc.

The tcgen05 builders deliberately expose semantic fields rather than every bit in the packed values. In particular, tcgen05_descriptor has no version keyword: PTX 9.3 §9.7.17.4.1 defines bits 46–48 as the fixed constant 0b001, not a version selector. Likewise, only tcgen05_instr_desc_i8 has a saturate keyword (the bit is integer-only), the i8 builder has no negate keywords and the mxf4 builders no transpose keywords (both ISA-unsupported for those kinds), and every architecture-gated encoding (the sm_107f-family K-dimension and scale-layout bits) is fixed at zero rather than exposed. An illegal descriptor is unrepresentable through the builders, not merely rejected. Matrix addresses and byte offsets must be 16-byte aligned and fit the 18-bit descriptor input window; invalid swizzle encodings and values that would spill into reserved fields fail with ArgumentError instead of being truncated.

Absolute leading-address mode (lbo_mode=1) is the PTX 9.3 §9.7.17.3.1.2 escape hatch for a 48-byte K dimension. The builder enforces its shared-descriptor restrictions: 128-byte swizzling with 16-byte atomicity (BlackwellLayout.B128, code 2) and base_offset=0. Because K and major axes live outside the shared descriptor, the caller must still pair it with K-major A and B (instruction-descriptor transpose bits 15 and 16 both zero), K=48B, and the architecture-specific sm_103a target.

wgmma_descriptor(smem_addr_u32; leading_byte_offset, stride_byte_offset,
                 swizzle = WgmmaSwizzle.NONE, base_offset = 0) -> UInt64

smem_addr_u32(p::Core.LLVMPtr{T, AS.Shared}) -> UInt32

step_desc(desc::UInt64, byte_offset::Integer) -> UInt64

tcgen05_descriptor(smem_addr_u32; leading_bytes, stride_bytes,
    swizzle=BlackwellLayout.NONE, base_offset=0, lbo_mode=0) -> UInt64

tcgen05_instr_desc_f16bf16_f32(; m, n, ab_dtype,
    a_major=:K, b_major=:K, scale_a=1, scale_b=1,
    sparse=false, max_shift=0) -> UInt32

tcgen05_instr_desc_i8(; m, n, a_dtype, b_dtype,
    a_major=:K, b_major=:K, saturate=false,
    sparse=false, max_shift=0) -> UInt32

tcgen05_instr_desc_f8f6f4(; m, n, a_dtype, b_dtype, d_dtype=:f32,
    a_major=:K, b_major=:K, scale_a=1, scale_b=1,
    sparse=false, max_shift=0) -> UInt32

tcgen05_instr_desc_mxf8f6f4(; m, n, a_dtype, b_dtype,
    scale_a_id, scale_b_id, a_major=:K, b_major=:K,
    scale_a=1, scale_b=1, sparse=false) -> UInt32

tcgen05_instr_desc_mxf4(; m, n, scale_a_id, scale_b_id,
    scale_a=1, scale_b=1, sparse=false, sparsity_version=0) -> UInt32

tcgen05_instr_desc_mxf4nvf4(; m, n, scale_dtype, scale_a_id, scale_b_id,
    scale_a=1, scale_b=1, sparse=false, sparsity_version=0) -> UInt32

GMMA layout helpers

For wgmma.mma_async SMEM operands, the descriptor's leading_byte_offset / stride_byte_offset / swizzle triple is fully determined by the tile geometry (dtype, M-or-N, K, major axis). The four canonical GMMA layout families (INTERLEAVE / B32 / B64 / B128) cover all wgmma-compatible SMEM tile widths.

pick_gmma_layout(; elem_bytes, m_or_n, k, major) -> GmmaLayout

layout_for_a(; dtype, m, k) -> GmmaLayout

layout_for_mn_major(; dtype, k, n) -> GmmaLayout

Host-side TMA descriptor encoder

Hopper TMA (cp.async.bulk.tensor.*) consumes a 128-byte CUtensorMap blob built host-side by the CUDA driver's cuTensorMapEncodeTiled. PTX.jl wraps the driver call so descriptors take Julia types / symbols instead of raw CUtensorMapDataType enums, with a thin convenience helper for the common 2D row-major case.

These methods live in ext/CUDACoreExt.jl and load automatically when CUDACore is in the environment.

tensor_map_encode_tiled(dtype, global_addr, global_dim, global_strides,
                        box_dim; kwargs...) -> CuTensorMap

tensor_map_tile_2d(dtype, global_addr, rows, cols, box_rows, box_cols;
                   swizzle=:B128, oob_fill=:NONE) -> CuTensorMap

When to extend

Most chain-default coverage is sufficient. Reach for a wrapper when:

• ptxas rejects the chain output ("Arguments mismatch", wrong constraint letter on a shared-AS pointer, missing brackets on a memory operand);
• the op returns multiple values;
• the operand layout has a brace group, tensor-coord vector, or other shape with no $N rendering;
• the dispatch needs to key on a fragment shape or count that the chain can't see.

The pattern: copy the closest existing wrapper file, adjust the table and asm template, add a for loop that emits one typed method per combination — spelling literal forms with optype"" and routing the generator's bookkeeping through register_wrapper! / wrapper_intrinsic_call (see src/wrappers/registry.jl). ~80 LOC for a new family in most cases.