SM100 Blackwell

The SM100 architecture (Blackwell, datacenter) represents the largest single-ISA definition embedded in nvlink v13.0.88. The instruction encoding infrastructure spans approximately 8 MB of .text across four major regions, defines 3,200+ encoding/decoding functions, 118 instruction families organized under 3 major opcodes, and introduces SM100-specific features including tcgen05 tensor intrinsics, unified tensor core (UTC) instructions, and a new ROT13-obfuscated mnemonic set internally codenamed "MERCURY" (ZREPHEL in ROT13).

This page documents the complete ISA encoding layer as reconstructed from the embedded ptxas backend, covering the 128-bit instruction word format, the opcode space partitioning, the encoder/decoder/descriptor infrastructure, and all SM100-specific instruction families.

Sub-variant Matrix

nvlink registers three SM100 sub-variants through the architecture dispatch at sub_15C0CE0:

All three sub-variants share the same encoding tables. The a suffix enables the full feature set; the f suffix restricts to the forward-compatible subset. The architecture dispatch registers 7 callbacks per sub-variant (nv.info emitter, resource usage table, instruction encoding table, compute capability array, perf-stats handler, cpf_optx handler, codegen options).

Instruction Word Format

All SM100 SASS instructions are 128 bits (16 bytes) wide. The instruction word is stored as two 64-bit halves at offsets +544 (bits 0-63) and +552 (bits 64-127) in the internal instruction representation structure.

Fixed Opcode Fields (Bits 0-31)

The first 32 bits encode the instruction identity through five fields set by the first five calls to the bitfield insertion primitive sub_4C28B0:

127                              64 63                                0
+----------------------------------+----------------------------------+
|     Modifier / Operand Fields    |  Fmt2 | Mod | SubOp | Minor | EF | MajOp |
+----------------------------------+----------------------------------+

Bit Range    Width   Name              Description
---------    -----   ----              -----------
  [3:0]      4 bits  Major Opcode      Instruction format class (1, 2, or 3)
  [6:4]      3 bits  Encoding Format    Sub-format / encoding variant (usually 0)
  [16:8]     9 bits  Minor Opcode      Instruction family (118 unique values)
  [24:17]    8 bits  Sub-Opcode        Specific instruction within family
  [31:25]    7 bits  Modifier           Data type, addressing mode, operation variant

Major Opcode Distribution

Encoding Formats

Operand Field Encoding

Bits 32-127 encode operand references, modifier flags, and immediates. The layout varies by instruction class but follows consistent patterns:

Register Operand Types

The register encoder sub_4C4D60 writes a structured operand field at the specified bit offset:

The decoder sub_4C60F0 uses a register class parameter: 2=GPR, 3=predicate, 9=constant buffer reference, 10=uniform register.

Encoding Infrastructure

The SM100 ISA definition is split across four binary regions, totaling ~8 MB of template-instantiated functions:

Region Map

Combined totals: 1,975 encoder functions, 648 decoder functions, 1,613 descriptor initializers = 4,236 template-instantiated functions.

Encoder Architecture

Every encoder follows an identical structural template:

__int64 __fastcall encode_XXX(__int64 buf, __int64 ir_instr)
{
    sub_4C2A60(buf);                          // Initialize encoding buffer
    sub_4C28B0(buf, 0, 4, major_opcode);      // Set bits[3:0] = major
    sub_4C28B0(buf, 4, 3, 0);                 // Set bits[6:4] = format ext
    sub_4C28B0(buf, 8, 9, minor_opcode);      // Set bits[16:8] = minor
    sub_4C28B0(buf, 17, 8, sub_opcode);       // Set bits[24:17] = sub
    sub_4C28B0(buf, 25, 7, modifier);         // Set bits[31:25] = modifier
    sub_4C2A90(buf, ir_instr, pred_idx);      // Encode predicate guard
    sub_4C4D60(buf, ir_instr, 0, 0x50);       // Encode register operand 0
    sub_4C4D60(buf, ir_instr, 1, 0x60);       // Encode register operand 1
    sub_4C52F0(buf, ir_instr, 2, 0x70);       // Encode immediate operand
    sub_4C5C30(buf, ir_instr, 3, 0x80);       // Encode predicate operand
    // ... instruction-specific modifier encoding via sub_A4xxxx/sub_A5xxxx ...
    sub_A50D10(lookup_table, opcode_value);    // Primary opcode mapping
}

Core Infrastructure Functions

Modifier Encoding Helpers

Decoder Architecture

Decoders mirror encoders -- each unpacks a 128-bit instruction word into the internal IR:

__int64 __fastcall decode_XXX(__int64 buf, __int64 ir_out)
{
    sub_4C5F90(buf, ir_out);                       // Finalize init
    sub_4C60F0(buf, ir_out, 0, 0x50, 2);           // Decode register (GPR)
    sub_4C60F0(buf, ir_out, 1, 0x60, 10);          // Decode register (uniform)
    sub_4C6380(buf, ir_out, 2, 0x70, 9);           // Decode constant buffer
    sub_4C6DC0(buf, ir_out, 3, 0x80, 3);           // Decode predicate
    sub_50C790(lookup, bit_val);                    // Read 1-bit flag
    sub_50BBA0(lookup, val);                        // Read 5-bit predicate field
    // ... field extraction via sub_50xxxx/sub_51xxxx ...
}

Dispatch Functions

Two functions serve as the entry points for the entire encoding/decoding pipeline:

Encoder dispatch (sub_E43C20, 92 lines): Reads the instruction opcode from *(a2+12) as a 16-bit value. Special-cases opcodes 120-121 (control flow) to sub_4C8810. Otherwise performs binary search in the 24-byte dispatch table at off_1EA4380. Each table entry contains {opcode_byte0, opcode_byte1, encode_func_ptr, this_offset}. Supports vtable-aware dispatch (checks LSB of function pointer).

Decoder dispatch (sub_EFE6C0, 93 lines): Reads the 128-bit instruction word from a1[68] (offset 544). Extracts the format field (bits>>4)&7 and immediate offset 16*(bits&0xF). Delegates format 2/3 to sub_4CB100. Extracts the 9-bit major opcode (bits>>8)&0x1FF. For format 1, adds 121 to create a separate opcode space. Binary search in dispatch table at off_1E957E0. Calls sub_A49B50 to finalize.

InstrDesc Initializers

The 1,613 descriptor initializers at 0x84DD70--0xA48290 populate instruction metadata objects defining:

• Operand count and types (register, immediate, memory, predicate)
• Instruction properties (flags, latencies, execution unit assignment)
• Scheduling hints and resource requirements

Each descriptor writes a class ID to *(a2+12). The class distribution spans 56 unique IDs (0-103), with ID 18 being the largest (78 descriptors, integer arithmetic class).

Complete Instruction Set Map

The following tables enumerate all 118 instruction families across the three major opcodes. Within each family, sub-opcodes define specific instruction variants (register-register, register-immediate, predicated, etc.).

Major Opcode 1: ALU/Scalar (558 encoders)

Major Opcode 2: Vector/Memory/Control (977 encoders)

Major Opcode 3: Special (1 encoder)

Instruction Families by Functional Class

The 118 instruction families group naturally into functional classes:

Integer Arithmetic (13 families, ~110 encoders)

Integer arithmetic dominates the ALU opcode space. The key families:

• IADD3 (minor 0x005 in the 0xDA region): 3-input integer add with carry. 26 encoding variants covering register-register, register-immediate, and different data widths. All variants include 1 immediate + 1 predicate + 3-4 register operands.
• IMAD (minors 0x016, 0x017): Integer multiply-add in standard (0x016, 7 variants) and wide (0x017, 6 variants) forms.
• ISCADD/LEA (minor 0x18): Integer scale-add / load effective address. 10 variants. Used extensively in address computation.
• ALU-MISC (minor 0x012 in 0xDA region): The largest single class with 63 distinct encodings. Includes BFE, BFI, FLO, POPC, LEA, PRMT, and many more.

Floating-Point Arithmetic (5 families, ~30 encoders)

• FADD (minor 0x007): 11 variants covering FP32 add, multiply, FMA.
• HADD2/HFMA2 (minors 0x06A, 0x06D): Half-precision packed arithmetic. 5 and 4 variants respectively. HFMA2 uses 2 predicate operands.
• HSETP2 (minor 0x06F): Half-precision set predicate with paired output. 7 variants including the sole format 3 instruction.
• DFMA (minor 0x020): Double-precision FMA with 4 variants.

Memory Operations (15 families, ~90 encoders)

• LDG (minor 0x025): Global memory load. 9 variants with up to 4 register operands.
• STG (minor 0x02E): Global memory store. 8 variants with up to 4 immediates.
• LDS/STS (minors 0x00D, 0x0E3, 0x0EC): Shared memory load and store. 6 variants each.
• LDL (minor 0x067): Local memory / stack. 13 variants with 2-4 immediates for complex addressing.
• LDC (minor 0x060): Constant buffer load. 6 variants, all with 4 immediate operands.
• Async Copy (minor 0x0B8): 6 variants for asynchronous data movement.

Texture and Surface (5 families, ~30 encoders)

• TEX (minor 0x05A): Texture fetch. 13 variants. The most operand-rich instructions in the entire ISA: up to 7 register + 1 immediate + 1 predicate operands. The largest encoder function (sub_DC6680, 8,794 bytes, 302 lines) belongs to this family.
• TXQ (minor 0x08B): Texture query. 4 variants with 4-7 register operands.
• SUST (minor 0x0CE): Surface store. 4 variants.
• Texture/Surface load (minor 0x03B): 6 variants.

Tensor Core / MMA (16 families, ~120 encoders)

The tensor core families constitute the second-largest functional class:

Control Flow (10 families, ~25 encoders)

• BRA (minor 0x034): Branch with 3 variants (conditional, unconditional, indirect).
• EXIT/BRK (minor 0x3E): Exit/break with 7-operand encoding (5 registers + 1 predicate + 1 memory).
• CALL/RET (minors 0x079, 0x07B): Function call and return.
• SSY/CAL/PCNT/PRET (minors 0x42-0x46): Structured control flow primitives, each with 7 operands.
• YIELD/NANOSLEEP (minor 0x07C): Thread yield with 2 variants.

Synchronization (6 families, ~25 encoders)

• BAR.SYNC (minor 0x085): Barrier synchronization.
• WARP-BARRIER (minor 0x56): Warp-level barrier with 4+ variants.
• DEPBAR (minor 0x096): Dependency barrier for scoreboard management.
• BARRIER-OP (minor 0x0A7): Extended barrier operations.
• FENCE (minor 0x08D): Memory fence with 4 variants.

SM100-Specific Instructions

ROT13 Mnemonic Table

The SM100 opcode table constructor at sub_1782540 (111,076 bytes, 3,227 lines) initializes ~400+ instruction mnemonics using ROT13 encoding. The ROT13 prefix ZREPHEL decodes to MERCURY, the internal codename for Blackwell. Key SM100-specific mnemonics decoded from the binary:

tcgen05 Intrinsics

SM100 introduces the tcgen05 (Tensor Core Generation 5) subsystem with dedicated PTX-level intrinsics. The code generation infrastructure resides at 0x16E0000--0x16E3AB0 and spans 14+ functions handling type classification, bounds checking, tensor memory address computation, MMA operand setup, and argument type mapping.

The two PTX instruction mnemonics are tcgen05.mma (standard) and tcgen05.mma.ws (warp-specialized). Modifier suffixes include _expand16bit, _pack16bit, _maxabs, _minabs, _fused, _blockscale, and _ashift.

Instruction Type Classifier

sub_16E0A70 (17,302 bytes, 322 lines) classifies a tcgen05 MMA instruction into one of 54 type IDs by chaining predicate functions against the instruction object's type field at *(*(a1+8)). Each predicate is a trivial equality check (*a1 == N); the classifier tests them in priority order and returns the first matching type ID, or 0 if no predicate matches.

The 54 predicate functions live at 0x12B5670--0x12B5C20 (16 bytes each, all identical in structure). The complete classifier chain and resulting type ID assignment:

Type IDs 9, 12, 14, 17, 20, 22, 57 are absent from the enum -- either reserved for future expansion or used by related subsystems not routed through this classifier.

The type ID groupings inferred from the classification order and the modifier suffix strings:

• IDs 1--8: Base MMA types (standard precision combinations -- the 8 fundamental tcgen05.mma configurations)
• IDs 10--16: Extended base types with non-standard precision or accumulator widths
• IDs 18--24: Extended type A variants (wider accumulator or non-standard rounding)
• IDs 25--29: Blockscale variants (_blockscale modifier -- block-level scaling for MX formats)
• IDs 30--33: Quantized MMA types (tcmma_*_q / tcmma_*_mxq internal names)
• IDs 34--37: Mixed-precision variants (asymmetric A/B input types)
• IDs 38--41: FP8/FP6/FP4 narrow-type variants (e4m3, e5m2, e3m2, e2m3, e2m1 combinations)
• IDs 42--43: Sparse MMA variants (structured sparsity -- 2:4 or 4:8 patterns)
• IDs 44--46: Ashift variants (_ashift modifier -- arithmetic shift on matrix A)
• IDs 47--48: Fused variants (_fused modifier -- fused accumulation)
• IDs 49--51: Pack/Expand variants (_pack16bit, _expand16bit modifiers)
• IDs 52--56: MXQ (Mixed-precision Quantized) types (_blockscale + quantization)
• IDs 58--59: Maxabs/Minabs variants (_maxabs, _minabs reduction modifiers)
• ID 55: Extended sparse variant
• ID 60: Extended sparse variant (last in chain)

Guardrails Code Generation

sub_16E1DB0 (10,365 bytes, 325 lines) generates inline PTX boundary-checking code that validates tensor memory accesses before tcgen05.mma execution. This is a compiler-inserted safety mechanism -- the guardrails are weak functions that can be overridden at link time.

Symbol selection (line 67--73): The function selects one of two guardrail symbol prefixes based on the instruction opcode at offset 776:

Both prefixes get nCols appended to form the final guardrail variable name (e.g., __cuda__sm10x_tcgen05_guardrails_in_physical_bounds_nCols).

nCols computation (line 82--318): The number of columns accessed by the MMA instruction is computed via a multi-level dispatch. The first dispatch key is BYTE2(v63) & 0x1C -- a 3-bit field extracted from the instruction descriptor at offset 632:

Instruction-opcode nCols table (lines 103--128): When none of the above field values match, the byte at offset 620 (v4, signed char interpreted as instruction opcode) determines nCols:

The nCols value is modified by two flags in the instruction descriptor at offset 624 (v62.m128i_i8[8]):

• & 0x03 (CTA group flag): If nonzero, the base nCols is doubled (nCols = nCols_wide instead of nCols_narrow)
• & 0x78 (scale factor flag): If nonzero, nCols is multiplied by sub_12AB660(v5) where v5 = (byte >> 3) & 0xF is a 4-bit index into the scale lookup table at dword_1F24300 (values indexed 0--7)

Dynamic descriptor dispatch (field value 0x14, lines 158--318): When the field value is 0x14, a secondary 3-bit dispatch on *(a1+634) >> 5 selects one of 6 nCols computation strategies:

Case 3 checks whether either matrix A (offset 611, bits[5:4]) or matrix B (offset 627, bits[5:4]) uses FP16 format (== 1). If so, nCols is fixed at 4; otherwise, a full bitfield extraction from the descriptor determines the value dynamically.

Case 4 is notable for using floating-point arithmetic (cvt.rp.f32.u32, div.rp.f32, cvt.rpi.u32.f32) to compute a ceiling division -- nCols = ceil(bits * 8 / 64) * 2. This computes the number of 64-byte-aligned column groups needed.

Tensor Memory Address Computation

Five functions generate the inline PTX for computing tensor memory (tmem) addresses used by tcgen05.mma operands:

A sixth function at sub_16E2D70 (93 lines) handles __cuda_sm10x_tcgen05_mma_spMetaTmem -- the sparsity metadata tensor memory address, reading from offset 672 (slot 3).

Each function follows the same two-path pattern:

1. 32-bit tmem address (when the address operand type is sub_12B5850 or sub_12B56C0, internal type 6): Emits a simple add.u32:

   add.u32 %s, <symbol>, %s, %s;
   

2. 64-bit tmem address (all other types): Wraps in a scoped register block with 64-to-32-bit conversion:

   {.reg .b32 __cuda_sm_100_tcgen05_tmem_addr_base;
   cvt.u32.u64 __cuda_sm_100_tcgen05_tmem_addr_base, %s;
   add.u32 %s, __cuda_sm_100_tcgen05_tmem_addr_base, %s;
   }
   

The scale operand functions (sub_16E2990, sub_16E2B80) dynamically select their operand slot: scaleTmemA uses slot 5 - (cta_group_flag == 0) and scaleTmemB uses slot 6 - (cta_group_flag == 0), where cta_group_flag = (*(a1+627) & 0x30).

Return Value / Input Operand Marshalling

Two functions handle parameter passing to/from the tcgen05 MMA helper functions:

Output marshalling (sub_16E1A80, 47 lines): Iterates over the operand list at offset 648, emitting ld.param.b32 %s, [%s + %d] for each register result. The return array symbol depends on the opcode:

• Opcode 143: __cuda_sm_100_tcgen05_ld_red_funcRetArr (load-reduce variant)
• Opcode 160: __cuda_sm_100_tcgen05_ld_funcRetArr (standard load variant)

For reduce operations (sub_12A9290 returns true), an additional operand is appended from offset 656.

Input marshalling (sub_16E1BC0, 30 lines): Iterates over the input operand list at offset 648 (slot determined by *(a1+596) - 1), emitting st.param.b32 [%s + %d], %s for each register input. Uses the symbol __cuda_sm_100_tcgen05_st_funcInputArr.

Arguments Type Mapper

sub_16E3AB0 (18,623 bytes, 337 lines) performs the same type classification as sub_16E0A70 but operates on an array of tcgen05 instruction arguments rather than a single instruction. It iterates over the argument array (count at offset 796, result stored at offset 944), calling the identical predicate chain to assign numeric type IDs.

Guardrail Trap Functions

The guardrails infrastructure defines five weak trap functions that are called when bounds violations are detected. These are declared as PTX prototypes embedded in the binary:

Two additional check functions return a result value:

All are declared .weak .func (overridable at link time) and use .FORCE_INLINE for the check variants, ensuring they are inlined at every call site.

Additional tcgen05 Symbols

Beyond MMA, the tcgen05 infrastructure references additional symbols for load, store, and load-reduce operations:

The EIATTR system tracks tcgen05 MMA usage through two compatibility attributes: EICOMPAT_ATTR_INST_TCGEN05_MMA (current) and EICOMPAT_ATTR_INST_TCGEN05_MMA_DEPRECATED (legacy).

Decoder Opcode Classes for Blackwell-New Instructions

The decoder region (0xE43DC0--0xF15A50) reveals Blackwell-specific instruction classes through high opcode IDs (>300) not present in earlier architectures:

SM100 Compiler Backend

The SM100-specific compiler backend extends from 0x1782540 through 0x17B9300 and includes:

Opcode Table Constructor

sub_1782540 (111,076 bytes, 3,227 lines) is the SM100 opcode table constructor. It calls the parent class constructor sub_19B11F0 and initializes all SM100 instruction mnemonics with name + length at offset +11,360. The vtable is at off_2415E98. This is over 3x larger than the SM70 equivalent (sub_1769B50, 24,230 bytes) reflecting the expanded ISA.

Instruction Property Initializer

sub_17884A0 (44,713 bytes, 1,603 lines) sets latencies, throughputs, and execution unit assignments for every SM100 opcode. Companion to the opcode table constructor.

Scheduling Table

sub_178AA00 (35,422 bytes, 1,205 lines) initializes the SM100-specific scheduling resource tables defining per-instruction throughputs and latencies for each execution unit type (ALU, SFU, LDST, TEX, MMA).

SM100-Specific Optimization Passes

Master Encoder Functions

Two monumental encoding functions handle the final SASS binary emission:

• sub_17F2670 (156,611 bytes, 4,858 lines): The master instruction encoder -- the largest function in the entire nvlink binary. Dispatches to individual encoding routines for all SASS instruction types. ~640 local variables, 0x2C8-byte stack frame.
• sub_17F9AE0 (61,531 bytes, 2,150 lines): Secondary encoder for less common instruction types.

Format Descriptor Tables

Each instruction references a 16-byte (128-bit) format descriptor loaded via SSE instructions. These descriptors define the operand layout template:

Encoder-Side Descriptors

Decoder-Side Descriptors

The universal operand descriptor at xmmword_1F460E0/xmmword_1F460F0 (32 bytes) is referenced by all 4,236 encoder+decoder+descriptor functions -- it defines the register class mapping for the SM100 architecture.

Internal Data Structures

Instruction Representation Object (~560 bytes)

The a1 parameter across all encoder/decoder functions points to:

Offset  Size   Description
------  ----   -----------
  0       4    Flags / instruction ID
  4       4    Scheduling control bits
  8      16    Format descriptor (128-bit SSE copy from xmmword table)
 12       4    Operand count metadata field 1
 16       4    Operand count metadata field 2
 24-60   40    Operand register indices (10 x 4 bytes)
 64-100  40    Operand type/class (10 x 4 bytes, -1 = unused)
104-140  40    Operand modifier flags (10 x 4 bytes)
144       4    Active operand count
148       4    Immediate encoding offset 1
152       4    Immediate encoding offset 2
156-276 120    Reserved / operand extension data
276-404 128    Decoded operand output buffer
404-468  64    Modifier flag output buffer
452       4    Modifier count field 1
456       4    Modifier count field 2
468-532  64    Operand remap table (16 DWORD slots)
532       4    Operand remap counter
536       8    Pointer to encoding/decoding lookup table
544       8    Instruction word 0 (bits 0-63 of 128-bit SASS)
552       8    Instruction word 1 (bits 64-127 of 128-bit SASS)
556       4    Decoded immediate / offset value

Encoder Dispatch Table Entry (24 bytes)

Offset  Size  Description
------  ----  -----------
  0       1   Opcode byte 0 (minor opcode low)
  1       1   Opcode byte 1 (minor opcode high / sub-opcode)
  8       8   Function pointer to encoder (LSB=1 indicates vtable indirection)
 16       8   this-pointer offset adjustment

Statistical Summary

Confidence Assessment

For general Blackwell architecture details, see the ptxas wiki: Blackwell and cicc wiki: SM100 Blackwell.

Cross-References

nvlink Internal

• Embedded ptxas Overview -- SM100+ codec at 0xDA0000--0xF16000 and 0x620000--0xA70000
• Architecture Dispatch -- SM100/100a/100f vtable registration and callbacks
• Architecture Profiles -- SM100 profile in linker database
• Compatibility -- SM100 finalization compatibility and capability bitmasks
• SM103/110/120/121 -- derived Blackwell-family architectures
• Mercury Overview -- Mercury encoding is the default for SM100+
• FNLZR -- finalizer for Mercury/capmerc SM100+ output
• Capsule Mercury Format -- capmerc output format for SM100+

Sibling Wikis

• ptxas: Blackwell -- standalone ptxas SM100+ target documentation
• ptxas: tcgen05 -- standalone ptxas tcgen05 tensor intrinsics
• ptxas: Mercury Encoder -- Mercury pipeline for SM100+ encoding
• ptxas: Capsule Mercury -- capmerc format in standalone ptxas
• cicc: SM100 Blackwell -- cicc compiler SM100 Blackwell target