SASS Execution Model — Pipelines, Scoreboards, Scheduling

Recovered from the decoded per-architecture SASS instruction tables in nvdisasm V13.1.115 (CUDA 13.1). The compiler-side hazard resolution (RAW/WAR/WAW) is the Dependency & Hazard Model; this page is the hardware's dynamic execution and warp-issue model.

The single most striking result of decoding the tables is that the warp-scheduler ISA contract has not changed since Volta. The control-word bit layout, the six scoreboards, the 6-bit wait mask, the usched_info (0–27) and batch_t enums are byte-identical from Turing (SM75) through Blackwell (SM121) — every table even self-identifies as ARCHITECTURE "Volta". Six GPU generations have added pipes and datapaths but left the scheduling model untouched. It is a single-issue, scoreboard-assisted, statically stall-scheduled machine: the compiler pre-computes every fixed stall and every variable-latency barrier; the hardware just evaluates the wait mask, honours the stall counter, and picks among ready warps.

The SM as a physical machine

The scheduling model runs on a fixed physical substrate, stable across the Volta→Hopper generations:

A warp is bound to one of the 4 sub-partitions, which owns its register-file slice, its dispatch unit, and the 6 scoreboards; the single-issue rule is per sub-partition. The 2-bank (even/odd) register file is exactly why the operand collector and the .reuse cache matter — two source operands drawn from the same bank in one cycle would conflict, so the allocator biases an accumulator into the opposite bank from its multiplicands and the reuse cache serves a repeated operand without a second read. (Banking and the allocator's bank bias are in Register Allocation.)

Coupled vs decoupled — the INSTRUCTION_TYPE split

Each instruction's INSTRUCTION_TYPE classifies how its latency is tracked. The enum grows from 9 codes (Turing–Hopper) to 12 (Blackwell adds scoreboard-elided emulatable variants):

Distribution by instruction form (PROPERTIES blocks; ~55–60 % COUPLED_MATH, ~40–45 % DECOUPLED on every arch):

Verified correlation: COUPLED_MATH forms carry no pipe (~99 %) — they are fixed-latency ALU ops dispatched into the back-to-back math pipe — while every DECOUPLED form has a VIRTUAL_QUEUE assignment (variable latency must route to a named functional-unit queue).

Pipe topology — the VIRTUAL_QUEUE functional units

VIRTUAL_QUEUE names the FU pipe an instruction dispatches into (39 codes). The pipe set is where generations diverge (distinct pipes: SM75=17, Ampere=20, SM90=26, Blackwell-DC=34, consumer Blackwell=27). Population by form:

Separate queues mean a long-latency texture or tensor op does not block an independent FMA — the scheduler tracks readiness per queue. The emulation→native migration is visible directly: Turing's dedicated REDIRECTABLE_FP16 folds away by SM80; FP64 stays on the emulated REDIRECTABLE queue until SM120 gains a native FP64 pipe; tensor goes emulated HMMA/IMMA → Hopper UMMA → Blackwell-DC TC_1CTA/TC_2CTA + TMEM. COUPLED_EMULATABLE reaches zero on SM120/121 — consumer Blackwell implements FP64 and its tensor variants natively.

The scheduling-control word

Bits 102–125 of every instruction carry the schedule ptxas computed (only bits 127–126 are reserved; layout identical on all arches):

opex = batch_t<<5 | usched_info

Verified across 4500 table rows, 0 mismatches: the 8-bit opex byte is the high-3-bit batch_t (122–124) over the low-5-bit usched_info (105–109).

usched_info (0–27) — the unified stall/yield field:

This is the Volta-family successor to the Maxwell/Pascal "stall count + yield bit" control word, packed into one field.

batch_t — the dual-issue/batch control: NOP(0) standalone · BATCH_START(1) · BATCH_START_TILE(2) · BATCH_END(4) · BARRIER_EXEMPT(5).

Dual-issue is in the instruction, not the scheduler

Every arch reports ISSUE_SLOTS 1 — one math issue per warp per cycle. Superscalar co-issue is expressed entirely through batch_t: BATCH_START … BATCH_END mark instructions that issue together into distinct pipes in one scheduler decision. The TABLES_opex_N(batch_t, usched_info) legality tables encode the co-issue rule: TABLES_opex_0 permits a stall (usched 17–27) in any batch position (~40 % of forms), but all other tables force a multi-cycle stall to batch_t=0 — an instruction inside a co-issue batch cannot also carry a stall. You may batch or stall, not both. The compiler, not the hardware, decides co-issue groups.

Operand-reuse cache

A separate 64-bit "OE" operand-encoding view exposes per-source reuse-cache flags OEReuseA/OEReuseB/OEReuseC (bits 46/45/44). When set, a source operand is taken from a small reuse cache instead of re-reading the register file — a power/bandwidth optimization the compiler controls per source slot. Reuse is mutually exclusive with ?DRAIN/?WAITn_END_GROUP scheduler tokens (a CONDITIONS rule enforces it).

Scoreboard / barrier model

Exactly six dependency scoreboards, SB0–SB5, on every arch SM75→SM121 (values 6/7 are the "no barrier" sentinel in the 3-bit producer fields). Resolution:

1. A variable-latency producer (DECOUPLED_*_SCBD) is assigned dst_wr_sb (set when its result is written back) and/or src_rel_sb (set when its sources have been read). The format syntax WR:wr = dst_wr_sb, RD:rd = src_rel_sb makes this explicit; a load (ldg) sets both.
2. The producer issues and the warp continues with independent instructions.
3. A consumer sets the matching bit in its 6-bit req_bit_set wait mask. The warp may not issue until every scoreboard in the mask has reached its target count — up to 6 outstanding dependencies waited on at once.

Fixed-latency (COUPLED_MATH) producers skip all of this: the RAW latency is a compile-time constant emitted as a usched stall, and MIN_WAIT_NEEDED gives the floor (0 for ~55 % of math, up to 6 for rpcmov/depbar).

DEPBAR.LE sb, cnt, {list} is an explicit scoreboard wait-down: block until the indexed scoreboard's outstanding count drops to ≤ cnt. This waits on partial completion of a batch of outstanding ops (e.g. "until ≤2 of 5 loads remain").

MEM_SCBD — the orthogonal memory-ordering axis. Separate from register scoreboards, MEM_SCBD (NONE/SOURCE_RD/SOURCE_WR/SINK/SOURCE_SINK_WR/…) + MEM_SCBD_TYPE (BARRIER_INST/MEM_INST/BB_ENDING_INST/ALL) sequence the memory subsystem: SINK/BARRIER_INST = membar/TMA drain the memory pipe; SOURCE_WR/ALL = fence produces global ordering; SOURCE_SINK_WR = async smem barriers. This is the fence/membar model layered on top of register scoreboarding.

The per-warp issue rule

Each cycle, for a warp's next decoded instruction I:

1. Wait-mask gate — if req_bit_set ≠ 0 and not all its scoreboards have reached target → WAIT (a variable-latency RAW/WAR dependency).
2. Dispatch-stall gate — if cycles_since_last_issue < MIN_WAIT_NEEDED(I) → STALL; honour the usched stall before the next issue (0 → yield-candidate; 1–15 → stall+end-batch; 17–27 → stall n−16).
3. Yield — usched == DRAIN(0) voluntarily yields the issue deck (greedy-then-oldest), avoiding SM monopolization.
4. Batch — BATCH_START … BATCH_END co-issue into distinct pipes in one decision (the only "dual issue"; mutually exclusive with a stall on most pipes).
5. Issue — into the instruction's VIRTUAL_QUEUE. COUPLED_MATH → fixed pipe, no scoreboard; DECOUPLED_* → FU queue, arm dst_wr_sb/src_rel_sb, move on. BRANCH_DELAY == 0 on all arches — no architected branch delay slot; convergence is the CBU pipe (BSSY/BSYNC) + control scoreboards.

Validation against ptxas-emitted SASS

The model above is reproduced exactly by the schedule ptxas computes. The check assembles PTX with ptxas, reads each instruction's 128-bit encoding with nvdisasm -c -hex, and decodes the control word (bits 105–124) back into usched/dst_wr_sb/src_rel_sb/req_bit_set/batch_t. A saxpy inner block on SM90 decodes as:

The two loads share one scoreboard (SB2); the consuming FFMA carries a single req_bit_set bit for SB2 and blocks until that scoreboard's outstanding count returns to its target — i.e. until both loads have written back. This is the producer→consumer pairing the model predicts, realized verbatim.

Across sm_75 sm_80 sm_86 sm_89 sm_90 sm_90a sm_100 sm_120, over a battery of math-chain, load, transcendental, tensor and cp.async probes (1912 decoded instructions): every consumer wait-bit pairs with an earlier producer that armed exactly that scoreboard — 221 pairs, 0 unmatched — and opex = batch_t<<5 | usched_info holds on every single instruction (0 violations, recomputed independently from the raw 128-bit words). A co-issue batch (batch_t ≠ 0) never carries a group-ending stall (usched 1–15); the only stall that rides a batch member is the 1–2-cycle no-group-end spacing of a .reuse co-issue pair — confirming "batch or stall, not both".

Observed dispatch stall by producer class

For fixed-latency math, the stall ptxas places between a producer and its immediate dependent consumer is small and stable across every generation:

The dispatch stall is the producer→consumer dependency latency for the coupled math pipe, which is narrower than the 6-cycle result band: a same-pipe pair emits 4 (the coupled base latency), a cross-pipe pair emits 5 (one extra cycle of inter-pipe hazard penalty). The 6-cycle band is the higher-level scalar that also covers the slower CC/predicate-write path — a coupled op writing a condition-code or predicate exposes 6 + 7 = 13, the control band, which is why an ISETP→predicate edge stalls 13. The stall always counts the cycles between the two issue slots (the wait the consumer owes after the producer issues), so it is the dependency edge weight directly, not a band with a constant subtracted. The MUFU pipe (transcendentals, F2*/I2F conversions, POPC) carries the band-13 latency but delivers it through a scoreboard, not a fixed stall — the band gives the magnitude, the INSTRUCTION_TYPE gives the mechanism.

cp.async (cp.async.commit_group → LDGDEPBAR; cp.async.wait_group k → DEPBAR.LE SBn, k) is the explicit partial-wait path: the count k is the number of outstanding copy groups the warp will tolerate, byte-identical from SM80 through SM120. WAR (read-then-overwrite) pairs carry no extra stall — ptxas renames the destination so the anti-dependency never reaches the schedule.

Reference simulator

decoded/sass-tools/sass_sched_sim.py is a single-warp issue/timing simulator that replays exactly this rule over a decoded stream — wait-mask gate, dispatch stall, yield-on-DRAIN, batch_t co-issue, scoreboard arm/wait — and emits a cycle-by-cycle trace plus a basic-block cycle estimate. Its --validate mode is the producer→barrier and opex self-check used above. sass_ctrl_decode.py is the control-word decoder it consumes.

Per-arch execution deltas

Cross-References

• SASS Instruction Encoding — the control-word bit placement.
• Dependency & Hazard Model — the ptxas-side RAW/WAR/WAW decision.
• Scoreboards & Dependency Barriers — ptxas scoreboard allocation.
• Architecture Evolution — the pipe additions per generation.
• Tooling: decoded/sass-tools/sass_ctrl_decode.py (control-word decoder) and sass_sched_sim.py (warp-issue / timing simulator + --validate self-check).