Reference

ptx"opcode.mod1.mod2..."

Construct an Operation{op, mods} singleton — op::Symbol is the opcode (first segment), mods::Tuple{Vararg{Symbol}} is the modifier chain. Splits on .; each segment becomes one Symbol verbatim, so :: (PTX sub-namespace separator), digit-leading tokens (3d, m16n8k32), and underscores in modifier names all flow through cleanly.

Supports $x / $(expr) interpolation. The macro expands to a chain of type-domain * compositions:

• Static segments fold into the initial Operation{op, mods}() and subsequent Chain{(...)}() constants.
• A glued interpolation (literal chars adjacent on either side, e.g. x$(N) or $(N)d) emits Symbol(pre, expr, post) and composes via Operation * Symbol. With a compile-time-constant interpolated value (e.g. N from a Val{N} unwrap) the whole site folds to the same singleton the literal form would produce — making it safe to use inside device kernels.
• A bare interpolation (between two .s) emits _ptx_dyn_seg(expr), which yields a Chain — supports String values containing . (split into multiple modifier segments) for host-side use.
• An interpolated opcode (first segment) falls back to _ptx_op_from_string(...). Device kernels should keep the opcode literal.

Examples:

ptx"add.f32"(a, b)
ptx"mma.sync.aligned.m16n8k16.row.col.f32.bf16.bf16.f32"(a, b, c)
ptx"cp.async.bulk.tensor.3d.shared::cta.global.tile.mbarrier::complete_tx::bytes"(...)
ptx"bar.sync"(Val(0))
ptx"mov.u32"(sreg"%tid.x")

dt = "u32"
ptx"mov.$dt"(x)              # ≡ ptx"mov.u32"(x)
ptx"st.$(space).b32"(p, v)   # $(...) for non-identifier exprs

# Glued — folds to a literal singleton when N is Val-known:
@inline f(p, ::Val{N}) where {N} =
    ptx"ldmatrix.sync.aligned.m8n8.x$(N).shared.b16"(p)

Empty literal parts (consecutive ., leading/trailing ., or empty string) error at expansion for the static path, and at runtime for the interp path.

See also: @mod_str for modifier-only chains usable on the right side of *.

optype"opcode.mod1.mod2..."

The method-definition companion of @ptx_str: expands to the annotation ::Operation{op, mods} for the same static spelling, so a typed wrapper is defined in ISA text instead of a hand-transcribed mods tuple:

@inline optype"add.f32"(a::Float32, b::Float32) = ...

is exactly

@inline (::Operation{:add, (:f32,)})(a::Float32, b::Float32) = ...

Both string macros share one parser, so the definition is dispatchable by the ptx"" spelling that reads back out of it, by construction — modifier transcription typos (which produce unreachable methods that only a count pin can catch) become impossible.

Definition-site only; no $ interpolation (a generated family should use an explicit @eval loop over its spec, which builds mods tuples directly).

sreg"name"

Construct a SpecialReg{Symbol("%name")} singleton — a compile-time literal for a PTX special register. Bakes the verbatim asm token, so underscore-bearing names (%cluster_ctarank, %lanemask_eq, %total_smem_size) round-trip losslessly. The legacy spelling %warpsize is the exception: PTX 9.3 defines WARP_SZ as an immediate, so it returns Val(32). Other names accept either form:

sreg"tid.x"            ≡ sreg"%tid.x"            → "%tid.x"
sreg"cluster_ctarank"  ≡ sreg"%cluster_ctarank"  → "%cluster_ctarank"
sreg"warpsize"         ≡ sreg"%warpsize"         → Val(32)

mod"mod1.mod2..."

Construct a Chain{mods} singleton — a sequence of PTX modifiers with no opcode. Splits on . like @ptx_str but does not accept $ interpolation: there is no opcode context to interpolate against, and the type-level * (below) already handles compile-time composition.

mod"" is the empty chain Chain{()}. Composing it via * is a no-op — useful as an "absent modifier" sentinel for conditional helpers.

Address{T}

An isbits role marker for a 32- or 64-bit integer that represents a bracketed PTX address operand. Address is not a pointer conversion and is removed before inline-assembly argument passing. Construct it through address.

Core.LLVMPtr values are deliberately not wrapped: their pointer type already preserves the address role, address space, and exact typed-wrapper dispatch.

address(value)

Preserve a bracketed PTX address role through Julia lowering. A 32- or 64-bit integer becomes an Address marker; an Address and a Core.LLVMPtr are returned unchanged. Other carriers are rejected.

reinterpret_addrspace(Val(AS′), p::Core.LLVMPtr{T,AS}) -> Core.LLVMPtr{T,AS′}

Reinterpret a pointer's raw 64-bit value in another address space — a ptrtoint/inttoptr pair, deliberately NOT an addrspacecast. NVPTX lowers addrspacecast to cvta window translation, which is wrong for several raw retypes that tier-2 wrappers need:

• Shared → SharedCluster: the ISA defines shared::cta addresses as valid shared::cluster addresses (own-CTA window), so the raw value is already correct; the cast's cvta.shared + cvta.to.shared::cluster round-trip is two wasted instructions.
• Const → Generic for TMA descriptors: the descriptor blob lives in global memory and is typed AS.Const by convention (TMADescriptorPtr), so its raw value is already a generic address — cvta.const would translate it as if it were a const-window offset and corrupt it.
• Global → Generic for a live tensor-map descriptor: global and generic use the same full virtual address for global storage. The tensor-map acquire intrinsic requires the latter carrier even when the preceding descriptor update correctly uses an explicitly global operand.

add_with_carry(a::NTuple{N,T}, b::NTuple{N,T}[, carry::Bool])
    -> (words::NTuple{N,T}, carry_out::Bool)

Add two little-endian limb tuples and return the same-width sum plus carry-out. An optional explicit carry-in is accepted. T may be UInt32, Int32, UInt64, or Int64; PTX defines identical carry behavior for signed and unsigned add. The complete straight-line sequence is emitted as one opaque, non-convergent inline-assembly call so CC.CF never crosses an LLVM boundary.

sub_with_borrow(a::NTuple{N,T}, b::NTuple{N,T}[, borrow::Bool])
    -> (words::NTuple{N,T}, borrow_out::Bool)

Subtract little-endian limb tuple b from a, returning the same-width difference and borrow-out. An optional explicit borrow-in is accepted. The entire sequence is one side-effecting, non-convergent asm block; see add_with_carry for supported limb types.

mul_wide(a::NTuple{2,T}, b::NTuple{2,T}) -> NTuple{4,T}

Multiply two little-endian, two-limb unsigned integers and return the full four-limb product. T may be UInt32 or UInt64. This embeds the PTX §9.7.2.6 UInt32 mad.cc/madc sequence, generalized to the legal UInt64 forms, in one side-effecting, non-convergent asm block. Signed multi-limb multiplication is intentionally not inferred from per-limb signed mad.hi semantics.

vector_load(ptx"ld....vN.type", address, Val(mask))

Emit a wide PTX vector load with an explicit destination-lane mask. true returns that lane and false emits PTX's _ sink, which guarantees the corresponding memory location is not read. Sink lanes are supported only for the ISA's wide ld.v8.{b32,u32,s32,f32} and ld.v4.{b64,u64,s64,f64} forms. The result is the homogeneous tuple of live lanes in source order. All-false masks are rejected: current ptxas cannot assemble an all-_ destination, and eliding a load would require a separate memory-model audit.

ptx_to_julia(source) -> String

Parse PTX source text and emit Julia source code. Returns a string of one or more function ... end definitions calling the ptx"..." macro after validating the complete module against the transpiler's deliberately closed subset. Lexical, parse, and transpiler-contract errors throw instead of returning a partial program. Acceptance is a lowering-contract guarantee, not a proof of semantic equivalence for every accepted PTX program.

ir_to_julia(mod::IR.Module) -> String

Validate a parsed IR.Module and convert it into Julia source code. Each accepted Function becomes one Julia function definition. Validation covers the complete module and finishes before emission, so an unsupported node, declaration, control-flow edge, or instruction form raises TranspilerError without returning partial Julia.

TranspilerError(path, category, detail)

Raised before PTX-to-Julia emission when parsed IR falls outside the closed semantic subset. path identifies the rejected IR node, category separates unsupported structure/operands from reviewed schema misses, and detail explains the missing contract. This is not evidence that the input PTX is invalid; parsing and lossless formatting intentionally support a broader IR.

tokenize(source::AbstractString) -> Vector{Token}

Tokenize PTX source text. The token stream includes NEWLINE / COMMENT tokens (needed for formatting preservation) and ends with an EOF token. Raises LexError on unrecognizable characters.

Token

A single token from the PTX source. leading_whitespace carries the spaces/tabs that preceded this token on the same line, used by the parser to reconstruct FormattingInfo for byte-identical round-trip.

TokenKind

Token kinds produced by tokenize. EnumX-nested so values are referenced as TokenKind.IDENTIFIER, TokenKind.NEWLINE, etc.

LexError

Raised by tokenize on unrecognizable input.

parse(source::AbstractString) -> IR.Module

Parse PTX source text into a Module. Module-header (.version / .target / .addresssize) and supported module/function/body statements are parsed into structured IR; blank lines and comments are preserved. Instructions do not require a known-opcode inventory. A post-header statement that fails the fault-tolerant statement parser is generally retained as an opaque RawLine, but lexical errors, invalid initial headers, malformed or misplaced .version/.target/`.addresssizedirectives, and target invariants still throw.RawLine` preserves text only and is not a semantic parse.

The original nonempty input is retained in Module.raw_source, so formatting an unmodified parsed module is byte-identical without exercising field-driven structural reconstruction.

ParseError

Raised by parse on input the recursive-descent parser cannot consume. Carries the source line and col of the failed token.

format(mod::IR.Module) -> String

Reconstruct PTX text from a parsed IR.Module. Returns mod.raw_source verbatim when set (the lossless fast path used by parser-produced IR); otherwise emits the module structurally — header, leading prelude, then each directive. A statement uses formatting.raw_line when present and field-driven reconstruction otherwise. An opaque RawLine is also emitted verbatim, but carries no reconstructed statement semantics.

Per-statement format(stmt) methods (one per IR.Statement kind) implement the structural fallback and can be called individually. Construction-time-only nodes without a PTX spelling, such as IntrinsicScope, error rather than discarding their metadata.