PTX Parser (Flex + Bison)

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

The ptxas front-end parses PTX assembly text into internal IR using a classic two-stage architecture: a Flex-generated DFA scanner (lexer) and a Bison-generated LALR(1) shift-reduce parser. Unlike most compiler front-ends, the parser does not construct an AST. Instead, Bison reduction actions directly build IR nodes, populate the instruction table, and emit validation calls -- the parse tree is consumed inline and never materialized as a data structure. A separate macro preprocessor handles .MACRO, .ELSE/.ELIF/.ENDIF, and .INCLUDE directives at the character level before tokens reach the Flex DFA. The instruction table builder (sub_46E000, 93 KB) registers all PTX opcodes with their legal type combinations during parser initialization, and an instruction lookup subsystem classifies operands into 12 categories at parse time.

Architecture

PTX source text
     │
     ▼
┌─────────────────────────────────────────────────────────┐
│  MACRO PREPROCESSOR (character-level, 0x71B000-0x720000)│
│  sub_71F630  dispatch: .MACRO / .ELSE / .INCLUDE        │
│  sub_71E2B0  conditional: .ELSE / .ELIF / .ENDIF (32KB) │
│  sub_71DCA0  macro definition handler                   │
│  sub_71C310  .INCLUDE file handler                      │
└────────────────────┬────────────────────────────────────┘
                     │ preprocessed character stream
                     ▼
┌─────────────────────────────────────────────────────────┐
│  FLEX DFA SCANNER  sub_720F00 (15.8KB, 552 rules)       │
│  off_203C020       DFA transition table                  │
│  Token codes:      258-422 (163 named emitted)           │
│  Helper:           sub_720410 (yy_get_next_buffer)       │
│                    sub_720630 (yy_get_previous_state)     │
│                    sub_720BA0 (yy_scan_string)            │
└────────────────────┬────────────────────────────────────┘
                     │ token stream (code + attribute)
                     ▼
┌─────────────────────────────────────────────────────────┐
│  BISON LALR(1) PARSER  sub_4CE6B0 (48KB, 512 prods)     │
│  5 LALR tables at 0x1D12xxx-0x1D15xxx                    │
│  443 reduction actions → direct IR construction           │
│  NO AST: reductions emit IR nodes inline                  │
└────────────────────┬────────────────────────────────────┘
                     │
          ┌──────────┴──────────┐
          ▼                     ▼
  INSTRUCTION TABLE         SEMANTIC VALIDATORS
  sub_46E000 (93KB)         sub_4B2F20 (52KB, general)
  sub_46BED0 (per-opcode)   sub_4C5FB0 (28KB, operands)
  sub_46C690 (lookup)       sub_4C2FD0 (12KB, WMMA/MMA)
  sub_46C6E0 (6.4KB match)  sub_4ABFD0 (11KB, async copy)
                            sub_4A73C0 (10KB, tensormap)
                            + 20 more validators

Flex DFA Scanner -- sub_720F00

The scanner is a standard Flex-generated DFA with the ptx prefix (all exported symbols use ptx instead of yy: ptxlex, ptxensure_buffer_stack, ptx_create_buffer, etc.). At 15.8 KB of core logic (64 KB including inlined buffer management), it is the largest single function in the lexer region. The DFA transition table lives at off_203C020 and is indexed by *(DWORD*)(state + 76) (the current start condition). The main loop structure follows the textbook Flex pattern:

// DFA transition core (reconstructed from sub_720F00)
while (1) {
    v10 = (DWORD*)(table_base + 8 * state);   // table[state]
    if (current_char == *v10) {                 // character match
        state = table_base + 8 * v10[1];       // goto next state
        action = *(unsigned int*)(state - 4);   // accept action (or 0)
    }
    if (action != 0) break;                     // matched a rule
}
// Giant switch on action number (0..~550)
switch (action) { ... }

The scanner returns integer token codes to the Bison parser. The value 550 is YY_NULL (end-of-input sentinel). Token attributes are communicated through the lexer state object, which the parser state carries as a pointer at offset +1096. The scanner receives this pointer as its a3 argument and dereferences it (e.g., *(_QWORD *)(a3 + 1096)) to reach the 2,528-byte lexer state.

Token Categories

The ≈552 Flex actions of ptxlex emit 163 distinct named terminal codes in the range [258..422], plus 25 one-character ASCII literals in [0x21..0x7E], plus the three Bison system tokens $end (0), error (256), and $undefined (257) -- a total of 191 concrete terminals. Bison reserves YYNTOKENS = 193 internal slots (see Grammar Parameters) because two of the 165 external slots Bison allocates between 258 and 422 are unused in v13.0.88 (gaps at external codes 321 and 398, confirmed absent from every return statement and every attribute-assignment branch in sub_720F00). The catalog below was reconstructed by walking the action switch linearly and recording every (action-case → token-code → attribute-value) triple from ptxas/decompiled/sub_720F00_0x720f00.c lines 389--2412.

Correction to earlier version

Previous revisions of this page stated "162 distinct token types." That figure was produced by counting entries in an incomplete switch-case enumeration and is wrong in two ways:

1. The actual number of distinct token codes emitted by the scanner is 163, not 162. The previously published table only surfaced 9 of the ≈160 category entries.
2. The Bison reserved-terminal count (YYNTOKENS) is 193, not 162 or 165 -- that constant is derived from yyr1[1] = 193 (the $accept non-terminal, whose symbol number equals YYNTOKENS by Bison convention) and cross-verified at decompiled line 2398 of sub_4CE6B0_0x4ce6b0.c: v31 = (unsigned __int16)word_1D15B80[v1136] - 193;. The three sub-totals decompose as 193 = 165 (external 258..422 slots) + 25 (ASCII literal tokens) + 3 (system: $end, error, $undefined).

Of the 165 external slots, 2 are dead (321, 398), giving the 163 codes actually reachable from the DFA.

1. System / internal (not surfaced as grammar terminals)

2. Identifiers, strings, and numeric literals

3. Structural multi-character punctuators

These have dedicated actions (not the . catch-all).

4. Top-level directives and linkage keywords (code 333--351, 339)

Lines 469--502.

5. Storage-class / state-space directives (code 271--287, 333--336, 383--385)

Lines 495--548.

6. Special-register families (code 264--270, 301--306, 311--319)

Lines 565--609.

7. Type qualifiers (unified under code 320)

All eighteen PTX type keywords are collapsed into token 320 with an integer attribute 1..18 written to *yylval. Decompiled lines 610--663 (cases 74--91).

Exact ordering of the upper bits depends on the PTX version; the decompiled binary only gives the numeric attribute. Attribute value 17 (.tf32-width) is also reachable via action 247 (→ code 404) and attribute 18 via action 248 (→ code 404), used for size-class contexts (line 1129).

8. Comparison operators and rounding modes (322, 323)

Lines 664--687.

9. Saturation / carry / precision modifiers (327, 410, 411, 324)

Lines 688--775.

10. Cache / memory-order modifier families (288--300, 310, 337--368, 407, 414--415)

A large family of fence / cache-operator / memory-ordering attributes, each with an attribute value encoding the specific variant. The block at decompiled lines 706--956 produces these:

11. Instruction-variant selectors (code 292--298, 325--331, 390, 401--413)

Lines 955--1128.

12. Typed-operand constructor tokens (code 275, 403, 404, 416)

These cases call helper constructors (sub_44BD30, sub_44BE60, sub_44C480, sub_44C660, sub_44BF70, sub_44C2F0, sub_44BB80, sub_44BAA0) that synthesize a typed operand descriptor in heap-allocated memory and return token 275. The helpers encode bitwidth and signedness in their call arguments:

Lines 1139--1906. These are not "token type" entries in the classical sense; they are scanner rules that short-circuit grammar reduction by constructing an IR node inline and passing it up as token 275 (the "generic typed-operand" terminal).

13. Opcode hash -- the 386 family (lines 1324--1853)

Code 386 is returned by 178 distinct scanner actions (cases 295--472), each assigning a different attribute value 1..178 to *yylval. This is the opcode identifier terminal: the scanner matches any recognized instruction mnemonic (add, sub, mul, mad, ld, st, mov, bra, call, ret, ...) and returns token 386 with an attribute that indexes into the opcode table built by sub_46E000. The two-level mapping is:

• Scanner: mnemonic-text → (386, opcode_id)
• Parser: opcode_tok(id) → reduces into a generic instruction production whose semantic action looks up opcode_id in the table built by sub_46E000 to find legal type combinations.

Attribute values 1..178 cover the full PTX ISA (v13.0.88). Attribute holes (0, plus values beyond 178) are not emitted. Case 386 at line 1594 returns result instead of the literal 386 -- a compiler oddity from switch-fallthrough optimization (both values are equal in that branch).

The presence of a single code 386 means the grammar does not have one rule per opcode. Instead, there is a single instruction : opcode_tok operand_list production and operand legality is checked at reduction time by cross-referencing the instruction descriptor. This is the classic "table-driven assembler" design and is why PTX can add new opcodes without grammar changes as long as the new opcode's operand shape matches an existing descriptor class.

14. Parameter / entry-function direction modifiers (code 387--396)

Lines 1855--1897. The 391--396 family is six single-valued keywords, likely .managed / .unified / .cluster_ctas / .maxnreg / .maxnctapersm / .maxntid / .reqntid class entry-function attributes.

15. Tensor / WMMA / async-copy modifiers (code 309, 417--422)

Lines 1910--2011. The 309 family's 15-way fanout matches the WMMA/MMA shape enumeration size, making it the most likely identification. Cases 490--504 are densely packed with sequential attribute values -- a strong signature of a generated table.

Coverage statistics

• 549 real Flex actions (numbered 1..549; 550/551 are sentinels)
• 178 actions emit opcode code 386 (32% of all actions)
• 163 distinct emitted named codes in [258..422]
• 2 gaps in the external range (321, 398) -- dead slots reserved for future expansion
• 25 ASCII literal codes reachable via the catch-all case 548
• 3 Bison system tokens ($end=0, error=256, $undefined=257)
• Total terminals (YYNTOKENS) = 193

Line and column tracking uses fields at *(state+48) (line number, ptxlineno) and *(state+52) (column, yycolumn), incremented on each \n character at decompiled lines 370--380 of sub_720F00. Each buffer in the lexer stack (lexer+40 + 8*buffer_idx) has its own line/column counters that are restored on buffer pop (in the YY_STATE_EOF handler at case 551, lines 2501--2502).

Buffer Management

The scanner uses the standard Flex buffer stack for nested input sources (includes, macros, inline strings). Key buffer management functions:

Notable: sub_720630 contains SSE2-optimized memmove using __m128i aligned 16-byte copies for buffer compaction -- a Flex optimization for large input buffers. The ptx_scan_bytes function (sub_724CC0) is called from the Bison parser actions (3 call sites in sub_4CCF30) to handle inline macro expansion during parsing.

Error strings in the buffer system:

• "out of dynamic memory in ptxensure_buffer_stack()"
• "out of dynamic memory in ptx_create_buffer()"
• "out of dynamic memory in yy_get_next_buffer()"
• "out of dynamic memory in ptx_scan_bytes()"
• "bad buffer in ptx_scan_bytes()"
• "out of dynamic memory in ptx_scan_buffer()"
• "fatal flex scanner internal error--no action found"
• "fatal flex scanner internal error--end of buffer missed"
• "unexpected EOF while scanning"

Macro Preprocessor

Before tokens reach the Flex DFA, a character-level macro preprocessor handles .MACRO/.ENDM, .ELSE/.ELIF/.ENDIF, and .INCLUDE directives. The preprocessor lives at 0x71B000--0x720000 (~20 KB) and operates on raw character streams, not tokens. This design is identical to C's preprocessor running before the lexer.

Preprocessor Dispatch -- sub_71F630

The top-level dispatcher (14 KB) is called from the Flex scanner's case 3 (directive detection). It examines the directive name and routes to the appropriate handler:

The dispatcher uses strstr for substring matching on directive names and returns token codes (e.g., 364 for end-of-directive).

Conditional Handler -- sub_71E2B0

At 32 KB, this is the largest preprocessor function. It handles .ELSE, .ELIF, and .ENDIF by scanning ahead through the input character stream, counting nesting levels, and skipping entire blocks of PTX text when conditions are false. It calls sub_4287D0 (the token reader) to evaluate conditional expressions and sub_428C40 (string compare) for keyword matching. Two nearly-duplicate code blocks handle .ELSE and .ELIF paths with identical scanning logic but different branch conditions.

Macro Definition -- sub_71DCA0

Handles .MACRO directives by recording the macro body text. The function is recursive to support nested .MACRO definitions. It delegates to sub_71D710 (macro body scanner, 7.5 KB) and sub_71D1B0 (macro argument scanner, 6.8 KB). The argument scanner uses strlen + strncmp for keyword matching against a delimiter string parameter.

Include Handler -- sub_71C310

Processes .INCLUDE by pushing a new file onto the lexer's input stack. The function is recursive (calls itself 4 times) for nested includes. It manages the include-stack pointers at offsets +2128, +2136, +2160, and +2168 of the lexer state object (the 2,528-byte struct pointed to by parser+1096), and uses the "pushback character" register at offset +2441 of the same lexer state. String reference: "ptxset_lineno called with no buffer".

Error Handling

Macro errors are reported through sub_71BF60 (fatal macro abort) which calls sub_71BF30 to print "out of dynamic memory..." messages, and sub_71C140 (format error) which calls sub_42CA60 (error output). Nesting depth is checked by sub_724CC0 which prints "macro nesting too deep!" on overflow.

Bison LALR(1) Parser -- sub_4CE6B0

The parser is a standard Bison-generated LALR(1) shift-reduce parser spanning 48 KB (addresses 0x4CE6B0--0x4DA337). It contains exactly 513 grammar productions (rules 1--513, plus the implicit $accept augmentation) with 443 reduction cases carrying non-default semantic actions; the remaining 70 rules use Bison's default action ($$ = $1, no code emitted). The function calls ptxlex (sub_720F00) to obtain tokens and uses nine LALR tables for state transitions, action lookup, and goto computation:

Table naming correction. Earlier versions of this page listed word_1D146A0 as yydefact, word_1D13360 as yypact, and word_1D150C0 as yypgoto. These labels were wrong: the decompiled indexing patterns (below) identify word_1D146A0 as yypact (indexed by state, values range [-921..2101] with the -921 sentinel), word_1D13360 as yytable (indexed by yypact[state]+token), and word_1D150C0 as yydefact (indexed by state, values are reduction rule numbers 0..513). The true yypgoto and yydefgoto (not previously documented) live at 0x1D14520 and 0x1D14F40.

Grammar Parameters

The following constants were recovered by cross-correlating the five hardcoded sentinels in the decompiled state machine with the .rodata table sizes and boundary signatures. File references use sub_4CE6B0_0x4ce6b0.c in /ptxas/decompiled/.

Terminal alphabet composition. The 193 terminals decompose as: YYEOF (code 0 → terminal 0), YYerror (code 256 → terminal 1), YYUNDEF (code 257 → terminal 2), 165 named tokens emitted by the Flex scanner (codes 258--422 → terminals 3--167), and 25 single-character literal tokens at ASCII codes ! % & ( ) * + , - / : ; < = > ? @ [ \ ] ^ { | } ~ (→ terminals 168--192). This explains the "162 token types (codes 258--422)" figure in the page header as an undercount: the true named-token range is 258--422, i.e. 165 distinct named tokens; with the 25 character literals and the three housekeeping tokens, the grammar's terminal count is 193.

Memory footprint of the tables. .rodata VA range 0x1D121A0..0x1D161E8 holds ≈19.9 KB of parser data: yycheck 4,538 B + yytable 4,538 B + yypgoto 364 B + yypact 2,198 B + yydefgoto 364 B + yydefact 2,198 B + yyr2 514 B + yyr1 1,028 B + yytranslate 423 B = 16,165 B logical, padded to 19,968 B including inter-table alignment. The 48 KB function body plus tables brings the total Bison footprint to ≈68 KB.

State Machine Walkthrough

The main parser loop (decompiled lines 1354--1507) follows the textbook Bison deterministic-LALR skeleton:

for ( state = 0; ; yypact_val = yypact[state] )  // line 1354
{
    if ( yypact_val == YYPACT_NINF )  goto yydefault;  // line 1357
    if ( lookahead == YYEMPTY )                        // line 1359
        lookahead = ptxlex();                          // line 1363 (sub_720F00)
    if ( lookahead <= 0 )      translated = 0;         // EOF
    else if ( lookahead > 422 ) translated = 2;         // YYUNDEF
    else                       translated = yytranslate[lookahead];  // line 1374
    idx = yypact_val + translated;                     // line 1375
    if ( idx > YYLAST )        goto yydefault;         // line 1376 (0x8DC)
    if ( yycheck[idx] != translated ) goto yydefault;  // line 1386
    // match: perform shift or reduce via yytable
    action = yytable[idx];                             // line 1497
    if ( action > 0 )                                  // shift
        push_state(action);
    else if ( action == 0 || action == YYTABLE_NINF )  // error
        goto yyerrlab;
    else
        reduce(-action);                               // negated = rule number
    continue;
  yydefault:
    action = yydefact[state];                          // line 1389
    if ( action != 0 ) reduce(action);
    else yyerrlab;
}

After a reduction the goto computation (lines 2392--2403) is:

nt_idx = yyr1[reduced_rule] - YYNTOKENS;  // line 2398 (- 193)
new_state = yypgoto[nt_idx] + current_state;
if ( new_state > YYLAST || yycheck[new_state] != current_state )
    new_state = yydefgoto[nt_idx];        // line 2403
else
    new_state = yytable[new_state];       // line 2401
push_state(new_state);

This is bit-for-bit identical to the Bison 3.x yacc.c skeleton (cf. Bison source tree data/yacc.c, functions yybackup: and yynewstate:).

Conflict Count (Not Recoverable)

Bison's deterministic parser tables preserve only the winning action for every (state, terminal) pair; shift-reduce and reduce-reduce conflicts resolved at generation time (via %left/%right/%nonassoc precedence or Bison's default shift-wins rule) leave no trace in yytable/yycheck. The tables cannot distinguish "state had a single action" from "state had several actions and the generator picked one". We can therefore characterize only the structural shape of conflict-resolution, not its count:

• 708 / 1,099 states (≈64%) have yypact[s] = YYPACT_NINF (no shift action); these states unconditionally reduce via yydefact[s].
• 255 states have yypact[s] ≠ NINF and yydefact[s] = 0: pure "shift-only" states, no default reduction.
• 136 states have both yypact[s] ≠ NINF and yydefact[s] ≠ 0: a shift set and a default reduction. This is the normal LALR configuration for states where the lookahead set partitions into "shiftable tokens → shift" and "everything else → reduce rule N". These 136 states are the candidates where a shift-reduce conflict could have been present in the source grammar and resolved by precedence, but the table shape alone does not prove a conflict occurred.
• 0 states have both yypact[s] = NINF and yydefact[s] = 0 ⇒ no state is "dead"; every state has at least one action.

The absence of any runtime conflict-handling code (no yyconfl/yyconflp tables, no GLR dispatch, no %glr-parser scaffolding) confirms a clean LALR(1) generation with zero runtime conflict-resolution logic. Any conflicts were resolved at generation time and baked into the tables; the generator's statistics (e.g. "N shift/reduce, M reduce/reduce") are gone. If an upper bound is needed for reimplementation purposes, 136 is a conservative overcount of the states that could have contained conflicts before resolution.

The parser is definitively not a GLR parser (no %glr-parser / yyconfl evidence), nor a push-pull API parser (the entry point at line 7490 takes parameters by value and returns a status code synchronously).

Direct IR Construction (No AST)

The critical architectural decision: Bison reduction actions directly construct IR nodes rather than building an intermediate AST. When a grammar rule is reduced, the semantic action immediately:

1. Allocates IR nodes via the pool allocator (sub_424070)
2. Populates instruction fields from token attributes
3. Calls instruction validators for semantic checking
4. Links nodes into the instruction stream
5. Registers symbols in the symbol table (via sub_426150, the hash map)

This means the parser is a single-pass translator from PTX text to IR. The trade-off is clear: no AST means no multi-pass source-level analysis, but it eliminates an entire allocation and traversal phase. For an assembler (as opposed to a high-level language compiler), this is the right choice -- PTX is already a linearized instruction stream with no complex scoping or overload resolution that would benefit from an AST.

Reduction Actions -- Semantic Processing

The 443 reduction cases in the parser body handle PTX constructs from simple register declarations to complex matrix instruction specifications. Diagnostic strings found in the parser tail (0x4D5000--0x4DA337) reveal the kinds of semantic checks performed during reduction:

Directive validation:

• "Defining labels in .section"
• "dwarf data" -- DWARF section processing
• "reqntid" / ".reqntid directive" -- required thread count
• ".minnctapersm directive" -- min CTAs per SM
• ".maxnctapersm" / ".maxnctapersm directive" -- max CTAs per SM (deprecated)
• ".maxntid and .reqntid cannot both be specified"
• ".maxnctapersm directive deprecated..."
• ".minnctapersm is ignored..."

Type and operand validation:

• "Vector Type not specified properly"
• ".f16x2 packed data-type" -- half-precision packed type
• "matrix shape" -- matrix instruction dimensions
• ".scale_vectorsize" -- vector scaling modifier
• "too many layout specifiers"

Resource limits:

• "Kernel parameter size larger than 4352 bytes"

Architecture gating:

• "sm_50", "sm_20", "sm_53" -- target architecture checks via sub_485520(ctx, sm_number)
• PTX version checks via sub_485570(ctx, major, minor)

Expression handling:

• "%s+%llu" / "%s-%s" -- label arithmetic in address expressions
• "Negative numbers in dwarf section" -- DWARF data validation

Symbol resolution:

• "unrecognized symbol" -- lexer/symbol table failure
• "syntax error" -- generic parse error
• ".extern" -- external declarations
• ".noreturn directive" -- function attributes
• "texmode_unified" / "texmode_raw" -- texture mode selection
• "cache eviction priority" / ".level::eviction_priority" -- cache policy

Error Recovery

Parse errors trigger sub_42FBA0 with "syntax error" as the message. The central diagnostic emitter (sub_42FBA0, 2,388 bytes, 2,350 callers) handles all severity levels:

The diagnostic system reads the source file to display context lines (prefixed with "# "), caching file offsets every 10 lines in a hash map for fast random-access seeking.

Parser Initialization -- sub_451730

Parser initialization (14 KB) builds the lexer's symbol table with all built-in PTX names before parsing begins. This function is called from the compilation driver (sub_446240) and performs three major tasks:

1. Special Register Registration

All PTX special registers are pre-registered in the symbol table with their internal identifiers:

2. Opcode Table Construction

Calls sub_46E000 -- the 93 KB instruction table builder -- to register all PTX opcodes with their legal type combinations. See the dedicated section below.

3. Context State Initialization

Allocates and initializes two objects: the parser state (1,128 bytes, sub_424070(pool, 1128)) and the lexer state (2,528 bytes, sub_424070(pool, 2528)). The parser state stores a pointer to the lexer state at offset +1096. The string "PTX parsing state" identifies the parser state allocation in memory dumps. The string "<builtin>" serves as the filename for built-in declarations. Both objects are zeroed via memset before field initialization.

Instruction Table Builder -- sub_46E000

This is the largest single function in the front-end region at 93 KB (disasm/sub_46E000_0x46e000.asm, 17,842 lines, 1,464,386 bytes). It is not a normal function body but a massive initialization sequence that calls sub_46BED0 exactly 1,141 times -- once per legal PTX instruction variant. Every call registers one (opcode name, operand encoding, type-suffix code, validator index) tuple.

Reconstruction note. sub_46E000 is intentionally not present in ptxas/decompiled/ for this build -- at 93 KB of straight-line calls it overflows the hand-decompilation budget. Everything in this section is reconstructed from three orthogonal sources that fully determine the answer: (1) the decompiled registration function sub_46BED0_0x46bed0.c (321 lines) which fixes the exact shape of each descriptor, (2) the decompiled matcher prologue sub_46C6E0_0x46c6e0.c (first 200 lines read for descriptor-field offsets and the 12-category operand classifier), and (3) the 1,080 .weak .func lowering targets in ptxas/extracted/embedded_ptx_intrinsics.json (entry_count: 1080, categorized as cuda_other: 549, sm70_intrinsics: 433, sm20_math: 70, redux_sync: 17, sanitizer: 7, sm80_intrinsics: 4) cross-referenced against the 322 SASS mnemonics in ptxas/extracted/opcode_master.json (opcode_master.count: 322) and the public PTX ISA 8.x reference. The per-family variant counts below were built by decoding the little-endian imm32 operands of every mov REG, imm32; ...; call sub_46BED0 call site in disasm/sub_46E000_0x46e000.asm and reading the pooled C-string at that RVA from .rodata; the sum of per-family counts equals exactly 1,141, matching the static call count observed in the disassembly.

368-byte Descriptor Layout

Every registration allocates one descriptor at sub_46BED0_0x46bed0.c:63 (v18 = sub_424070(v15, 368)). The layout below is recovered by mapping every field write in sub_46BED0 to a byte offset, then cross-checking each offset against the reads in sub_46C6E0 (the matcher) where the same descriptor is walked operand-by-operand. All offsets are cited back to the decompiled line where the write occurs.

Notes on the layout:

• The type_class[] and width_mask[] arrays are indexed by v21, the running operand-slot counter (initialized to -1 at line 67 so the first ++v21 writes slot 0). v21 is bumped either pre-switch (N, O, P, T, E cases, which use ++v21) or deferred into v48 and committed at LABEL_26 (F, H, I, B, Q, R cases, which allow a [width|width|...] group to follow). Both paths write the same logical slot, but the offsets differ in the decompiled source: v18+v21+9 in the pre-increment path vs v18+v21+10 in the deferred path -- that is because at the deferred write-site v21 is still the previous operand's index, so the +10 reaches one slot forward. Either way the class lands at byte_offset = 36 + 4*slot.
• The width_mask[] array holds pointers to bitmask objects, not bitmasks themselves. sub_1CB0790() allocates an empty mask; sub_1CB0850(mask, w) adds bit w to it. The matcher tests a candidate operand width against the mask via sub_1CB06F0/sub_1CB0700 (not decompiled here; referenced from sub_46C6E0).
• The type_subclass[] / type_digit[] split lets one suffix character encode either a symbolic tag (upper- or lowercase letter, 22 distinct values) or a small integer (0..9). A digit-class slot has type_subclass[i] = 0 and type_digit[i] = digit; the matcher looks at type_digit[i] only when type_subclass[i] == 0.
• The 16-slot cap on per-operand arrays is the hard upper bound on operand count for any PTX instruction: no registered opcode has more than 16 operands (the biggest real examples are _mma.warpgroup at 12 and wmma.mma at 10).

Each call passes four distinguishing arguments in the System V AMD64 calling convention (verified against the hand-decoded prologue in sub_46BED0_0x46bed0.c lines 4--12, which declares a1..a8):

Concretely, an add.f32 registration looks like this in the raw disassembly (addresses from disasm/sub_46E000_0x46e000.asm):

0x472fe9: b9 13 92 ce 01   mov     ecx, (offset a10000+3); "000"     ; a4 = type suffixes "000"
0x472fee: ba 39 81 d0 01   mov     edx, offset aAdd_0; "add"          ; a3 = opcode name "add"
0x472ff3: be 35 b8 02 02   mov     esi, (offset aNanNotAllowedW+18h); "F32"  ; a2 = encoding "F32"
0x472ff8: 48 89 df         mov     rdi, rbx                           ; a1 = lexer state
0x472ffb: 41 b8 2f 00 00 00 mov     r8d, 2Fh                          ; a5 = 47  (semantic index)
        ... xmm0 / stack-slot setup for a7/a8 ...
0x47302d: e8 9e 8e ff ff   call    sub_46BED0

The string "F32" is itself a substring of a longer pooled string "NaN not allowed with W..." -- IDA represents it as (offset aNanNotAllowedW+18h) because ptxas deduplicates its read-only strings aggressively, so a 3-character suffix can be reused across many callers. IDA also truncates long symbol comments: names such as _tcgen05.guardrails.sp_consistency_across_idesc_mod_scale appear with "..." in the comment, but the full string is recoverable by decoding the little-endian imm32 in the mov edx, imm32 instruction bytes and reading the null-terminated C-string at that address from .rodata (range 0x1CE2E00..0x240BF90). The catalog below was built by exactly that procedure.

Operand Encoding Alphabet

The encoding string is consumed character-by-character by the first switch in sub_46BED0 (lines 90--229). The original wiki listed six codes; the decompiled switch actually handles eleven type-class codes plus structural punctuation:

The loop accumulates decimal width digits into v46 at line 217 (v46 = v26 + 10*v46 - 48), then commits v46 to the current operand's width mask on |, ], or boundary via sub_1CB0850(descriptor, width) (line 223). Structural punctuation:

• [ opens a width-set group (line 210, continue, no state change)
• | commits the pending digit as an additional legal width (LABEL_12 at line 220)
• ] commits the final width and closes the group (same LABEL_12)

So B[8|16|32|64] is parsed as: open group, push width 8, push width 16, push width 32, push width 64, close. The result is a B-class descriptor whose "allowed widths" mask contains exactly those four bits. F[32|64] restates the default set explicitly.

Encoding-string first-character frequencies observed across all 1,141 registrations:

The remaining 143 of the 1,141 registrations pass an empty string for the encoding (""). These are the control-flow, barrier, async-bulk, and tcgen05.mma opcodes that take no typed-operand sequence -- the encoding is determined entirely by state-space modifiers on the opcode itself, checked by the semantic validator rather than by descriptor-list matching.

Type-Suffix Code String (a4)

The second string argument drives the loop at lines 234--316 of sub_46BED0_0x46bed0.c. Each character maps the i-th operand's fine-grained type flavor into *((_DWORD *)v18 + v27 + 59) (the descriptor's per-operand type-subclass array, sixteen DWORDs starting at byte offset +236):

Digit characters (0--9) write 0 to the class slot and the digit value to *((_DWORD *)v18 + v27 + 75) (default case at line 309). Common suffix strings from the catalog:

• "000" (49 full-string matches, 877 '0' characters in total) -- "three operands, all default type class" (e.g. scalar add, sub, min, max)
• "0000" -- four operands all default (e.g. fma, mad)
• "M0" / "0M" -- memory operand paired with a typed register (e.g. ld, st)
• "hhhhdC", "fhhddC", "sddsdC" -- the MMA per-operand tag sequence (h=half, f=float32, s=int, d=int32 accumulator, C=.cc flag)
• "0i1" / "0i1s" -- texture sequence (sampler index, texref, coord, optional shadow)

Two Hash Tables: +2472 and +2480

sub_46E000 populates two separate hash tables on the lexer state object, both allocated by back-to-back sub_425CA0 calls in the prologue (disasm lines 8--24). The two tables correspond to different string namespaces:

• lexer_state+2472 (offset 0x9A8) -- the primary opcode table (128 buckets, sub_425CA0(sub_427630, sub_4277B0, 0x80)); every user-visible PTX opcode lives here.
• lexer_state+2480 (offset 0x9B0) -- a smaller 16-bucket table (sub_425CA0(…, 0x10)) used by the matcher as a secondary probe.

sub_46BED0 always inserts into +2472 (line 317: v29 = *(char**)(a1 + 2472)). The matcher side (sub_46C690_0x46c690.c:11) probes +2472 first, then falls back to +2480; sub_46C6E0:258 adds a special-case fast path that probes only +2480 when the opcode starts with ASCII _ (byte value 95). That fast path is used for the 18 _-prefixed compiler-private pseudo-opcodes (_mma, _mma.warpgroup, _ldsm, _warpgroup.*, _tcgen05.guardrails.*, _movm, _gen_proto, _jcall, _match, _checkfp.divide, _sulea.*, _createpolicy.*, _ldldu, _warpsync), which are emitted by earlier passes and are not part of the public PTX ISA. The +2480 table is therefore the internal-opcodes index and is populated by code paths inside sub_46E000 that are not visible as sub_46BED0 calls (they use sub_425CA0 infrastructure directly); identifying those sites is a follow-up item.

Registration Function -- sub_46BED0 Pseudocode

// Reimplementation derived from sub_46BED0_0x46bed0.c, lines 4--321.
//   a1 : LexerState*  — primary hash table at a1+2472
//   a2 : const char*  — encoding string (iterated by first switch)
//   a3 : const char*  — opcode name (interned; becomes hash key)
//   a4 : const char*  — type suffix string (iterated by second switch)
//   a5 : int          — semantic index / validator selector (stored at +8)
//   a6 : __int64      — extra payload (stack arg, unused in hot path)
//   a7 : __m128i      — 16-byte validator function pointer pair (stored at +12)
//   a8 : int          — validator flags (stored at +28)
char *register_opcode(LexerState *a1, const char *a2, const char *a3,
                      const char *a4, int a5, __int64 a6, __m128i a7, int a8)
{
    // Count "tokens" in the encoding string: every non-space char advances v12.
    int token_count = 0;
    for (const char *p = a2; *p; p++)
        token_count += ((*__ctype_b_loc())[*p] & 0x400) ? 0 : 1;   // 0x400 = _ISspace

    // Allocate a 368-byte descriptor from the compilation pool (line 63).
    Pool *pool = /* pool handle cached behind sub_4280C0 */;
    Descriptor *d = pool_alloc(pool, 368);             // sub_424070
    memset(d, 0, 368);
    d->name       = a3;                                // *v18 = a3                (line 72)
    d->token_cnt  = token_count;                       // ((int*)d)[8]              (line 73)
    d->sem_index  = a5;                                // ((int*)d)[2]              (line 76)
    d->validator  = a7;                                // 16-byte pair at byte +12  (line 79)
    d->flags      = a8;                                // ((int*)d)[7]              (line 78)
    d->suffix_len = strlen(a4);                        // ((int*)d)[58]             (line 80)

    // --- Parse encoding string a2 into per-operand (type_class, width_mask) tuples
    int     operand_idx    = -1;
    unsigned pending_digits = 0;
    for (size_t j = 0, n = strlen(a2); j < n; /* advanced inside */) {
        char c = a2[j++];
        switch (c) {
        case 'F':   ((int*)d)[++operand_idx + 9] = 1;  // float                  (line 107)
                    d->width_mask[operand_idx] = mask_new();
                    if (end_of_token(a2, j)) { mask_add(..., 32); mask_add(..., 64); }
                    break;
        case 'H':   /* class=2, defaults {32,64} */           break;    // line 111
        case 'I':   ((int*)d)[++operand_idx + 9] = 4;  // integer                (line 124)
                    d->width_mask[operand_idx] = mask_new();
                    if (end_of_token(a2, j)) { mask_add(..., 16); mask_add(..., 32); mask_add(..., 64); }
                    break;
        case 'B':   /* class=5, defaults {1,16,32,64}   */   break;    // line 92
        case 'N':   /* class=3, defaults {32}            */  break;    // line 139
        case 'P':   /* class=6, defaults {32}            */  break;    // line 156
        case 'E':   /* class=8, widths always explicit   */  break;    // line 103
        case 'O':   /* class=7, no width                 */  break;    // line 150
        case 'Q':   /* class=10, defaults {8,16,32}      */  break;    // line 165
        case 'R':   /* class=11, defaults {4,8,16}       */  break;    // line 178
        case 'T':   /* class=9, no width                 */  break;    // line 202
        case '[':   continue;                                          // line 210
        case ']':
        case '|':   goto commit_digit;                                 // line 213
        default:
            if ((unsigned)(c - '0') <= 9u) {                           // line 215
                pending_digits = pending_digits * 10 + (c - '0');
                if (!end_of_token(a2, j)) continue;
commit_digit:
                if (operand_idx < 0) operand_idx = 0;
                mask_add(d->width_mask[operand_idx], pending_digits);   // sub_1CB0850 (line 223)
                pending_digits = 0;
            }
            break;
        }
    }

    // --- Parse suffix string a4 into per-operand type-subclass array (lines 234--316)
    for (size_t k = 0; k < d->suffix_len; k++) {
        switch (a4[k]) {
        case 'A': d->type_sub[k] = 20; break;  case 'b': d->type_sub[k] = 8;  break;
        case 'C': d->type_sub[k] = 13; break;  case 'c': d->type_sub[k] = 9;  break;
        case 'D': d->type_sub[k] = 14; break;  case 'd': d->type_sub[k] = 10; break;
        case 'L': d->type_sub[k] = 22; break;  case 'e': d->type_sub[k] = 11; break;
        case 'M': d->type_sub[k] = 17; break;  case 'f': d->type_sub[k] = 5;  break;
        case 'P': d->type_sub[k] = 15; break;  case 'h': d->type_sub[k] = 6;  break;
        case 'Q': d->type_sub[k] = 16; break;  case 'i': d->type_sub[k] = 12; break;
        case 'S': d->type_sub[k] = 18; break;  case 'l': d->type_sub[k] = 7;  break;
        case 'T': d->type_sub[k] = 19; break;  case 's': d->type_sub[k] = 4;  break;
        case 'U': d->type_sub[k] = 3;  break;  case 'u': d->type_sub[k] = 2;  break;
        case 'V': d->type_sub[k] = 21; break;  case 'x': d->type_sub[k] = 1;  break;
        default:                               // digit                     (line 309)
            d->type_sub[k]        = 0;
            d->type_sub_digit[k]  = a4[k] - '0';
            break;
        }
    }

    // --- Append into the hash bucket at LexerState+2472 ---
    HashTable *tbl = *(HashTable**)((char*)a1 + 2472);
    HashBucket *bucket = sub_426D60(tbl, d->name);       // bucket lookup (line 318)
    sub_42CA00(d, bucket);                               // link into bucket list (line 319)
    return sub_426150(tbl, d->name, d);                  // commit (line 320)
}

The hash insert uses d->name (the interned opcode pointer, deduped across the entire binary) as the key. Because PTX opcode names are pooled, the string pointer equality is the hash key equality -- no strcmp is needed on insert or lookup.

Catalog -- 1,141 Registrations by PTX Family

The tables below give every registered opcode, with its variant count and the distinct operand encoding strings it accepts. Methodology: all 1,141 mov edx, <opcode>; mov esi, <encoding>; mov ecx, <suffix>; ...; call sub_46BED0 call sites were extracted from disasm/sub_46E000_0x46e000.asm by decoding the little-endian imm32 in each mov REG, imm32 instruction's byte stream and reading the null-terminated C-string at that address from ptxas_rodata.bin. This bypasses IDA symbol-comment truncation ("_tcgen05.guardrails.sp_consistency_acro"...) and substring-reference opacity (1D07286h -> inner offset of the pooled string "F32F16" -> the 3-byte substring "F16").

Totals by category:

The three dominant families -- tensor cores (360 = 31.6%), texture (222 = 19.5%), and arithmetic (96 = 8.4%) -- account for 60% of all registrations. The tensor-core bloat comes almost entirely from _mma.warpgroup (135 variants) and the public _mma/mma families (41+38). Texture bloat is structural: each of tex, tex.base, tex.level, tex.grad registers the identical 8-encoding x 6-mode cross-product = 48 variants, giving 192 before adding tld4, txq, sust, sured, _sulea, suq, suld.b.

Arithmetic (96 variants / 32 opcodes)

Full expansion for add (all 11 variants with suffix codes):

And the full fma expansion (9 variants):

Comparison and select (80 variants / 5 opcodes)

set blows up to 54 variants because it registers every setp-to-data type pair, i.e. it emits an integer/bitwise destination rather than a predicate. All 16 setp variants come from a (6 type classes) x (with/without predicate-merge) cross product plus two forms for E16/E32. Full setp expansion:

Logic and bitwise (16 variants / 10 opcodes)

Bit manipulation (10 variants / 10 opcodes, one apiece)

Transcendental (19 variants / 8 opcodes)

Note: rcp.approx.ftz.f32, rcp.approx.f64, etc. are not separate registrations. Approximation mode, ftz, saturation, and rounding mode are all carried by modifier tokens on the opcode and rejected by the semantic validator rather than by a per-mode descriptor. This is why rcp has only 2 variants despite several user-visible modes.

Data movement (44 variants / 8 opcodes)

cvt has the most per-opcode diversity of any family: 28 type-pair variants. Full list:

cvt.pack adds 3 more for cvt.pack.sat.<T>.<Tsrc> widening-pack forms.

Memory load/store (55 variants / 23 opcodes)

ld has 12 variants = 4 encodings x 3 suffix-string shapes (regular, vector-packed, .U-suffix uniform). B128 appears three times with different suffix flags (0M, 0M, 0MU) to separate non-vector 128-bit loads, vector-packed 128-bit loads, and uniform 128-bit loads.

Atomic / reduction (41 variants / 3 opcodes)

The 21 atom variants span 10 encodings x up to 3 suffix variants each. The suffix tail "0M0U" marks forms where a .uni (uniform) post-modifier is legal. atom with B128 appears three times: atom.cas.b128, vector-packed B128, and the uniform form.

Barrier / fence / mbarrier (56 variants / 35 opcodes)

The bar.red/barrier.red fourfold expansion is (I32 integer arg, P predicate arg) x (cta, non-cta scope). The 13 distinct mbarrier.* opcodes (init, inval, complete_tx, expect_tx, pending_count, test_wait, test_wait.parity, try_wait, try_wait.parity, arrive, arrive.expect_tx, arrive_drop, arrive_drop.expect_tx) sum to 17 total variants because 4 of them register both "with-count" and "without-count" forms.

Control flow (17 variants / 9 opcodes)

call registering 8 variants is a surprise: each one corresponds to a different argument-list shape that the Bison grammar produces (call (retvals), fn, (args); versus call.uni, prototype reference, no prototype, direct target, indirect target, etc.). They all share the empty encoding string because the arguments are validated semantically after parsing, not through the descriptor type matcher.

Texture / surface (222 variants / 15 opcodes)

Each of the four tex* opcodes registers the exact same 8 encodings (F16F32, F16I32, F32F32, F32I32, H32F32, H32I32, I32F32, I32I32) six times, once per texture-mode suffix ("0i1", "0i1s", and four more for 1D / 2D / 3D / cube / array / shadow-array combinations). That is 4 x 8 x 6 = 192 registrations on its own -- 17% of the entire instruction table -- which is why texture is the single largest non-MMA category.

Warp-level collectives (17 variants / 11 opcodes)

vote has 4 variants: vote.all.pred, vote.any.pred, vote.uni.pred, and vote.ballot.b32.

Cooperative copy / tensor memory (75 variants / 21 opcodes)

cp.async.bulk.tensor registers 10 variants (one per tensor map dimensionality x load/store x multicast) without any per-variant type checking -- the 10 come from different suffix strings using the M and C tags.

Tensor cores -- MMA / wgmma / wmma / tcgen05 (360 variants / 37 opcodes)

This is the largest single category, accounting for 31.6% of the entire instruction table. _mma.warpgroup alone registers 135 variants, making it the single largest opcode in the binary:

The 8 encoding shapes shared by _mma.warpgroup and wgmma.mma_async are the story of SM 90 / SM 100 tensor cores:

Each of those eight encodings is registered up to 18 times for _mma.warpgroup via different suffix strings, covering (dense vs sparse) x (scale vs no-scale) x (stride A vs B) x (predicate fill) = up to 16+ modes. The P-tailed suffixes (hUUhP, fUUfP, ...) are the predicate-returning forms used to signal MMA completion back to the warpgroup scheduler.

Video SIMD (31 variants / 23 opcodes)

Every video/SIMD opcode has exactly 1 or 2 variants and uses the same I32I32I32 encoding (three-operand int32 SIMD); differences are all carried in the suffix string:

(vmad is technically not a packed form but registers with the same encoding.) These are the PTX 5.0 video SIMD intrinsics -- effectively thin wrappers around the .s32/.u32 variants of each underlying SASS instruction, so the semantic difference lives entirely in the suffix string (e.g. "sddsdC" vs "uuuddC").

Verification / debug (2 variants / 2 opcodes)

Instruction Lookup -- sub_46C690 and sub_46C6E0

At parse time, when the parser reduces an instruction production, it calls sub_46C690 to look up the instruction name in the hash table built by sub_46E000. The lookup returns a descriptor list, and sub_46C6E0 (6.4 KB, the descriptor matcher) walks the list to find the variant matching the actual operands present in the source.

sub_46C690 (lines 4--16 of sub_46C690_0x46c690.c) is a trivial wrapper: it probes the two opcode hash tables at lexer-state offsets +2472 and +2480 with sub_426D60 and returns the first nonzero bucket's *(_DWORD*)(entry+8+8) (the descriptor head pointer). The real work happens in sub_46C6E0, which is called directly from the Bison reduction actions with the raw token list.

Operand Classification -- 12 Categories

The descriptor matcher classifies every operand into one of twelve category codes before walking the candidate descriptor list. The classifier is the leading loop in sub_46C6E0 (lines 142--249 of sub_46C6E0_0x46c6e0.c): it iterates a8 times (a8 = parsed operand count), reads each 8-byte operand-token pointer v14 = *(_DWORD **)(a6 + 8*i) (note: the source uses 2 * v13 with v13 stepped by 4, which is a byte stride of 8), and dispatches on *v14 (the first DWORD, a lexer token-kind enum, distinct from the AST-node 6-bit tag of IR-08). The switch writes two parallel slots of the stack array v133:

• v133[i] -- the category code (0--11), occupying [0..15]
• v133[i + 16] -- the operand's bit width obtained from sub_44B390(v14) (which walks the type token and folds *= 2/4/8/16/32/64/128 or *= arraylen for aggregates)

Category 0 is the implicit default (any token that hits the default: break; at line 244 leaves v133[i] unwritten, so it is effectively the sentinel "unclassified"). That produces twelve distinct states numbered 0--11. Every classification is a pure table lookup on *v14; no flag bits, no uniform-register 0x6000000 mask, no (field>>28)&7 test -- those checks live in the lexer (where the token kind itself was assigned), not in the matcher. By the time sub_46C6E0 sees the operand, the distinction between R5 vs UR5 vs %r5 is already baked into the numeric value of *v14.

The 12 Category Codes

The distinction is that the classifier runs once per operand in a tight switch, while the matcher then walks a list of candidate descriptors and rejects each one using a series of descriptor-bit-against-category-code filters before finally doing a full per-operand check. The 11 explicit categories plus category 0 (default/unclassified) give 12 states, which fits in 4 bits -- but the binary stores them as full _DWORDs (v133 is _DWORD v133[32], line 135) because the compare at line 950 (if ( v50[9] != v133[0] )) is a direct int compare against the descriptor's pre-serialized category sequence.

Classifier Pseudocode

// Extracted from sub_46C6E0 lines 142-249. Pure function over token array.
void classify_operands(
    const OperandToken *tokens[],  // a6: array of 8-byte token pointers
    size_t count,                  // a8: number of operands
    uint32_t cat[32],              // v133[0..15]  -- output category codes
    uint32_t width[32])            // v133[16..31] -- output bit widths
{
    // v133 is zero-initialized only implicitly; default: leaves cat[i] at
    // its prior (unset) value, which effectively acts as "category 0".
    for (size_t i = 0; i < count; i++) {
        const uint32_t *tok = (const uint32_t *)tokens[i];
        switch (tok[0]) {                       // first DWORD = token-kind enum
            case 0x01: case 0x02: case 0x03: case 0x04:
            case 0x05: case 0x06: case 0x07: case 0x08:
                cat[i] = 5; break;              // type-width / qualifier
            case 0x09: case 0x0A: case 0x0B: case 0x0D:
            case 0x0F: case 0x10: case 0x11: case 0x12:
            case 0x13: case 0x15: case 0x17: case 0x18:
                cat[i] = 4; break;              // small/medium integer register
            case 0x0C: case 0x0E: case 0x14: case 0x16:
                cat[i] = 3; break;              // 32b+ float / packed vector reg
            case 0x19: case 0x1A: case 0x1B: case 0x1C:
            case 0x1D: case 0x1E: case 0x1F: case 0x22:
            case 0x25: case 0x26: case 0x27: case 0x28:
            case 0x29: case 0x2A: case 0x2B: case 0x2C:
            case 0x2D: case 0x2E: case 0x2F: case 0x30:
            case 0x33:
                cat[i] = 10; break;             // typed integer immediate
            case 0x20: case 0x21: case 0x23: case 0x24:
            case 0x31: case 0x32:
                cat[i] = 11; break;             // typed float immediate
            case 0x34: case 0x3A: case 0x3B:
                cat[i] = 1;  break;             // label
            case 0x35: case 0x36:
                cat[i] = 8;  break;             // .const / .param addr
            case 0x37:
                cat[i] = 9;  break;             // .global addr
            case 0x38: case 0x39:
                cat[i] = 2;  break;             // integer data reg
            case 0x3C:
                cat[i] = 6;  break;             // predicate reg
            case 0x40:
                cat[i] = 7;  break;             // aggregate constant
            default:
                /* cat[i] left unset -> effective category 0 */
                break;
        }
        width[i] = bit_width_of(tok);           // sub_44B390
    }
}

Matcher Pseudocode

Once the category array is built, the matcher (same function, lines 250--1473) walks the descriptor candidates returned by the dual hash lookup against lexer_state+2472 and lexer_state+2480. The candidates are copied into a local v135[326] buffer with a running count v20, then filtered in phases: each phase tests a descriptor bit against a category predicate and zeroes non-matching entries in place, decrementing v22 (the live-candidate count). Surviving descriptors are compared operand-by-operand in a final pass.

Descriptor *match_instruction(
    LexerState *lex,        // a1
    const char *opcode,     // a2  -- opcode name for the hash probe
    /* ... */,
    OperandToken *ops[],    // a6
    int op_count,           // a8
    uint64_t *diag_lock)    // a10
{
    InsnTableCtx *ctx = lex->ctx;         // v12 = *(a1 + 1096)
    uint32_t cat[32], width[32];
    classify_operands(ops, op_count, cat, width);

    // --- hash lookup: both tables, linked-list concat into v135 ---
    Descriptor *list1 = hash_lookup(lex->tbl_2472, opcode);  // sub_426D60
    Descriptor *list2 = hash_lookup(lex->tbl_2480, opcode);
    if (!list1 && !list2) {
        emit_diag(dword_29FB550, diag_lock, parser_state);   // "unknown opcode"
        return NULL;
    }
    Descriptor *cand[326];
    int n = 0;
    for (Descriptor *p = list1; p; p = p->next) cand[n++] = p->payload;
    for (Descriptor *p = list2; p; p = p->next) cand[n++] = p->payload;
    int live = n;

    // --- Phase 1: opcode-class gate (v12+617 & 0x20) ---
    // Only descriptors whose byte at +22 has bit 1 set survive (line 348).
    if (ctx->byte_617 & 0x20) {
        for (int i = 0; i < n; i++)
            if (cand[i] && !(cand[i]->flags_22 & 2)) { cand[i] = NULL; --live; }
        if (!live) { emit_diag(sub_708200(ctx), diag_lock, ...); return NULL; }
    }

    // --- Phase 2: modifier-bit gates (v12+644, .ftz/.sat/.rnd set) ---
    uint32_t m = ctx->dword_644;
    if (m) {                                        // lines 390-505
        // per-descriptor byte-flag filter selected by (m & 2) and (m & 4),
        // reading either desc->byte_21 & 1, desc->byte_20 >> 7, or desc->byte_12 & 1
        // depending on whether ctx->dword_640 == 27 (FMA family special-case).
        filter_by_modifier_bits(cand, &live, m, ctx);
        if (!live) { emit_diag(sub_707530(ctx), diag_lock, ...); return NULL; }
    }

    // --- Phase 3..N: one filter per feature bit in ctx+600..+630 ---
    // Each phase maps 1:1 to a PTX modifier class. The complete list from
    // sub_46C6E0 (lines 510-930):
    //   ctx+628 & 0x40 -> desc->byte_25 & 1    (predicated)
    //   ctx+600 & 0x02 -> desc->byte_13 & 1    (wide/no-wide)
    //   ctx+600 & 0x80 -> desc->byte_13 & 0x10 or sub_4CE100()  (vector form)
    //   ctx+621 & 0x70 -> desc->byte_23 & 0x08 (cache-op variant)
    //   ctx+630 & 0x02 -> desc->byte_26 & 0x08 (async-copy group)
    //   ctx+629 & 0x80 -> desc->byte_26 & 0x02 (TMA / tensor-mem)
    //   ctx+612 & 0x08 -> desc->byte_18 & 0x80 (level qualifier)
    //   ctx+612 & 0x70 -> desc->byte_19 & 0x01 (scope qualifier)
    //   ctx+610 & 0x3C0 ->desc->byte_18 & 0x02 (ordering .relaxed/.acq/.rel)
    //   ctx+612 & 0x04 -> desc->byte_18 & 0x40 (mmu / tex level)
    //   ctx+620 & 0x38000->desc->byte_23 & 0x10(shared-memory variant)
    //   ctx+629 & 0x40 -> desc->byte_26 & 0x01 (dst-predicate)
    //   ctx+627 & 0x30 -> desc->byte_24 & 0x10 when AST kind 13 (reserved)
    //   ctx+611 & 0x30 -> desc->byte_18 & 0x08 (wmma/mma layout)
    //   ctx+612 & 0x80 -> desc->byte_19 & 0x02 (half-precision lane)
    //   ctx+613 & 0x03 -> desc->byte_19 & 0x04 (tensor-core accumulator)
    // Each phase that drops live to 0 emits a distinct diagnostic
    // (sub_707610, sub_707CE0, sub_70A180, sub_708860, sub_707B60, sub_70AFA0,
    //  sub_70B080, sub_70AAD0, sub_70AEF0, sub_70A0D0, sub_707AB0, sub_709860,
    //  sub_70ACC0, sub_70AB30, sub_70ABA0) so the user sees exactly which
    //  modifier family disqualified every candidate.

    // --- Phase N+1: operand-category comparison (line 940-1560) ---
    // v50 = cand[j].  v50[8] = descriptor's op_count.
    // v50[9..9+op_count-1] = expected category sequence.
    for (int j = 0; j < n; j++) {
        Descriptor *d = cand[j];
        if (!d) continue;
        if (d->op_count != op_count)      goto fail;
        if (!op_count)                    break;        // opcode-only match ok
        if (d->cat[0] != cat[0])          goto fail;    // line 950

        for (int k = 1; k < op_count; k++) {
            if (ops[k]->kind == 64) continue;           // skip aggregate wrapper
            // Per-slot detailed check: the descriptor slot at v50[2*k+24]
            // stores an "operand-check selector" (0..16). sub_1CB0820 dispatches
            // on it and on cat[k]/width[k], and also consults the type bits via
            // a big switch at lines 1009-1469:
            //   0: width compare (with optional %tid/%ntid/%ctaid special-case,
            //      line 1232: recognizes "%gridid" -> 32->64 widening)
            //   1: integer-only (via sub_457610 / sub_457490)
            //   2..7: fp / bit / vector subclass checks
            //   8,9,A,B: exact-width (8/16/32/64) constraints
            //   C: %tid / %laneid / %warpid / %smid special-register whitelist
            //   D,F: width == 2 (i.e. half-word) gate
            //   E: AST kind == 3 (register-triplet)
            //   10: integer-or-half via sub_457A00 guard
            //   11: only accept non-"standard" widths if ptx_major > 1 or (2, m1)
            //   12: AST kind == 4 required
            //   13: AST kind == 61 (identifier) or sub_457B60/sub_457B80 pass
            //   14: AST kind == 15 required (predicate)
            //   16: AST kind == 4 and sub_44A220()[0] == 0 (symbol undef)
            // Any failure "goto LABEL_153" zeros cand[j] and --live.
            if (!check_op_slot(d->slot[k], cat[k], width[k], ops[k]))
                goto fail;
            // Also require d->cat[k] == cat[k] (category sequence identity,
            // line 964).
            if (d->cat[k+1] != cat[k]) goto fail;
        }
        // Bonus suffix-check (lines 1474-1553): d->cat[9] onward is a
        // 16-slot ragged "trailing-modifier sequence". Ordered pair-scan --
        // if any (expected, present) disagrees and expected is not zero, drop.
        if (!trailing_modifier_match(d)) goto fail;
        continue;
    fail:
        cand[j] = NULL;
        --live;
    }

    // --- Ambiguity / failure reporting ---
    if (!live) {
        // None survived the operand-category pass.
        emit_diag(dword_29FB630, diag_lock, parser_state); // "no matching variant"
        return NULL;
    }

    // If a7 (the expected sm_target_code) is 0 and exactly one candidate
    // survives, return it -- the first non-null one (line 994: `if (!a7) return *v107;`).
    // Otherwise, iterate survivors and keep the one whose d->sm_target (at
    // offset +232) equals a7.  The PTX version guard at (a1+160 > 1 || (a1+164 > 2
    // && a1+160 == 1)) additionally excludes pre-ISA-1.x descriptors.
    Descriptor *hit = NULL;
    for (int j = 0; j < n; j++) {
        if (!cand[j]) continue;
        if (cand[j]->sm_target == a7) { hit = cand[j]; break; }
    }
    if (hit) return hit;

    // Survivors exist but none match the active SM target.
    emit_diag(dword_29FB640, diag_lock, parser_state); // "no variant for this sm_XX"
    return NULL;
}

Ambiguity Resolution

The matcher is not a pure "first match wins" scheme. When multiple descriptors survive every filter, disambiguation is by the SM target code stored at descriptor byte +232 (compared against a7, the compile target). If a7 == 0 (target-independent lookup, as during macro expansion), the matcher returns the first surviving descriptor unconditionally (line 994). If more than one descriptor survives and more than one has matching sm_target, the matcher still returns the first one encountered in list order -- there is no tie-breaking heuristic, so instruction-table registration order (which is deterministic because sub_46E000 registers in source order) is the silent arbiter. Truly-ambiguous encodings are prevented at table-build time rather than at parse time.

Failure Diagnostics

sub_46C6E0 emits five distinct diagnostic message IDs via sub_42FBA0(id, lock, parser_state):

Each per-phase emitter reports exactly which modifier family cost the match, which is why ptxas produces targeted messages like ".wide not valid for this instruction" rather than a generic "operand mismatch".

The classification examines token attributes set by the lexer. The bit tests mentioned in older wiki drafts ((field >> 28) & 7, 0x1000000, 0x6000000) live in the lexer (sub_44F2A0 and its callees) where the token-kind value is assigned -- the matcher itself only reads the already-encoded token kind at *v14.

Parser State Object (1,128 bytes)

The parser passes a state object through all phases. This 1,128-byte structure (sub_424070(pool, 1128)) carries compilation context and pointers to sub-systems. It is indexed as _QWORD* (8-byte slots), so QWORD index [N] = byte offset N*8. The highest accessed byte is +1120 (index [140]), fitting exactly within the 1,128-byte allocation.

Lexer State Object (2,528 bytes)

The lexer state is a separate heap-allocated object (sub_424070(pool, 2528)) pointed to by parser_state+1096. It is the primary state carrier for the Flex DFA scanner and the instruction table subsystem. All functions that need scanner state (the Bison parser, the Flex scanner, the include handler, and the instruction table builder) access this object through the pointer at +1096.

Version checks use sub_485520(ctx, sm_number) (SM architecture >= N) and sub_485570(ctx, major, minor) (PTX version >= major.minor). For example, the address-space attribute setter (sub_4035D3) checks sm_90 and PTX 7.8:

if (!sub_485520(ctx, 90))
    sub_42FBA0(&err, loc, "sm_90", ...);   // Error: requires sm_90
if (!sub_485570(ctx, 7, 8))
    sub_42FBA0(&err, loc, "7.8", ...);     // Error: requires PTX 7.8
*(byte*)(v15 + 632) = (old & 0xFC) | (a2 & 3);   // Set address space bits

Semantic Validators

The parser's reduction actions dispatch to specialized validator functions for each instruction category. These functions live in 0x460000--0x4D5000 and check SM architecture requirements, type compatibility, operand constraints, and instruction-specific invariants.

Validators follow a uniform pattern: they receive the parser context and instruction data, check constraints against the current SM architecture and PTX version, and call sub_42FBA0 with descriptive error messages when violations are found. The general validator (sub_4B2F20, 52.6 KB) is the second-largest function in the front-end and covers the broadest range of PTX instructions.

ROT13 Opcode Name Obfuscation

PTX opcode names stored in the binary are ROT13-encoded as an obfuscation measure. The static constructor ctor_003 at 0x4095D0 (17 KB, ~1,700 lines) decodes and populates the opcode name table at 0x29FE300 during program startup. Each entry is a (string_ptr, length) pair. Decoded examples:

The table covers the entire PTX ISA vocabulary -- hundreds of opcodes. A separate ROT13 table in ctor_005 (0x40D860, 80 KB) encodes 2,000+ internal Mercury/OCG tuning knob names (see Knobs System).

Compilation Pipeline Integration

The parser is invoked from the top-level compilation driver sub_446240 (11 KB), which orchestrates the full pipeline:

Parse  →  CompileUnitSetup  →  DAGgen  →  OCG  →  ELF  →  DebugInfo

The driver reports timing for each phase:

• "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%%)"

The parse phase encompasses the Flex scanner, macro preprocessor, Bison parser, instruction table lookup, and all semantic validation. Since the parser directly builds IR, the output of the parse phase is a populated instruction stream ready for the DAG generation phase.

PTX Text Generation (Reverse Direction)

The inverse of parsing -- converting IR back to PTX text -- lives in 0x4DA340--0x5A8E40 (580 formatter functions). Each handles one PTX opcode. A dispatcher at sub_5D4190 (12.9 KB) routes by opcode name using 81 direct string comparisons plus a 473-entry hash switch. Every formatter follows an identical allocation pattern:

pool = sub_4280C0(ctx)[3];              // Get allocator pool
buf = sub_424070(pool, 50000);          // 50KB temp buffer
// ... sprintf() operands into buf ...
len = strlen(buf);
result = sub_424070(pool, len + 1);     // Exact-size allocation
strcpy(result, buf);
sub_4248B0(buf);                        // Free temp buffer
return result;

A monolithic format string table (~1.8 MB) at the a2 parameter contains pre-assembled PTX text templates with %s/%llu/%d placeholders. This trades memory for speed: instead of building instruction text dynamically, ptxas simply fills in operand names at runtime.

Function Map

Cross-References

• Pipeline Overview -- where the parser fits in the compilation flow
• PTX Directive Handling -- detailed directive processing after parsing
• PTX-to-Ori Lowering -- what happens to the IR the parser builds
• Knobs System -- ROT13-encoded knob names from ctor_005
• Memory Pool Allocator -- sub_424070/sub_4248B0 pool system
• Hash Tables & Bitvectors -- sub_426150/sub_426D60 hash map
• PTX Instruction Table -- full opcode catalog
• CLI Options -- sub_432A00/sub_434320 option handling