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 @ptx_str spelling as the method head (@inline ptx"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

Mixed-precision packed arithmetic

PTX ISA 9.4 adds 36 packed mixed-precision add, sub, mul, and fma forms for sm_107f (or higher in the same family, including sm_107a). Their type tokens specify the destination first, then the sources. f32x2 uses a UInt64 bit carrier; f16x2 and bf16x2 use UInt32 bit carriers.

sum = ptx"add.rn.f32x2.f16x2.f32x2"(half_pair_bits, float_pair_bits) # UInt64
result = ptx"add.rz.bf16x2.f32x2.f32x2"(sum, float_pair_bits)        # UInt32

The packed mixed fma forms require an explicit rounding mode. Their second and third inputs are both f32x2 (UInt64). Narrowing to f16x2 uses .rz.ftz for add/sub and .ftz.rz for mul; these modifier orders follow the instruction grammars. Unsupported spellings and incorrect operand widths fail before LLVM compilation, including through ptx"..."raw.

The assembler tests cover every form on sm_107f and sm_107a when the managed compiler supports PTX ISA 9.4. Older compilers explicitly skip that partition. Execution on sm_107 hardware remains unverified.

Alternate packed arithmetic

PTX ISA 9.4 alternate floating-point add, sub, mul, and fma support four-lane FP8/FP6/FP4 arithmetic on sm_100a and sm_103a only. These forms are unavailable on sm_100f, sm_103f, and sm_107a.

The result is always packed e4m3x4 or e5m2x4 in a UInt32 bit carrier. Source carriers are fixed by the instruction schema:

compact = UInt16(0x2222)     # four e2m1 values
padded = UInt32(0x02020202)  # the same values in e2m1p4x4
acc = UInt32(0x38383838)    # four e4m3 values
sum = ptx"add.rn.e4m3x4.e2m1x4"(compact, acc)
product = ptx"mul.rn.satfinite.e5m2x4.e2m1x4.e2m1p4x4"(compact, padded)
result = ptx"fma.e4m3x4.e2m1x4.e2m1x4"(compact, compact, sum)

add/sub use a shared destination/last-source type followed by the first source type. mul specifies destination, first source, and second source types; fma shares its destination type with the accumulator and then specifies its two multiplicand types. Thus the final type token alone cannot determine the result or all input widths.

Rounding defaults to .rn; .satfinite optionally clamps overflow to the largest finite value of the destination format, preserving the sign. The canonical order puts .rn and .satfinite before the type tokens. The ledger admits all 896 canonical forms plus seven exact specification examples with postfix .satfinite, marked :ptxas_compat. Other modifier permutations and incorrect carrier widths are rejected, including through ptx"..."raw.

The assembler suite checks every form with stored results on both supported targets when PTX ISA 9.4 is available. Execution on B200/B300 hardware remains unverified for these new forms.

set comparison results

set compares two values and returns a numeric result. Its closed grammar covers 1,152 general scalar forms, 1,960 half/BF16 forms, and the 32 packed integer forms introduced in PTX ISA 9.4. An optional .and, .or, or .xor combines the comparison with a trailing Bool input.

mask = ptx"set.lt.u32.s64"(Int64(-2), Int64(1))       # UInt32(0xffffffff)
value = ptx"set.lt.f32.s64"(Int64(-2), Int64(1))      # 1.0f0
gated = ptx"set.lt.and.f32.s64"(a, b, gate)           # gate::Bool
lanes = ptx"set.ge.s16x2"(packed_a, packed_b)         # UInt32 lane masks

For scalar forms, the destination type precedes the source type. Integer destinations return an all-ones mask for true and zero for false; floating destinations return one or zero. .f16 results use Float16, while .bf16 results use BFloat16. Integer result types preserve their signedness and width.

Scalar .f16 results are formed from a 32-bit integer comparison mask, then converted to the half zero/one encoding. This avoids a ptxas widening defect where native set.f16.f16 can retain a duplicate high half through cvt.u32.u16. The public result remains Float16, including on the raw path. For .f16 sources, this lowering uses PTX ISA 6.5 instructions.

Half/BF16 x2 comparisons return two packed results. Integer destinations pack 16-bit masks; .f16x2/.bf16x2 destinations pack the corresponding floating one/zero encodings. Packed integer set uses a single type token and returns one 8-bit or 16-bit mask per lane in UInt32, including signed input formats.

Ordered floating comparisons return false when either input is NaN; unordered comparisons and .num/.nan follow their PTX definitions. .ftz is restricted to the reviewed f32/f16 comparison forms. Incorrect modifier order, comparison operators, operand widths, or predicate inputs fail before LLVM compilation, including through ptx"..."raw.

All admitted spellings have offline assembly coverage. Packed integer forms require PTX ISA 9.4 and the sm_107f family (also sm_107a); older compilers explicitly skip that partition. Their hardware execution remains unverified.

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 and ti16 builders no transpose keywords (ISA-unsupported, or ISA-contradictory, for those kinds). The architecture-gated encodings (the sm_107f-family K-dimension and scale-layout bits, the 15-bit shared-descriptor fields, the .decompress::lut::b segment bit) default to zero and are set only through their explicit keywords (k, scale_a_layout, wide, lut_k_offset). 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 (19-bit with wide=true); 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,
    lut_k_offset=0, wide=false) -> 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_ti16(; m, n, scale_a=1, scale_b=1,
    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, k=nothing, lut_b=false) -> 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,
    k=nothing, scale_a_layout=32, lut_b=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,
    k=nothing, scale_a_layout=32) -> UInt32

spcompress_desc(; op, dtype) -> UInt32

spcompress_elemsize(dtype) -> Symbol

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

Register dependencies on waits

PTX.wait_registers attaches register values to a wait's input/output operands. Consumers must use the returned values; calling a void wait and returning the original values does not establish that compiler dependency. The wait keeps its existing memory clobber and collective semantics.

The initial surface covers tcgen05.wait::ld.sync.aligned and tcgen05.wait::st.sync.aligned, with scalar or homogeneous tuple values. Adjacent 32-bit tuple elements use paired 64-bit operands. The ptxas/wait_registers test compiles six attention configurations with and without register bindings and compares machine-instruction encodings under the same toolchain, including any spill instructions. Debug metadata is stripped before backend code generation so it cannot perturb the instruction comparison.

Register operands and memory clobbers describe different compiler effects; see NVIDIA's inline-PTX documentation and LLVM's tied input constraints. This is an API addition for expressing register dependencies; no wrong-result bug has been demonstrated in the existing no-argument wait.

wait_registers(wait_instruction, values)

values = ptx"tcgen05.ld.sync.aligned.32x32b.x32.b32"(tmem_addr)
# Independent work can run while the load is in flight.
values = PTX.wait_registers(ptx"tcgen05.wait::ld.sync.aligned", values)
# Consume the returned values here.

Host-side TMA descriptors

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.

PTX.upload_tma_descriptor uploads the encoded descriptor and returns (; ptr, blob): blob owns the allocation, and ptr is the borrowed PTX.TMADescriptorPtr passed to kernels. Its AS.Const address-space type is the TMA wrapper carrier convention; the allocation lives in device global memory. Pointer conversion happens on the host.

Preserve the returned owner and the source tensor until GPU work completes, including synchronization after an asynchronous launch. The uploader's example shows the GC.@preserve pattern. Uploading copies the descriptor; it does not take ownership of the tensor whose address it contains.

Encoding and upload live in ext/CUDACoreExt.jl and become available when CUDA.jl or CUDACore is loaded alongside PTX.

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

upload_tma_descriptor(tmap::CuTensorMap) -> (; ptr, blob)

using PTX, CUDA

src = CuArray(reshape(UInt16.(1:64), 8, 8))
GC.@preserve src begin
    tmap = PTX.tensor_map_tile_2d(:u16, pointer(src), 8, 8, 8, 8;
                                 swizzle=:NONE)
    descriptor = PTX.upload_tma_descriptor(tmap)
    GC.@preserve descriptor begin
        @cuda kernel!(descriptor.ptr)  # kernel accepts PTX.TMADescriptorPtr
        CUDA.synchronize()
    end
end

TMADescriptorPtr = Core.LLVMPtr{UInt8, AS.Const}

Pitfalls in warp-specialized kernels

Things ptxas and LLVM do silently that cost hours to rediscover:

• setmaxnreg is ignored by ptxas unless the kernel declares .reqntid; pass minthreads=N through @cuda or the compiler config, and check EIATTR_REG_COUNT in the cubin.
• @inbounds CuDynamicSharedArray(...): the @boundscheck emits a gpu_report_exception call in the entry region.
• Indexing an NTuple with a runtime index demotes the tuple to local memory; make the index a Val type parameter. Large ntuple(...) do closures become real device calls; write explicit tuples or check code_llvm for call .*julia_.
• UInt32(x) on a loop-carried Int leaves live InexactError branches; use % UInt32 when the value is exact by construction.
• Debug loop: code_llvm and grep call .*julia_|report_exception|alloca; nvdisasm -c and grep LDL|STL|CALL; ptxas -v for spill counts.

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 as ptx"..."(...) = ... heads 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.

PTX 9.4 readonly loads and proxy fences

The readonly-load spelling places the proxy after the result type:

value = ptx"ld.global.u32.proxy::readonly"(global_ptr)
value = ptx"ld.u32.proxy::readonly"(generic_ptr)

The address must refer to global memory, or to global/constant memory through generic addressing; its contents must remain read-only for the kernel's lifetime. Scalar result types preserve their width and signedness. The finite transpiler grammar does not yet include these loads.

The generic fence chains include:

ptx"fence.proxy.alias.release.sys"()
ptx"fence.proxy.alias.acquire.sys"()
ptx"fence.proxy.async::generic.release.sync_restrict::shared::cluster::read.cluster"()

These require PTX ISA 9.4 and sm_90 or higher. The cluster-read fence orders shared-cluster reads across the specified proxies and is non-cumulative: do not rely on it to transitively order other accesses. The assembly tests check the emitted instructions and target floors; concurrent ordering still depends on the surrounding synchronization protocol.

Stochastic conversion (cvt.rs) does not permit .pzo. Other rounding and saturation modifier combinations continue to be checked by ptxas.