Internals
Machinery documented for contributors and the curious — not part of the public API. Everything here may change without notice; the public surface is the Reference page.
Form registry
The blessing boundary's data: every promise the chain default makes to the optimizer (purity, memory effects, convergence) lives in one auditable table in src/ledgers/forms.jl.
PTX.form_contract — Function
form_contract(op, mods) -> Union{FormContract, Nothing}The registry's contract for a chain form: the opcode's default, refined by the longest matching mods-prefix override. Returns nothing for an unregistered opcode or for a typed-wrapper-only form whose semantics cannot use the generic/raw chain (currently the implicit-CC.CF family). Callers distinguish the fail-loud semantic boundary from an ordinary unregistered opcode before considering the explicit raw tier.
Wrapper registry
Every typed wrapper method registers a record (family, op, mods, tier, intrinsic) through one entry point; the accessors below are what the conformance and coverage tests query instead of per-family bookkeeping constants.
PTX.register_wrapper! — Function
register_wrapper!(family, op, mods, tier[, intrinsic])Record one wrapper registration in WRAPPER_REGISTRY. Idempotent: a key already recorded is skipped, so registration helpers may be re-invoked (coverage testsets do) without inflating the inventories the count pins protect. tier must be :intrinsic, :asm, :core_ir, or :missing; intrinsic is required for :intrinsic/:missing records and must be absent otherwise.
PTX.wrapper_records — Function
wrapper_records(families...) -> Vector{WrapperRecord}All registry records for the given families, in registration order.
PTX.wrapper_intrinsic_names — Function
wrapper_intrinsic_names(families...) -> Vector{String}The distinct tier-2 intrinsic names the given families stand on, in registration order (the conformance replays pin these sets and counts).
PTX.wrapper_asm_forms — Function
wrapper_asm_forms(families...) -> Vector{Tuple{Vararg{Symbol}}}The mods tuples of the given families' asm-tier registrations, in registration order (the forms with no intrinsic at the pinned backend).
PTX.wrapper_missing_intrinsics — Function
wrapper_missing_intrinsics(families...) -> Vector{String}Intrinsic names a generator expected but the pinned NVVM registry lacks (no method was defined). Non-empty means a form silently lost its lowering — the conformance suite asserts emptiness per family.
PTX.wrapper_intrinsic_call — Function
wrapper_intrinsic_call(family, op, mods, name; missing_ok = false)The shared tier-2 shell of every generated wrapper family: verify name against the pinned NVVM registry, record the registration, and return the NVVM.IntrinsicCall singleton for the method body. Generated names bypass the nvvm"" macro's expansion-time existence check; this keeps its can't-survive-load property. With missing_ok = true an unregistered name is recorded as a :missing record and nothing is returned — the caller skips the method and the conformance counts fail loudly (used by families whose forms have no asm fallback).
PTX.ceiled — Function
ceiled(call::NVVM.IntrinsicCall, op_singleton::Operation) -> NVVM.IntrinsicCallThread the reviewed form contract for the PTX spelling onto a bare nvvm"..." singleton, so the effect ceiling is checked at every use. The spelling is passed as a ptx"..." singleton — the (op, mods) the wrapper is lexically implementing — and resolved to its contract at generation time. Every nvvm"" literal in src/ must flow through this or wrapper_intrinsic_call (enforced by test/host/effect_ceiling.jl).
Lowering reflection
Ask any operation what it will actually become — which lowering tier a call binds to, and to which intrinsics — without compiling for a device.
PTX.lowering — Function
lowering(op::Operation, argtypes) -> NamedTupleReflect on how calling op with arguments of the given types will lower, without a device. argtypes is a tuple of types or a Tuple{...} type. Returns (; tier, method, rettype, intrinsics, asm):
tier = :intrinsic— a wrapper routes tollvm.nvvm.*intrinsics (tier 2); their names are inintrinsics, andNVVM.intrinsic(name)has each record.tier = :core— a wrapper emits target-independent LLVM IR (tier 1): a realfence,load/store, etc. No intrinsic, no asm.tier = :asm— a hand-written wrapper embeds inline PTX asm (no intrinsic or core-IR spelling exists at the pinned backend).tier = :chain_asm— no wrapper: the chain default renders inline asm from the mods and argument types under the form registry's contract (RAW_CONTRACTfor aRawOperation); the text is inasm.tier = :unregistered— no wrapper and the opcode is not in the form registry: the call errors at the blessing boundary.tier = :forbidden— the selected generic path is unsafe: a spelling accesses implicit architectural state that cannot cross the call boundary, a typed-wrapper-only form missed its exact method, or a closed structured-, vector-, scalar-result, or immediate grammar island missed its audited ABI or constant domain. Explicit raw remains available for the typed-wrapper structural case, but not hidden state, an unknown result ABI, or an invalid ISA-required immediate.
Binding is not selectability: an :intrinsic form can still fail ISel below its capability floor — that gate lives in the backend, not the registry.
NVVM intrinsic registry
Generated from the backend LLVM's IntrinsicsNVVM.td (see gen/) and committed as source, so a backend bump is a reviewable diff. The nvvm"..." tier emits calls through this registry with explicit attribute groups — the in-process LLVM doesn't know these names, so the attributes here are the only effect/convergence information the optimizer gets.
The attributes are bounded by the package's own review: the FORMS/ledger contract for the PTX form being emitted is the ceiling on optimizer permissiveness for both lowering tiers. Table attributes may refine below it (speculatable, return ranges, param attrs, narrower memory location classes); a permissive-direction divergence — the table claiming purity, reorderability, or non-convergence the reviewed contract does not grant — fails at generation time (check_ceiling), and is resolved only by a reviewed overlay (CONVERGENT_OVERLAY_PREFIXES, MEMORY_WIDEN_OVERLAY) or a reviewed contract change. Every package call site carries its ceiling (ceiled(nvvm"...", ptx"..."), wrapper_intrinsic_call, the sreg fast path); the bare nvvm"" macro is unceiled plumbing, and test/host/effect_ceiling.jl keeps it out of src/.
One asm-tier honesty note: the :observable effects class (never deletable, touches no tracked memory) currently renders with the same conservative barrier as :clobbers — LLVM treats sideeffect inline asm without call-site memory(...) attributes as unknown memory, so the class documents reviewed semantics without yet granting the optimizer freedom. Attaching call-site memory attributes through the handwritten-IR path is the known lever; measurement showed motion windows in real pipelined kernels are bounded by their barrier waits regardless, so the lever stays unpulled until a profile demands it.
PTX.NVVM — Module
NVVMThe registry of llvm.nvvm.* intrinsics understood by the pinned external NVPTX backend. This module holds the hand-written half: the record type, the type-token vocabulary, and the query API. The data itself lives in table.jl, machine-generated from IntrinsicsNVVM.td by gen/generate_registry.jl — never edited by hand.
Names and signatures here are stable only within a backend JLL major, by explicit policy: this vocabulary is upstream's, surfaced honestly.
PTX.NVVM.Ceiling — Type
Ceiling(rank, convergent)The two ceiling-audited axes of a reviewed form contract, in the isbits shape a type parameter requires (FormContract carries a Symbol field and cannot be a type parameter itself). rank is the observability class: 3 = :pure, 2 = :observable, 1 = :clobbers. Constructed PTX-side from a FormContract (see _ceiling in ledgers/forms.jl).
PTX.NVVM.Intrinsic — Type
One intrinsic record, as extracted from the backend's IntrinsicsNVVM.td.
name: the LLVM intrinsic name (base name — overloaded intrinsics take a mangling suffix at the call site).ret,params: type-token tuples. Multipleretentries mean a literal struct return.props: intrinsic-level properties, lowercased from tblgen (:convergent,:nomem,:readmem,:writemem,:argmemonly,:inaccessiblememonly,:inaccessiblemem_or_argmemonly,:speculatable,:sideeffects,:noreturn,:nocallback,:nofree,:willreturn,:nocreateundefpoison,:commutative). These drive the attribute groups the emission layer attaches;:convergentis the load-bearing one.immargs: parameter positions (1-based) that must be immediate constants in the emitted IR, not SSA values.ranges:(pos, lo, hi)triples — value is in[lo, hi); position 0 is the return value.argnames:(pos, name)pairs — operand names from the .td, where present (documentation metadata).argattrs:(pos, attr)pairs for per-position attributes (:nocapture,:noalias,:noundef,:readonly,:writeonly); position 0 is the return value.
PTX.NVVM.abityp — Method
llvmcall ABI spelling of an accepted Julia type (Bool lowers as i8).
PTX.NVVM.accepts — Method
Whether Julia type T is accepted at a position with concrete token tok.
PTX.NVVM.callsite_ceiling — Method
Reviewed contract ceiling threaded onto a call site (nothing = none).
PTX.NVVM.canonical — Method
Canonical Julia type for a concrete token (used for returns).
PTX.NVVM.check_ceiling — Method
check_ceiling(name, ceiling::Union{Ceiling, Nothing}) -> nothingError unless the named intrinsic's effective table attributes stay at or below the reviewed contract ceiling on both audited axes. nothing means the call site carries no ceiling (the bare nvvm"" macro — package-internal sites are swept onto ceiled by test/host/effect_ceiling.jl).
PTX.NVVM.describe — Method
Human description of what a token accepts, for error messages.
PTX.NVVM.intrinsic — Method
intrinsic(name) -> IntrinsicLook up an intrinsic by its base name. On a miss, the error suggests registered names sharing the longest matching dotted prefix — a renamed intrinsic after a backend bump typically keeps most of its segments.
PTX.NVVM.isintrinsic — Method
isintrinsic(name) -> BoolWhether name is in the backend's intrinsic table (by base name, without any mangling suffix).
PTX.NVVM.llvmtype — Method
llvmtype(tok) -> StringLLVM IR spelling of a concrete type token. Overload tokens (AnyTok, SlotTok) have no spelling of their own — binding them to concrete types is the emission layer's job — so they error here.
PTX.NVVM.mangle — Method
Mangling suffix contributed by the Julia type bound to an overload slot.
PTX.NVVM.matching — Method
matching(prefix) -> Vector{String}All registered intrinsic names starting with prefix, sorted.
PTX.NVVM.observability_class — Method
observability_class(i::Intrinsic) -> SymbolThe intrinsic's effective (post-overlay) table effects on the contract's three-class resolution: :pure (deletable), :observable (undeletable, touches no memory), :memory (any location claim, or unspecified).
PTX.NVVM.overloaded — Method
overloaded(i::Intrinsic) -> BoolWhether the intrinsic has overload slots, i.e. needs a mangled callsite name.
PTX.NVVM.synthesize — Method
synthesize(name, argtypes) -> (; ir, rettype, tupletype, runtime)Build the Base.llvmcall IR for calling intrinsic name with arguments of the given Julia types. runtime is the subset of (1-based) argument positions passed at run time — immediate-operand positions take Val(x) and are spliced into the IR as constants instead, validated against the registry's legal ranges. Overload slots are bound from the concrete argument types and produce the canonical mangled callsite name. Every check errors at compile time, naming the intrinsic and position.
PTX.NVVM.@nvvm_str — Macro
nvvm"llvm.nvvm.name" # or nvvm"name" — the prefix is impliedA callable for a backend intrinsic, validated against the registry at macro expansion. Calling it emits attribute-correct Base.llvmcall IR:
- Pointers are
Core.LLVMPtr{T,AS}; the address space participates in overload resolution (mangled names are synthesized canonically). - Immediate operands (
immargsin the registry) are passed asVal(x)and range-checked at compile time. - Vector operands are
NTuple{N,VecElement{T}}. - Integer returns are unsigned (
UInt32, ...); multi-result intrinsics return tuples;i1maps toBool.
This is tier-2 plumbing: names and signatures are stable only within a backend JLL major, surfaced as-is from the LLVM table. The PTX-vocabulary ptx"..." notation is the stable surface above this.
Golden-harness IR
The structural PTX comparison behind the golden tests: parsed modules are normalized and canonicalized (registers, labels, and symbols renamed to a stable scheme) so goldens pin modeled instruction structure, not formatting or register-allocation accidents. Opaque RawLine nodes are not treated as structural coverage; strict goldens reject them recursively before comparison. Neither canonicalization nor an empty module diff is PTX ISA validation.
PTX.IR.normalize — Function
normalize(m::IR.Module) -> IR.ModuleReturn the structural comparison form of m. Source/header snapshots, the leading prelude, comments, blank lines, and formatting metadata are removed. Semantic headers and ordered statements remain; opaque RawLine nodes remain and compare by their exact text, while Block and IntrinsicScope retain their lexical boundaries.
PTX.IR.canonicalize — Function
canonicalize(m::Module) -> ModuleCanonical form for comparing PTX structure: normalized, with virtual registers, labels, parameters, declared variables, and function names renamed to position-stable canonical names. Scalar .reg declarations are dropped, while semantic vector declarations remain with canonicalized roots. Lexical scope nodes remain, with their nested operands renamed. For parser-produced PTX IR, format-identical canonical results have the same structure modulo allocator-style naming — the golden-harness equivalence. Construction-time IntrinsicScope nodes retain metadata but intentionally cannot be formatted; compare their canonical IR directly instead.
PTX.IR.diff — Function
diff(a::IR.Module, b::IR.Module; entry_only=false) -> Vector{String}Compare all normalized module nodes and return human-readable differences. With entry_only=true, every non-entry .func directive is excluded in full (declaration, signature, linkage, function directives, and body), while other module directives remain. Empty output means identical normalized IR structure, with opaque RawLine text compared verbatim. It does not canonicalize names and is not a substitute for PTX ISA, ptxas, or runtime semantic validation.
PTX.IR._sym — Function
Rename a symbol that may be a known name, a virtual register, or a label.
Transpiler contract
The transpiler validates the complete IR module against its deliberately narrow semantic boundary before emitting any Julia source.
PTX.Codegen.validate_transpilable — Function
validate_transpilable(mod::IR.Module)Validate the complete, deliberately narrow IR subset that ir_to_julia can lower without dropping declarations or guessing operand/result roles. The visitor runs before emission and rejects every unsupported node with its tree path. Parsing/formatting remain broader; rejection here is not a claim that the input PTX itself is invalid.
Evidence archaeology
One-off validation scripts lived in spikes/ until their findings were baked into code, tests, and goldens; the directory was then deleted. Comments citing a spikes/*.jl script refer to these — view any of them with git show ccdfb8a~1:spikes/. The load-bearing ones:
spikes/convergence.jl— reproduces the divergent-duplication miscompile class that motivatesconvergenton collective ops.spikes/raw_asm_attrs.jl— proves aconvergentattribute group on an inline-asm call site parses throughBase.llvmcalland survives the optimized module (theconvergent_asm_irmechanism).spikes/aggregate_return.jl— hardware validation of the ldmatrix aggregate-return repack and mangled overloaded-callsite names.