PTXAS v13.0 — Reverse Engineering Reference

Purpose: reimplementation-grade documentation of NVIDIA's PTX-to-SASS assembler, recovered entirely from static analysis of the stripped x86-64 binary.

PTX (Parallel Thread Execution) is NVIDIA's virtual ISA for GPU compute. SASS (Shader Assembly) is the native machine code executed by GPU hardware. PTXAS is the binary that transforms PTX into SASS. At 37.7 MB stripped, it is a fully proprietary compiler with no LLVM code, no EDG frontend, and no third-party optimizer components. Every pass, every data structure, and every encoding table was built in-house by NVIDIA. This wiki documents its internal architecture using IDA Pro 8.x and Hex-Rays decompilation.

Version note: All addresses and binary offsets in this wiki apply to ptxas v13.0.88 (CUDA Toolkit 13.0). Other versions will have different addresses.

Glossary

Three Subsystems

PTXAS is not a monolithic assembler. It decomposes into three largely independent subsystems with distinct coding conventions, data structures, and lineages:

1. PTX Frontend (~3 MB, 0x400000--0x5AA000) — A Flex-generated DFA scanner (sub_720F00, 64 KB, ~552 rules) feeds tokens into a Bison-generated LALR(1) parser (sub_4CE6B0, 48 KB). The parser is driven from sub_446240 (the real main, 11 KB), which orchestrates the full pipeline: parse, DAGgen, OCG, ELF, DebugInfo. The frontend also contains 1,141 instruction descriptors registered via sub_46E000 (93 KB) that define accepted type combinations for every PTX opcode, 608 CUDA runtime intrinsics registered in sub_5D1660 (46 KB), and a suite of per-instruction semantic validators (0x460000--0x4D5000) that check architecture requirements, type compatibility, and operand constraints before lowering. See PTX Parser and Entry Point & CLI.

2. Ori Optimizer (~8 MB, 0x5AA000--0xC52000) — A proprietary 159-phase optimization pipeline managed by the PhaseManager (sub_C62720). The phase factory at sub_C60D30 is a 159-case switch that allocates polymorphic phase objects from a vtable table at off_22BD5C8. Each phase has virtual methods for execute(), isNoOp(), and getName(). Major subsystems include: a fatpoint-based register allocator (sub_957160 core, sub_95DC10 driver, sub_926A30 interference graph builder), a 3-phase instruction scheduler (sub_688DD0 with ReduceReg/DynBatch modes and 9 register pressure counters), copy propagation, strength reduction, predication (if-conversion), rematerialization, and GMMA/WGMMA pipelining. The pipeline reads its default phase ordering from a 159-entry table at 0x22BEEA0. See Optimization Pipeline and Phase Manager.

3. SASS Backend (~14 MB, 0xC52000--0x1CE3000) — The Mercury encoder generates native SASS binary code. Instruction encoding is handled by ~4,000 per-variant handler functions (683 + 678 = 1,361 in the SM100 Blackwell encoding tables alone at 0xED1000--0x107B000, with additional tables for other SM generations). Each handler follows a rigid template: set opcode ID, load a 128-bit encoding format descriptor via SIMD, initialize a 10-slot register class map, register operand descriptors via sub_7BD3C0/sub_7BD650/sub_7BE090, finalize with sub_7BD260, then extract bitfields from the packed instruction word. The backend also contains 3 peephole optimizers (the PeepholeOptimizer class at 0x7A5D10 with Init, RunOnFunction, RunOnBB, RunPatterns, SpecialPatterns, ComplexPatterns, and SchedulingAwarePatterns methods), a capsule Mercury ELF embedder for debug metadata (sub_1CB53A0, section .nv.capmerc), and a custom ELF emitter (sub_1C9F280, 97 KB) that builds the final CUBIN output. See SASS Code Generation, Mercury Encoder, and Peephole Optimization.

Additionally, the binary embeds a custom pool allocator (sub_424070, 3,809 callers), MurmurHash3-based hash maps (sub_426150 insert / sub_426D60 lookup), a thread pool with pthread-based parallel compilation support, and a GNU Make jobserver client for integration with build systems.

Compilation Pipeline

Both standalone and library-mode invocations converge on the same pipeline, visible in the timing strings emitted by sub_446240:

PTX text (.ptx file or string)
  |
  +-- Flex Scanner (sub_720F00, 64KB)
  |     552-rule DFA, off_203C020 transition table
  |     Tokens: 340+ terminal symbols for Bison grammar
  |
  +-- Bison LALR(1) Parser (sub_4CE6B0, 48KB)
  |     Semantic validators: 0x460000-0x4D5000
  |     1,141 instruction descriptors via sub_46E000
  |
  +-- Ori IR Construction (DAGgen phase)
  |     Internal representation: basic blocks + instruction DAG
  |     608 CUDA runtime intrinsics (sub_5D1660)
  |
  +-- 159-Phase Optimization Pipeline (PhaseManager, sub_C62720)
  |     Phase factory: sub_C60D30 (159-case switch)
  |     Fatpoint register allocator (sub_957160)
  |     3-phase instruction scheduler (sub_688DD0)
  |     Copy propagation, CSE, strength reduction, predication,
  |     rematerialization, GMMA pipelining, late legalization
  |
  +-- Mercury SASS Encoder
  |     Instruction encoding: ~4000 per-variant handlers
  |     3 peephole optimizers (PeepholeOptimizer at 0x7A5D10)
  |     WAR hazard resolution (sub_6FC240)
  |     Operand expansion (Opex pipeline)
  |
  +-- ELF/CUBIN Output (sub_1C9F280, 97KB)
        Sections: .text, .nv.constant0, .nv.info, .symtab
        Capsule Mercury: .nv.capmerc (debug metadata)
        DWARF: .debug_line, .debug_info, .debug_frame

The driver at sub_446240 reports per-stage timing: Parse-time, CompileUnitSetup-time, DAGgen-time, OCG-time, ELF-time, DebugInfo-time, plus PeakMemoryUsage in KB. For multi-entry PTX files, each compile unit is processed independently with the header "\nCompile-unit with entry %s".

Dual Compilation Modes

PTXAS operates in two modes selected at invocation:

The main function (0x409460, 84 bytes) is a thin wrapper: it stores argv[0], sets stdout/stderr to unbuffered via setvbuf, and delegates to sub_446240. The --input-as-string flag enables accepting PTX source directly as a CLI argument rather than reading from a file.

Configuration

PTXAS exposes three layers of configuration:

CLI Options (~100 flags) — Registered in sub_432A00 and parsed by sub_434320. Key options include --gpu-name (target SM), --maxrregcount (register limit), --opt-level (0--4), --verbose, --warn-on-spills, --warn-on-local-memory-usage, --fast-compile, --fdevice-time-trace (Chrome trace JSON output), --compile-as-tools-patch (sanitizer mode), and --extensible-whole-program. Help is printed by sub_403588 which calls sub_1C97640 to enumerate all registered options.

Internal Knobs (1,294 ROT13-encoded entries) — A separate configuration system implemented in generic_knobs_impl.h (source path recovered: /dvs/p4/build/sw/rel/gpgpu/toolkit/r13.0/compiler/drivers/common/utils/generic/impl/generic_knobs_impl.h). The knob table is populated by two massive static constructors: ctor_005 at 0x40D860 (80 KB, ~2,000 general OCG knobs) and ctor_007 at 0x421290 (8 KB, 98 Mercury scheduler knobs). All knob names are ROT13-obfuscated in the binary. Examples after decoding: MercuryUseActiveThreadCollectiveInsts, MercuryTrackMultiReadsWarLatency, MercuryPresumeXblockWaitBeneficial, ScavInlineExpansion, ScavDisableSpilling. Knobs are read from environment variables and knob files via ReadKnobsFile (sub_79D070) which parses [knobs]-header INI files. Lookup is performed by GetKnobIndex (sub_79B240) with inline ROT13 decoding and case-insensitive comparison. See Knobs System.

SM Profile Tables — Per-architecture capability maps initialized by sub_607DB0 (14 KB) which creates 7 hash maps indexing sm_XX / compute_XX strings to handler functions. Profile objects are constructed by sub_6765E0 (54 KB) with architecture-to-family mappings (sm_75 -> Turing, sm_80/86/87/88 -> Ampere, sm_89 -> Ada Lovelace, sm_90/90a -> Hopper, sm_100/100a/100f -> Blackwell, sm_103/103a/103f -> Blackwell Ultra, sm_110/110a/110f -> Jetson Thor, sm_120/120a/120f -> RTX 50xx, sm_121/121a/121f -> DGX Spark). See SM Architecture Map.

Reading This Wiki

The wiki is organized around the compilation pipeline. Every page is written at reimplementation-grade depth for an audience of senior C++ developers with GPU compiler experience.

Section Index

Overview

• Function Map — Address-to-identity lookup for key functions with confidence levels.
• Binary Layout — Subsystem address map at pass granularity.
• Methodology — How this analysis was performed.
• Version Tracking — Cross-version address deltas.

Compilation Pipeline

• Pipeline Overview — End-to-end PTX-to-SASS flow diagram with links to every stage.
• Entry Point & CLI — CLI parsing, main at 0x409460, the real driver at sub_446240.
• PTX Parser (Flex + Bison) — 552-rule Flex DFA scanner, Bison LALR(1) parser, instruction descriptor table.
• PTX Directive Handling — .version, .target, .entry, .func, .reg, .shared, .const processing.
• PTX-to-Ori Lowering — How parsed PTX is lowered into the Ori internal representation.
• Optimization Pipeline (159 Phases) — PhaseManager, phase factory, default phase ordering, per-phase timing.
• SASS Code Generation — Mercury encoder, instruction selection, operand expansion.
• ELF/Cubin Output — Custom ELF emitter, section layout, capsule Mercury, DWARF generation.

Ori IR — Internal Representation

• IR Overview & Design — Instruction DAG, basic blocks, typed virtual registers.
• Instructions & Opcodes — Ori opcode set and instruction encoding.
• Basic Blocks & CFG — Control flow graph construction and manipulation.
• Register Model (R/UR/P/UP) — Four register classes and their constraints.
• Data Structure Layouts — Memory layout of key IR objects.

Optimization Passes

• Pass Inventory & Ordering — All 159 phases with names, addresses, and pipeline positions.
• Phase Manager Infrastructure — Phase factory, vtable dispatch, execute/isNoOp/getName.
• GeneralOptimize Bundles — Mega-pass bundles that group related sub-passes.
• Loop Passes — Unrolling, LICM, induction variable optimization, strength reduction.
• Copy Propagation & CSE — Value forwarding and common subexpression elimination.
• Predication — If-conversion for GPU divergence control.
• Rematerialization — Recomputing values to reduce register pressure.
• Synchronization & Barriers — Barrier insertion and dead barrier elimination.
• Late Expansion & Legalization — Final lowering before codegen.

Register Allocation

• Allocator Architecture — Fatpoint algorithm, interference graph, spilling, ABI constraints.
• Fatpoint Algorithm — Core allocation loop and heuristics.
• Spilling — Spill cost model and spill code generation.
• GPU ABI & Calling Convention — Register assignment rules and caller/callee contracts.

Instruction Scheduling

• Scheduler Architecture — 3-phase scheduler, ReduceReg/DynBatch modes.
• Scheduling Algorithm — Priority list scheduling with register pressure tracking.
• Latency Model & HW Profiles — Per-SM instruction latency tables.
• Scoreboards & Dependency Barriers — WAR hazard resolution, barrier allocation.

SASS Code Generation

• Code Generation Overview — Instruction selection, encoding, peephole, Mercury.
• Instruction Selection — Pattern-based DAG-to-SASS lowering.
• SASS Instruction Encoding — 128-bit instruction word format and bitfield packing.
• Peephole Optimization — Three peephole dispatchers with SM-variant patterns.
• Mercury Encoder — Per-variant handler architecture, encoding tables.
• Capsule Mercury & Finalization — .nv.capmerc section, debug metadata embedding.
• SASS Text Generation — Disassembly-format printing for --verbose output.

GPU Architecture Targets

• SM Architecture Map — SM feature gates from sm_75 through sm_121f.
• Turing & Ampere (SM 75--88) — Feature delta between generations.
• Ada & Hopper (SM 89--90a) — Async copy, TMA, distributed shared memory.
• Blackwell (SM 100--121) — TCGen05, fifth-gen tensor cores, new SM variants.
• TCGen05 — 5th Gen Tensor Cores — Blackwell tensor core instruction set.

CUDA Intrinsics

• Intrinsic Table (608 Entries) — Math, tensor, sync, warp intrinsics.
• Math Intrinsics — Fast-math, Newton-Raphson, special functions.
• Tensor Core Intrinsics — WMMA, GMMA, WGMMA instruction families.
• Sync & Warp Intrinsics — Barrier, vote, shuffle, match.

ELF/Cubin Output

• Custom ELF Emitter — ELFW internals, section construction, symbol table.
• Section Catalog & EIATTR — .nv.info attribute encoding, per-kernel metadata.
• Debug Information — DWARF generation for GPU debugging.
• Relocations & Symbols — CUBIN relocation types and symbol conventions.

Configuration

• CLI Options — ~100 flags registered in sub_432A00.
• Knobs System (1,294 Knobs) — ROT13 knob table, environment variables, INI files.
• Optimization Levels — O-level to phase mapping, --fast-compile tiers.
• DUMPIR & NamedPhases — Dumping IR at specific pipeline points.

Infrastructure

• Memory Pool Allocator — sub_424070, 3,809 callers, arena-style allocation.
• Hash Tables & Bitvectors — MurmurHash3-based maps, bitvector liveness sets.
• Thread Pool & Concurrency — pthread pool, GNU Make jobserver client.

Reference

• SASS Opcode Catalog — Complete SASS opcode enumeration.
• PTX Instruction Table — All PTX instructions with type signatures.
• EIATTR Attribute Catalog — Tag-length-value format for .nv.info attributes.

Reading Path 1: End-to-End Pipeline Understanding

Goal: understand how PTX text becomes SASS binary, what each stage does, and how control flows between subsystems.

1. Pipeline Overview — The complete flow diagram. Establishes all stages and their address ranges.
2. Entry Point & CLI — How ptxas is invoked, the ~100 CLI flags, and the sub_446240 driver function.
3. PTX Parser — The Flex scanner and Bison parser. How PTX text becomes an internal parse tree.
4. PTX-to-Ori Lowering — How the parse tree is lowered to Ori IR (basic blocks + instruction DAG).
5. Optimization Pipeline — The 159-phase PhaseManager. Phase factory, ordering, timing infrastructure.
6. SASS Code Generation — Mercury encoder, instruction selection, operand expansion, peephole.
7. ELF/Cubin Output — Custom ELF emitter, section layout, DWARF debug info, capsule Mercury.

Reading Path 2: Reimplementing a Specific Pass

Goal: reproduce the exact behavior of one optimization phase deeply enough to write a compatible replacement.

1. Pass Inventory & Ordering — Locate the phase in the 159-entry table. Note its index, vtable address, and pipeline position.
2. The phase's dedicated page (e.g., Copy Propagation & CSE, Predication). Every dedicated page contains the function address, decompiled algorithm, data flow, and controlling knobs.
3. Knobs System — Find which ROT13 knobs control the phase's behavior (enable/disable toggles, thresholds).
4. Ori IR Overview — Understand the IR data structures the phase operates on.
5. Register Model — The R/UR/P/UP register classes and their constraints.
6. Function Map — Cross-reference internal function addresses with the master function map.

Reading Path 3: Debugging Correctness

Goal: diagnose a miscompilation, crash, or incorrect SASS output by tracing the problem to a specific phase.

1. DUMPIR & NamedPhases — How to dump IR at specific pipeline points. Use DUMPIR to observe the IR before and after each phase.
2. Optimization Levels — Compare phase pipelines at different O-levels. If a bug appears at -O2 but not -O1, the diff identifies suspect phases.
3. Pipeline Overview — The pipeline is linear: Parse -> DAGgen -> OCG (159 phases) -> Mercury -> ELF. The stage where output first goes wrong narrows the search.
4. Knobs System — Check whether the suspect phase has enable/disable knobs. Toggle them to confirm or rule out the phase.
5. Instruction Scheduling and Scoreboards & Dependency Barriers — If the generated SASS hangs or produces wrong results under specific warp configurations, the scheduler or barrier insertion may be at fault.

Reading Path 4: Tuning Performance

Goal: understand what ptxas does at each optimization level and what knobs control aggressiveness.

1. Optimization Levels — The O-level to phase mapping, including --fast-compile tiers.
2. Knobs System — The 1,294 ROT13-encoded internal tuning parameters. The primary mechanism for fine-grained control.
3. Register Allocation — The fatpoint allocator directly determines register count, which determines maximum occupancy.
4. Instruction Scheduling — The scheduler's ReduceReg and DynBatch modes, WAR hazard resolution, and interaction with register pressure.
5. Peephole Optimization — The 3 peephole dispatchers that perform late SASS-level rewrites.
6. SM Architecture Map — Per-SM feature gates that influence code generation decisions.

Function Map

Binary: ptxas v13.0.88, 37.7 MB stripped ELF, ~40,000 functions
Documented: 2,063 unique functions across 70 wiki pages
This page: Top ~100 most cross-referenced functions, plus routing tables
Complete listings: Each wiki page has its own Function Map section with full details

This page is the central lookup index for identified functions in ptxas. It lists the functions that appear most frequently across the wiki (cross-cutting infrastructure and major entry points), and provides routing tables to find any function by address range or subsystem.

Confidence levels: CERTAIN = named in symbols or strings. HIGH = strong evidence from strings and call patterns (>90%). MEDIUM = structural analysis with partial string evidence (70-90%).

Core Infrastructure

These functions appear in 10+ wiki pages -- they are the universal building blocks called by nearly every subsystem.

Details: Memory Pools, Hash & Bitvector, Threading

Compilation Driver & CLI

Details: Pipeline Entry, Pipeline Overview, CLI Options

PTX Front End

Details: PTX Parser, PTX Directives, PTX to ORI

Static Initialization

Details: Pipeline Entry, Binary Layout

Phase Manager & Optimization Framework

Details: Phase Manager, Pass Inventory, Optimizer Pipeline

ORI IR & Instruction Access

Details: Instructions, Registers, Data Structures, CFG

Intrinsic Infrastructure

Details: Intrinsics Index, Math Intrinsics, Tensor Intrinsics, Sync & Warp

Register Allocator

Details: RegAlloc Overview, RegAlloc Algorithm, Spilling, ABI

Instruction Scheduling

Details: Scheduling Overview, Scheduling Algorithm, Latency Model, Scoreboards

Codegen & ISel

Details: ISel, Encoding, Peephole, Mercury, Templates

Bitfield Encoding

Details: Encoding, SASS Printing

ELF / CUBIN Output

Details: ELF Emitter, Sections, Relocations, Debug Info, Capsule Mercury

Knobs System

Details: Knobs, Opt Levels

Target-Specific Code

Details: Targets Index, Turing-Ampere, Ada-Hopper, Blackwell, tcgen05

Subsystem Routing Table

To find a specific function, locate it by address range or subsystem topic in this table. Each page contains a detailed Function Map section with complete listings.

By Subsystem Topic

By Address Range

Functions in the binary are clustered by subsystem. This table maps address ranges to the pages that document them.

Statistics

Top 10 Most-Called Functions

Top 5 Largest Functions

Top 10 Most Cross-Referenced (by wiki page count)

Documentation Coverage

Binary Layout

All addresses in this page apply to ptxas v13.0.88 (CUDA 13.0). Other versions will differ.

PTXAS v13.0.88 is a 37,741,528-byte stripped x86-64 ELF executable. Its .text section spans 26.2 MB (0x403520--0x1CE2DE2) containing 40,185 functions. This page maps every byte of the binary to the subsystem that owns it, derived from all 40 sweep reports covering the complete address range.

ELF Section Map

Total file composition:

Program Headers

Entry point: 0x42333C (ELF e_entry), which is inside .text (the CRT startup stub _start). The actual main is at 0x409460.

Three Subsystems

The .text section decomposes into three subsystems with distinct coding styles, data structures, and origins:

  .text linear address map (26.2 MB)
  0x403520                 0x67F000        0xC52000                          0x1CE2DE2
  |--- PTX Frontend 2.9 MB ---|-- Ori Optimizer 5.8 MB --|---- SASS Backend 17.6 MB ----|
  |          11%               |          22%             |              67%              |
  |  parsers, validators,      | passes, regalloc,        | encoding handlers, ISel,      |
  |  intrinsics, formatters    | scheduling, CFG analysis  | peephole, codecs, ABI, ELF    |

The backend dominates the binary because SASS instruction encoding is template-generated code: each of the ~4,000 encoding handler functions is a standalone vtable entry, never called directly. The optimizer has the highest function density (many small pass helpers), while the frontend has the largest average function size (complex validators and parsers).

Complete .text Address Map

The table below maps every address range in the .text section to its subsystem, function count, and key entry points. Data is aggregated from the 30 sweep partitions (p1.01 through p1.30).

PTX Frontend (0x403520--0x67F000, 2.9 MB)

Note on the 0x400000--0x403520 gap. The LOAD segment begins at 0x400000, but the first 13.6 KB before .text contains the ELF header (64 B at 0x400000), program headers (7 entries, 392 B), .interp (28 B, path to ld-linux-x86-64.so.2), .hash / .gnu.hash (symbol hash tables), .dynsym / .dynstr (dynamic symbol table, 146 entries), .gnu.version / .gnu.version_r (symbol versioning), .rela.plt (PLT relocations, 146 entries), and the .plt stub table (2,336 B, 146 stubs at 0x402C00--0x403520). These are standard ELF infrastructure, not ptxas application code. The first ptxas function begins at 0x403520.

Ori Optimizer (0x67F000--0xC52000, 5.8 MB)

SASS Backend (0xC52000--0x1CE2DE2, 17.6 MB)

.rodata Contents (7.5 MB)

The .rodata section at 0x1CE2E00--0x240BF8F is 29% of the binary by size. Its dominant consumers:

.bss Contents (84 KB)

.data Contents (14 KB)

Static Constructors

The .ctors section holds 12 entries executed before main. The four largest are:

The remaining 8 constructors handle memory allocator pool initialization, hash map infrastructure setup, diagnostic system initialization, and architecture vtable factory registration (sub_1CCD900).

Mega-Functions (>50 KB binary)

These eight functions account for 1.2 MB of code (4.8% of .text) but only 0.02% of the function count.

Most-Called Functions

The top five functions are all in the runtime infrastructure region (0x403520--0x42F000). Together they represent the core allocation, error handling, and data structure layer that the rest of the binary depends on.

Binary Composition by Purpose

Estimated from function classification across 30 sweep reports (p1.01--p1.30). Each function was assigned to a single purpose category based on its dominant behavior; functions straddling categories (e.g., a scheduling pass that also emits SASS) are attributed to the category consuming the larger share of their code.

The single largest consumer of code space is SASS instruction encoding. Each SM architecture generation requires its own set of per-opcode encoding/decoding handler functions. With support for SM75 through SM121 (six major generations), this yields approximately 4,000 encoding handlers, each a standalone function averaging 1,400 bytes.

Methodology

All addresses in this page apply to ptxas v13.0.88 (CUDA 13.0). Other versions will differ.

This page documents how the reverse engineering of ptxas v13.0.88 was performed. It serves as a transparency record so readers can assess the confidence of any claim in this wiki, and as a practical guide for anyone who wants to reproduce or extend the analysis.

Scope and Scale

PTXAS is a 37.7 MB stripped x86-64 ELF binary with no debug symbols, no DWARF information, and no export table beyond 146 libc/libpthread PLT stubs. Unlike NVIDIA's cicc (which is an LLVM fork), ptxas contains no LLVM code, no EDG frontend, and no third-party optimizer components. Every pass, data structure, and encoding table is proprietary NVIDIA code. This makes the analysis harder than LLVM-derived binaries -- there is no upstream source to compare against.

The 304 functions that Hex-Rays could not decompile are predominantly PLT stubs, computed-jump trampolines in the Flex DFA scanner, and the four mega-dispatch functions exceeding 200 KB (too large for Hex-Rays to handle within default limits). None are in critical analysis paths -- the dispatch functions are understood from their callee lists and the PLT stubs from their import names.

Why PTXAS Is Harder Than LLVM-Based Binaries

Reverse engineering cicc (NVIDIA's LLVM-based CUDA compiler) benefits from extensive prior art: LLVM's open-source codebase provides structural templates, pass names are registered in predictable patterns, and cl::opt strings directly name their global variables. PTXAS offers none of these advantages:

• No upstream source. Every identified function is identified from first principles -- string evidence, callgraph position, structural fingerprinting, or decompiled algorithm analysis. There is no reference implementation to compare against.
• ROT13 obfuscation. Internal names for tuning knobs and PTX opcode mnemonics are ROT13-encoded in the binary, requiring decoding before they become useful anchors.
• Vtable-heavy architecture. 39.6% of functions have zero static callers because they are dispatched through vtable pointers or function pointer tables. The call graph alone cannot reach them.
• Template-generated code. The SASS backend contains approximately 4,000 encoding handler functions generated from templates, each structurally near-identical. These dominate the function count but carry almost no unique identifying features.
• No pass registration infrastructure. LLVM passes register themselves via PassInfo objects with name strings. PTXAS phases are allocated by a factory switch (sub_C60D30) and their names are only visible through the NamedPhases registry and AdvancedPhase* timing strings -- far fewer anchors than LLVM's registration system.

Toolchain

All analysis was performed with IDA Pro 8.x and the Hex-Rays x86-64 decompiler. The entire effort is static analysis of the binary at rest -- no dynamic analysis (debugging, tracing, instrumentation) was used for function identification. Runtime tools (ptxas --stat, DUMPIR knob, --keep) were used only for validation and cross-referencing.

IDA Pro Setup and Initial Analysis

Loading the Binary

PTXAS is a dynamically-linked ELF with 146 PLT imports but no symbol table beyond those imports. IDA auto-analysis settings:

1. Processor: Meta PC (x86-64)
2. Analysis options: default. IDA correctly identifies the Flex DFA scanner tables, Bison parser tables, and the .ctors/.dtors sections.
3. Auto-analysis time: approximately 8-10 minutes on a modern machine for the 37.7 MB binary.
4. Compiler detection: IDA identifies GCC as the compiler. The binary uses the Itanium C++ ABI (confirmed by the embedded C++ name demangler at sub_1CDC780, 93 KB).

Post-Auto-Analysis Steps

After auto-analysis completes:

1. Run string extraction. IDA's auto-analysis finds 30,632 strings. All are exported via the analyze_ptxas.py IDA Python script.
2. Force function creation. Some address ranges, particularly the template-generated encoding handlers, are not automatically recognized as functions. IDA's "Create function" (P key) was applied selectively in the 0xD27000--0x1579000 range where encoding handler stubs are tightly packed.
3. Batch decompile. The IDA Python script iterates all 40,185 detected functions and calls ida_hexrays.decompile() on each, saving per-function .c files. 39,881 succeeded; 304 failed (PLT stubs, computed-jump trampolines, and 4 mega-functions exceeding decompiler limits).
4. Export control flow graphs. For each function, the script extracts the FlowChart (basic blocks, edges, per-instruction disassembly) as JSON. 80,078 graph files were produced.

Type Recovery

PTXAS uses no C++ RTTI (no typeid, no dynamic_cast -- the binary has no .data.rel.ro RTTI structures). Type recovery relies on:

• Vtable layout analysis. Each vtable is a contiguous array of function pointers in .data.rel.ro (4,256 bytes total). The vtable at off_22BD5C8 contains 159 entries, one per optimization phase. Each entry points to the phase's constructor function.
• Structure offset patterns. The pool allocator struct has free-list bins at offset +2128 and a mutex at +7128. The thread-local context is a 280-byte struct accessed via pthread_getspecific. These offsets were recovered from the decompiled code of sub_424070 (pool alloc, 3,809 callers) and sub_4280C0 (TLS accessor, 3,928 callers).
• Parameter/return type propagation. Once a function's signature is established (e.g., pool_alloc(pool*, size_t) -> void*), Hex-Rays propagates types to all 3,809 call sites, improving decompilation quality throughout the binary.

String-Driven Analysis

Strings are the single most productive source of function identification in ptxas. Of the 30,632 strings extracted, several categories are particularly valuable.

ROT13-Encoded Knob Names (2,000+ entries)

PTXAS uses ROT13 encoding as a light obfuscation layer on internal configuration names. Two massive static constructors populate these tables at startup:

• ctor_005 at 0x40D860 (80 KB) registers approximately 2,000 general OCG tuning knobs
• ctor_007 at 0x421290 (8 KB) registers 98 Mercury scheduler knobs

Each entry pairs a ROT13-encoded name with a hex-encoded default value. Decoding examples:

The knob names directly reveal subsystem organization. Names prefixed with Mercury* belong to the SASS encoder. Names prefixed with Scav* belong to the register allocator's scavenger. Names like XBlockWait* and WarDeploy* belong to the instruction scheduler. The knob lookup function GetKnobIndex at sub_79B240 performs inline ROT13 decoding and case-insensitive comparison, which was itself identified by tracing the xrefs from the ROT13-encoded strings.

ROT13-Encoded PTX Opcode Names (~900 entries)

A third static constructor, ctor_003 at 0x4095D0 (17 KB), populates a table of ~900 ROT13-encoded PTX opcode mnemonics. Decoding examples:

These strings are used by the PTX parser to match instruction mnemonics. Each xref from one of these strings leads to a parser action or instruction validator function.

Timing and Phase Name Strings

The compilation driver at sub_446240 emits per-stage timing via format strings:

Parse-time            : %.3f ms (%.2f%%)
CompileUnitSetup-time : %.3f ms (%.2f%%)
DAGgen-time           : %.3f ms (%.2f%%)
OCG-time              : %.3f ms (%.2f%%)
ELF-time              : %.3f ms (%.2f%%)
DebugInfo-time        : %.3f ms (%.2f%%)
PeakMemoryUsage = %.3lf KB

Tracing the xrefs from these format strings identifies the code that brackets each pipeline stage, revealing the stage boundaries within sub_446240.

The NamedPhases registry (string at 0x21B64C8, xrefs to sub_9F4040) and the AdvancedPhase* timing strings provide phase-level anchors within the 159-phase optimization pipeline:

• AdvancedPhaseBeforeConvUnSup, AdvancedPhaseAfterConvUnSup
• AdvancedPhaseEarlyEnforceArgs, AdvancedPhaseLateConvUnSup
• AdvancedPhasePreSched, AdvancedPhaseAllocReg, AdvancedPhasePostSched
• AdvancedPhaseOriPhaseEncoding, AdvancedPhasePostFixUp
• GeneralOptimizeEarly, GeneralOptimize, GeneralOptimizeMid, GeneralOptimizeMid2
• GeneralOptimizeLate, GeneralOptimizeLate2
• OriPerformLiveDead, OriPerformLiveDeadFirst through OriPerformLiveDeadFourth

Each AdvancedPhase* string xrefs to exactly one call site, which is a boundary marker in the phase pipeline. These 15 markers divide the 159-phase pipeline into named segments whose boundaries were used to identify the phases between each pair of markers.

Error and Diagnostic Strings

The central diagnostic emitter sub_42FBA0 (2,350 callers) prints error messages whose text reveals the calling function's purpose. Examples:

• "Please use -knob DUMPIR=AllocateRegisters for debugging" -- identifies the register allocator failure path at sub_9714E0
• "SM does not support LDCU" -- identifies SM capability checking in the instruction legalizer
• "Invalid knob identifier", "Invalid knob specified (%s)" -- identifies the knob parsing infrastructure around sub_79D070
• "fseek() error knobsfile %s", "[knobs]" -- identifies ReadKnobsFile at sub_79D070

Source File Path

One recovered source path provides a structural anchor:

/dvs/p4/build/sw/rel/gpgpu/toolkit/r13.0/compiler/drivers/common/utils/generic/impl/generic_knobs_impl.h

This string (at 0x202D4D8, 66 xrefs) is referenced from assertion checks throughout the knobs infrastructure, confirming that the knob system is a shared utility component (generic_knobs_impl.h) used across NVIDIA's compiler drivers.

Build and Version Strings

Cuda compilation tools, release 13.0, V13.0.88
Build cuda_13.0.r13.0/compiler.36424714_0

The version string at sub_612DE0 identifies both the exact build and the version reporting function. The Usage : string at 0x1CE3666 identifies the usage printer. The "\nCompile-unit with entry %s" string identifies the per-kernel compilation loop within the driver.

Vtable-Driven Discovery

The Phase Vtable Table

The most productive vtable discovery was the phase vtable table at off_22BD5C8 in .rodata. This is an array of 159 pointers, each pointing to a vtable for one optimization phase class. The phase factory function at sub_C60D30 is a 159-case switch statement that allocates a 16-byte phase object and assigns the corresponding vtable from this table:

// Simplified from decompiled sub_C60D30
switch (phase_index) {
    case 0:  obj->vtable = off_22BD5C8[0];  break;
    case 1:  obj->vtable = off_22BD5C8[1];  break;
    ...
    case 158: obj->vtable = off_22BD5C8[158]; break;
}
return obj;

Each vtable contains pointers to the phase's virtual methods. The virtual method at slot 0 is execute() (the phase body). The virtual method at slot 1 is isNoOp() (returns whether the phase should be skipped). The virtual method at slot 2 is getName() (returns the phase name string).

By following each of the 159 vtable entries to their execute() slot, every optimization phase's main function was identified. The getName() slot provided the phase name for phases that implement it. For phases that return a constant empty string, the name was inferred from the NamedPhases registry or from the AdvancedPhase* timing strings that bracket the phase in the pipeline.

Encoding Handler Vtables

The SASS backend uses vtable dispatch for instruction encoding. Each SASS opcode variant has its own encoding handler function, registered in dispatch tables rather than called directly. This explains why 15,907 functions (39.6%) have zero static callers -- they are reached exclusively through indirect calls via function pointer tables.

The encoding handler vtables were identified by their structural uniformity: every handler in the 0xD27000--0x1579000 range follows an identical template:

1. Set opcode ID via bitfield insert into the instruction word at a1+544
2. Load a 128-bit format descriptor from .rodata via SSE (movaps xmm0, xmmword_XXXXXX)
3. Initialize a 10-slot register class map
4. Register operand descriptors via sub_7BD3C0 / sub_7BD650 / sub_7BE090
5. Finalize encoding via sub_7BD260
6. Extract bitfields from the packed instruction word

The uniformity of this template allowed batch identification: once the template was recognized in a few handlers, the remaining ~4,000 were identified by structural matching alone.

Peephole Optimizer Vtable

The PeepholeOptimizer class at 0x7A5D10 has a reconstructed vtable with 7 virtual methods:

The three peephole dispatch mega-functions (sub_143C440 at 233 KB, sub_18A2CA0 at 231 KB, sub_198BCD0 at 239 KB) each serve a different SM generation family and call 1,100--1,336 pattern matcher functions. These dispatchers were identified by their enormous callee counts and their position in the pipeline after instruction encoding.

Callgraph Analysis

The 548,693-edge call graph, exported from IDA, reveals the binary's module structure and function relationships. Several callgraph properties were systematically exploited.

Hub Function Identification

Functions with extreme callee or caller counts serve as structural anchors:

Top callees (hub functions -- "fan-out" nodes):

Top callers (utility functions -- "fan-in" nodes):

The fan-out nodes identify the mega-dispatch functions: ISel, peephole, and dataflow. The fan-in nodes identify the shared infrastructure layer: memory allocation, encoding primitives, string formatting, and error reporting.

Module Boundary Detection

The call graph reveals clear module boundaries. Functions in the 0x400000--0x67F000 range (PTX frontend) rarely call functions in 0xC52000--0x1CE3000 (SASS backend) directly, and vice versa. The optimizer region (0x67F000--0xC52000) bridges the two, calling into both the frontend (for IR construction) and the backend (for encoding).

The call graph was used to validate the three-subsystem decomposition:

Propagation from Known Functions

Once a high-confidence function is identified, its callees and callers gain contextual identity. The most productive propagation chains:

1. sub_446240 (real main, CERTAIN) -> calls stage entry points for Parse, DAGgen, OCG, ELF, DebugInfo. Each stage's entry point was identified by following the timing format string pattern.

2. sub_C62720 (PhaseManager constructor) -> allocates 159 phase objects via sub_C60D30 (factory). The factory's 159 case targets are the phase constructors. Each constructor installs a vtable whose slot 0 points to the phase's execute() method.

3. sub_79B240 (GetKnobIndex) -> called from every function that reads a tuning knob. The first argument to GetKnobIndex is the ROT13-encoded knob name, so every call site reveals which knob a function checks.

4. sub_42FBA0 (diagnostic emitter) -> the format string argument at each of the 2,350 call sites reveals the error context. A call with "Cannot take address of texture/surface variable (%s)" identifies a PTX semantic checker.

Pattern Recognition

16-Byte Phase Objects

All 159 optimization phases share a uniform object layout:

Offset 0: vtable pointer (8 bytes) -- points to phase-specific vtable
Offset 8: phase data pointer or inline data (8 bytes)

The phase factory (sub_C60D30) allocates each phase as a 16-byte object from the pool allocator, sets the vtable pointer from the vtable table at off_22BD5C8, and returns the object. The PhaseManager stores these 159 objects in its internal array and iterates them to execute the pipeline.

Pool Allocator Usage Pattern

The custom pool allocator (sub_424070, 3,809 callers) is the dominant allocation mechanism. Its usage pattern is recognizable throughout the binary:

ptr = sub_424070(pool, size);   // Allocate
if (!ptr) sub_42BDB0();         // Fatal OOM -- never returns
// ... use ptr ...
sub_4248B0(ptr);                // Free (1,215 callers)

The OOM handler sub_42BDB0 (14 bytes, 3,825 callers) is a tiny wrapper that calls sub_42F590 (fatal internal error). Because every allocation site checks for failure and calls the same handler, the allocator usage pattern is a reliable structural marker. Finding sub_42BDB0 in a function's callee list confirms that function performs heap allocation.

SASS Encoding Handler Template

Every encoding handler in the backend follows a rigid 6-step template (described in the vtable section above). The key identification markers:

• Calls to sub_7B9B80 (bitfield insert, 18,347 callers)
• SSE movaps loading a 128-bit constant from .rodata
• Calls to sub_7BD3C0, sub_7BD650, or sub_7BE090 (operand registrars)
• Final call to sub_7BD260 (encoding finalize)

Any function matching this pattern is a SASS encoding handler. This template recognition identified approximately 4,000 handlers spanning 6 SM architecture generations.

Hash Map Infrastructure Pattern

The MurmurHash3-based hash map infrastructure (sub_426150 insert, sub_426D60 lookup, sub_427630 MurmurHash3) appears throughout the binary with a consistent usage pattern:

map = sub_425CA0(hash_fn, cmp_fn, initial_capacity);  // Create
sub_426150(map, key, value);                           // Insert (2,800 callers)
result = sub_426D60(map, key);                         // Lookup (422 callers)
sub_425D20(map);                                       // Destroy

The MurmurHash3 constants (0xcc9e2d51, 0x1b873593) in sub_427630 confirmed the hash algorithm. The hash map supports three modes (custom function pointers, pointer hash, integer hash) selected by flags at struct offset 84.

Data Artifacts

The complete IDA database was exported via analyze_ptxas.py into 8 JSON artifacts. These artifacts are the foundation for all subsequent analysis.

Total artifact storage: 1.14 GB (dominated by the 978 MB xref database).

What Each Artifact Reveals

Functions (ptxas_functions.json): The master index. Every function's address, size, instruction count, caller list, and callee list. The caller/callee lists are the basis for callgraph analysis. The is_thunk flag identifies PLT stubs (exclude from analysis). The is_library flag identifies functions IDA tagged as library code (CRT startup, jemalloc-like allocator internals).

Strings (ptxas_strings.json): The primary identification tool. Each string's xref list shows which functions reference it. Searching for "AdvancedPhase" returns 15 strings, each xref pointing to a pipeline boundary in the PhaseManager. Searching for strings starting with "Z" (ROT13 "M" for "Mercury") returns the Mercury subsystem's knob names. The 2,035 hex-encoded default value strings ("0k..." / "0x...") are paired 1:1 with knob name strings in the constructors.

Call graph (ptxas_callgraph.json): The structural backbone. Each edge records a direct call from one function to another. Indirect calls (vtable dispatch, function pointer callbacks) are not captured, which is the primary limitation -- the 15,907 zero-caller functions are almost all vtable-dispatched. The call graph is used for module boundary detection, propagation from known functions, and entry/exit point analysis.

Cross-references (ptxas_xrefs.json): The most comprehensive artifact. Contains all code-to-code, code-to-data, and data-to-data references detected by IDA. At 7.4 million entries, it is too large to load into memory on machines with less than 16 GB RAM. Used for deep analysis of specific functions: finding all references to a particular .rodata constant, tracing data flow through global variables, and identifying vtable consumers.

Comments (ptxas_comments.json): IDA's auto-generated comments (e.g., "File format: \\x7FELF") plus analyst-added annotations. The auto-comments on function prologues identify calling conventions and stack frame layouts. Analyst comments record identification rationale for reviewed functions.

Names (ptxas_names.json): IDA's auto-generated names for data and code addresses. Of 16,019 entries, approximately 9,670 are auto-generated string reference names (aLib64LdLinuxX8, aGnu, etc.) and ~6,349 are analyst-assigned or IDA-recovered names (PLT stubs, constructors, etc.). These names appear in the callgraph edges as from/to identifiers.

Imports (ptxas_imports.json): The 146 PLT imports. Key imports include pthread_* (13 functions), malloc/free/realloc, _setjmp/longjmp (used by the error recovery system), select/fcntl (used by the GNU Make jobserver client), and clock (used by the timing infrastructure).

Segments (ptxas_segments.json): The 24 ELF segments/sections. Used to establish the address space layout and map code/data boundaries. The .ctors section (104 bytes, 12 entries) is particularly important -- it lists the static constructors that initialize the ROT13 tables and the knob registry.

The 30-Region Sweep Approach

The primary analysis was conducted as a systematic address-range sweep of the entire .text section, divided into 30 contiguous regions. Each region was analyzed independently in a single session, producing a raw sweep report. The 40 report files (including sub-region splits) total 34,880 lines of working notes.

Region Partitioning

The .text section (0x403520--0x1CE2DE2, 26.2 MB) was divided into approximately 870 KB regions. The partitioning was not arbitrary -- region boundaries were chosen to align with subsystem boundaries where possible, so that each sweep report covers a coherent functional area.

Several regions were further split into sub-reports (p1.04a/b, p1.05a/b, p1.06a/b, p1.07a/b, p1.08a/b) when the initial analysis revealed that a region contained multiple distinct subsystems requiring separate treatment.

Sweep Report Structure

Each sweep report follows a consistent format:

================================================================================
P1.XX SWEEP: Functions in address range 0xAAAA000 - 0xBBBB000
================================================================================
Range: 0xAAAA000 - 0xBBBB000
Files found: NNN decompiled .c files (of which ~MMM are > 1KB)
Total decompiled size: X,XXX,XXX bytes
Functions in range (from DB): NNN
Named functions: NNN (or 0 if all are sub_XXXXXX)
Functions with identified callers: NNN

CONTEXT: [1-paragraph summary of the region's purpose]

================================================================================
SECTION 1: [Subsystem name]
================================================================================

### 0xAAAAAA -- sub_AAAAAA (NNNN bytes / NNN lines)
**Identity**: [Function identification]
**Confidence**: [CERTAIN / HIGH / MEDIUM]
**Evidence**:
  - [String evidence]
  - [Structural evidence]
  - [Callgraph evidence]
**Key code**:
  [Relevant decompiled excerpts]
**Note**: [Additional observations]

Each function entry records the address, size, decompiled line count, proposed identity, confidence level, evidence citations, and key code excerpts. The reports are raw working notes -- they contain false starts, corrections, and evolving hypotheses that were resolved as more context became available.

Analysis Ordering

The sweep was not performed in address order. The analysis followed an information-maximizing sequence:

1. p1.01 (infrastructure + CLI) first -- establishes the allocator, hash map, TLS, and diagnostic patterns that appear throughout the binary.
2. p1.11 (PhaseManager) second -- identifies all 159 phases and their vtable entries, providing the skeleton of the optimization pipeline.
3. p1.07 (register allocator) and p1.06 (scheduler) third -- these are the highest-complexity subsystems with the richest string evidence.
4. p1.12--p1.15 (SASS encoders) in batch -- once the encoding template was recognized, all encoder regions were swept rapidly with template matching.
5. p1.30 (library layer) late -- identifies shared infrastructure (ELF emitter, demangler, thread pool) referenced by earlier regions.
6. Remaining regions filled in by decreasing information density.

Cross-Referencing with PTXAS CLI

Several ptxas command-line features and internal mechanisms provide runtime validation of static analysis findings.

--stat and --verbose

Running ptxas --stat input.ptx prints per-kernel resource usage (register count, shared memory, stack frame size). This output is generated by sub_A3A7E0 (the IR statistics printer), which was identified from the format strings:

ptxas info    : Used %d registers, %d bytes smem, %d bytes cmem[0]

Comparing the --stat output against the decompiled statistics printer confirms the register counting and resource tracking logic.

--compiler-stats

Enables the timing output (Parse-time, DAGgen-time, OCG-time, etc.) from sub_446240. This confirms the pipeline stage ordering and the stage boundary functions identified by string xrefs.

--fdevice-time-trace

Generates Chrome trace JSON output showing per-phase timing. The trace parser at sub_439880 and the ftracePhaseAfter string at 0x1CE383F confirm the per-phase instrumentation infrastructure. The trace output lists phase names that can be cross-referenced against the 159-entry phase table.

DUMPIR Knob

The internal DUMPIR knob (accessed via -knob DUMPIR=<phase_name>) dumps the Ori IR at specified pipeline points. The string "Please use -knob DUMPIR=AllocateRegisters for debugging" at 0x21EFBD0 confirms this mechanism. The NamedPhases registry at sub_9F4040 maps phase names to pipeline positions. Available DUMPIR points include:

• OriPerformLiveDead, OriPerformLiveDeadFirst through OriPerformLiveDeadFourth
• AllocateRegisters (the register allocation phase)
• swap1 through swap6 (swap elimination phases)
• shuffle (instruction scheduling)

The DUMPIR output format reveals the IR structure: basic block headers, instruction opcodes, register names (R0--R255, UR0--UR63, P0--P7, UP0--UP7), and operand encodings. This runtime output was used to validate the IR format reconstructed from static analysis.

--keep Flag

The --keep flag preserves intermediate files. While ptxas does not emit intermediate text files in the same way as nvcc, the --keep behavior in the overall CUDA compilation pipeline (nvcc -> cicc -> ptxas) allows inspecting the PTX input that reaches ptxas, confirming the PTX grammar and instruction format expectations.

Confidence Levels

Every function identification in this wiki carries one of three confidence levels:

The distribution across the ~200 key identified functions in the Function Map:

• CERTAIN: ~30 functions (PLT imports, main, functions with unique identifying strings)
• HIGH: ~130 functions (string evidence + structural confirmation)
• MEDIUM: ~40 functions (inferred from callgraph context or structural similarity)

The remaining ~39,985 functions are either unidentified (template-generated encoding handlers, small utility stubs) or identified at subsystem level only (e.g., "this is an SM100 SASS encoding handler" without knowing which specific opcode it encodes).

Reproducing the Analysis

To reproduce this analysis from scratch:

1. Obtain the binary. Install CUDA Toolkit 13.0. The binary is at <cuda>/bin/ptxas. Verify: ptxas --version should report V13.0.88 and the binary should be 37,741,528 bytes. Build string: cuda_13.0.r13.0/compiler.36424714_0.

2. Run IDA auto-analysis. Open ptxas in IDA Pro 8.x with default x86-64 settings. Allow auto-analysis to complete (8-10 minutes). Accept GCC as the detected compiler.

3. Run the extraction script. Load analyze_ptxas.py in IDA's Python console. The script exports all 8 JSON artifacts plus per-function decompiled C files, disassembly files, and control flow graph JSON files. Expected runtime: 4-8 hours for the full export (the xref export dominates).

4. Decode ROT13 strings. Apply codecs.decode(s, "rot_13") to all strings in the knob constructors (ctor_003, ctor_005, ctor_007). This decodes ~3,000 obfuscated names into readable English identifiers.

5. Identify anchor functions. Start with the highest-confidence identifications:

   • main at 0x409460 (named in symbol table)
   • sub_446240 (real main -- called from main, contains timing format strings)
   • sub_C60D30 (phase factory -- 159-case switch)
   • sub_C62720 (PhaseManager constructor -- references phase vtable table)
   • sub_79B240 (GetKnobIndex -- inline ROT13 decoding)
   • sub_42FBA0 (diagnostic emitter -- 2,350 callers, severity dispatch)

6. Sweep the address space. Work through the .text section in regions of ~870 KB. For each region:

   • Count functions and decompiled file sizes
   • Identify string anchors (search for region-specific strings)
   • Classify functions by structural template (encoding handler, phase body, utility, etc.)
   • Propagate identities from known callers/callees
   • Record findings in the sweep report format

7. Cross-reference with runtime. Compile a simple CUDA kernel and run ptxas --stat --verbose --compiler-stats to observe runtime behavior. Use -knob DUMPIR=<phase> to dump IR at specific pipeline points. Compare the dumped IR format against the IR structure reconstructed from decompiled code.

Dependencies

The extraction script (analyze_ptxas.py) requires IDA Pro 8.x with Hex-Rays decompiler and Python 3.x. No external Python packages are needed -- only the IDA Python API (idautils, idc, idaapi, ida_bytes, ida_funcs, ida_segment, ida_nalt, ida_gdl, ida_hexrays).

Post-export analysis requires only the Python 3.8+ standard library (json, codecs, collections).

Debug Infrastructure: bugspec.txt

ptxas contains an internal fault injection framework that deliberately corrupts the Mercury IR to test compiler verification passes. The mechanism is entirely file-driven: if a file named ./bugspec.txt exists in the current working directory when ptxas runs, the function sub_A83AC0 reads it and injects controlled mutations into the post-register-allocation instruction stream. No CLI flag activates this -- file presence alone is sufficient. If the file is absent, a diagnostic is printed to stdout (Cannot open file with bug specification) and compilation proceeds normally.

File Format

The file contains a single line of six integers:

COUNT0,COUNT1,COUNT2,COUNT3 COUNT4 COUNT5

The first four are comma-separated; then a space; then two space-separated values. Each integer specifies the number of faults to inject for that bug category. Zero or negative disables the category.

Example: 3,2,1,0 0 1 injects 3 register bugs, 2 predicate bugs, 1 offset bug, and 1 bit-spill bug.

Bug Kind String Table

Each injected fault record carries a kind code (1--10) mapped to a string table at 0x21F0500:

Injection Algorithm

The injection proceeds in four phases:

1. Candidate collection. The function walks the Mercury IR instruction linked list (from context[0]+272). For each instruction, it checks which bug categories are active and whether the instruction qualifies:

• Register bugs (field0): Scans operands for type-tag 1 (register) with register class 6 (general) or 3 (predicate), excluding opcodes 41--44. Eligible instructions are collected into a candidate list.
• Predicate bugs (field1): Checks flag byte at instruction+73 for bit 0x10 (predicated). Eligible instructions are collected separately.
• Offset/spill bugs (field2): Calls sub_A56DE0 / sub_A56CE0 against the register allocator state (context[133]) to identify spill/refill instructions.
• Remat bugs (field3): Queries the rematerialization hash table (context+21 via sub_A54200) for instructions with remat entries.
• R2P/P2R bugs (field4): Checks instruction opcode (offset +72) for values 268, 155, 267, 173 (the R2P and P2R conversion opcodes, with bit-masked variants).
• Bit-spill bugs (field5): Checks operand count > 2, flag bit 0x10 at offset +28, and calls sub_A53DB0 / sub_A53C40 / sub_A56880 for bit-spill eligibility.

2. Random selection. Seeds the RNG with time(0) via srand(). For each active category, sub_A83490 randomly selects N instruction indices from the candidate list, where N is the count from bugspec.txt. The selector uses FNV-1a hashing on instruction addresses for collision avoidance, re-rolling duplicates.

3. Mutation application. For register and predicate categories, sub_A5EC40 iterates over selected instructions and calls sub_A5E9E0, which finds the last register operand, allocates a new register of the same class via sub_91BF30, and replaces the operand value. For offset bugs, the mutation adds +16 to the signed 24-bit offset field directly: *operand = (sign_extend_24(*operand) + 16) & 0xFFFFFF | (*operand & 0xFF000000).

4. Reporting. Prints to stdout:

Num forced bugs N
Created a bug at index I : kind K inst # ID [OFF] in operand OP correct val V replaced with W

Fault Record Structure (40 bytes)

Records are stored in a dynamic array at context[135].

Function Map

Significance

This is NVIDIA's internal compiler testing infrastructure for stochastic fault injection. It targets specific vulnerability surfaces in the register allocator and post-allocation pipeline: wrong-register assignments, address calculation errors, predicate propagation failures, rematerialization correctness, spill code integrity, and register-predicate conversion accuracy. The time(0)-seeded RNG produces different fault patterns on each run for the same bugspec.txt, enabling randomized stress testing of verification passes.

Embedded C++ Name Demangler

PTXAS statically embeds an Itanium ABI C++ name demangler rather than linking libc++abi or libstdc++. The demangler is a self-contained 41-function cluster spanning 0x1CD8B00--0x1CE1E60 in .text, with a single external entry point. The core recursive-descent parser at sub_1CDC780 (93 KB decompiled, 3,442 lines) handles the full Itanium mangling grammar: nested names, template arguments, substitutions, function types, and special names.

API and Integration

The public-facing function is sub_1CE23F0, whose signature matches __cxa_demangle exactly: it takes a mangled name string, an optional output buffer with length pointer, and a status pointer; it returns a malloc-allocated demangled string or NULL with a status code (-1 = memory failure, -3 = invalid arguments). The only caller of this function is the embedded terminate handler at sub_1CD7850, which prints the standard "terminate called after throwing an instance of '...'" diagnostic to stderr, demangling the exception type name before display.

Why Embedded

PTXAS imports only libc, libpthread, libm, and libgcc_s (146 PLT stubs total). It has no dependency on any C++ runtime library. The only C++ ABI symbol in the PLT is __cxa_atexit (at 0x401989), used to register the terminate handler. By embedding the demangler and terminate handler directly, NVIDIA avoids a runtime dependency on libstdc++ or libc++abi, which would otherwise be required solely for exception type name display in fatal error messages. This is consistent with the binary's overall strategy of minimizing external dependencies.

Function Map

Version Update Procedure

All addresses, function counts, and structural offsets in this wiki are specific to ptxas v13.0.88 (build cuda_13.0.r13.0/compiler.36424714_0, 37,741,528 bytes). When a new CUDA toolkit ships a different ptxas binary, the wiki must be updated. This section documents the procedure.

Version-Stable vs Version-Fragile Findings

Not everything changes between versions. Understanding what is stable dramatically reduces update effort.

Version-stable (survives across minor and most major releases unchanged):

Version-fragile (changes with every recompilation):

Identifying Function Address Changes

When loading a new ptxas version into IDA:

1. Extract the same 8 JSON artifacts using analyze_ptxas.py (or equivalent). The critical artifacts for diffing are ptxas_functions.json (address, size, callee list) and ptxas_strings.json (string content, xref locations).

2. Match functions by invariant properties. Functions cannot be matched by address alone. Use these matching criteria in priority order:

   • String anchors. Functions containing unique string references (e.g., the function referencing "Please use -knob DUMPIR=AllocateRegisters") can be matched across versions by searching for the same string in the new binary. This is the highest-confidence matching method.
   • Size + callee signature. For functions without string anchors, match by (approximate size, sorted callee list). A function of ~2,100 bytes calling the pool allocator, OOM handler, and hash map insert is almost certainly the same function even if its address shifted by megabytes.
   • Callgraph position. Functions identified by their caller/callee topology: the phase factory is the function called from the PhaseManager constructor with 159+ case targets. The diagnostic emitter is the function with 2,000+ callers that calls vfprintf.
   • Vtable slot position. Phase execute() methods are at vtable slot 0. If the vtable table address changes but still contains 159 entries, the slot positions identify each phase.
   • Template fingerprinting. Encoding handlers matching the 6-step template (bitfield insert via the highest-caller utility, movaps from .rodata, operand registrars, finalize call) are encoding handlers in any version.

3. Diff the function lists. Produce a mapping {old_addr -> new_addr} for all matched functions. Functions present in the new binary but absent in the old are new (likely new SM target support). Functions absent in the new binary are removed (dropped legacy SM support) or merged.

Updating Sweep Reports

The 30-region sweep reports in ptxas/raw/ are version-locked historical records -- they document the analysis of v13.0.88 and should not be overwritten. For a new version:

1. Re-run the sweep with new address ranges derived from the new binary's function list. The region partitioning should follow the same subsystem-aligned strategy: infrastructure first, then PhaseManager, then high-complexity subsystems, then batch encoding handlers.

2. Name new reports with a version suffix: p2.01-sweep-v13.1-0xNNN-0xMMM.txt (or whatever scheme distinguishes the version).

3. Cross-reference against old reports. For each region, note which functions moved, which are new, and which disappeared. The old sweep reports provide the expected function identities; the new sweep validates whether those identities still hold at the new addresses.

Pages Most Sensitive to Version Changes

These wiki pages require immediate updates when the binary changes:

Recommended Update Workflow

The update follows a five-step sequence. Steps 1-2 are mechanical; steps 3-5 require analyst judgment.

Step 1: Extract new IDA artifacts.

Load the new ptxas binary into IDA Pro 8.x. Run analyze_ptxas.py to produce the 8 JSON artifacts and per-function decompiled .c files. Store them in a version-specific directory (e.g., ptxas/ida-v13.1/ or alongside the existing artifacts with clear version labeling).

Step 2: Diff against the old artifacts.

Write or use a diff script that:

• Compares ptxas_functions.json (old vs new) by matching on string anchors, size+callee signature, and callgraph position.
• Produces a {old_addr -> new_addr} mapping for matched functions.
• Lists unmatched functions in both directions (new functions, removed functions).
• Compares ptxas_strings.json to detect new strings, removed strings, and strings whose xref functions changed.
• Reports total function count delta, binary size delta, and new section addresses.

Step 3: Update address-sensitive pages.

Using the address mapping from Step 2:

• Update every sub_XXXXXX reference in function-map.md, binary-layout.md, and all pages listed in the sensitivity table above.
• Update the "Scope and Scale" table in methodology.md with new function counts, string counts, binary size, and build string.
• Update VERSIONS.md with the new binary metadata.
• For pages with address ranges (sweep boundaries, subsystem regions), recompute the ranges from the new function list.

Step 4: Verify key struct layouts.

Struct offset changes are the most dangerous kind of version drift because they silently invalidate decompiled code analysis. For each documented struct:

• Re-decompile the struct's primary accessor function (e.g., sub_424070 for the pool allocator, sub_4280C0 for the TLS context).
• Compare field offsets against the documented layout.
• If offsets shifted, update the struct documentation and propagate the change to all pages that reference those offsets.

Priority structs to verify: pool allocator (free-list bins at +2128, mutex at +7128), TLS context (280 bytes), Mercury instruction node (112 bytes), scheduler context (~1000 bytes), allocator state (1590+ bytes), phase objects (16 bytes).

Step 5: Validate phase pipeline.

• Re-extract the phase vtable table (find the new address of the 159-entry pointer array in .data.rel.ro).
• Verify all 159 phases are present and in the expected order.
• Check for new phases (count > 159) or removed phases (count < 159).
• Re-run ptxas --fdevice-time-trace on a test kernel and cross-reference the phase names in the trace output against the wiki's phase list.

Raw Data Locations

All raw analysis artifacts for the current version (v13.0.88) live in the repository under ptxas/:

When updating to a new version, preserve the existing artifacts for v13.0.88 (rename or move to a versioned subdirectory) and store new artifacts alongside them. The sweep reports in ptxas/raw/ are historical records and should never be overwritten.

Limitations and Known Gaps

• No dynamic validation of optimization correctness. All findings are from static analysis. The identified phase algorithms have not been tested against runtime inputs to verify they produce correct output for all corner cases.

• 39.6% of functions are vtable-dispatched. Functions with zero static callers can only be reached by finding the vtable or function pointer table that references them. Some vtables in deep .rodata may have been missed, leaving some functions orphaned.

• No upstream reference for any code. Unlike cicc (LLVM fork) or nvcc (EDG frontend), ptxas has no open-source analog. Every identification is from first principles. This limits confidence for functions where string evidence is absent and structural analysis is the only basis.

• Template-generated code is indistinguishable. The ~4,000 SASS encoding handlers are generated from internal templates. Without the template source, mapping individual handlers to specific opcodes requires tracing the dispatch table entries, which has only been done for select handlers.

• Mega-functions are partially opaque. The four functions exceeding 200 KB (sub_169B190 at 280 KB, sub_143C440 at 233 KB, sub_198BCD0 at 239 KB, sub_18A2CA0 at 231 KB) could not be decompiled by Hex-Rays. Their behavior is understood from their callee lists (13,000--15,870 callees each) and their position in the pipeline, but the internal dispatch logic is known only at the disassembly level.

• ROT13 decoding is necessary but not sufficient. Decoding the 2,000+ knob names reveals the existence of tuning parameters but not their semantics. A knob named MercuryPresumeXblockWaitBeneficial can be decoded from ROT13, but understanding what "xblock wait beneficial" means requires analyzing the code paths that read the knob.

• Version-specific addresses. All addresses in this wiki apply to ptxas v13.0.88 (build cuda_13.0.r13.0/compiler.36424714_0). Other CUDA toolkit versions will have different addresses, different function counts, and potentially different phase orderings. However, the analysis methodology (string-driven, vtable-driven, callgraph propagation) applies to any version.

• Indirect calls are undercounted. The 548,693-edge call graph captures only direct call instructions resolved by IDA. Virtual calls through vtable pointers, function pointer callbacks, and computed jumps are not fully captured. The true call graph is significantly denser than what is recorded.

Corrections Log

This section documents every factual error discovered and corrected during the wiki improvement pass. Each entry records the error, the correction, affected pages, and the agent task that performed the fix. The full detail for each correction is in ptxas/raw/P5_11_corrections_log_report.txt.

Summary

Corrections by Severity

Systematic errors (affected 5+ pages each)

Identity misattributions

Inverted semantics

Wrong numeric values

Phantom data and scope errors

Minor corrections

Error Categories

Lessons Learned

1. Behavioral inference is unreliable for opcode identity. Observing that an opcode appears in branch contexts does not make it BRA. Always check the authoritative ROT13 name table.

2. Cross-page consistency checks catch conflicting speculations. Five pages independently naming the same field (context+1584) is a strong signal that at least four are wrong.

3. Counts from partial analysis are systematically low. The "3 values" for context+1552 and "139 named phases" both resulted from stopping the search too early. Exhaustive binary sweeps consistently reveal more entries.

4. Function size is not a reliable identity signal. sub_83EF00 (29KB) was large enough to seem like a major driver, but size alone does not distinguish a peephole optimizer from a loop unroller.

5. ROT13 decoding + binary cross-validation is the gold standard. Every correction that replaced speculative labels with ROT13-decoded names has held up under subsequent audits.

Version Tracking

All addresses in this page apply to ptxas v13.0.88 (CUDA 13.0). Other versions will differ.

This page documents the exact ptxas binary under analysis and the version-related metadata recovered from the stripped ELF.

Binary Under Analysis

Embedded Version Strings

sub_432A00 (0x432A00, CLI option registration) contains the self-identification strings that ptxas prints for --version / --list-version:

The "ptxocg.0.0" tag also appears in sub_43A400 (compilation setup) and at address 0x1CE74AB in the .rodata section, identifying the backend optimizer component embedded inside ptxas.

Default Target Architecture

sub_6784B0 returns sm_75 (Turing) as the default compilation target when no --gpu-name flag is supplied. This is consistent with the CUDA 13.0 toolkit defaulting to a Turing-class GPU.

The full set of architecture strings referenced in the front-end validators (addresses 0x460000-0x4D5000) includes:

sm_20  sm_30  sm_35  sm_50  sm_60  sm_75  sm_80  sm_86  sm_89  sm_90

with sm_%d format-string patterns covering all supported SM codes.

Output ELF Format

Cubins emitted by ptxas use the ELF standard with:

The SM-version-to-code-object mapping lives in the ELF emitter at sub_1C9F280. Example encodings recovered from sub_A3D000 range:

Build System Metadata

The source path leaked through __FILE__ macros in the knobs infrastructure reveals the NVIDIA internal build tree layout:

/dvs/p4/build/sw/rel/gpgpu/toolkit/r13.0/compiler/
  drivers/common/utils/generic/impl/generic_knobs_impl.h

Key observations:

• /dvs/p4/ -- Perforce depot root on the DVS (Driver Verification System) build farm.
• sw/rel/gpgpu/toolkit/r13.0/ -- Release branch for CUDA toolkit 13.0.
• compiler/drivers/common/ -- Shared compiler driver code (used by both ptxas and cicc).
• generic_knobs_impl.h -- The knob system implementation header; the __FILE__ macro at lines 395-1090 of this file is embedded in ptxas error metadata.

Evidence Index

Key Functions

Compilation Pipeline Overview

All addresses in this page apply to ptxas v13.0.88 (CUDA 13.0). Other versions will differ.

This page maps the complete end-to-end flow of a PTX assembly through ptxas v13.0.88, from the initial CLI invocation to the final ELF/cubin binary output. Each stage is a self-contained subsystem with its own address range, data structures, and failure modes. The links below lead to dedicated pages with reimplementation-grade detail for every stage.

Pipeline Diagram

nvcc / cicc
  |  (PTX text file or --input-as-string)
  v
+================================================================+
| ptxas v13.0.88 (37.7 MB, ~40,000 functions)                   |
|                                                                |
|  1. Entry & CLI Parsing ----------> [entry.md]                 |
|     |  main -> sub_446240 -> sub_434320                        |
|     |  target arch, opt level, --maxrregcount, knobs           |
|     v                                                          |
|  2. PTX Lexer + Parser -----------> [ptx-parser.md]            |
|     |  sub_451730: Flex scanner, Bison grammar                 |
|     |  ROT13-decoded opcode table (900+ mnemonics)             |
|     |  30+ per-instruction semantic validators                 |
|     v                                                          |
|  3. PTX Directive Handling --------> [ptx-directives.md]       |
|     |  .version, .target, .entry, .func, .reg, .shared         |
|     |  register constraints, ABI configuration                 |
|     v                                                          |
|  4. PTX-to-Ori Lowering ----------> [ptx-to-ori.md]           |
|     |  PTX AST -> Ori IR (basic blocks, virtual registers)     |
|     |  address space annotation, special register mapping      |
|     v                                                          |
|  5. 159-Phase Optimization -------> [optimizer.md]             |
|     |  PhaseManager: sub_C62720 (constructor),                 |
|     |                sub_C64F70 (executor)                     |
|     |  10 stages, 17 AdvancedPhase hooks,                     |
|     |  8-phase Mercury encoding sub-pipeline                   |
|     |  per-kernel via sub_7FBB70 -> sub_7FB6C0                 |
|     v                                                          |
|  6. Register Allocation ----------> [../regalloc/overview.md]  |
|     |  Fatpoint algorithm, phase 101 (AdvancedPhaseAllocReg)   |
|     |  spill/fill insertion, ABI register reservations          |
|     v                                                          |
|  7. Instruction Scheduling -------> [../scheduling/overview.md]|
|     |  3-phase: pre-schedule (97), post-schedule (106),        |
|     |           post-fixup (111)                               |
|     |  scoreboard generation, dependency barriers              |
|     v                                                          |
|  8. SASS Encoding ----------------> [../codegen/encoding.md]   |
|     |  530 instruction encoding handlers (vtable dispatch)     |
|     |  Mercury format: phases 113-122                          |
|     |  Capsule Mercury (default on sm_100+)                    |
|     v                                                          |
|  9. ELF/Cubin Output -------------> [output.md]               |
|     |  sub_612DE0 (finalizer) -> sub_1C9F280 (ELF emitter)    |
|     |  section layout, symbol table, relocations               |
|     |  DWARF debug info, EIATTR attributes                     |
|     v                                                          |
|  OUTPUT: .cubin / .o (ELF)                                    |
+================================================================+

Side paths:
  * Capsule Mercury (--cap-merc) -----> [../codegen/capmerc.md]
  * Debug info (all stages) ----------> [../output/debug-info.md]
  * SASS text (--verbose) ------------> [../codegen/sass-printing.md]

Narrative Walk-Through: One Kernel, Start to Finish

A concrete trace of a single-kernel PTX module compiled for sm_100 at -O2:

1. PTX text arrives (~2--200 KB). Either read from a .ptx file or received in-memory via --input-as-string from nvcc. The driver sub_446240 establishes a setjmp recovery point, parses CLI options into the 1,352-byte options block, and allocates the "Top level ptxas memory pool".

2. Lexer + Parser (sub_451730). A Flex-generated scanner tokenizes the PTX text into a token stream. Tokens flow into a Bison-generated LALR parser that builds an AST. The opcode dispatch table (sub_46E000, 93 KB, 1,168 callees) routes each instruction mnemonic through ROT13 decoding, type resolution, and 30+ per-instruction semantic validators. For a 5 KB PTX kernel, the parser typically produces ~200--500 AST nodes with ~50 virtual register declarations. The "PTX parsing state" pool holds all AST memory.

3. Directive processing and CompileUnitSetup. .version/.target directives configure the SM profile via sub_6765E0 (54 KB profile constructor). .entry/.func directives establish the kernel boundary. .reg/.shared/.const directives declare resources. sub_43B660 computes the physical register budget from .maxnreg, --maxrregcount, and .maxntid constraints. The 1,936-byte profile object is now populated with codegen factory value (36864 for sm_100), scheduling parameters, and capability flags.

4. PTX-to-Ori lowering (DAGgen). sub_6273E0 (44 KB) converts each AST instruction into an Ori IR node: a basic block with virtual registers, control flow edges, and memory space annotations. Special registers (%ntid, %laneid, %smid) map to internal IDs. Address computation uses a 6-bit operand type encoding. A 500-instruction PTX kernel typically produces ~600--1,200 Ori instructions (expansion from pseudo-ops, address calculations, and predicate materialization). The "Permanent OCG memory pool" is created here to hold all IR state.

5. 159-phase OCG pipeline (sub_C62720 constructs, sub_C64F70 executes). Each phase is a 16-byte polymorphic object with execute(), isNoOp(), and getName() vtable methods. The PhaseManager iterates the phase table at 0x22BEEA0, skipping any phase whose isNoOp() returns true. At -O2, roughly 80--100 of the 159 phases are active. Typical expansion factors: the initial 1,000 Ori instructions may grow to 1,200--1,500 after unrolling and intrinsic expansion, then shrink to 800--1,000 after CSE/DCE, then re-expand to 1,500--2,500 after register allocation spill/fill insertion. The PhaseManager logs "Before <phase>" / "After <phase>" strings (visible in the sub_C64F70 decompile) for DUMPIR.

6. Register allocation (phase 101, sub_971A90). The Fatpoint algorithm attempts NOSPILL allocation first. If pressure exceeds the register budget, the spill guidance engine (sub_96D940, 84 KB) computes spill candidates across 7 register classes, and the retry loop makes up to N attempts (knob 638/639) with progressively more aggressive spilling. Physical register assignments are committed; spill/fill instructions are inserted into the Ori IR.

7. Instruction scheduling (phases 97, 106, 111). Three scheduling passes assign dependency barriers and reorder instructions for pipeline throughput. The scoreboard generator tracks 6 dependency barriers per warp. For a 1,500-instruction kernel, scheduling typically produces a ~2,000--3,000-entry instruction stream after barrier insertion and NOP padding.

8. SASS encoding (phases 113--122). Each Ori instruction is lowered to a 128-bit SASS binary instruction via the 530-handler vtable dispatch. The 1,280-bit (160-byte) encoding workspace at instruction+544 is filled by sub_7B9B80 (bitfield insert, 18,347 callers). A 2,000-instruction kernel produces ~32 KB of raw SASS binary. On sm_100+, Capsule Mercury (capmerc) is the default format, embedding PTX source alongside the SASS.

9. ELF/cubin emission (sub_612DE0, 47 KB). The finalizer assembles the cubin: .text.FUNCNAME (SASS binary), .nv.info.FUNCNAME (EIATTR attributes), .nv.shared.FUNCNAME (shared memory layout), .nv.constant0.FUNCNAME (constant bank), plus global sections (.shstrtab, .strtab, .symtab). Section layout (sub_1CABD60, 67 KB) assigns addresses; the master ELF emitter (sub_1C9F280, 97 KB) writes headers, section tables, and program headers. A single-kernel cubin for a medium-complexity kernel is typically 40--120 KB.

Approximate data sizes at each stage (medium kernel, sm_100, -O2):

Timed Phases

The compilation driver sub_446240 measures six timed phases per compile unit and reports them when --compiler-stats is enabled. The format strings are embedded directly in the binary:

Additional aggregate stats:

CompileTime = %f ms (100%)
PeakMemoryUsage = %.3lf KB

The per-unit header prints "\nCompile-unit with entry %s" before each kernel's phase breakdown.

Per-Kernel Parallelism

ptxas supports two compilation modes for multi-kernel PTX modules:

Single-Threaded Mode (Default)

The compilation driver sub_446240 iterates over compile units sequentially. For each kernel entry:

1. sub_43CC70 -- per-entry compilation unit processor, skips __cuda_dummy_entry__
2. sub_7FBB70 -- per-kernel entry point, prints "\nFunction name: " + kernel name
3. sub_7FB6C0 -- pipeline orchestrator: builds phases via sub_C62720, executes via sub_C64F70
4. Cleanup: destroys 17 analysis data structures (live ranges, register maps, scheduling state)

Each kernel runs through the entire 159-phase pipeline independently. Cross-kernel state is limited to shared memory layout and the global symbol table.

Thread Pool Mode (--split-compile)

When --allow-expensive-optimizations or --split-compile is active, ptxas uses a pthread-based thread pool for per-kernel parallelism:

• Pool constructor (sub_1CB18B0): allocates a 184-byte pool struct (0xB8), spawns N detached worker threads via pthread_create, initializes mutex at +24, two condition variables at +64 and +112
• Task submit (sub_1CB1A50): allocates a 24-byte task node {func_ptr, arg, next}, enqueues via linked list, broadcasts on cond_work
• Jobserver integration (sub_1CC7300): reads MAKEFLAGS environment variable, parses --jobserver-auth= for either fifo: named pipes or pipe-based file descriptors, throttles thread count to respect GNU Make's -j slot limit

The thread pool is used throughout the OCG and ELF phases (stages 5-9 in the diagram). Each worker thread receives its own thread-local context (sub_4280C0, 280-byte TLS struct with per-thread error flags, diagnostic suppression state, Werror flag, and synchronization primitives).

Thread-Local Context Layout

struct ThreadLocalContext {  // 280 bytes (0x118), per-thread via pthread_getspecific
    uint64_t error_flags;          // +0:   error/warning state flags
    uint64_t has_error;            // +8:   nonzero if error occurred
    // +16..+48: internal fields (jmp_buf pointer, pool pointer, counters)
    uint8_t  diag_suppress;        // +49:  diagnostic suppression flag
    uint8_t  werror_flag;          // +50:  --Werror promotion flag
    // +51..+127: reserved / internal state
    pthread_cond_t  cond;          // +128: condition variable (48 bytes)
    pthread_mutex_t mutex;         // +176: per-thread mutex (40 bytes)
    sem_t           sem;           // +216: semaphore (32 bytes)
    // +248..+279: linked-list pointers (global thread list at +256/+264)
};

Accessed by sub_4280C0 (3,928 callers -- the single most-called function in the binary). On first call in a new thread, allocates and initializes via malloc(0x118) + memset + pthread_cond_init + pthread_mutex_init + sem_init. The decompiled code confirms the 280-byte size: v5 = malloc(0x118u), followed by memset(v5, 0, 0x118u), pthread_cond_init(v5 + 128), pthread_mutex_init(v5 + 176), sem_init(v5 + 216). After initialization, the struct is inserted into a global doubly-linked list (offsets +256 and +264 hold prev/next pointers, protected by a global mutex). The pthread_setspecific(key, v5) call stores the pointer for subsequent pthread_getspecific retrieval.

Key Function Call Chain

The top-level control flow from program entry to ELF output:

main (0x409460, 84 bytes)
  |  setvbuf(stdout/stderr, unbuffered)
  v
sub_446240 (0x446240, 11KB) ---- "Top-level compilation driver"
  |
  |-- sub_434320 (0x434320, 10KB) -- Parse CLI options, validate flags
  |     reads: --gpu-name, --maxrregcount, --opt-level, --verbose,
  |            --compiler-stats, --split-compile, --fast-compile
  |
  |-- [allocate "Top level ptxas memory pool"]
  |-- [allocate "Command option parser" pool]
  |
  |-- sub_445EB0 (setup) ----------- Target configuration, texturing mode
  |     sub_43A400 --------------- SM-specific defaults ("ptxocg.0.0")
  |     sub_43B660 --------------- Register/resource constraint calculation
  |
  |-- sub_451730 (0x451730, 14KB) -- Parser initialization
  |     |  "PTX parsing state" pool allocation
  |     |  Builtin symbol table: %ntid, %laneid, %smid, %clock64, ...
  |     |  sub_46E000 (93KB) ---- Opcode-to-handler dispatch table (1168 callees)
  |     v
  |   [Flex lexer + Bison parser: PTX text -> AST]
  |
  |-- for each compile unit:
  |     sub_4428E0 (0x4428E0, 14KB) -- PTX input validation
  |     |  .version/.target checks, ABI mode selection
  |     |  --extensible-whole-program, --compile-only handling
  |     |
  |     sub_43CC70 (5.4KB) --------- Per-entry unit processor
  |     |  skip __cuda_dummy_entry__
  |     |  generate .sass and .ucode sections
  |     |
  |     sub_7FBB70 (198 bytes) ----- Per-kernel entry point
  |       |
  |       sub_7FB6C0 (1.2KB) ------- Pipeline orchestrator
  |         |  check knob 298 (NamedPhases mode)
  |         |  if NamedPhases: delegate to sub_9F63D0
  |         |  else:
  |         |    sub_C62720 -- PhaseManager constructor (159 phases)
  |         |    sub_C60D20 -- get default phase table (at 0x22BEEA0)
  |         |    sub_C64F70 -- execute all phases
  |         |  cleanup: destroy 17 analysis data structures
  |         v
  |       [159-phase pipeline: optimization -> regalloc -> scheduling -> encoding]
  |
  |-- sub_612DE0 (0x612DE0, 47KB) -- Kernel finalizer / ELF builder
  |     |  "Finalizer fastpath optimization"
  |     |  version: "Cuda compilation tools, release 13.0, V13.0.88"
  |     |  build:   "Build cuda_13.0.r13.0/compiler.36424714_0"
  |     |
  |     sub_1CB53A0 (13KB) ------- ELF world initializer (672-byte object)
  |     |  "elfw memory space", .shstrtab, .strtab, .symtab
  |     |
  |     sub_1CB3570 (10KB) ------- Add .text.FUNCNAME sections (44 callers)
  |     sub_1CB68D0 (49KB) ------- Symbol table builder
  |     sub_1CABD60 (67KB) ------- Section layout & memory allocation
  |     sub_1CD48C0 (22KB) ------- Relocation resolver
  |     sub_1C9B110 (23KB) ------- Mercury capsule builder (capmerc)
  |     sub_1C9F280 (97KB) ------- Master ELF emitter (largest in range)
  |     sub_1CD13A0 (11KB) ------- Final file writer
  |
  v
[report CompileTime, PeakMemoryUsage, per-phase breakdown]

Memory Pools

ptxas uses a custom hierarchical pool allocator (sub_424070 / sub_4248B0, the most-called allocation functions with 3,809 and 1,215 callers respectively) instead of the system malloc/free. Three named pools are created during the top-level driver:

Additional per-subsystem pools exist:

• "PTX parsing state" -- created by sub_451730, holds the lexer/parser symbol tables and AST nodes
• "elfw memory space" -- created by sub_1CB53A0, holds the ELF world object (672 bytes) and section data

Pool Allocator Internals

The allocator at sub_424070 implements a dual-path design:

• Small allocations (up to 4,999 bytes / 0x1387): 8-byte-aligned, size-class binned free lists at pool struct offset +2128. Pop from free list head on alloc, push back on free.
• Large allocations (above 4,999 bytes): boundary-tag allocator with coalescing of adjacent free blocks.
• Thread safety: pthread_mutex_lock/unlock around all pool operations, mutex at pool struct offset +7128.
• OOM handling: calls sub_42BDB0 (3,825 callers) which triggers longjmp-based fatal abort via sub_42F590.

Pipeline Stage Breakdown

Terminology note. The 6 stages below (Parse, CompileUnitSetup, DAGgen, OCG, ELF, DebugInfo) correspond to the 6 timed phases measured by --compiler-stats. They cover the entire program lifecycle. The OCG stage (Stage 4 here) is itself subdivided into 10 internal stages in the Pass Inventory, numbered OCG-Stage 1--10. To avoid confusion, cross-references use "timed phase" for the 6 whole-program stages and "OCG stage" for the 10 optimizer sub-stages.

Stage 1: Parse (Parse-time)

The Flex-generated scanner and Bison-generated parser consume PTX text and produce an internal AST. The opcode dispatch table at sub_46E000 (93KB, 1,168 callees) registers type-checking rules for every PTX instruction. Thirty separate validator functions (in 0x460000-0x4D5000) enforce SM architecture requirements, PTX version constraints, operand types, and state space compatibility. See PTX Parser.

Stage 2: CompileUnitSetup (CompileUnitSetup-time)

Target configuration via sub_43A400: sets SM-specific defaults (texturing mode, cache policies, def-load-cache, force-load-cache), applies --fast-compile shortcuts, configures ABI (parameter registers, return address register, scratch registers). Register constraints computed by sub_43B660 from .maxnreg, --maxrregcount, .minnctapersm, and .maxntid directives. See Entry Point & CLI.

Stage 3: DAGgen (DAGgen-time)

Lowers the validated PTX AST into the Ori intermediate representation: basic blocks with a control flow graph, virtual registers, and memory space annotations. Special PTX registers (%ntid, %laneid, %smid, %ctaid, etc.) are mapped to internal identifiers. Operand processing at sub_6273E0 (44KB) handles address computation with a 6-bit operand type encoding. See PTX-to-Ori Lowering.

Stage 4: OCG (OCG-time)

The core of ptxas: the 159-phase Optimized Code Generation pipeline. This single timed phase encompasses:

• Early optimization (phases 13-36): general optimization, branch/switch, loop simplification, strength reduction, unrolling, pipelining, barrier removal
• Mid-level optimization (phases 37-58): GVN/CSE, reassociation, commoning, late expansion, speculative hoisting
• Late optimization (phases 59-95): loop fusion, predication, GMMA propagation, legalization
• Register allocation (phase 101): Fatpoint algorithm
• Instruction scheduling (phases 97, 106, 111): pre-schedule, post-schedule, post-fixup
• Mercury encoding (phases 113-122): SASS binary format generation

The PhaseManager (sub_C62720) instantiates phases via a 159-case factory switch (sub_C60D30), each phase a 16-byte polymorphic object with a vtable providing execute(), isNoOp(), and getName() methods. See Optimization Pipeline.

Stage 5: ELF (ELF-time)

The finalizer sub_612DE0 (47KB) assembles the NVIDIA ELF/cubin from the compiled SASS. Section layout (sub_1CABD60, 67KB) assigns addresses to shared memory, constant banks (with OCG deduplication), local memory, and reserved shared memory (.nv.reservedSmem.begin/cap/offset0). The master ELF emitter sub_1C9F280 (97KB) constructs headers, section tables, and program headers. Three binary output modes exist:

1. mercury -- traditional SASS binary format
2. capmerc -- Capsule Mercury (default on sm_100+), embeds PTX source in .nv.merc.* sections
3. sass -- direct SASS output

See ELF/Cubin Output.

Stage 6: DebugInfo (DebugInfo-time)

DWARF debug information generation: .debug_info, .debug_line, .debug_frame sections. The LEB128 encoder at sub_45A870 handles all variable-length integer encoding. Source location tracking uses the location map (hash map at sub_426150/sub_426D60) with file offset caching every 10 lines for fast random access. Labels follow the pattern .L__$locationLabel$__%d. See Debug Information.

Error Paths and Recovery

ptxas uses setjmp/longjmp as its sole error recovery mechanism -- there are no C++ exceptions (the binary is compiled as C). Three nested recovery points exist, each catching progressively more localized failures.

Recovery Point Hierarchy

sub_446240 (top-level driver)
  setjmp(jmp_buf_global)         // Level 1: catches any fatal anywhere
    |
    sub_43A400 (per-kernel worker)
      setjmp(jmp_buf_kernel)     // Level 2: catches per-kernel fatals
        |
        sub_432500 (finalization bridge)
          setjmp(jmp_buf_local)  // Level 3: catches OCG pipeline fatals
            |
            [159-phase pipeline, regalloc, encoding, ELF]

Level 1 (global). Established by sub_446240 at function entry. If any code path anywhere in ptxas calls sub_42FBA0 with severity >= 6 (fatal), execution longjmps back here. The handler cleans up global resources and returns a non-zero exit code. This is the last-resort handler.

Level 2 (per-kernel). Established by sub_43A400 before the OCG pipeline runs. On longjmp, the handler destroys the partially-compiled kernel's state, clears the error flags in the TLS context, and continues to the next kernel. This allows multi-kernel compilations to survive a single kernel's failure.

Level 3 (finalization). Established by sub_432500, which saves and replaces the TLS jmp_buf pointer for nested recovery. On longjmp: restores the previous jmp_buf, sets error_flags = 1, releases output buffers, and calls report_internal_error(). Execution returns false to the Level 2 handler.

Parse Error Recovery

Parse errors in sub_451730 (the Flex/Bison parser) invoke sub_42FBA0 with the message "syntax error":

• Severity 4--5 (non-fatal error): The error is printed with file:line location, and the parser attempts to continue via Bison's error recovery rules. Multiple non-fatal parse errors can accumulate. After parsing completes, if the error count > 0, the compilation is aborted before entering the OCG pipeline.
• Severity 6 (fatal): Triggers longjmp to the Level 1 handler immediately. The parser state pool is leaked (accepted trade-off since the process is about to exit).

Bison error recovery operates through the error token in the grammar. When the parser encounters a token that matches no production, it discards tokens until it finds one that allows the error production to reduce, then resumes parsing. This provides reasonable error recovery for common mistakes (missing semicolons, misspelled opcodes) but can cascade badly for structural errors (mismatched braces, corrupt PTX).

Phase Failure in PhaseManager

The phase executor sub_C64F70 runs each phase by calling its vtable execute() method. There is no explicit per-phase error check -- phases that detect internal errors call the diagnostic emitter sub_42FBA0 directly. The error handling cascade:

1. Non-fatal phase error (severity 3--5): The error is printed and the error flag is set in the TLS context. The PhaseManager continues executing subsequent phases. This allows multiple diagnostics to be collected in a single run.
2. Fatal phase error (severity 6): Triggers longjmp to Level 2 or Level 3. The current kernel's compilation is aborted. The PhaseManager's loop is unwound non-locally -- no cleanup of intermediate phase state occurs. Resources are reclaimed when the per-kernel memory pool is destroyed.
3. OOM during phase execution: Any allocation failure calls sub_42BDB0 (3,825 callers), which forwards to sub_42F590 with a severity-6 descriptor at unk_29FA530. This always triggers longjmp.

The PhaseManager logs phase transitions using "Before <phase>" and "After <phase>" string construction (visible in sub_C64F70). When DUMPIR is set to a phase name, the IR is dumped to a file after that phase completes. This enables bisection of phase failures: --knob DUMPIR=<phase_name> isolates which phase corrupted the IR.

Register Allocation Failure and Retry

The register allocator has its own retry mechanism that operates within the normal pipeline (not via longjmp). The retry driver sub_971A90 (355 lines) wraps the Fatpoint allocator in a two-phase strategy:

Phase 1 -- NOSPILL. Attempt allocation without spilling. If the allocator fits within the register budget, proceed directly to finalization.

Phase 2 -- SPILL retry loop. If NOSPILL fails:

1. The spill guidance engine sub_96D940 (84 KB) computes per-register-class spill candidates
2. The allocator retries with progressively more aggressive spilling, up to N attempts (controlled by knobs 638/639)
3. Each attempt prints: "-CLASS NOSPILL REGALLOC: attemp %d, used %d, target %d" (note: the typo "attemp" is in the binary)
4. The best result across all attempts is tracked by sub_93D070
5. The finalization function sub_9714E0 (290 lines) commits the best result or emits a fatal error

On allocation failure (all retry attempts exhausted):

Register allocation failed with register count of '%d'.
Compile the program with a higher register target

This error is emitted by sub_9714E0 through two paths: with source location (via sub_895530, including function name and PTX line number) or without source location (via sub_7EEFA0, generic). After this error, sub_9714E0 returns with HIBYTE(status) set, causing the retry driver to clear all register assignments to -1 and propagate the failure.

A dedicated DUMPIR hook exists: "Please use -knob DUMPIR=AllocateRegisters for debugging" -- this string (found at sub_9714E0's error path) directs users to dump the IR state before the allocator runs.

Fatal Error Handler Chain

The complete chain from any error site to process termination:

[any function, 2,350 call sites]
  sub_42FBA0(descriptor, location, ...)   // central diagnostic emitter
    |  checks descriptor[0] for severity
    |  severity 0: silently ignored
    |  severity 1-2: prints "info    " message
    |  severity 3: prints "warning " (or "error   " if TLS[50] Werror flag set)
    |  severity 4-5: prints "error   " / "error*  " (non-fatal)
    |  severity 6: prints "fatal   " then:
    v
  longjmp(tls->jmp_buf, 1)
    |  unwinds to nearest setjmp recovery point
    v
  [Level 3] sub_432500 -> restore jmp_buf, set error_flags, return false
  [Level 2] sub_43A400 -> cleanup kernel state, continue to next kernel
  [Level 1] sub_446240 -> cleanup global state, exit(non-zero)

Resource leak note. Because longjmp bypasses normal stack unwinding, all heap allocations made between the setjmp and the fatal error are leaked unless tracked in a pool. This is why ptxas uses pool allocators -- the per-kernel pool can be destroyed wholesale at the Level 2 recovery point, reclaiming all leaked memory without tracking individual allocations.

Architecture Dispatch

An architecture vtable factory at sub_1CCEEE0 (17KB, 244 callees) constructs a 632-byte vtable object (79 function pointers) based on the target SM version. The version dispatch ranges:

The distinction between "validation only" and "active" is critical: the base validation table at unk_1D16220 contains 32 entries including all legacy SMs back to sm_20, allowing ptxas to parse PTX files that declare .target sm_30 without immediately rejecting them. However, the capability dispatch initializer sub_607DB0 only registers handler functions for sm_75 through sm_121 (13 base targets). Attempting to compile code for an unregistered SM produces a fatal error during codegen factory lookup -- the architecture vtable factory sub_1CCEEE0 cannot construct a backend object for these targets.

The legacy codegen factory values (12288 for sm_30, 16385/20481 for sm_50, 24576 for sm_60) survive as comparison constants in feature-gating checks throughout the backend (e.g., if (factory_value > 28673) gates sm_90+ features), but the code paths they would activate no longer exist.

Each vtable entry is a function pointer to an SM-specific implementation of a codegen or emission primitive. This is the central dispatch mechanism for all architecture-dependent behavior in the backend.

Obfuscation: ROT13 Encoding

All internal identifiers in ptxas's static initializers are ROT13-encoded:

• Opcode table (ctor_003 at 0x4095D0, 17KB): 900+ PTX opcode mnemonics. Example: NPDOHYX decodes to ACQBULK, SZN decodes to FMA, RKVG decodes to EXIT.
• General knob table (ctor_005 at 0x40D860, 80KB): 2,000+ Mercury/OCG tuning knob names with hex default values. Example: ZrephelHfrNpgvirGuernqPbyyrpgvirVafgf decodes to MercuryUseActiveThreadCollectiveInsts.
• Scheduler knob table (ctor_007 at 0x421290, 8KB): 98 scheduler-specific knob names. Example: XBlockWaitOut, ScavInlineExpansion.

The ROT13 decoding is performed inline during lookup (in sub_79B240, GetKnobIndex) using character-range detection: bytes in A-M get +13, bytes in N-Z get -13, with case-insensitive comparison via tolower().

Cross-References

• Binary Layout -- address ranges for every subsystem
• Function Map -- master index of recovered function addresses
• CLI Options -- complete flag catalog
• Knobs System -- 1,294 internal tuning parameters
• Optimization Levels -- what changes at -O0/-O1/-O2/-O3
• Phase Manager -- PhaseManager object layout and dispatch
• Memory Pool Allocator -- pool struct layout and allocation algorithm
• Thread Pool & Concurrency -- thread pool struct, task submit, jobserver

Function Map

Entry Point & CLI

All addresses in this page apply to ptxas v13.0.88 (CUDA 13.0). Other versions will differ.

The ptxas binary has a deceptively simple entry point. The exported main at 0x409460 is an 84-byte wrapper that sets up unbuffered I/O and immediately tail-calls sub_446240 -- the real compilation driver. This driver is a monolithic 11 KB function that allocates a 1,352-byte master options block on the stack, establishes setjmp-based error recovery, parses all command-line options through a generic framework, reads PTX input, and then loops over compile units running the full Parse -> CompileUnitSetup -> DAGgen -> OCG -> ELF -> DebugInfo pipeline for each. The entire error-handling strategy is non-local: any of the 2,350 call sites to the central diagnostic emitter sub_42FBA0 can trigger a longjmp back to the driver's recovery point on fatal errors.

The same binary doubles as an in-process library. When nvcc loads ptxas as a shared object rather than spawning a subprocess, three extra arguments to the driver carry an output buffer pointer, an extra option count, and an extra options array. Callback function pointers at fixed offsets in the options block allow the host process to receive diagnostics and progress notifications without going through stderr.

Architecture

main (0x409460, 84B)
  │
  ├─ nullsub_1(*argv)          // store program name (no-op)
  ├─ setvbuf(stdout, _IONBF)
  ├─ setvbuf(stderr, _IONBF)
  │
  └─ sub_446240(argc, argv, 0, 0, 0)   // REAL MAIN
       │
       ├─ setjmp(jmp_buf)               // fatal error recovery point
       │
       ├─ sub_434320(opts_block, ...)    // OPTION PARSER
       │    └─ sub_432A00(...)           // register ~100 options via sub_1C97210
       │
       ├─ sub_4428E0(...)                // PTX INPUT SETUP
       │    ├─ validate .version / .target
       │    ├─ handle --input-as-string
       │    └─ generate __cuda_dummy_entry__ if --compile-only
       │
       ├─ sub_43A400(...)                // TARGET CONFIGURATION
       │    └─ set cache defaults, texmode, arch-specific flags
       │
       ├─ FOR EACH compile unit:
       │    ├─ sub_451730(...)           // parser/lexer init + special regs
       │    ├─ sub_43B660(...)           // register constraint calculator
       │    ├─ sub_43F400(...)           // function/ABI setup
       │    └─ sub_43CC70(...)           // per-entry: DAGgen → OCG → ELF → DebugInfo
       │
       ├─ timing / memory stats output (--compiler-stats)
       │
       └─ cleanup + return exit code

Pre-main Static Constructors

Before main executes, four static constructors run as part of the ELF .init_array. Three of them populate ROT13-obfuscated lookup tables that are foundational to the rest of the binary. This obfuscation is deliberate -- it prevents trivial string searching for internal opcode names and tuning knob identifiers in the stripped binary.

ctor_001 -- Thread Infrastructure (0x4094C0, 204 bytes)

Initializes the POSIX threading foundation used throughout ptxas:

pthread_key_create(&key, destr_function);       // TLS key for sub_4280C0
pthread_mutexattr_settype(&attr, PTHREAD_MUTEX_RECURSIVE);
pthread_mutex_init(&mutex, &attr);
dword_29FE0F4 = sched_get_priority_max(SCHED_RR);
dword_29FE0F0 = sched_get_priority_min(SCHED_RR);
__cxa_atexit(cleanup_func, ...);                // registered destruction

The TLS key created here is the one used by sub_4280C0 (3,928 callers), making it the single most important piece of global state in the binary.

ctor_003 -- PTX Opcode Name Table (0x4095D0, 17,007 bytes)

Populates a table at 0x29FE300+ with approximately 900 ROT13-encoded PTX opcode mnemonic strings. Each entry is a (string_ptr, length) pair. The ROT13 encoding maps A-Z to N-Z,A-M and a-z to n-z,a-m, leaving digits and punctuation unchanged.

These decoded names are the canonical PTX opcode mnemonics used during parsing and validation. The table is consumed by the PTX lexer initialization at sub_451730 and the opcode-to-handler dispatch table at sub_46E000 (93 KB, the largest function in the front-end range).

ctor_005 -- Mercury Tuning Knob Registry (0x40D860, 80,397 bytes)

The single largest function in the front-end address range. Registers 2,000+ ROT13-encoded internal tuning knob names, each paired with a hexadecimal default value string. These are the "Mercury" (OCG) backend tuning parameters that control every aspect of code generation, scheduling, and register allocation.

The knob system is documented in detail in the Knobs System page. The ROT13 encoding applies identically to all knob name strings in all four constructors.

ctor_007 -- Scheduler Knob Registry (0x421290, 7,921 bytes)

A smaller companion to ctor_005 that registers 98 scheduler-specific knobs. These control the instruction scheduler (Mercury/OCG) behavior at a finer granularity than the general knobs:

Decoded examples: XBlockWaitOut, XBlockWaitInOut, XBlockWaitInOnTarget, WarDeploySyncsFlush_SW4397903, WaitToForceCTASwitch, VoltageWar_SW4981360PredicateOffDummies, TrackMultiReadsWarLatency, ScavInlineExpansion, ScavDisableSpilling.

Knob names containing _SW followed by a number (e.g., _SW4397903) indicate workarounds for specific hardware bugs identified by NVIDIA's internal bug tracking system.

Real Main -- sub_446240

The exported main() tail-calls sub_446240 with three zero arguments appended. This function is the complete compilation orchestrator: it owns the options block, the error recovery, the compilation loop, and the statistics output.

Signature and Library Mode

int sub_446240(int argc, char **argv,
               void *output_buf,        // a3: cubin output buffer (NULL for standalone)
               int   extra_opt_count,   // a4: count of extra options from nvcc
               char **extra_opts);      // a5: array of extra option strings

When main calls this, a3/a4/a5 are all zero -- standalone mode. When nvcc loads ptxas as a shared library and calls the entry point directly, these arguments carry non-null values:

• a3 (output_buf): Pointer to a memory buffer where the compiled cubin is written. Eliminates the need for temporary files and filesystem round-trips, which matters for large CUDA compilations where nvcc may invoke ptxas hundreds of times.
• a4 (extra_opt_count): Number of additional option strings injected by nvcc beyond what appears on the command line.
• a5 (extra_opts): Array of those extra option strings.

Additionally, callback function pointers at offsets 37--39 of the 1,352-byte options block (byte offsets ~296, ~304, ~312) allow the host process to receive progress notifications and diagnostic messages in-process rather than through stderr.

Error Recovery with setjmp/longjmp

The first significant action in sub_446240 is establishing a setjmp recovery point:

if (setjmp(jmp_buf) != 0) {
    // Fatal error occurred somewhere in the pipeline.
    // Clean up and return non-zero exit code.
    goto cleanup;
}

This is the only error recovery mechanism in ptxas -- there are no C++ exceptions (the binary is compiled as C, not C++). Any function anywhere in the call tree that encounters an unrecoverable error calls sub_42FBA0 with severity >= 6, which internally calls longjmp(jmp_buf, 1) to unwind directly back to this point. The approach is simple but has a critical implication: all resources allocated between the setjmp and the fatal error are leaked unless explicitly tracked and cleaned up at the recovery site.

The 1,352-Byte Options Block

The master options block lives on the stack and accumulates all compilation state during option parsing. It is passed by pointer to virtually every subsystem. Key fields (approximate offsets based on access patterns):

Compilation Loop

After option parsing and PTX input setup, the driver enters a loop over compile units. Each unit corresponds to one entry function (or device function in --compile-only mode). The per-entry processing is handled by sub_43CC70, which prints a separator:

printf("\n# ============== entry %s ==============\n", entry_name);

and then sequences: DAGgen (PTX-to-Ori lowering), OCG (optimization and code generation), ELF (binary emission), and DebugInfo (DWARF generation). The special entry __cuda_dummy_entry__ is silently skipped.

Timing and Memory Statistics

When --compiler-stats is active, sub_446240 prints per-phase timing and peak memory after all compile units complete:

CompileTime = 42.3 ms (100%)
Parse-time            : 12.1 ms (28.61%)
CompileUnitSetup-time :  1.4 ms ( 3.31%)
DAGgen-time           :  8.7 ms (20.57%)
OCG-time              : 15.2 ms (35.93%)
ELF-time              :  3.8 ms ( 8.98%)
DebugInfo-time        :  1.1 ms ( 2.60%)
PeakMemoryUsage = 2048.000 KB

When --compiler-stats-file is specified, the same data is written in JSON format using the shared JSON builder (sub_1CBA950). When --fdevice-time-trace is active, sub_439880 parses Chrome DevTools trace format JSON and merges ptxas timing events into the trace.

Option Parser -- sub_434320 and sub_432A00

Option parsing is split into two phases: registration and processing.

Option Registration -- sub_432A00

This 6,427-byte function calls sub_1C97210 approximately 100 times, once per recognized option. Each call provides the option's long name, short name, value type, help text, and default value to the generic option framework (implemented in the 0x1C96xxx--0x1C97xxx range, shared with other NVIDIA tools).

Option Processing -- sub_434320

The 10,289-byte parser iterates over argv (and any extra options from library mode), matches each argument against registered options via the framework, and populates fields in the 1,352-byte options block. Special handling exists for:

• --version: Prints the identification string "Ptx optimizing assembler" followed by the version (e.g., "Cuda compilation tools, release 13.0, V13.0.88") and exits.
• --help: Delegates to sub_403588, which prints "Usage : %s [options] <ptx file>,...\n" followed by all registered options, then exits.
• --fast-compile: Validated against conflicting optimization options.
• -cloning=yes/-cloning=no: Inline cloning control parsed as an equality option.

Generic Option Framework

The option parsing library lives in the 0x1C96000--0x1C97FFF range and is shared with other NVIDIA tools (nvlink, fatbinary, etc.):

Diagnostic System -- sub_42FBA0

The central diagnostic emitter is the most important error-reporting function in ptxas. With 2,350 call sites, it handles every warning, error, and fatal message in the entire binary.

Signature

void sub_42FBA0(
    int *descriptor,    // a1: points to severity level at *a1
    void *location,     // a2: source location context
    ...                 // variadic: printf-style format args
);

Severity Levels

The machine-readable tags (@E@, @W@, @O@, @I@) allow nvcc and other tools to parse ptxas output programmatically, extracting severity without parsing the human-readable text.

Warning-to-Error Promotion

Severity level 3 has context-dependent behavior controlled by two flags in the thread-local storage:

v5 = *a1;   // severity
if (v5 == 3) {
    if (sub_4280C0()[49])   // TLS offset 49: suppression flag
        return;             // silently discard
    if (sub_4280C0()[50])   // TLS offset 50: Werror flag
        prefix = "error   ";
    else
        prefix = "warning ";
}

This implements the --Werror equivalent: when the Werror flag is active in the TLS context, all warnings become errors.

Output Format

<filename>, line <N>; <severity>: <message>

When source is available, the diagnostic emitter reads the PTX input file, seeks to line N, and prints the source line prefixed with "# ". To avoid O(n) seeking through large files on every diagnostic, it maintains a hash map (sub_426150/sub_426D60) that caches file byte offsets every 10 lines for fast random access to arbitrary line numbers.

Fatal Error Handler -- sub_42BDB0

A 14-byte wrapper called from 3,825 sites (nearly every allocation in ptxas). It fires whenever the pool allocator sub_424070 returns NULL:

void sub_42BDB0(...) {
    return sub_42F590(&unk_29FA530, ...);   // internal error descriptor
}

The descriptor at unk_29FA530 has severity 6 (fatal), so this always triggers longjmp back to the driver's recovery point.

Thread-Local Storage -- sub_4280C0

The most-called function in the entire binary (3,928 callers). Returns a pointer to a 280-byte per-thread context struct, allocating and initializing it on first access.

void *sub_4280C0(void) {
    void *ctx = pthread_getspecific(key);
    if (ctx) return ctx;

    ctx = malloc(0x118);        // 280 bytes
    memset(ctx, 0, 0x118);
    pthread_cond_init(ctx + 128, NULL);
    pthread_mutex_init(ctx + 176, NULL);
    sem_init(ctx + 216, 0, 0);
    pthread_setspecific(key, ctx);
    return ctx;
}

TLS Context Layout (280 bytes)

The TLS key is created by ctor_001 before main runs, and a destructor function registered via pthread_key_create frees the 280-byte struct when a thread terminates. This per-thread context enables concurrent compilation of multiple compile units (when the thread pool is active), with each thread maintaining independent error state and diagnostic suppression flags.

PTX Input Setup -- sub_4428E0

After options are parsed, this 13,774-byte function reads and preprocesses the PTX input:

1. Version and target validation. Checks .version and .target directives in the input. Emits synthetic headers ("\t.version %s\n", "\t.target %s\n") when needed.

2. Compile-only mode. When --compile-only is active and no real entries exist, generates a dummy entry: "\t.entry %s { ret; }\n" with name __cuda_dummy_entry__.

3. Input-as-string mode. When --input-as-string is active, PTX is read from memory (passed as a string argument) rather than from a file. This is used by the library-mode interface.

4. Whole-program mode. --extensible-whole-program enables inter-function optimization across all entries in the compilation unit.

5. Cache and debug configuration. Applies --def-load-cache, --def-store-cache, --force-load-cache, --force-store-cache, and suppress-debug-info settings.

6. Tools-patch mode. --compile-as-tools-patch activates the CUDA sanitizer compilation path, checking for __cuda_sanitizer symbols.

Key diagnostic strings from this function:

• "'--fast-compile'"
• "calls without ABI"
• "compilation without ABI"
• "device-debug or lineinfo"
• "unified Functions"

Target Configuration -- sub_43A400

A 4,696-byte function that configures target-specific defaults after option parsing completes. It reads the SM architecture number from the options block and sets:

• Texturing mode: texmode_unified vs raw texture mode.
• Cache defaults: Based on architecture capabilities.
• Feature flags: Hardware-specific workaround flags (e.g., --sw4575628).
• Indirect function support: "Indirect Functions or Extern Functions" validation.

The function references "NVIDIA" and "ptxocg.0.0" (the internal name for the OCG optimization pass), suggesting it also initializes the pass pipeline configuration for the target architecture.

Register Constraint Calculator -- sub_43B660

A 3,843-byte function that resolves potentially conflicting register limit specifications into a single register budget per function. Register constraints come from four sources with different priorities:

The occupancy-derived limit is computed from .minnctapersm and .maxntid: given a minimum number of CTAs per SM and a maximum thread count per CTA, the function calculates the maximum register count that allows the requested occupancy level, accounting for per-SM register file size.

Diagnostic strings indicate the resolution process:

• "computed using thread count" -- derived from .maxntid
• "of .maxnreg" -- explicit per-function limit
• "of maxrregcount option" -- CLI override
• "global register limit specified" -- global cap applied

Per-Entry Compilation -- sub_43CC70

A 5,425-byte function that processes each entry function through the complete backend pipeline. For each entry:

1. Skips __cuda_dummy_entry__ (generated by compile-only mode).
2. Prints the entry separator: "\n# ============== entry %s ==============\n".
3. Runs DAGgen (PTX-to-Ori lowering).
4. Runs OCG (the 159-phase optimization pipeline + SASS code generation).
5. Generates .sass and .ucode ELF sections.
6. Generates DWARF debug information if requested.

The function also handles reg-fatpoint configuration (the register allocation algorithm, documented in the Fatpoint Algorithm page).

Function/ABI Setup -- sub_43F400

A 9,078-byte function that configures the calling convention for each function before compilation. This includes:

The function handles both entry functions (kernels launched from the host) and device functions (callable from other device code), with different ABI requirements for each. Entry functions use a simplified ABI where parameters come from constant memory, while device functions use register-based parameter passing.

The --compile-as-tools-patch and --sw200428197 flags activate a special ABI variant for CUDA sanitizer instrumentation, which inserts additional scratch registers for sanitizer state.

Function Map