A case study of the B200 pipeline model

A note on methodology: Everything in this article is based on my analysis of microbenchmarks executed directly on B200 silicon. Nvidia does not publish instruction latencies, pipeline depths, or scoreboard encoding details for its GPUs. The numbers and mechanisms described here represent my best empirical understanding. Readers should do their own due diligence and verify against their own hardware.

When working with modern, deep-pipeline GPUs like the Nvidia B200, static analysis is necessary but insufficient for validating instruction schedules. It is a humbling experience to see a scheduler report 100% test coverage on dependency tracking, only to watch the emitted code fail silently on actual silicon.

Why does this happen? The hardware pipeline itself is the final arbiter of correctness.

When a scheduler under-stalls a dependency, it allows a consumer instruction to issue into the pipeline before the producer’s result is firmly committed to the register file. The hardware does not raise an exception. Instead, it executes the schedule, reading stale state, and propagates incorrect values through the rest of the computation.

These are not defects in the silicon. They are schedule violations where the hardware exposes the compiler’s incorrect assumptions. In compiler backends, compiler engineers generally adhere to the rule: over-stalling is a performance bug, but under-stalling is a silent correctness bug.

To catch these issues, I constructed a registry of hardware hazards ^{backed by minimal, reproducible on-silicon tests.}

Prerequisites & Terminology

Before diving into specific B200 hazards, it helps to establish some baseline context:

• Instruction Scheduling: A phase in the compiler backend that reorders instructions to maximize hardware utilization. It must explicitly encode delays (stalls) or synchronization (scoreboards) between dependent instructions.
• Pipeline Depth: The number of stages an instruction passes through (fetch, decode, execute, writeback). Deeper pipelines take longer to complete an instruction.
• RAW (Read-After-Write) Hazard: A scenario where an instruction tries to read a register before a previous instruction has finished writing to it.
• Variable-Latency Operations: Operations whose execution time is not fixed. This includes global memory loads (LDG), shared memory operations (LDS), atomic operations (ATOM), and multi-function units (MUFU).

Silicon Doesn’t Lie

Modern GPU streaming multiprocessors (SMs) are designed for extreme throughput. To achieve this, the pipeline is deep. The hardware relies on the compiler to explicitly encode dependency information^.

Consider a simple dataflow path where Instruction A produces a value that Instruction B consumes.

flowchart TD
    subgraph "Producer (Instruction A)"
        FetchA[Fetch] --> DecodeA[Decode]
        DecodeA --> IssueA[Issue]
        IssueA --> ExecA1[Execute Stage 1]
        ExecA1 --> ExecA2[Execute Stage N]
        ExecA2 --> WriteBack[Writeback to Register]
    end

    subgraph "Consumer (Instruction B)"
        FetchB[Fetch] --> DecodeB[Decode]
        DecodeB --> IssueB{Wait on Stall / Scoreboard}
        IssueB --> ReadReg[Read Register]
        ReadReg --> ExecB1[Execute Stage 1]
    end

    WriteBack -->|Data Forwarding / Register File| ReadReg
    IssueA -.->|Static Latency L| IssueB

If Instruction B issues too early, its Read Register phase fetches the register’s old contents before WriteBack completes.

On CPUs, sophisticated out-of-order execution engines mask these latencies dynamically. On GPUs, the philosophy is to maximize die area for ALUs. This pushes the complexity of instruction scheduling onto the compiler^{. This is reminiscent of VLIW architectures, which share the same philosophy of offloading scheduling decisions from hardware to the compiler.}

This architectural tradeoff means that compiler engineers must be pedantic about low-level constraints like pipeline depths and barrier encodings.

The Predicate-Consumer Under-Stall

The most difficult bugs slip through rigorous static checks. Recently, while hacking on the B200, I discovered a critical bug involving predicate evaluation in an instruction scheduler. This occurred despite static metrics claiming full RAW coverage across the test suite.

The pattern involves an integer set-predicate instruction (ISETP) that computes a condition. It writes it to a predicate register, which is then read by a branch instruction. This is classically seen in a back-edge branch defining a loop.

// 1. Produce the predicate P1 based on some condition.
// R0 and R1 are compared; the boolean result is written to P1.
ISETP.GE.AND P1, PT, R0, R1, PT;

// 2. Consume P1 as the branch target condition.
// P0 is the execution guard (is the thread active?), P1 is the branch condition.
@!P0 BRA P1, target;

The Mechanism of Failure

The bug surfaced while I was hacking on the B200’s predicate handling. The compiler correctly recorded the guard predicate P0 as a use for the branch, but it missed the branch condition operand P1.

Consequently, the ISETP $\rightarrow$ BRA RAW dependency was missed entirely. The scheduler failed to insert the required predicate-latency stall.

sequenceDiagram
    participant HW as B200 Pipeline
    participant Reg as Predicate Register File

    Note over HW: Under-stalled Schedule Execution
    HW->>HW: Issue ISETP (Computes P1)
    Note right of HW: ISETP latency is ~13 cycles.
    HW->>HW: Wait 4 cycles (Arbitrary delay, not a dependent stall)
    HW->>HW: Issue BRA (Reads P1)
    HW->>Reg: Read Predicate P1
    Reg-->>HW: Returns STALE state (0 instead of 1)
    HW->>HW: Takes wrong execution edge
    Note right of HW: 9 cycles later...
    HW->>Reg: ISETP Writeback to P1 (Too late)

The branch issued roughly 4 cycles after the ISETP. This was well before the predicate’s actual modeled latency of 13 cycles had elapsed. The branch instruction read a stale value, took the wrong execution edge, and resulted in a silent miscomputation.

Ground-Truth Mitigation

To prevent this, a scheduler’s operand analysis must correctly identify both P0 and P1 as uses of @!P0 BRA P1.

However, the true defense is an on-silicon probe. I sweep the stall cycles between the ISETP and the branch, verifying the minimum latency required for correct execution.

The baseline predicate latency must be dynamically probed on-device because architectural models are often approximations. On the B200, microbenchmarking probes confirmed the divergence between the modeled 13 cycles and the actual pipeline depths, where the physical predicate latency floor sits at approximately 4 cycles.

Fixed-Latency RAW Under-Stalls

Fixed-latency arithmetic instructions form the backbone of matrix multiplication and tensor core workloads. They require precise, fixed cycle delays before their destination registers can be safely read. Examples include:

• FFMA (Single-precision Fused Multiply-Add)
• DFMA (Double-precision Fused Multiply-Add)

If a scheduler emits a stall with a cycle count strictly below the hardware’s fixed latency, the consumer reads the destination register early.

Latency Measurement and Tradeoffs

Through direct hardware probing on the B200, I measured the exact latency floors where execution transitions from incorrect (stale read) to correct (valid read)^.

Notice the tradeoff here: higher precision arithmetic naturally requires deeper pipelines. The FP64 unit requires exactly twice the latency of the FP32 unit.

When building latency validation tests, it is critical to construct floating-point recurrence chains rather than integer linear chains. Integer chains can be folded or bypassed via pipeline forwarding networks in hardware, which masks under-stalls^{. Floating-point chains, due to strict execution pipeline stages and rounding, make dependency latencies visible.}

To validate these latencies, I wrote a probe kernel that builds a long dependent FFMA chain: a = a*1 + 1, repeated 64 times, so the expected result is exactly seed + 64. Every FFMA reads the immediately preceding FFMA’s result, creating a pure RAW hazard chain. A post-processing script then rewrites the stall field of every FFMA in the compiled SASS binary to a forced value and runs each variant on the B200.

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29

// probe_ffma.cu -- Tier-2 "minimum correct stall" probe for FFMA.
//
// A single dependent FFMA chain: a = a*1 + 1, CHAIN times.
// The exact result is seed + CHAIN (integer-valued float, bit-exact on CPU).
// Every FFMA reads the immediately-preceding FFMA's result -> a RAW hazard
// on the FMA pipe.
//
// We compile this once, then rewrite the stall subfield of every FFMA
// to a forced value N and run each variant on the B200:
//   * if N < true FFMA latency -> consumer reads stale register -> WRONG
//   * smallest N that stays CORRECT == the FFMA latency floor

#ifndef CHAIN
#define CHAIN 64
#endif

extern "C" __global__ void probe(float *out, unsigned *cyc, float seed) {
  float a = seed;
  unsigned t0, t1;
  asm volatile("mov.u32 %0, %%clock;" : "=r"(t0));
#pragma unroll
  for (int i = 0; i < CHAIN; i++)
    asm volatile("fma.rn.f32 %0,%0,%1,%2;" : "+f"(a) : "f"(1.0f), "f"(1.0f));
  asm volatile("mov.u32 %0, %%clock;" : "=r"(t1));
  if (threadIdx.x == 0) {
    out[0] = a;          // expected == seed + CHAIN if no hazard
    cyc[0] = t1 - t0;    // timed cycles for the chain
  }
}

Running this probe across stall values 0–15 on the B200 produces the following output:

1
2
3
4
5
6
7
8
9
10

# stall  cycles  result   (chain=64)
     0    2205  CORRECT
     1     225  WRONG
     2     227  WRONG
     3     291  WRONG
     4     355  CORRECT
     5     419  CORRECT
     6     483  CORRECT
     ...
    15    1075  CORRECT

The boundary is unambiguous. Stall 3 yields WRONG (the consumer reads a stale register), stall 4 yields CORRECT. Stall 0 is a special encoding that defaults to a large wait, which is why it reports CORRECT but at a much higher cycle count. I used tools like cuobjdump -sass to verify the assembled control codes before running the resulting binaries directly on the GPU.

The “Friendliest Dependency” Trap

A word of caution: the FFMA floor of 4 came from the friendliest possible dependency shape — an FFMA→FFMA chain where the consumer reads a single live register and two folded-immediate 1.0f operands. This is the easiest case for the hardware to service.

The true latency floor is not a property of the producer alone. It is a function of stall(producer_op, consumer_op, operand_layout). Several mechanisms can push the required stall higher for real workloads:

• Consumer op-class read stage. Different consumer ops read their source operands at different pipeline stages. An FFMA result may need to be ready earlier for an FADD consumer than for another FFMA.
• Operand-collector and register-bank pressure. A consumer reading three distinct live registers in conflicting banks needs an extra operand-gather cycle. The probe above reads one register and two immediates, so it never pays this cost.
• Operand .reuse flags. A homogeneous same-slot chain is exactly what ptxas tags with .reuse. A served-from-latch operand has different timing than a cold gather, so the probe chain can see an artificially low floor that a real (non-reuse) instruction stream would not tolerate.

A scheduler is correctness-first: if any real consumer needs stall 5 and the model emits 4, the result is a silent wrong-answer bug. The safe approach is to sweep across consumer op-classes and operand layouts, then set the model’s scalar latency to the maximum observed floor. Do not trust a single producer→consumer microbenchmark in isolation.

Variable Latency and Uncovered Scoreboards

Fixed latencies apply only to deterministic ALU operations. However, a vast portion of a GPU’s workload involves variable-latency operations:

• LDG (Global memory loads)
• LDS (Shared memory operations)
• ATOM (Atomic memory operations)
• MUFU (Multi-function units)

These instructions have execution times that vary based on cache hits, TLB state, and structural hazards (resource conflicts in the hardware pipeline).

For these operations, static stall counts are entirely inadequate. Instead, the architecture utilizes a scoreboard.

The Scoreboard Mechanism

When a variable-latency instruction is issued, it allocates a slot in a hardware scoreboard. The compiler must explicitly encode a barrier identifier (e.g., scoreboard indices 0 through 5) in the control fields of the instruction. The consumer instruction must then be encoded to wait on that specific scoreboard barrier before issuing.

stateDiagram-v2
    [*] --> IssueProducer

    state IssueProducer {
        direction LR
        Decode --> AllocateBarrier : "e.g., Barrier 0"
        AllocateBarrier --> DispatchToMemory
    }

    DispatchToMemory --> MemorySubsystem : Cache Miss

    state MemorySubsystem {
        direction LR
        L1 --> L2
        L2 --> HBM
    }

    MemorySubsystem --> Writeback : Data Returns
    Writeback --> ClearBarrier : "Clear Barrier 0"

    state ConsumerWait {
        direction LR
        DecodeConsumer --> WaitBarrier : "Wait on Barrier 0"
        WaitBarrier --> IssueConsumer : Barrier Cleared
    }

    IssueProducer --> ConsumerWait : Compiler encodes dependency
    ClearBarrier --> WaitBarrier : HW Signal

If an assembler strips these control barriers, the pipeline coherence breaks down. Consumers read destination registers before the variable-latency memory result has landed.

This leads to incorrect mathematical results. More critically, if the stale data is used as a memory address in a subsequent access, it triggers a CUDA_ERROR_ILLEGAL_ADDRESS.

Validation requires generating stripped binaries for various kernels and asserting that they must either compute the wrong result or crash. This proves that the scoreboard barriers in the fully assembled binaries are the sole mechanism guaranteeing correctness.

Crash-Amplified Load-Use Hazards

The fixed-latency under-stalls described above cause silent data corruption by reading a nearby valid—but mathematically wrong—value. Load-use hazards, however, can be intentionally amplified to provide a deterministic, loud failure^.

This is highly desirable for Continuous Integration (CI) systems, where binary pass/fail crash signals are easier to triage than heuristic output differencing.

A load-use hazard occurs when a value intended to be used as a memory address is read before the load producing it has landed.

To amplify this into a guaranteed crash, I explicitly poisoned the index register. I loaded a wild constant into the register (e.g., 0x40000000, translating to a +4 GiB offset in memory) before the actual load occurred. I maintained this poison value’s liveness via a runtime-unknown guard so the compiler’s dead-code elimination (DCE) pass could not optimize it away.

1
2
3
4
5
6
7
8
9
10

// 1. Initialize the target register with a known poison value pointing to unmapped memory space.
uint32_t addr_reg = 0x40000000;

// 2. Issue a global memory load (variable latency) to overwrite the register with a valid address.
// Under correct execution, this load must complete and update the register before any subsequent read.
addr_reg = load_actual_address();

// 3. Immediately consume the register as a pointer or address operand.
// If the compiler fails to emit a scoreboard wait, this instruction reads the stale register state (0x40000000), triggering an MMU fault.
execute_memory_operation(addr_reg);

The Amplification Tradeoff

This testing methodology provides significant value:

1. If the schedule is correct, the load lands, the valid pointer is used, and execution succeeds cleanly (ok).
2. If the schedule under-covers the latency, the hardware reads the poisoned register, resulting in a deterministic MMU fault (CUDA_ERROR_ILLEGAL_ADDRESS).

It is crucial to note that this is a recoverable software fault. The MMU and the CUDA driver safely contain the illegal memory access. It does not cause physical hardware damage; at worst, it results in a dead CUDA context that requires process restart or an nvidia-smi --gpu-reset.

The primary advantage is an unambiguous, self-contained, minimal form of the scoreboard hazard described earlier that leaves no room for debate in the CI logs.

The Path Forward: Trusting Silicon

Building reliable compilers and instruction schedulers requires treating the hardware as the ultimate source of truth. Software models are theoretical; silicon is absolute.

As we scale into more complex architectures, the abstraction gap between the high-level language and the physical pipeline deepens. Relying solely on static analysis or internal graph coverage metrics is fundamentally insufficient.

By systematically building a registry of targeted, minimal hardware hazards and executing them continually on actual silicon, compiler engineers can ensure that their scheduling logic remains sound against the unyielding reality of the pipeline.

References

Disclaimer: This article was generated using the Gemini 3.1 Pro and Claude Opus 4.8 models.