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_contractFunction
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.

source

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.

source
PTX.wrapper_recordsFunction
wrapper_records(families...) -> Vector{WrapperRecord}

All registry records for the given families, in registration order.

source
PTX.wrapper_intrinsic_namesFunction
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).

source
PTX.wrapper_asm_formsFunction
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).

source
PTX.wrapper_missing_intrinsicsFunction
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.

source
PTX.wrapper_intrinsic_callFunction
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).

source
PTX.ceiledFunction
ceiled(call::NVVM.IntrinsicCall, op_singleton::Operation) -> NVVM.IntrinsicCall

Thread 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).

source

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.loweringFunction
lowering(op::Operation, argtypes) -> NamedTuple

Reflect 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 to llvm.nvvm.* intrinsics (tier 2); their names are in intrinsics, and NVVM.intrinsic(name) has each record.
  • tier = :core — a wrapper emits target-independent LLVM IR (tier 1): a real fence, 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_CONTRACT for a RawOperation); the text is in asm.
  • 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.

source

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.NVVMModule
NVVM

The 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.

source
PTX.NVVM.CeilingType
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).

source
PTX.NVVM.IntrinsicType

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. Multiple ret entries 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; :convergent is 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.
source
PTX.NVVM.check_ceilingMethod
check_ceiling(name, ceiling::Union{Ceiling, Nothing}) -> nothing

Error 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).

source
PTX.NVVM.intrinsicMethod
intrinsic(name) -> Intrinsic

Look 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.

source
PTX.NVVM.isintrinsicMethod
isintrinsic(name) -> Bool

Whether name is in the backend's intrinsic table (by base name, without any mangling suffix).

source
PTX.NVVM.llvmtypeMethod
llvmtype(tok) -> String

LLVM 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.

source
PTX.NVVM.matchingMethod
matching(prefix) -> Vector{String}

All registered intrinsic names starting with prefix, sorted.

source
PTX.NVVM.observability_classMethod
observability_class(i::Intrinsic) -> Symbol

The 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).

source
PTX.NVVM.overloadedMethod
overloaded(i::Intrinsic) -> Bool

Whether the intrinsic has overload slots, i.e. needs a mangled callsite name.

source
PTX.NVVM.synthesizeMethod
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.

source
PTX.NVVM.@nvvm_strMacro
nvvm"llvm.nvvm.name"           # or nvvm"name" — the prefix is implied

A 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 (immargs in the registry) are passed as Val(x) and range-checked at compile time.
  • Vector operands are NTuple{N,VecElement{T}}.
  • Integer returns are unsigned (UInt32, ...); multi-result intrinsics return tuples; i1 maps to Bool.

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.

source

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.normalizeFunction
normalize(m::IR.Module) -> IR.Module

Return 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.

source
PTX.IR.canonicalizeFunction
canonicalize(m::Module) -> Module

Canonical 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.

source
PTX.IR.diffFunction
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.

source
PTX.IR._symFunction

Rename a symbol that may be a known name, a virtual register, or a label.

source

Transpiler contract

The transpiler validates the complete IR module against its deliberately narrow semantic boundary before emitting any Julia source.

PTX.Codegen.validate_transpilableFunction
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.

source

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 motivates convergent on collective ops.
  • spikes/raw_asm_attrs.jl — proves a convergent attribute group on an inline-asm call site parses through Base.llvmcall and survives the optimized module (the convergent_asm_ir mechanism).
  • spikes/aggregate_return.jl — hardware validation of the ldmatrix aggregate-return repack and mangled overloaded-callsite names.