SM103a Instruction Set Architecture

NVIDIA Blackwell datacenter (B300 SXM6) — Unofficial Reference
163 opcodes · 397 instruction forms · 1291 encoding variants · 397 with scheduling data
Sources: nvdisasm pairing on a 31.8M-instruction vendor corpus, tcgen05 PTX probe corpus (CPU-exact encoding pairs), hardware validation on B300

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

1. Architecture Overview

SM103a is the streaming multiprocessor architecture used in NVIDIA's Blackwell datacenter GPUs (B300 SXM6), the refresh of SM100a (B200). Across all 119 probe pairs shared with SM100a in the tcgen05 corpus, the 128-bit instruction words are byte-identical between the two targets.

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 SM103a 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

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

7. Barrier and Scoreboard System

SM103a 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 SM103a:

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