Reference

ptx"opcode.mod1.mod2..."
ptx"%name"

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"(ptx"%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.

A spelling that starts with % is a special register rather than an instruction: ptx"%tid.x" constructs the SpecialReg{Symbol("%tid.x")} operand, which renders as the verbatim PTX token. The name must be a scalar special register of the PTX inventory (%tid alone is a vector; read %tid.x). Interpolation folds the same way as a glued modifier:

@inline block_index(axis::Symbol) = ptx"mov.u32"(ptx"%ctaid.$axis")

A static spelling is also a method-definition head. The same ptx"" that calls a wrapper defines it, so the method is dispatchable by its spelling by construction:

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

Interpolated spellings are call-only.

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

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.

PTX.on_device() -> Bool

false in host code and true in device code. Every PTX.jl call that emits inline PTX asm or an NVVM intrinsic is guarded by it, so calling one on the host throws a DeviceOnlyError instead of handing NVPTX code to the host backend.

The device answer is a method overlay, which PTX.jl's CUDACore extension installs for CUDA.jl's method table. A GPUCompiler client that compiles for NVPTX under a method table of its own opts in by overlaying it there:

Base.Experimental.@overlay MyMethodTable PTX.on_device() = true

The overlay must be a plain @overlay, not @consistent_overlay (which CUDACore.@device_override expands to on Julia 1.11+): a consistent overlay licenses the compiler to evaluate the host method in its place.

DeviceOnlyError(what)

Thrown when PTX device code is called from the host. what names the inline asm or intrinsic the call would have emitted.

PTX.device_code_typed(f, argtypes; optimize = true) -> Vector{Pair{CodeInfo, Type}}

Base.code_typed for device code: f applied to argtypes, inferred under a device compiler's method table, where the host gate resolves to the device body. Needs no GPU. Provided by PTX.jl's CUDACore extension.

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.