Embedded ptxas: Architecture Overview

Note: This page is the entry point to the nvlink-internal documentation of the ptxas backend statically linked into nvlink v13.0.88. The address map, mega-hub layout, ROT13 details, and sub_1112F30 per-module compilation driver below are recovered from nvlink's binary -- not from the standalone ptxas binary. For the standalone ptxas reverse-engineering reference (159-phase pipeline, full Ori IR architecture, target catalog), see the ptxas wiki -- in particular Pipeline Overview, Codegen Overview, and Targets.

The single most important structural fact about nvlink v13.0.88 is that approximately 95% of its 25.2 MB .text section is not linker code -- it is a complete, statically embedded copy of the ptxas assembler/compiler backend. The actual device linker (ELF merge, symbol resolution, relocation, layout, output) occupies roughly 1.2 MB in the address range 0x400000--0x530000. Everything from 0x530000 through the end of .text at 0x1D32172 (~24 MB, ~38,000 functions) is the ptxas compiler backend: IR primitives, instruction selection, register allocation, instruction scheduling, SASS binary encoding, PTX parsing, and ELF/cubin output generation.

This page documents the evidence for this claim, the complete address map of the embedded ptxas subsystems, the five mega-hub instruction selector dispatch functions, and the ROT13 obfuscation applied to SASS mnemonics.

Evidence for Embedded ptxas

The embedded compiler is not a stripped-down stub -- it is a full-featured PTX-to-SASS compilation pipeline identical in capability to the standalone ptxas binary shipped in the CUDA toolkit. Key evidence:

1. Named memory pools. The linker creates "nvlink option parser" and "nvlink memory space" arenas at startup. The embedded compiler creates its own arenas with ptxas-specific names. Memory pool diagnostics at 0x1AEE070 report pool usage statistics (total, freeable, leaked) for the compiler's internal allocations.

2. Full option parser. sub_1103030 (29,803 bytes) registers the complete ptxas command-line option set via sub_42F130: --maxrregcount, --opt-level, --gpu-name, --device-debug, --fast-compile, --register-usage-level, --compile-only, --minnctapersm, --warn-spills, --lineinfo, --sp-bounds-check, --device-stack-protector, --sanitize, --position-independent-code, and approximately 50 more. These are forwarded from nvlink's LTO pipeline into the embedded compiler.

3. Full compilation pipeline. sub_1112F30 (65,018 bytes) at 0x1112F30 is the top-level per-module compilation driver. It writes PTX headers (.version, .target, .entry __cuda_dummy_entry__ { ret; }), selects codegen callbacks based on mode flags (--compile-as-tools-patch, --extensible-whole-program, --compile-only), validates SM version compatibility, and dispatches to per-function codegen initialization.

4. Multi-architecture support. sub_15C0CE0 (14,517 bytes) initializes 7 dispatch hash maps covering sm_75, sm_80, sm_86, sm_87, sm_88, sm_89, sm_90/90a, sm_100/100a/100f, sm_103/103a/103f, sm_110/110a/110f, sm_120/120a/120f, and sm_121/121a/121f. Each architecture gets 7 registered callbacks (nv.info emitter, resource usage table, instruction encoding table, compute capability array, perf-stats handler, cpf_optx handler, codegen options).

5. Register allocation and instruction scheduling. The range 0x1850000--0x1A00000 contains the full backend compiler core: ScheduleInstructions (85 KB), ScheduleInstructionsReduceReg, DynBatch, HoistInvariants, ConvertMemoryToRegister, spilling regalloc, SMEM spilling, multi-class register allocation (R-regs, UR-regs, predicates), setmaxnreg CTA-reconfig for Blackwell+, and codegen verification passes.

6. ISel mega-hubs. Five functions exceed 160 KB each. These are the top-level instruction selector dispatch functions, too large for Hex-Rays to decompile. Each calls hundreds of pattern matchers, selects the highest-priority match, and dispatches to the corresponding emitter.

Relationship to Standalone ptxas

The standalone ptxas binary in the CUDA toolkit and the compiler backend embedded in nvlink share the same codebase. They differ in how they are invoked:

• Standalone ptxas: Invoked as a separate process by nvcc. Reads .ptx files from disk, writes .cubin files.
• Embedded ptxas in nvlink: Invoked in-process during LTO (-lto) and PTX JIT compilation. The entry point is sub_4BD760 (called from main() when a PTX input file is detected) or sub_4BC6F0 (called for LTO IR compilation after libnvvm produces PTX output). Options are forwarded programmatically rather than via argc/argv.

The embedded copy supports thread-pool parallelism for split compilation (sub_43FDB0 creates the pool, sub_4264B0 dispatches per-function work items). This is the same --split-compile-extended feature available in standalone ptxas.

Embedded ptxas Address Map

The following table maps the full address range of the embedded ptxas backend. All addresses are within the .text section of nvlink v13.0.88.

IR Primitives (0x530000 -- 0x620000, ~960 KB)

The IR node structure is accessed through 22 leaf functions that constitute the most-called code in the entire binary. sub_530FB0 (get operand by index) at 31,399 callers and sub_A49150 (get instruction attribute) at 30,768 callers form the universal accessor layer. The IR node layout:

Offset  Size   Field
  0     1B     operand type tag (1=immediate, 2=register, 6=memref, ...)
  4     4B     register class / encoding field (1023 = wildcard "any")
 14     1B     flag A
 15     1B     flag B
 20     4B     data type / secondary encoding
 28     2B     IR opcode
 32     8B     pointer to operand array (each operand = 32 bytes)
 40     4B     total operand count
 92     4B     first source operand index

Number of source operands = *(off+40) + 1 - *(off+92). Number of destination operands = *(off+92).

ISA Encoding Tables (0x620000 -- 0xA70000, ~4.3 MB)

This is the largest contiguous subsystem -- 4.3 MB of template-instantiated functions defining the complete NVIDIA GPU instruction set encoding and metadata.

The 1,537 SM100+ encoders each translate one instruction variant into a 128-bit SASS instruction word via the core primitive sub_4C28B0(buf, bit_offset, width, value). Opcode breakdown: major=1 (ALU/Scalar) 37.2%, major=2 (Vector/Memory/Control) 62.7%, major=3 (Special) 0.1%, across 118 instruction families.

The 1,613 InstrDesc initializers populate per-instruction metadata: operand count, operand types/constraints, scheduling hints, latency estimates, and execution unit assignments. Combined, the encoder + descriptor tables define the complete NVIDIA GPU ISA from SM50 through SM121.

Instruction Codecs (0xA70000 -- 0xCA0000, ~2.2 MB)

Multi-architecture instruction encoding and decoding, organized per-SM.

Each encoder packs IR operands into a 128-bit SASS instruction word at *(a1+40). Each decoder unpacks a 128-bit word back into IR form. The sentinel value 1023 (register field) maps to RZ (zero register), and 31 (predicate field) maps to PT (true predicate). Architecture-specific encoder variants are differentiated by the helper functions they call: sub_A5A000 (SM70 Volta), sub_A5AB30 (SM75 Turing), sub_A59D80 (SM80 Ampere), etc.

Per-Arch ISel Backends

Instruction selection is implemented as a linear-scan architecture: for each IR instruction, every pattern matcher is called in sequence, and the match with the highest priority wins. Each backend has its own set of pattern matchers, a mega-hub dispatch function (too large for Hex-Rays), and instruction emitters.

SM80 (Ampere) ISel Backend (0xCA0000 -- 0xDA0000, ~1 MB)

Three-phase pipeline: (1) pattern match on IR attributes/operand types, (2) operand emission into instruction descriptor, (3) binary encoding into 128-bit SASS word.

SM100+ (Blackwell) SASS Codec -- Second Table (0xDA0000 -- 0xF16000, ~1.5 MB)

Format 1 instructions: 147. Format 2 (extended with modifiers): 290. Format 3 (special wide): 1. Every encoder has a mirror decoder; the decoder count exceeds encoders because decoders also handle architecture-variant forms.

SM75 (Turing) ISel Backend (0xF16000 -- 0x100C000, ~984 KB)

This is the largest single-architecture ISel backend. sub_FBB810 at 280 KB is the largest function in the binary.

SM89/90 (Ada/Hopper) Backend (0x100C000 -- 0x11EA000, ~1.9 MB)

PTX Frontend (0x11EA000 -- 0x15C0000, ~3.5 MB)

The PTX frontend parses PTX assembly text, validates instructions against SM version constraints, and lowers them to the internal IR consumed by the per-architecture ISel backends.

The cuda_builtin_prototype_generator is the second-largest function in the binary at 345 KB. It maps builtin index numbers to PTX prototype strings of the form .weak .func (...) __cuda_smXX_foo (...). Function families include div, rem, rcp, sqrt, dsqrt, barrier, wmma, shfl, vote, matchsync, warpsync, reduxsync, sanitizer_memcheck, tcgen05, bulk_copy, and cp_async_bulk_tensor.

Compilation Pipeline (0x15C0000 -- 0x1A00000, ~4.2 MB)

This region contains the per-function compilation pipeline from SM dispatch through code generation to backend verification.

SASS Emission (0x1A00000 -- 0x1D32172, ~3.2 MB)

The final segment of .text handles SASS instruction lowering, ABI enforcement, ELF/cubin output, name demangling, and DWARF debug info.

The Five Mega-Hub Functions

Five functions exceed 160 KB each. They are the top-level instruction selector dispatch functions for different SM architecture generations. Each contains a massive jump table that calls hundreds of ISel pattern matchers in sequence, selects the highest-priority match, then dispatches to the corresponding emitter. All five are too large for Hex-Rays to decompile.

The ISel protocol is the standard ptxas linear-scan pattern matcher: every matcher is invoked with (ctx, ir_node, &pattern_id, &priority), the highest-priority match wins, the emitter table dispatches by pattern_id. Matchers query IR through sub_A49150 (attribute), sub_530FD0/sub_530FC0 (operand count), and sub_530FB0 (operand by index). For the algorithm in full detail see ptxas: Instruction Selection; the table above lists the nvlink-binary addresses of the five mega-hub dispatch functions that implement it.

ROT13 Obfuscation of SASS Mnemonics

NVIDIA applies ROT13 encoding to SASS instruction mnemonic strings stored in the binary. The decoder function sub_1A40AC0 uses SSE/SIMD vectorization for bulk ROT13 processing (loading 16 bytes at a time via _mm_load_si128). The SASS opcode table initializer at 0x1A85E40 stores all mnemonics in ROT13-encoded form; they are decoded at runtime before use.

Known decoded mnemonics:

The "MERCURY" prefix (ZREPHEL in ROT13) corresponds to sm_100+ (Blackwell) and appears throughout the compilation pipeline as a codename. ROT13 is also applied to some internal ELF section names: .sync_restrict::shared::read::mma::a is stored as its ROT13 equivalent, .acc::f16 as .npp::s16, and .sp::2to4 as .fc::2gb4.

Size Summary

Cross-Reference: Key Functions

Compilation Pipeline: sub_1112F30

sub_1112F30 (65,018 bytes at 0x1112F30, ~2,164 decompiled lines) is the top-level per-module compilation driver inside nvlink's embedded ptxas. It receives a module context a1 and a PTX module descriptor a2, then orchestrates the full PTX-to-SASS compilation across 26 phases before returning. Confidence: HIGH — derived directly from Hex-Rays output of this function.

This driver corresponds structurally to the entry/dispatch path in standalone ptxas — see ptxas: Pipeline Overview for the generic 159-phase pipeline narrative and ptxas: Entry Point for option-parser behavior. The table and notes below preserve the binary-specific phase order, callback choices, and helper addresses that are unique to the nvlink-embedded copy.

Phase Table

Per-Function Inner Pipeline (Phase 23)

For each function in the compile list, Phase 23 (either sequential or thread-pool worker) runs the following sub-stages:

In parallel mode each worker additionally allocates three local sorted maps (cap 8), copies a 15×16-byte snapshot of driver state into the per-function state, and allocates a 216-B per-function DWARF state via dwarf_register. After the barrier, the main thread merges per-thread maps, restores DWARF and pipeline snapshots, and runs codegen_finalize sequentially for each function so that register-budget propagation observes a deterministic order.

Key Subroutine Reference

Compilation Mode Matrix

The mode flag dispatch at Phase 6 selects one of four codegen pathways. The choice is determined by command-line flags forwarded into the embedded compiler:

The 48-Pass Codegen Pipeline (sub_1655A60)

The per-function codegen entry point sub_1655A60 runs a 48-pass pipeline (passes 0--47). Each pass is enable-gated by the SM dispatch vtable at a1[3757] (registered by sub_15C0CE0); enable flags occupy a1[160..207]. The pass numbering, vtable offsets, and binary slot allocation are the nvlink-embedded copy's own and do not match the standalone ptxas pass numbering (~159 phases) — for the corresponding standalone passes see ptxas: Passes Index, ptxas: Instruction Selection, and ptxas: Scheduling Algorithm.

After the pass loop, sub_1655A60 registers additional IR lowering callbacks (sub_161F1C0, sub_161F800, sub_1620460) on the function's basic block list, sets up UDT/UFT relocations for Blackwell+ (SM ordinal > 26 = sm_100+), and processes the function's call graph for register pressure analysis.

Sequential vs. Parallel Compilation

Selected by ctx->thread_count at a1+668. Sequential (count = 0) runs codegen_init -> codegen_compile -> codegen_finalize on the main thread with timing recorded between stages. Parallel (count > 0) creates a thread pool via sub_43FDB0, builds 48-byte work items containing the per-function state snapshot (360 B + 3 local sorted maps + 216-B DWARF state + pipeline snapshot), enqueues each via sub_43FF50, barriers on sub_43FFE0, then destroys via sub_43FE70. Workers run sub_1107420, which delegates to sub_1102B30 (setjmp-wrapped compile) and records timing + peak memory. After the barrier the main thread merges per-thread maps back, restores DWARF and pipeline snapshots, and runs codegen_finalize sequentially so register-budget propagation is deterministic. If qword_2A64430 is non-null, each worker error-checks via sub_1D1E060/sub_1D1E300 after its work item. This is the same --split-compile-extended mechanism available in standalone ptxas.

Cross-References

nvlink Internal

• IR Nodes -- IR node structure and universal accessor functions
• ISel Hubs -- the five mega-hub instruction selector dispatch functions
• Peephole -- peephole optimization passes (ORI, scheduling-phase, linker-level)
• PTX Parsing -- the embedded PTX assembler frontend
• Register Allocation -- graph-coloring register allocator with spilling
• Scheduling -- pre-RA and tepid (post-RA) instruction schedulers
• Architecture Dispatch -- per-SM vtable dispatch system
• Mercury Overview -- Mercury ISA encoding pipeline
• FNLZR -- post-link binary rewriter for Mercury targets
• LTO Overview -- how the LTO pipeline invokes the embedded compiler

Sibling Wikis

• ptxas: Pipeline Overview -- standalone ptxas 159-phase compilation pipeline
• ptxas: Entry Point -- standalone ptxas main() and option parsing
• ptxas: Optimizer -- standalone ptxas optimization passes
• ptxas: Codegen Overview -- standalone ptxas code generation