SM100a Instruction Set Architecture

NVIDIA Blackwell datacenter (B200) — Unofficial Reference
163 opcodes · 400 instruction forms · 1628 encoding variants · 400 with scheduling data
Sources: canonical SM103a encoding layer (byte-identical on all 119 shared tcgen05 probe pairs), offline decoder/encoder replay on an archived 20.0M-instruction sm_100 corpus; no B200 re-validation — scheduling and pipeline sections are inherited from SM103a and unvalidated for B200 timing

Contents

1. Architecture Overview
2. Execution Model
3. Register Files
4. Instruction Encoding
5. Control Word and Scheduling
6. Pipeline Units
7. Barrier and Scoreboard System
8. Memory Hierarchy
9. Instruction Reference
10. Derivation and Validation (vs SM103a)

1. Architecture Overview

SM100a is the streaming multiprocessor architecture of NVIDIA's Blackwell datacenter GPUs (B200, GB100). SM103a (B300) is its refresh: across all 119 probe pairs shared between the two targets in the tcgen05 corpus, the 128-bit instruction words are byte-identical, so the encoding layer documented here is the SM103a canonical table carried over 1:1, with the tcgen05 acceptance deltas listed below.

1.1 Architecture Notes

2. Execution Model

Each SM executes warps of 32 threads in lockstep (SIMT). The warp scheduler selects a ready warp each cycle and dispatches one instruction.

Divergence Handling

Structured divergence uses convergence barriers: BSSY sets a reconvergence point, BSYNC waits for all threads in the region. SM120 has 6 barrier slots (B0–B5). For unstructured control flow, WARPSYNC forces reconvergence.

Occupancy

Register usage determines how many warps can co-reside on an SM. With 65,536 registers and a maximum of 255 per thread, the register pressure is the primary occupancy limiter. NVIDIA’s compiler uses a sophisticated register target selection algorithm (947 decoded configuration knobs) that balances register count against warp occupancy, including shared memory spilling as a first-class strategy (RegAllocCudaSmemSpillEnable).

3. Register Files

4. Instruction Encoding

Every SM100a instruction is exactly 128 bits (16 bytes), stored little-endian as two 64-bit words:

Instruction Word [63:0]

Encoding Formula

Each instruction variant is defined by:

• and_base — 128-bit constant containing all fixed bits (opcode, modifiers, default scheduling, constant operands)
• fields[] — list of variable bitfields, each with {shift, bits, extraction}
• variable_mask — OR of all field masks

To encode: instruction = and_base | (field_value & mask) << shift for each field.

Extraction Types

5. Control Word and Scheduling

The upper 64 bits of each instruction contain scheduling information that controls instruction issue timing, dependency tracking, and pipeline utilization. These bits are set by NVIDIA’s compiler or by a custom barrier insertion pass.

6. Pipeline Units

SM100a has 37 distinct pipeline classes identified from scheduling records. Each instruction is dispatched to exactly one pipeline unit.

7. Barrier and Scoreboard System

SM100a uses a scoreboard-based dependency tracking system with 6 barriers (B0–B5). The barrier insertion algorithm works as follows:

1. Build dependency graph from register definitions and uses within each basic block.
2. Classify dependencies as short-latency (≤ threshold) or long-latency (> threshold):
   • Short: handled by stall count in control word bits [67:64], range 0–15 cycles.
   • Long: requires write barrier assignment (e.g., LDG, TEX, S2R).
3. Assign write barriers. The producer instruction sets write_barrier = Bn in its control word. Uses interval graph coloring (like register allocation) to map live barrier ranges to B0–B5.
4. Set wait masks on consumers. When an instruction reads a register produced by a long-latency op, set bit N in the consumer’s wait_mask.
5. Compute stall counts. For short-latency dependencies: stall = max(0, latency − distance_in_instructions).

Key scheduling knobs: ShortStallThreshold (boundary between short/long), FirstLongLatencySB (which barrier index starts long-latency allocation), ScoreboardUseMultSBLimit (max barriers per instruction). Cross-block dependencies handled by SbXBlock, CrossBlockDepbar.

8. Memory Hierarchy

Address Formats

Global/shared/local loads use descriptor-based addressing on SM100a:

LDG.E Rd, desc[UR4][Ra + imm]    # Global: uniform descriptor + GPR offset + immediate
LDS   Rd, [Ra + imm]              # Shared: GPR base + immediate offset
LDC   Rd, c[bank][Ra + imm]       # Constant: bank select + GPR base + immediate

The .E modifier on LDG/STG indicates “extended” 64-bit addressing.

9. Instruction Reference