Intrinsic Table Architecture (607 Registered Entries)

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

ptxas maintains two separate intrinsic subsystems that together cover every CUDA runtime helper function, every PTX opcode requiring inline code generation, and every Blackwell+ OCG builtin operation. The first subsystem (sub_5D1660 + sub_5D4190 + sub_5D7430 + sub_5FF700) handles 607 classical CUDA intrinsics and PTX opcode dispatch through a name-to-ID hash map, a body template name table, and a giant prototype generator. The second subsystem (sub_6C9EB0 and its handler cluster at 0x6C0000--0x6CC000) handles OCG (Optimized Code Generation) builtins for SM100+ targets. Both subsystems use the same hash map infrastructure (sub_425CA0 / sub_426150 / sub_426D60) documented in Hash Tables & Bitvectors.

Per-Family Deep Dives:

• OCG Intrinsic System -- SM100+ OCG builtins (44 operations), lowering pipeline, SASS handler map
• Math Intrinsics -- IEEE math software emulation (div, rcp, sqrt, rem)
• Tensor Core Intrinsics -- WMMA, MMA, WGMMA, tcgen05 lowering
• Sync & Warp Intrinsics -- Barriers, vote, shuffle, match, redux

System Overview

sub_451730 (intrinsic lowering context constructor)
  │
  ├── sub_5D4190(ctx)  ── register PTX opcode & MMA handlers ──────────────┐
  │     │ (1) Calls sub_5D1660 to populate intrinsic ID table (607 entries) │
  │     │ (2) Registers ~120 PTX opcode -> codegen handler mappings         │
  │     │ (3) Registers ~400 MMA hash -> codegen handler mappings           │
  │     │                                                                   │
  │     ├─ Hash map at +808  ── PTX opcode name -> codegen function ptr     │
  │     │    "div"     -> sub_5B76D0  (64KB)                                │
  │     │    "sqrt"    -> sub_5B4040  (49KB)                                │
  │     │    "wmma.mma"-> sub_5C7A50  (173KB)                               │
  │     │    "mma"     -> sub_5C10A0  (120KB)                               │
  │     │    ... ~116 more                                                  │
  │     │                                                                   │
  │     ├─ Hash map at +816  ── numeric MMA hash -> codegen handler ptr     │
  │     │    "2644314910" -> sub_4DDB80                                     │
  │     │    ... ~399 more (shape/type/layout combinations)                 │
  │     │                                                                   │
  │     └─ ID table at +1056 ── 9728-byte array (memcpy from unk_1D4D940)  │
  │        Hash map at +1064 ── name -> integer ID (sub_5D1660, 607)        │
  │        Count at +1072 = 608 (includes null ID 0 slot)                   │
  │                                                                         │
  ├── sub_4CE230(ctx)  ── register modifier keywords (GUARD, PRED, ...)     │
  │                                                                         │
  ├── sub_5D7430(ctx, sregs)  ── body template name table (161KB) ──────────┤
  │     │ 1,079 entries, each constructed from:                             │
  │     │   16-byte .rodata prefix (e.g. "__cuda_sm20_div_")               │
  │     │ + 4-byte type suffix (e.g. "s16\0", "u64\0", "rn_f")            │
  │     │ → registered into hash map at +824 with sequential integer IDs    │
  │     │                                                                   │
  │     └─ Hash map at +824  ── intrinsic name -> body template ID          │
  │          "__cuda_sm20_div_s16" -> 0                                     │
  │          "__cuda_sm20_div_u16" -> 1                                     │
  │          ... 1,079 total entries                                        │
  │                                                                         │
  └── sub_451330("<fermi macros>", ...)  ── load Fermi macro library        │
                                                                            │
sub_5FF700 (354KB) ─────────────────────────────────────────────────────────┘
  │ switch(body_template_id) with hundreds of cases
  │ Each case: allocate buffer via sub_4DA340, strcpy() PTX prototype
  │
  │ case 0:  ".weak .func (.reg .s32 %d) __cuda_sm20_div_s16
  │           (.reg .s32 %a0, .reg .s32 %a1)"
  │ case 4:  ".weak .func (.reg .u64 %rdv1) __cuda_sm20_div_u64
  │           (.reg .u64 %rda1, .reg .u64 %rda2)"
  │ case 9:  ".weak .func (.reg .f32 %fv1) __cuda_sm20_div_rn_f32
  │           (.reg .f32 %fa1, .reg .f32 %fa2)"
  │ case 25: ".weak .func (.reg .f64 %fdv1) __cuda_sm20_div_rn_f64_full
  │           (...)"
  │ ... hundreds more for rcp, sqrt, dsqrt, barrier, wmma, mma, etc.
  v
Emitted into PTX output as .weak .func declarations
(linker resolves calls to runtime helper functions)

Master Registration -- sub_5D1660

This 46KB function is the master catalog. It allocates a 9728-byte table (memcpy from unk_1D4D940, 0x2600 bytes = 608 x 16B slots), creates a hash map with initial capacity 0x80 via sub_425CA0, then calls sub_426150(hashmap, "name", (char*)ID) exactly 607 times to register every CUDA runtime helper function with an integer ID (IDs 1--607, contiguous). The hash map is stored at a1+1064, the table at a1+1056, and the count 608 at a1+1072 (includes the unused null ID 0 slot).

Complete ID Allocation

607 intrinsics are registered with contiguous IDs from 0x01 through 0x25F. The binary stores count=608 at a1+1072 because the pre-built 9,728-byte table (608 x 16B slots) includes a null ID 0 sentinel. The ID ranges partition cleanly by SM generation and functional category.

Total: 607 registered intrinsics across 13 prefix groups. Table has 608 slots (ID 0 unused).

sm_70 Intrinsic Breakdown (IDs 0x89--0x211)

The sm_70 block is by far the largest at 393 entries. It covers every Volta-era warp synchronous intrinsic plus the complete WMMA API. The explosion in count comes from the combinatorial product of shapes, layouts, data types, address spaces, and predicate/satfinite variants.

The WMMA entries dominate the count. Each combination of shape (m16n16k16/m32n8k16/m8n32k16), operation (load_a/load_b/load_c/store_d/mma), layout (row/col for each matrix), data type (f16/f32), address space (generic/global/shared), and optional satfinite flag produces a separate intrinsic registration.

Opcode Dispatch -- sub_5D4190

This 41KB function first calls sub_5D1660(a1) to populate the intrinsic ID table, then builds two more hash maps for PTX opcode dispatch.

Named Opcode Table (at a1+808)

~120 PTX instruction names mapped to codegen handler function pointers. Each handler allocates a 50,000-byte buffer, queries instruction properties through accessor functions on the instruction object at a1+1096, and generates inline PTX code via sequential sprintf() calls.

Numeric MMA Hash Table (at a1+816)

~400 entries where the key is a numeric string representation of a hash value (e.g., "2644314910") that encodes a specific MMA shape/type/layout combination. The hash encodes the instruction variant completely: matrix dimensions (m16n8k16, m16n8k32, etc.), data type (f16, bf16, tf32, f32, f64, s8, u8, s4, u4, b1), and layout (row/col combinations). Each entry maps to a codegen handler function pointer. This avoids a multi-dimensional lookup by collapsing the full variant space into a single hash probe.

Body Template Name Table -- sub_5D7430

At 161KB of machine code (0x5D7430--0x5FF700), this is the largest function in the intrinsic infrastructure by code size and the 6th largest function in the entire ptxas binary. IDA failed to decompile it; all analysis comes from raw x86-64 disassembly. The function constructs a third hash map (at context offset +824 / 0x338) containing 1,079 entries that map dynamically constructed __cuda_* intrinsic names to sequential body template IDs (0--1078).

Why 1,079 Body Templates for 607 Logical Intrinsics

The master registration table (sub_5D1660) maps 607 intrinsic names to logical IDs. The body template table (sub_5D7430) maps 1,079 variant-specific names to prototype generator case numbers. The 1.78x expansion has one dominant cause: WMMA template proliferation across GPU generations.

The 204 logical WMMA entries in sub_5D1660 cover only the original sm_70 Volta shapes (m16n16k16/m32n8k16/m8n32k16 with f16/f32 types). But the body template table includes all later-generation WMMA variants -- sm7x sub-byte/bit, sm72 integer, sm8x tf32/bf16/f64 -- that were added as hardware evolved. These ~416 extra WMMA templates have no matching entry in the 607 logical ID table; they exist only in the body template hash map and the prototype generator switch.

Non-WMMA intrinsics map approximately 1:1 between logical IDs and body templates. The math operations (div, rcp, sqrt) are already fully type-specialized at the logical level -- each rounding-mode/type combination is a separate logical intrinsic.

Three sources of expansion beyond the 607 logical entries:

1. Later-generation WMMA variants (~416 template-only entries):
   • sm7x sub-byte WMMA (s4/u4 m8n8k32) + bit WMMA (m8n8k128): ~231 templates
   • sm72 integer WMMA (m16n16k16/m32n8k16/m8n32k16 integer types): ~105 templates
   • sm8x tf32 WMMA (m16n16k8) + bf16/f64 WMMA: ~80 templates
2. Aligned warp sync variants (~13 extra templates): matchsync_aligned, votesync_aligned, votesync_ballot_groupwise, query_activemask/query_activemask_groupwise for cooperative group support
3. Additional SM100 specializations (~8 extra templates): tcgen05_alloc_two_sm, extra guardrails check variants, get_warp_rank

Conversely, 18 sm1xx bulk copy intrinsics have logical IDs but zero body templates -- they bypass the template/prototype mechanism entirely and are lowered directly to inline PTX by the opcode dispatch handlers (sub_593210, sub_5AB460).

Template Distribution Table

WMMA subtotal: 204 logical entries expand to 665 body templates (3.3x). Non-WMMA: 403 logical entries map to 408 templates (~1.0x). The remaining 7 templates (1,080 prototype switch cases minus 1,073 classified) are sanitizer/cache variants where IDA produced qmemcpy instead of strcpy, preventing exact name extraction.

The sm70 WMMA group itself expands from 204 to 249 templates because the prototype generator includes update_ptr and desc (descriptor-based addressing) variants of certain load/store operations that the logical table does not separate.

The three "tmpl-only" WMMA rows (sm7x/sm72/sm8x) are the single largest contributor to the expansion. They represent ~416 templates with zero logical ID counterparts. These families use .FORCE_INLINE .func linkage in their prototypes instead of the .weak .func used by the original sm70 WMMA entries:

sm70 (original):   .weak .func (...) __cuda_sm70_wmma_m16n16k16_load_a_col (...)
sm72 (integer):    .FORCE_INLINE .func (...) __cuda_sm72_Integer_wmma_m16n16k16_load_a_row (...)
sm7x (sub-byte):   .FORCE_INLINE .func (...) __cuda_sm7x_sub_byte_wmma_m8n8k32_load_a_row (...)
sm8x (tf32):       .FORCE_INLINE .func (...) __cuda_sm8x_tf32_wmma_m16n16k8_load_a_row (...)

The .FORCE_INLINE directive forces inlining at every call site rather than emitting a separate callable function. The later-gen WMMA implementations are more complex and performance-sensitive, making per-call-site specialization profitable.

Name Construction Algorithm

The function contains zero string references because it constructs all 1,079 names at runtime. For each entry:

1. Allocate a 20-byte buffer via sub_424070(allocator, 20)
2. Copy prefix (16 bytes) from .rodata via SSE movdqa + movups (e.g., "__cuda_sm20_div_")
3. Append suffix (4 bytes) via movl immediate at offset +16 (e.g., "s16\0", "u64\0", "rn_f")
4. Register via sub_426150(context+824, buffer, template_id) with sequential integer IDs

The 533 unique .rodata prefix addresses fan out through multiple suffixes per prefix:

.rodata prefix (16B)       suffix (4B)     result (20B buffer)
───────────────────────    ───────────     ──────────────────────
"__cuda_sm20_div_"    +    "s16\0"    =   "__cuda_sm20_div_s16"
"__cuda_sm20_div_"    +    "u16\0"    =   "__cuda_sm20_div_u16"
"__cuda_sm20_div_"    +    "u64\0"    =   "__cuda_sm20_div_u64"
"__cuda_sm20_div_"    +    "s64\0"    =   "__cuda_sm20_div_s64"
"__cuda_sm20_div_"    +    "rn_f"     =   "__cuda_sm20_div_rn_f" (truncated)
"__cuda_sm20_rem_"    +    "s16\0"    =   "__cuda_sm20_rem_s16"
"__cuda_sm20_rcp_"    +    "rn_f"     =   "__cuda_sm20_rcp_rn_f" (truncated)
"__cuda_sm70_barr"    +    "ier_"     =   "__cuda_sm70_barrier_" (prefix chain)

Names truncated at the 20-byte buffer limit are still sufficient for hash map lookup -- the full untruncated name appears only inside the prototype string in sub_5FF700.

Worked Example: Division (Cases 0--26)

The __cuda_sm20_div operation illustrates the template-to-prototype mapping. Division has 19 logical IDs and 19 body templates (1:1 ratio) because each type/rounding/precision variant is already a separate logical intrinsic. The suffix encodes the type specialization:

Cases 2--3 (rem s16/u16), 6--7 (rem s64/u64) are interleaved between the division entries. Cases 8, 13 are _slowpath variants that implement Newton-Raphson refinement fallbacks. Cases 18--21 are the sm3x-optimized division variants with the same suffix scheme. Note: s16/u16 division uses .s32/.u32 register types because PTX has no 16-bit register class; the 16-bit operation is performed by 32-bit hardware with appropriate sign/zero extension.

Statistics

Context Hash Map Summary

The intrinsic lowering context object holds five hash maps and one flat table:

Instruction Property Accessors

All codegen handlers query instruction properties through accessor functions on the instruction object at a1+1096. These are the same accessors used by WMMA, MMA, and tcgen05 codegen.

Prototype Generator -- sub_5FF700

At 354KB, this is the single largest function in the intrinsic infrastructure and the 2nd largest function in the entire ptxas binary. It takes a body template ID (a1, range 0--1079) and an allocator context (a2), allocates a buffer via sub_4DA340(size, a2), fills it with a PTX prototype string via strcpy(), and returns the result. The output is a complete .weak .func or .FORCE_INLINE .func PTX declaration that gets emitted into the PTX output stream so the linker can resolve calls to CUDA runtime helper functions.

The function is a single switch(a1) with 1,080 case labels (0--1079) plus a default case that returns an empty string "". Each case allocates an exact-sized buffer (72--1,200 bytes), copies a hardcoded PTX prototype string into it, and returns the pointer.

Prototype Generator Architecture

sub_5FF700(template_id, allocator)
  │
  │  switch(template_id)     ← 1,080 cases, 0--1079
  │
  ├── case N:
  │     buf = sub_4DA340(byte_count, allocator)    ← allocate exact-fit buffer
  │     strcpy(buf, ".weak .func (...) name (...)")  ← copy PTX prototype
  │     return buf
  │
  ├── case M:  (45 large WMMA mma cases)
  │     buf = sub_4DA340(byte_count, allocator)    ← up to 1,200 bytes
  │     *(u64*)buf = 0x662E206B6165772E            ← ".weak .f" (inline store)
  │     *(u64*)(buf+N-8) = <trailer>               ← last 8 bytes inline
  │     qmemcpy(buf+8, .rodata_addr, size-16)      ← bulk copy middle
  │     return buf
  │
  └── default:
        return ""

Three copy strategies appear in the decompilation, all producing the same result:

The qmemcpy cases are the 45 WMMA mma operations with the largest parameter lists (3--4 fragment matrices of 8 elements each). IDA stores the first and last 8 bytes as inline immediates (0x662E206B6165772E = ".weak .f", trailer varies per case) and bulk-copies the middle from .rodata. The prototype content is structurally identical to the strcpy cases.

Linkage Directives

Two PTX linkage types are emitted, controlling how the linker handles the declared function:

The .weak linkage supports user-supplied replacements: if the user provides their own implementation of __cuda_sm20_div_s16, the linker will use that instead of the built-in runtime version. The .FORCE_INLINE directive forces per-call-site specialization -- the later-generation WMMA implementations are more complex and performance-sensitive, making inlining profitable.

A subset of .weak prototypes (~410) carry the .unique qualifier:

.weak .func (.reg .b32 dst) __cuda_sm70_barrier_sync (.reg .b32 arg0) .unique ;

.unique instructs the PTX linker to keep exactly one copy of the function body even if multiple compilation units reference it. All barriers, redux sync, warpsync, non-aligned vote/match/shuffle use .unique.

Prototype Format

Every emitted prototype follows one of these structural patterns:

<linkage> .func (<return_params>) <name> (<input_params>) [.unique] ;

Parameter Passing Conventions

Five distinct parameter-passing ABIs appear across the 1,080 prototypes:

Convention A -- Register-only (.reg): Used by math operations, barriers, warp sync, redux sync, video emulation. Return and input parameters are individual .reg declarations with typed names. This is the simplest and most common convention.

.weak .func (.reg .f32 %fv1) __cuda_sm20_div_rn_f32 (.reg .f32 %fa1, .reg .f32 %fa2) ;

Convention B -- Param-array with alignment (.param .align N .b32 name[K]): Used by WMMA load/mma, MMA, Hopper sub-byte MMA, Blackwell MMA. Returns an aligned array of .b32 elements. Array sizes: dst[2], dst[3], dst[4], dst[5], dst[8], mma_dst[2], mma_dst[4], mma_dst[8], ret_dst[3], ret_dst[5]. 326 prototypes use .align 16; 1 prototype (mma_shfl_f16) uses .align 8.

.weak .func (.param .align 16 .b32 d[8]) __cuda_sm70_wmma_m16n16k16_mma_row_col_f32_f32
  (.param .align 16 .b32 a[8], .param .align 16 .b32 b[8], .param .align 16 .b32 c[8]) ;

Convention C -- Param-scalar (.param .b64): Used exclusively by the 7 compute-sanitizer hooks. Parameters use fully-qualified names (__cuda_sanitizer_memcheck_malloc_param_0).

.weak .func (.param .b64 func_retval0) __cuda_sanitizer_memcheck_malloc
  (.param .b64 __cuda_sanitizer_memcheck_malloc_param_0,
   .param .b64 __cuda_sanitizer_memcheck_malloc_param_1) ;

Convention D -- Void return (): Used by WMMA store_d, tcgen05 guardrail traps, sanitizer_free. ~140 prototypes (45 .weak + 95 .FORCE_INLINE).

.weak .func () __cuda_sm70_wmma_m16n16k16_store_d_row_f32
  (.reg .b64 ptr, .reg .b32 ldm, .reg .b32 sreg0, .reg .b32 sreg1, ...) ;

Convention E -- Multi-register return (.FORCE_INLINE only): Used by extended WMMA load operations (SM7x/SM72/SM8x). Returns 1--4 registers in the return position (never 8 -- 8-element returns use Convention B's .param arrays instead).

.FORCE_INLINE .func (.reg .b32 dst0, .reg .b32 dst1, .reg .b32 dst2, .reg .b32 dst3)
  __cuda_sm8x_tf32_wmma_m16n16k8_load_a_row (.reg .u64 ptr, .reg .u32 ldm) ;

PTX Register Types

Eight PTX register types appear across the prototypes:

Note: .b32 is used instead of .s32/.u32/.f32 for operations where the type interpretation is determined by the instruction rather than the register declaration (WMMA fragments, MMA accumulators, barrier IDs). The .s32 type appears only in the 4 oldest SM20 div/rem_s16/u16 prototypes (cases 0--3).

Register Naming Convention

The prototype register names encode the data type and role:

Buffer Allocation Sizes

sub_4DA340(size, allocator) allocates an exact-fit buffer per prototype:

The 45 qmemcpy cases have the largest buffers: 386--1,200 bytes. These are WMMA mma operations whose prototypes enumerate all 3--4 fragment matrices (a, b, c, d) with 4--8 elements each, producing prototype strings that exceed 900 bytes.

Case Range Layout

The 1,080 cases follow the body template registration order from sub_5D7430, roughly grouped by SM generation:

Statistics

Major Codegen Handlers

The four largest codegen handlers together represent ~500KB of code and cover the tensor core instruction families.

sub_5C7A50 -- WMMA.MMA Codegen (173KB)

The largest codegen handler. Generates inline PTX code for wmma.mma instructions across all variant combinations.

• Allocates a 50,000-byte buffer for code generation
• Covers shapes: m16n16k16, m32n8k16, m8n32k16
• Data types: f16, f32, bf16, tf32, s8, u8, s4, u4, b1
• Layouts: row/col for each of the A, B, C, D matrices (4 layout combinations)
• Satfinite variants for each configuration
• Address spaces: generic, global, shared

sub_5C10A0 -- MMA Codegen (120KB)

Handles the newer mma.sync API (non-WMMA). Covers the post-Volta PTX MMA instructions.

• Shapes: m8n8k4, m16n8k8, m16n8k16, m16n8k32, m16n8k64, m16n8k128, m16n8k256
• Types: f16, bf16, tf32, f32, f64, s8, u8, s4, u4, b1
• Sparse variants for sm_80+ and sm_90+ (structured sparsity 2:4)

sub_5BBC30 -- TCGen05.MMA Codegen (90KB)

Blackwell 5th-generation tensor core MMA code generation. Handles the tcgen05.mma instruction family introduced in sm_100.

• Allocates a 50,000-byte buffer
• Queries sub_70FA00(*, 29) to validate tcgen05 capability
• Handles standard, sparse, and warp-shared (.ws) variants
• Uses sub_70F0A0 for sparse metadata parameter extraction
• Generates code for tcgen05-specific tensor memory addressing

sub_5B76D0 -- Division Codegen (64KB)

Generates inline PTX code for all div variants.

• Integer division: s16, s64, u16, u64
• Floating-point division: f32, f64 with all rounding modes (rn, rd, ru, rz)
• Flush-to-zero (ftz) variants for f32
• Checks operand type via sub_70CA60(*(_QWORD *)(a1+1096), 0) == 21
• Emits both fastpath and slowpath (Newton-Raphson) code sequences

OCG Intrinsic System -- sub_6C9EB0

The OCG (Optimized Code Generation) intrinsic subsystem is a separate, parallel dispatch mechanism for SM100+ builtin operations. While the classical system at sub_5D1660 maps CUDA runtime helper names to integer IDs and generates inline PTX, the OCG system maps __nv_ptx_builtin_ocg_* function names to type-specific handlers that validate parameters and emit SASS instructions directly -- bypassing PTX entirely. The OCG table contains 44 operations across 9 categories: arithmetic, packed float, vector integer, async copy/TMA, load/store/cache, reduction/fence, tensor core, tensor memory, and synchronization.

See OCG Intrinsic System (44 Operations) for the complete builtin name table, handler functions, validation strings, SASS-level handlers, and the full five-stage lowering pipeline with operand buffer layout.

Intrinsic Families by SM Generation

Each SM generation introduces new intrinsic families while preserving all earlier ones. The per-SM intrinsic table initializer functions (sub_60AXXX cluster, registered in Map 3 of the capability dispatch) control which intrinsics are available on each target.

sm_20 -- Software IEEE Math (70 entries)

The foundation layer. 70 intrinsics providing IEEE-754-compliant software implementations of math operations that either lack hardware support or need exact rounding guarantees. All later SM targets inherit these.

• Division: div_s16, div_u64, div_rn_f32, div_rn_f64_full, etc. -- all rounding modes (rn/rd/ru/rz) and types (s16/s64/u16/u64/f32/f64)
• Reciprocal: rcp_rn_f32, rcp_rn_f64, etc. -- all rounding modes
• Square root: sqrt_rn_f32, sqrt_rn_f64, etc. -- all rounding modes
• Double-precision sqrt: dsqrt_rn, dsqrt_rd, dsqrt_ru, dsqrt_rz
• Double-precision reciprocal sqrt: drsqrt_rn
• Bit extract/insert: bfe (bit field extract), bfi (bit field insert)
• Remainder: rem_s32, rem_u32, rem_s64, rem_u64

Codegen handlers: sub_5B76D0 (div, 64KB), sub_5B0CD0 (rcp, 44KB), sub_5B4040 (sqrt, 49KB).

sm_3x -- Optimized Division (4 entries)

Four optimized division variants introduced on Kepler to improve throughput on common division patterns.

sm_62 -- Integer Dot Product (2 entries)

dp2a and dp4a integer dot product intrinsics introduced on Pascal (GP10x). Software emulation of the hardware instructions added in sm_61/sm_62.

sm_70 -- Volta Warp-Synchronous + WMMA (393 entries)

The largest single block. Volta introduced mandatory warp-synchronous programming with explicit sync masks and the first generation of tensor core (WMMA) instructions.

Synchronization primitives:

• barrier_arrive / barrier_sync / barrier_red (0--15, with/without count)
• matchsync_all/any_b32/b64 with predicate variants
• shflsync_bfly/down/idx/up with predicate variants
• votesync_all/any/ballot/uni
• warpsync

WMMA (Warp Matrix Multiply-Accumulate):

• Shapes: m16n16k16, m32n8k16, m8n32k16
• Operations per shape: load_a, load_b, load_c, store_d, mma
• Layouts: row/col combinations for A and B matrices
• Types: f16, f32 (with satfinite optional)
• Address spaces: generic, global, shared

sm_80 -- Ampere Cache Policy (3 entries)

Three createpolicy intrinsics for L2 cache management: createpolicy_fractional, createpolicy_fractional_encode, createpolicy_range_encode.

sm_10x -- Blackwell MMA Metadata + Bit MMA (10 entries)

10 hmma_mdata/imma_mdata + bit MMA intrinsics for sm_100: metadata variants at m16n8k16/k32/k64 shapes, and 1-bit AND/XOR MMA at m8n8k128/m16n8k128/m16n8k256.

sm_8x -- Direct MMA (14 entries)

14 mma_* intrinsics for sm_8x: 12 direct MMA operations (f16/f32 accumulator x 4 layout combinations of col/row for A and B) plus mma_shfl_f16 and mma_shfl_f32 for register-to-register MMA shuffle.

sm_9x -- Sub-Byte + Bit MMA (51 entries)

51 Hopper-era intrinsics: 3 bit-XOR MMA (m8n8k128/m16n8k128/m16n8k256), 24 dense sub-byte MMA (s4/u4 at m16n8k32/m16n8k64/m8n8k32, with satfinite), 8 sparse m16n8k128, and 16 sparse m16n8k64 (with _0/_1 split variants and satfinite).

sm_10x (via __cuda_sm10x_*) -- Blackwell Tensor Memory + Guardrails (11 entries)

• 1 create_mask_from_bit_idx_and_alloc_size_v1 helper
• 10 tcgen05_guardrail_trap_* intrinsics for debug validation of tensor memory operations

sm_1xx -- Bulk Copy (18 entries)

18 bulk copy and cp.async.bulk.tensor intrinsics covering 1D through 5D tensor copies with tile and im2col addressing modes, both unicast and multicast variants.

Intrinsic Lookup Flow

The lookup path from a function call in PTX source to the codegen handler follows this sequence:

PTX source: call.uni __cuda_sm70_warpsync, (%mask);
                    |
                    v
            sub_5D1660 hash map (a1+1064)
            key: "__cuda_sm70_warpsync"
            value: integer ID (within 0x89..0x211 range)
                    |
                    v
            sub_5FF700 switch(ID)
            Emits: .weak .func __cuda_sm70_warpsync (.reg .u32 %a0)
                    |
                    v
            sub_5D4190 named opcode hash map (a1+808)
            key: PTX opcode (e.g., "shfl", "vote", "barrier")
            value: codegen handler function pointer
                    |
                    v
            Codegen handler (e.g., sub_5801D0 for "shfl")
            Queries instruction properties via sub_70XXXX accessors
            Generates inline PTX code into 50KB buffer

For OCG intrinsics on SM100+, the path bypasses PTX entirely: sub_6A97B0 matches call nodes to SASS instructions via an RB-tree, sub_6C9BC0 parses the __nv_ptx_builtin_ocg_* name into an operation enum + sub-op array, sub_6CC690 routes to type-specific handlers and assembles operands, and sub_6CB8A0 emits the final SASS instruction. See the OCG Intrinsic System page for the full five-stage pipeline breakdown with operand buffer layout and internal SASS opcode enum values.

Per-SM Intrinsic Initializers

Each SM target has its own intrinsic table initializer function registered in Map 3 of the capability dispatch (sub_607DB0). These functions control which subset of the 607 intrinsics are available on each target.

Sub-variants (e.g., sm_100a, sm_100f) share the same initializer as their base SM since they represent the same silicon with different feature exposure levels.

Instruction Description Loader -- sub_9EE390

sub_9EE390 (3,584 bytes, 0x9EE390--0x9EF190) is the constructor for an instruction description object that feeds the register allocator's pre-coloring pass. Despite the diagnostic string "IntrinsicDescrFile=%s", the function loads instruction descriptions broadly -- not just intrinsic operations. It determines which instructions exist for the target SM, what register classes they use, and what scheduling properties apply. The sole caller is sub_991790 (pre-coloring pass, 12KB).

Invocation pattern: The pre-coloring pass checks context+1936 before calling sub_9EE390. If the descriptor for the current SM class already exists, it is reused. This means the expensive initialization happens once per SM architecture per ptxas process lifetime.

Initialization Sequence

1. Extract target properties. Read the target descriptor from context+1584. Compute the SM architecture class: v111 = target_descriptor[+372] >> 12. Read resource descriptors from option interface slots 41--44.

2. Check option 404 (IntrinsicDescrFile). Query the option interface at context[208]. If option 404 is set, extract the file path and log " IntrinsicDescrFile=%s". This CI-internal mechanism supplies an external description file that overrides or extends the built-in instruction table. When absent, the built-in database is used.

3. Determine instruction format class. Call sub_7DDB50(context) (GetSmVersionIndex), subtract 1, index into dword_21E6330:

   sub_7DDB50 return	v114	Format
   1	0	basic 64-bit
   2	1	128-bit
   3	3	extended
   4	2	192-bit
   5+	3	extended (default)

4. Determine SM generation class. Read context+12 (sm_version_id), subtract 1, index into dword_21E5C80. The table is an identity mapping (1--11), one entry per SM generation.

5. Construct instruction table (648 bytes). Call sub_10AFF80 with 32 parameters including memory pool, register count, format class, description file path, architecture descriptor (16 bytes from context+1888), SM generation class, instruction count limits, and context flags. Follow with sub_10B1A90 (init pass 2) and sub_10AEF10 (finalization).

6. Apply option overrides. Options 497, 738, 739 from the option interface set register limits and allocation budget values on the instruction table sub-object at +312.

7. Select SM-specific instruction set descriptor. Based on v111 (SM architecture class):

   v111	SM range (inferred)	Alloc size	Constructor	Vtable
   5	sm_50--sm_62	200 B	sub_9CDF90	off_23F3B00
   6	sm_70--sm_75	216 B	sub_9CE030	off_22BB738
   7	sm_80--sm_89	232 B	sub_9CE120	off_22B5150
   8+	sm_90--sm_121	240 B	sub_9CE190	off_22AD230
   <5	(reuse existing)	--	--	--

   Each successor inherits the previous class and extends it with generation-specific instructions. The descriptor is stored at context+1936 and this+48.

Object Layout

Diagnostic Strings

Function Map

Cross-References

• OCG Intrinsic System -- SM100+ OCG builtin table (44 operations), lowering pipeline, SASS handler map
• SM Architecture Map -- Per-SM capability dispatch tables and intrinsic initializer assignments
• Math Intrinsics -- Detailed coverage of sm_20 IEEE math intrinsic codegen (div, rcp, sqrt, rem)
• Tensor Core Intrinsics -- WMMA, MMA, WGMMA, tcgen05 instruction lowering
• Sync & Warp Intrinsics -- Barrier, vote, shuffle, match, redux intrinsics
• Newton-Raphson Templates -- Software math slowpath sequences used by div/rcp/sqrt
• TCGen05 -- 5th Gen Tensor Cores -- Blackwell tensor core ISA detail
• Hash Tables & Bitvectors -- Hash map infrastructure (sub_425CA0 / sub_426150 / sub_426D60)
• Mercury Encoder -- Master SASS encoder sub_6D9690 (94KB) that encodes validated intrinsics
• SASS Instruction Encoding -- Instruction encoding infrastructure
• Pipeline Overview -- OCG-time measurement covers intrinsic lowering

Appendix: Complete Intrinsic Name Catalog (607 Entries)

Every intrinsic registered by sub_5D1660, extracted from the decompiled source. IDs are contiguous 1--607 (0x001--0x25F). The suffix after stripping the prefix encodes the operation, data type, rounding mode, address space, and optional modifiers.

__cuda_reduxsync_* -- Redux sync (17 entries, 0x001--0x011, sm_70)

__cuda_sanitizer_memcheck_* -- Compute-sanitizer hooks (7 entries, 0x012--0x018, --)

__cuda_scalar_video_emulation_* -- Video emulation (7 entries, 0x019--0x01F, sm_20)

__cuda_sm10x_* -- Blackwell tcgen05 guardrails + mask (11 entries, 0x020--0x02A, sm_100)

__cuda_sm1xx_* -- Bulk copy + cp.async.bulk.tensor (18 entries, 0x02B--0x03C, sm_100+)

__cuda_sm20_* -- IEEE math (70 entries, 0x03D--0x082, sm_20)

__cuda_sm3x_* -- Optimized division (4 entries, 0x083--0x086, sm_30)

__cuda_sm62_* -- Integer dot product (2 entries, 0x087--0x088, sm_62)

__cuda_sm70_* -- Volta sync/warp/WMMA (393 entries, 0x089--0x211, sm_70)

__cuda_sm80_* -- Ampere createpolicy (3 entries, 0x212--0x214, sm_80)

__cuda_sm_10x_* -- Blackwell hmma/imma/bit MMA (10 entries, 0x215--0x21E, sm_100)

__cuda_sm_8x_* -- Direct MMA + shfl (14 entries, 0x21F--0x22C, sm_80+)

__cuda_sm_9x_* -- Hopper sub-byte/bit MMA (51 entries, 0x22D--0x25F, sm_90)