PTX Parsing

The embedded ptxas compiler in nvlink v13.0.88 contains a complete PTX assembler frontend: a flex-generated lexer, a bison/yacc-generated parser, a module initialization system that defines every PTX special register and builtin type, an instruction handler registry covering the full PTX ISA 9.0, a CUDA builtin function database of 608 entries, and a recursive expression printer. This page documents each subsystem at the function level, reconstructed from decompiled binary analysis across the 0x12AF000--0x12C0000, 0x1430000--0x15C0000, and 0x16E0000--0x16F6000 address ranges.

Architecture Overview

PTX parsing is organized into seven subsystems that form a pipeline from raw text to internal IR:

PTX source text
    |
    v
[1] Flex Lexer (sub_16EAF60, 65KB)
    |  DFA state table at off_22788C0
    |  550+ token actions
    |  Macro preprocessor (sub_16E8310)
    v
[2] Bison/Yacc Parser (sub_16F0E60, 31KB)
    |  1172 lines of parser actions
    |  Operand list builders (sub_16E4D60..sub_16E7D00)
    v
[3] Module Initializer (sub_12AF950, 60KB)
    |  Special registers, builtin types, hash tables
    |  PTX version 9.0, ~20 symbol caches
    v
[4] Instruction Handler Registry (sub_158D130, 41KB)
    |  115 named mnemonic handlers + 473 hash-encoded handlers (588 total)
    |  Two-level dispatch: name table *(ctx+808), hash table *(ctx+816)
    v
[5] Semantic Validators (sub_147EF50, 288KB + sub_146BEC0, 206KB + ...)
    |  SM version gates, PTX version gates, operand type checks
    |  Relaxed mode bypass (sub_12B3090)
    v
[6] Code Template Generators (~250 functions, 3--50KB each)
    |  50KB temp buffer, sprintf-based PTX emission
    v
[7] Expression Printer (sub_12B33A0, 9KB)
       17 expression kinds, recursive descent

Module Initialization

sub_12AF950 -- ptx_init_compilation_state (59,874 bytes, 1,440 lines)

The master initialization function. Creates the entire PTX compilation state in a single call, including the symbol table infrastructure, all PTX special registers, all builtin aggregate types, and the version/target configuration. Takes 11 parameters including a function pointer callback.

Allocation. Creates two major structures: a 1,128-byte compilation context and a 2,528-byte instruction state block. Both are arena-allocated via sub_4307C0.

PTX version. Sets PTX_MAJOR_VERSION = 9, PTX_MINOR_VERSION = 0 via sub_448E70 hash table inserts. This corresponds to CUDA Toolkit 13.0 / PTX ISA 9.0.

Special register initialization. Registers every PTX special register in declaration order, using sub_448E70 to insert each name into the symbol table:

Builtin type initialization. Creates three aggregate type descriptors via sub_12AD9E0:

.texref fields: width, height, depth, channel_data_type, channel_order, normalized_coords, filter_mode, addr_mode_0, addr_mode_1, addr_mode_2, array_size, num_mipmap_levels, num_samples (13 fields).

.samplerref fields: force_unnormalized_coords, filter_mode, addr_mode_0, addr_mode_1, addr_mode_2 (5 fields).

.surfref fields: width, height, depth, channel_data_type, channel_order, array_size, memory_layout (7 fields).

Hash table allocation. Initializes approximately 20 hash tables (via sub_4489C0) for symbol caches, type caches, instruction lookup, and builtin registration. Each table uses the standard pattern: sub_4489C0(hash_fn, eq_fn, bucket_count).

Target configuration. Calls sub_1426EE0 to initialize target-specific configuration and sub_16E3AA0 to set SM version info (field 30 = SM major, field 31 = SM minor).

sub_12AF200 -- ptx_init_builtin_macros (4,556 bytes, 144 lines)

Initializes Fermi-era backward-compatibility macro definitions. Referenced as "<fermi macros>" in calling context.

sub_12AF550 -- ptx_init_target_features (4,625 bytes, 144 lines)

Initializes the target feature set based on the SM version. Companion to the main initializer.

Lexer (Flex-Generated)

sub_16EAF60 -- ptx_flex_scanner_main_loop (64,758 bytes, 2,530 lines)

The main scanner loop, generated by flex. This is the lexical analysis engine that tokenizes PTX source text.

DFA structure. The scanner uses a DFA state table at off_22788C0, indexed by (current_state, input_character). The action dispatch table at dword_2278020 maps accepted token types to action categories. The action switch has 550+ cases covering the entire PTX token vocabulary: keywords, directives, register names, literals, operators, and punctuation.

EOF handling. Case 550 handles end-of-file. The scanner tracks line numbers by counting newline characters (ASCII 10).

Buffer management. Three helper functions manage scanner input:

sub_16E8310 -- ptx_macro_preprocessor (33,290 bytes, 1,068 lines)

The PTX macro preprocessor. Runs before lexical analysis to expand .MACRO, .ELSE, .ELIF directives.

Input handling. Processes input character by character, checking for whitespace (ASCII 9 = tab, 32 = space). Manages macro nesting with a depth limit -- exceeding it triggers "macro nesting too deep!". Uses sub_44FB20 for character classification.

State tracking. Key offsets in the preprocessor context:

• +48: current line number
• +52: current column number
• +2144: input buffer pointer
• +2441: lookahead character
• +8: stdin fallback (when no input buffer is active)

sub_16EF680 -- ptx_macro_expansion (4,010 bytes, 142 lines)

Handles individual macro expansion with nesting depth checking. Emits "macro nesting too deep!" on overflow.

Parser (Bison/Yacc-Generated)

sub_16F0E60 -- ptx_parser_yacc_main (30,684 bytes, 1,172 lines)

The main parser driver, generated by bison/yacc. Works in conjunction with the flex scanner (sub_16EAF60). At 1,172 lines, this encodes the full PTX grammar reduction rules and semantic actions.

sub_16E9690 -- ptx_directive_parser (14,777 bytes, 525 lines)

Parses PTX directives including .surfref, .samplerref, .texref, and section directives. The 525-line function implements a complex state machine that handles multiple directive types and their parameter lists.

sub_1442040 -- section_directive_validator (3,300 bytes, 82 lines)

Validates .section and @@DWARF directives. The @@DWARF handler processes inline DWARF debug data embedded in PTX source, which is the mechanism by which nvcc passes debug information through PTX to the assembler.

Instruction Construction

The parser builds instruction IR nodes through a family of builder functions in the 0x16E4D60--0x16E7D00 range. All work with dynamic arrays of 48-byte records (one per operand), using a growth factor of 1.5x with a minimum of 16 elements. The block allocation unit is 768 bytes (16 elements times 48 bytes). Memory allocation uses vtable callbacks at **obj+24.

Instruction Handler Registration

sub_158D130 -- ptx_instruction_handler_table_init (40,900 bytes, 595 lines)

The master instruction handler registration function. First calls sub_158A600 to initialize the CUDA builtin name table, then populates two dispatch hash tables with a combined 588 unique handler functions spanning 988.7 KB of code.

Named Mnemonic Table (115 entries)

Stored at *(context+808). Maps PTX instruction mnemonic strings to handler function pointers via sub_448E70. Every entry below is extracted verbatim from the decompiled registration sequence -- no entries are omitted.

Arithmetic and Math (17 entries)

Bit Manipulation (6 entries)

Warp-Level Operations (5 entries)

Barrier and Synchronization (12 entries)

Texture and Surface (6 entries)

Matrix (Tensor Core) Operations (16 entries)

tcgen05 (5th-Gen Tensor Core) Management (9 entries)

tcgen05 Guardrail Intrinsics (8 entries)

These are internal pseudo-instructions (prefixed with _) used for runtime validation of tensor core operations. They emit guard traps expanded by sub_15B86A0.

Async Copy and Bulk Operations (6 entries)

Multi-Memory (CTA-level Multicast) (3 entries)

Cache Policy (3 entries)

SIMD Video Instructions (23 entries)

Internal / Special (1 entry)

Named handler statistics. 115 total entries, 115 unique function addresses. Total code size: 426,096 bytes (416.1 KB). Handler sizes range from 304 bytes (wgmma.fence) to 39,952 bytes (wmma.mma). The five largest handlers account for 99,632 bytes (23.4% of all named handler code): wmma.mma (39.9 KB), mma (27.1 KB), tcgen05.mma (21.6 KB), div (17.8 KB), sqrt (14.0 KB).

Hash-Encoded Opcode Table (473 entries)

Stored at *(context+816). Maps 32-bit hash keys (stored as decimal string keys like "2644314910") to handler function pointers. These 473 entries cover instruction variants that are dispatched by opcode encoding hash rather than mnemonic string. Each handler follows the same code template pattern as the named handlers: allocate 50 KB temp buffer, build PTX text via sprintf from string table offsets, shrink-allocate, return.

Structure. The hash keys are 32-bit unsigned integers representing pre-computed hashes of instruction encoding descriptors. The hash function is not applied at lookup time to a mnemonic string; instead, the parser pre-computes the hash from the instruction's opcode + type + modifier combination and looks it up directly. This covers complex variant combinations (e.g., add.s32, add.f32, mad.lo.s32, setp.eq.f64) where encoding a separate named mnemonic for each variant would be impractical.

Statistics. 473 entries, 473 unique function addresses, zero overlap with the named mnemonic table. Total code size: 586,352 bytes (572.6 KB). Handler sizes range from 736 to 8,080 bytes, with median 1,136 bytes and mean 1,239 bytes. Size distribution:

Address distribution. Hash handlers cluster in the 0x1490000--0x1520000 range (92% of entries), while named handlers dominate 0x1520000--0x1580000. This reflects a compilation unit boundary: the hash handlers are generated from table-driven templates, while the named handlers are handwritten for semantically complex instructions.

Complete hash-encoded dispatch table. All 473 entries in registration order:

Two-Level Dispatch Protocol

When the parser encounters an instruction:

1. Look up the mnemonic in the named table at *(context+808) via sub_449A80.
2. If not found, compute the instruction encoding hash and look up in *(context+816).
3. The resolved handler receives (context, instruction_node) and generates PTX code template text.

Combined scale. The two tables register 588 unique handler functions covering 988.7 KB of code (416.1 KB named + 572.6 KB hash-encoded). This is the complete PTX ISA 9.0 instruction set as implemented by nvlink's embedded assembler. The handler address range spans 0x1493320--0x15809F0 (949.7 KB of .text).

Code Template Generators

Approximately 250 functions in the 0x14A0000--0x15C0000 range follow an identical code generation pattern. Each handles one specific instruction variant (a particular MMA shape, a barrier variant, an arithmetic operation with specific type combination).

Common pattern (verified across sub_14A4770, sub_14AAA30, sub_14C3C10):

1. arena = sub_44F410(a1, a2) + offset[3]     // get memory arena
2. buf   = sub_4307C0(arena, 50000)            // allocate 50KB temp buffer
3. // Build PTX text via successive sprintf() calls:
   //   - Base template from read-only string table (a2 + constant_offset)
   //   - Conditional sections based on sub_16DF3B0 (has descriptor),
   //     sub_16DF3D0 (has scale), sub_16E4530 (matrix dimension)
   //   - Register names, immediates, matrix shapes substituted in
4. len   = strlen(buf)
5. out   = sub_4307C0(arena, len + 1)          // shrink-allocate final string
6. strcpy(out, buf)
7. sub_431000(buf)                             // free temp buffer
8. return out

The larger handlers deviate from this pattern by size but not structure. The six largest handlers each exceed 100KB:

Instruction Property Query Functions

Code template generators query instruction properties through a family of accessor functions in the 0x16DF000--0x16E4000 range:

CUDA Builtin Registration

sub_158A600 -- cuda_builtin_name_table_init (45,700 bytes, 630 lines)

Registers 608 CUDA runtime builtin function names into a hash table for the instruction lowering pipeline. These builtins are PTX pseudo-functions that the assembler expands into instruction sequences during compilation.

Data structure. Creates a hash table via sub_4489C0 with capacity 0x80. Allocates a 9,728-byte builtin data array from unk_1F8E0C0. Stores: *(context+1056) = data array, *(context+1064) = hash table, *(context+1072) = 608 (entry count).

Each builtin is registered with a sequential index (1--608) mapping the name to an index used by sub_15B86A0 (the 345KB prototype generator) to produce the corresponding .weak .func (...) __cuda_smXX_foo (...) ; declaration.

Builtin families (608 total, organized by SM generation and functional category):

sub_15B86A0 -- cuda_builtin_prototype_generator (345,200 bytes, 9,666 lines)

A giant switch on builtin index (608 cases). For each case, allocates a string via sub_14932E0(length, arena) and writes the PTX prototype. Returns a string of the form:

.weak .func (.param .b32 retval) __cuda_sm20_div_s16 (.param .b32 a, .param .b32 b) ;

Covers all SM generations (sm20 through sm10x) and all function families: div, rem, rcp, sqrt, dsqrt, drsqrt, barrier, wmma, shfl, vote, matchsync, warpsync, reduxsync, sanitizer_memcheck, scalar_video_emulation, tcgen05, bulk_copy, cp_async_bulk_tensor.

Semantic Validation

Instruction validation follows a consistent protocol across all validators:

1. sub_12B3090(context)       // Check "relaxed mode" -- if true, skip all checks
2. sub_12B3CA0(major, minor)  // Check PTX version requirement
3. sub_1441FB0(context, ...)  // Emit error if PTX version too low
4. sub_12A8530(sm_version)    // Check SM target capability
5. sub_14422F0(context, ...)  // Emit error if SM version too low
6. (operand type/count checks)
7. sub_467A70(error_code)     // Emit diagnostic with specific error dword

Error codes are global dword_ variables:

sub_147EF50 -- ptx_instruction_semantic_analyzer (287,900 bytes, 7,803 lines)

The master instruction validator. Takes (context, char* name, char mode, error_sink). Validates every aspect of a PTX instruction against the current SM target and PTX version:

• Dimensions, texture modes (texmode_independent, texmode_raw, texmode_unified)
• Cache policies, state spaces, vector types
• Scoping (.cta, .cluster, .sys), ordering (.acquire, .release)
• Proxy fences, comparison operators, eviction priorities
• Directives: .branchtargets, .calltargets, .callprototype, .weak, .common, .noreturn, .abi_preserve, .abi_preserve_control, .unified, .allocno, .scratch, .retaddr_allocno, .ptr

SM version gates referenced: sm_20, sm_30, sm_32, sm_50, sm_60, sm_61, sm_70, sm_72, sm_75, sm_80, sm_89, sm_90.

sub_146BEC0 -- ptx_load_store_validator (206,100 bytes, 5,546 lines)

Validates all memory operations: ld/st, atomics, reductions, fence, membar, cp.async, st.async/red.async. Checks:

• Vector types (.v2, .v4, .v8), data widths (.b32, .b64, .b128)
• Memory ordering (.acquire/.release), scope (.sys)
• Cache eviction priority for L2
• Address space modifiers (.shared::cluster, .shared::cta)
• .sync_restrict, .op_restrict, .cluster scope on atomics
• SM gates: sm_11, sm_12, sm_20, sm_30, sm_50, sm_60, sm_80, sm_90
• PTX version gates: 1.1, 1.2, 2.0, 6.0, 7.8, 8.0, 8.1, 8.4, 8.6

References the string "igonre-src on cp.async" -- a genuine typo ("ignore") in NVIDIA's source code at approximately offset +1100.

sub_1487650 -- ptx_statement_processor (240,300 bytes, 6,428 lines)

Top-level PTX statement handler. Processes kernel/function declarations including:

• Performance directives: .maxnctapersm, .minnctapersm, .maxntid, .reqntid, .reqnctapercluster, .blocksareclusters
• Kernel parameter size validation (4,352 byte limit)
• Function attributes: .noreturn, .extern, .FORCE_INLINE, .unique
• Start/end label creation: "__$startLabel$__%s", "__$endLabel$__%s"
• Parameter registers, scratch registers, return address registers
• PTX version gates: 2.0, 2.1, 2.2, 2.3, 3.0, 3.1, 4.3, 5.0, 6.0, 6.3, 6.4, 7.8, 8.0, 8.1, 8.8, 9.0

Instruction Modifier Groups

sub_14871D0 (3,400 bytes, 54 lines) initializes the modifier group enum. Registers bidirectional group mappings via sub_465720 and sub_448E70:

GUARD  <-> PRED     (bidirectional)
TYPES, POSTOP, COMPARE, APRX, FTZ, SAT, SHAMT, ORDER, NC, ROUND,
TESTP, FLOW, TEXTURE, QUERY, CLAMP, SHR, VMAD, PRMT, SHFL,
ENDIS, UNIFORM, VECTOR, VOTE

These 25 modifier groups classify the possible instruction suffixes. The GUARD <-> PRED bidirectional mapping reflects that guard predicates and predicate operands share the same register file.

Expression Printer

sub_12B33A0 -- ptx_print_expression (8,692 bytes, 291 lines)

A recursive pretty-printer that serializes expression AST nodes to text output. Dispatches on (*node & 0x3F), a 6-bit expression kind discriminant:

The binary operator table at off_1F24200 has 30 entries (indices 0 through 0x1D), covering arithmetic, bitwise, shift, logical, and comparison operators.

Symbol Management

sub_12ADB80 -- ptx_register_function_symbol (5,093 bytes, 171 lines)

Registers a function in the PTX module's symbol table. Allocates an 88-byte function descriptor (type = 4). Checks for __cudart_ prefix to classify CUDA runtime functions, which are stored in a separate linked list at module offset +168 (versus offset +56 for user functions). Parameter lists are linked at offsets +32 and +40.

sub_12ADFE0 -- ptx_register_global_symbol (5,025 bytes, 153 lines)

Registers a global/shared/const variable. Takes 15 parameters covering all variable attributes. Allocates three structures: 88-byte descriptor (type = 5), 248-byte extended info (at offset +80), and 80-byte type descriptor (at extended info offset +136).

sub_12AEA80 -- ptx_resolve_struct_member (7,282 bytes, 278 lines)

Resolves member access on struct types. Walks member name strings character by character, uses ctype_b_loc() for digit detection to extract numeric array index suffixes. Temporarily replaces characters with '<' during lookup. Recursive: calls sub_12AE6F0 for base lookup, then recurses on the result for chained member access.

sub_12B8790 -- ptx_define_global_symbol (6,937 bytes, 264 lines)

Defines a new global symbol with allocation. Takes 10 parameters: (module, name, storage_class, alignment, size, ...). Storage class dispatch:

Magic relocation base: 0x70000064 (= R_CUDA relocation type constant).

Bindless Texture Lowering

sub_12B9660 -- ptx_process_bindless_textures (7,505 bytes, 220 lines)

Processes bindless texture/surface operand pairs from the module. Iterates over three lists: samplers (offset +144), textures (offset +152), surfaces (offset +160). For each (texture, sampler) pair:

1. Query maximum resource limits via vtable calls at offsets +24+56, +24+40, +24+48.
2. Compare actual counts against limits; warn via sub_467460 if exceeded.
3. Create a merged, sorted list of (texture, sampler) combinations.
4. For each pair within the address space limit, create a "$BINDLESS$%s$%s$%s" symbol (module name, texture name, sampler name). Allocate a 112-byte descriptor with type = 4, alignment = 4.
5. Set relocation offset = vtable[304]() - 0x70000064.
6. Register relocation entries via sub_12B8650.

Unpaired textures get "$BINDLESS$%s$%s" (two-part, no sampler) via sub_12B9320. Surfaces are processed in a final pass via sub_465390.

DWARF Debug Data Processing

Several functions handle DWARF debug information embedded in PTX via @@DWARF directives:

SM100 tcgen05 Intrinsic Codegen

A cluster of functions in the 0x16E0A70--0x16E3AB0 range generates PTX inline code for Blackwell (SM100+) tensor core gen05 operations. These are not parser functions per se but are invoked during instruction lowering to produce PTX helper sequences.

sub_16E0A70 -- tcgen05_instruction_type_classifier (17,302 bytes, 322 lines)

Chains approximately 50 type-check predicates (sub_12B5670--sub_12B5950) against an instruction object at *(a1+8). Returns an integer type ID (1--54) corresponding to specific tcgen05 MMA variants. Type IDs 1--8 map to base types, 18--29 to extended variants. Modifiers include _expand16bit, _pack16bit, _maxabs, _minabs, _fused, _blockscale, _ashift.

sub_16E1DB0 -- tcgen05_guardrails_codegen (10,365 bytes, 325 lines)

Generates boundary-checking PTX code for tcgen05 tensor memory operations. Emits "mov.u32 %s, %s;" and "mov.u32 %s, %d;" sequences. References: "__cuda__sm10x_tcgen05_guardrails_in_physical_bounds_", "__cuda__sm10x_tcgen05_guardrails_are_columns_allocated_". Checks SM100/SM10x variants.

sub_16E2410 -- tcgen05_tmem_addr_compute (3,411 bytes, 111 lines)

Emits "add.u32 %s, %s, %s;" for tensor memory address computation. References "__cuda_sm_100_tcgen05_tmem_addr".

sub_16E3AB0 -- tcgen05_instruction_args_type_mapper (18,623 bytes, 337 lines)

Maps arrays of instruction argument pointers to arrays of numeric type IDs, using the same predicate chain as sub_16E0A70. Stores results at a1+944 per argument, with argument count at base struct offset +796.

PTX Type System

sub_12A7B90 -- ptx_init_version_table (17,596 bytes, 430 lines)

Initializes the PTX ISA version-to-feature mapping table. The entire body is SIMD divmod-by-10 computation on constant arrays at xmmword_1F24440--1F244E0. The pattern q = (x * 0xCCCCCCCD) >> 34; r = x - 10*q decomposes version numbers into major/minor pairs. Writes 30+ (quotient, remainder) pairs. String reference "%d.%d\n" confirms version formatting.

sub_12AA190 -- ptx_get_cvt_opcode (6,332 bytes, 171 lines)

Resolves conversion instruction opcodes from source/destination type pairs. Takes (src_type, dst_type, rounding_mode, context). Returns PTX opcode enum values (25--59). Examples:

• src=8, dst=7 -> opcode 53 or 55 (f16 to f32 or similar)
• src=2, dst=3 -> opcode 29 (s32 to s64)
• src=3, dst=2 -> opcode 28 (s64 to s32)

Calls sub_12A72D0 to validate that the conversion is legal. Calls sub_16E08F0/sub_16E60D0 for illegal conversion warnings.

sub_12AA9C0 / sub_12AAD20 -- ptx_canonicalize_opcode (5,110--6,944 bytes)

Opcode renumbering for stable serialization or cross-version compatibility. Maps old opcode IDs to canonical form via a giant if-else chain. Most values map to themselves, but some are reordered: 33->31, 32->30, 49->50, 50->49, 53->58, 54->59, 55->57, 56->54, 57->55, 58->53, 59->56.

Shared Utility Functions

The entire PTX frontend is built on a consistent set of utility primitives:

Context Structure (Partial Reconstruction)

The main compiler context is accessed via *(a1 + 1096) throughout the frontend. Key fields in the instruction structure pointed to by this offset:

Offset  Size  Field
  +88    8B   scope/function pointer
 +104    8B   modifier enum table
 +112    8B   modifier group table
 +220    4B   operand count
 +224    4B   primary operand type enum
 +228    4B   secondary operand type enum
 +244    4B   instruction opcode / shape identifier
 +256   40B   operand type descriptors (8 bytes each, 5 operands max)
 +560    4B   parameter stride
 +596    4B   operand count (second field)
 +601    1B   flags (bit 1: has_scale, bits 2-5: type code)
 +602    1B   flags (bits 7-8: block_scale_kind)
 +603    1B   flags (bits 1-2: scale variant)
 +620    1B   instruction subtype byte
 +627    1B   flags (bits 4-5: layout variant)
 +648+   8B   per-operand tree pointers (8 bytes each)
 +776    4B   instruction class (11/12 = MMA class)
 +808    8B   named instruction handler hash table
 +816    8B   hash-encoded instruction handler hash table
+1056    8B   CUDA builtin data array
+1064    8B   CUDA builtin hash table
+1072    4B   CUDA builtin count (608)
+2488    8B   pragma definition table

SM Version Gate Reference

The validators enforce feature availability across GPU generations. This table summarizes when major PTX features became available, as encoded in the validation logic:

Cross-References

nvlink Internal

• Embedded ptxas Overview -- PTX frontend address map at 0x11EA000--0x15C0000
• IR Nodes -- the IR node structures produced by PTX parsing
• Architecture Dispatch -- SM version dispatch tables consulted during parsing
• LTO Overview -- how LTO pipeline feeds PTX to the embedded parser

Sibling Wikis

• ptxas: PTX Parser -- standalone ptxas PTX parsing infrastructure
• ptxas: PTX-to-Ori -- standalone ptxas PTX-to-Ori IR lowering
• ptxas: PTX Directives -- directive handling in standalone ptxas