Skip to content
Markdown

Register bank conflicts and the instruction control word

Scope: the register file's physical bank layout and the per-instruction control word (register-reuse flags, dependency-barrier indices, stall cycles, yield flag) that the compiler and assembler emit beneath CUDA C++, PTX, and even most hand-written SASS. Covers register bank conflicts, a silent, correctness-preserving throughput cost distinct from the shared-memory bank conflicts in Shared Memory, Bank Conflicts, and Tiling; how the control-word encoding evolved from Kepler's 8-bit code shared across an instruction bundle to Volta's combined 128-bit instruction-plus-control word; and the reverse-engineering method (disassembly plus targeted microbenchmarks) the field uses to discover any of this, since NVIDIA does not document it.

flowchart LR
  subgraph RF["Register file, per SM"]
    direction TB
    REGS["32-bit registers R0...Rn"]
    BANKS["Physical banks: N banks, W bits wide<br/>Kepler/Maxwell/Pascal: 4 banks x 32-bit<br/>Volta: 2 banks x 64-bit"]
    REGS --> BANKS
  end
  BANKS -->|"one operand per bank per cycle (narrow) or two 32-bit operands per bank per cycle (wide)"| OC["Operand collector"]
  OC --> ALU["FFMA / ALU pipe"]

  subgraph CW["Per-instruction control word"]
    direction TB
    REUSE["Register reuse-cache flags"]
    BARR["Wait / read / write dependency-barrier index"]
    STALL["Stall-cycle count + yield flag"]
  end
  CW -->|"one word per instruction on Volta; shared across a 7- or 3-instruction bundle pre-Volta"| SCHED["Warp scheduler / processing block"]
  SCHED --> OC

What it is

Below the SASS you read in Inline PTX and SASS-level Tuning, two more hardware structures decide a kernel's real throughput, and neither is documented by NVIDIA. Both were discovered the same way: independent researchers disassembled compiled kernels, flipped bits, and re-disassembled to see what changed.

Register bank conflicts. The register file is not a flat, uniformly fast array; it is physically divided into banks, and each bank can serve only a limited number of operands per clock. Zhang et al. measured this on an NVIDIA Kepler K20m: the register file has 4 banks, 32 bits wide, and an FFMA instruction (d = a*b+c, three source registers) that draws all three sources from the same bank cannot fetch them simultaneously; the accesses serialize.1 Jia et al. found Volta redesigned this to 2 banks, 64 bits wide, so one bank now serves two 32-bit operands per cycle; a Volta FFMA only conflicts when all three source registers land in the same bank, since the wider port already absorbs any two of them for free.2 The conflict is invisible at the CUDA C++, PTX, or even algorithm level: it produces correct results, just slower ones, and no compiler diagnostic names it.

The instruction control word. Alongside each instruction (or bundle of instructions, pre-Volta), the assembler emits scheduling metadata the hardware consumes to avoid data hazards without extra runtime logic. Jia et al. document six fields, in a 21-bit section on Volta, Pascal, and Maxwell:3

Width (bits) 4 6 3 3 1 4
Field Reuse flags Wait barrier mask Read barrier index Write barrier index Yield flag Stall cycles
  • Reuse flags: 4 bits, one per source-operand slot. Volta, Pascal, and Maxwell each have 4 register reuse caches; setting a flag stashes that operand's value in its slot's cache for the next instruction to consume, sidestepping a second bank fetch (and a possible conflict) entirely. This is the .reuse suffix visible in nvdisasm/cuobjdump output.
  • Dependency barriers: most instructions have fixed, staticaly-schedulable latency, but memory and shared-resource ops do not. The assembler assigns a variable-latency instruction's result to one of 6 barriers (its write barrier index); a later instruction that consumes that result sets the matching bit in its wait barrier mask and stalls until the hardware signals completion. Read barrier index protects the same hazard in the other direction, for instructions that read a register that a later write must not clobber before the read completes.
  • Stall cycles + yield flag: a 4-bit count (0-15) telling the scheduler how long to wait before issuing the next instruction from this warp, plus a 1-bit hint for whether the scheduler should keep favoring this warp or switch to another.

On Volta, each 128-bit instruction word carries one instruction plus its own control section: at least 91 bits of instruction, a 2-bit zero header plus the 21-bit, six-field section above (23 bits total), and roughly 14 bits unused.4 Pre-Volta, the layout is different: a 64-bit control word (6 zero high bits, 2 zero low bits, 7 sections of 8 bits each on Kepler) precedes a bundle of instructions and governs all of them, rather than each instruction carrying its own.3 Zhang et al.'s independent, from-scratch reverse engineering of that Kepler control code found the same shape inside its 8-bit-per-instruction section: bits 0-3 are a 4-bit stall-cycle count with the identical 0-15 semantics, and bits 4, 5, and 7 flag shared-memory, global-memory, and texture-cache dependency barriers.5 Two research groups, six years and two architectures apart, independently rediscovered the same shape: a small stall-cycle counter governing issue timing, plus barrier bits guarding variable-latency dependencies.

Kepler's control code additionally carried an explicit, compiler-scheduled dual-issue mode: a warp scheduler could issue two instructions from the same warp in one cycle by setting specific stall/control-code combinations, and Zhang et al. reverse engineered which combinations triggered it.6 Jia et al. checked for the same mechanism on Volta and did not find it: "on Volta there is only one dispatcher in a processing block, and we do not observe dual issue in the generated code."7 Volta's SM is partitioned into 4 processing blocks, with scheduler_id = warp_id % 4; Jia et al. confirmed this by running FFMA sequences on different warp-index combinations and measuring per-block throughput, finding at least 128 resident threads are needed to keep all four processing blocks fed.8 This corroborates, from an independent, earlier-generation measurement, the four-scheduling-sub-partition structure GPU Execution Model: SMs, Warps, and SIMT describes for current hardware.

Why use it

Both papers open with the same justification: the gap between a naive CUDA kernel and NVIDIA's own cuBLAS/cuDNN is not explained by algorithm or occupancy tuning alone. Jia et al. put it directly: "there is consensus that the high degree of optimization found in NVIDIA's libraries such as cuBlas and cuDNN is inaccessible to authors of CUDA code, even when they use inline PTX assembly. We believe that manufacturers write optimized libraries as low-level SASS code, possibly with assemblers which they do not make available to the public."9 Register bank layout and the control word are exactly the layer where that gap lives.

The measured numbers back this up:

  • Kepler (Zhang et al.): starting from a conventionally tuned SGEMM baseline (register blocking, double buffering, unrolling, but no assembly-level tuning), eliminating register bank conflicts alone contributed about 10% of the total gain; wide LD.128 loads contributed 27-35%; texture-cached LDG.128 contributed 5-12%; and forcing dual-issue via control codes contributed 84-106%, the single largest lever. Combined, this took SGEMM from the baseline to 2.6x faster, reaching 3.1 Tflop/s at 88% efficiency on a Tesla K20m, 15% faster than cuBLAS 7.0's 2.7 Tflop/s at 76%.10 The same optimizations applied to a convolution kernel beat cuDNN 4.0 by 39-62% across three CNN configurations (AlexNet, VGG-A, OverFeat).11
  • Volta (Jia et al.): hand-remapping only the register allocation of an 8x8 register-blocked matrix-multiply inner loop, with no algorithmic change, raised measured throughput from 132.05 to 152.43 GFlops/s per SMX, a 15.4% gain, purely by moving operands into different physical banks.2

Neither result is a compiler flag or a documented tuning knob; both required disassembling the compiled kernel to see the conflict at all.

When to use it (and when not)

Reach for register-bank and control-word-level reasoning when:

  • You are hand-authoring a from-scratch, register-blocked, compute-bound kernel (a GEMM- or convolution-like inner loop) and a profiler already confirms it is issue-bound at high occupancy (see the Not Selected-dominant regime in Instruction-Level Parallelism and Warp Stall Analysis), yet throughput still falls short of the vendor library or the roofline compute bound.
  • You have already exhausted tiling, coalescing, occupancy, and ILP tuning, and the remaining gap is specifically in per-instruction issue efficiency, not memory bandwidth or latency.
  • You are reading nvdisasm/cuobjdump output (per Inline PTX and SASS-level Tuning) and need to interpret the .reuse suffixes, stall annotations, or dependency-barrier markers you see there.

Skip it when:

  • You call cuBLAS, cuDNN, CUTLASS, Triton, or torch.compile. These already fold register allocation and instruction scheduling in, generated or hand-tuned by people with access to information this page had to reconstruct by disassembly. Re-deriving it in application code duplicates work library authors already did better.
  • You have not profiled. Both source papers describe substantial, architecture-specific manual effort for a percentage-level gain on a single kernel; Jia et al. state plainly that this effort "may not, in general, be worth the gains it produces, except for tight computational kernels that vastly dominate an application's execution time."12
  • You are tempted to port Zhang et al.'s exact Kepler dual-issue control-code pattern forward. Jia et al.'s own Volta measurement shows that specific mechanism gone: one dispatcher per processing block, no software-visible dual issue.7 What generalizes is the underlying lesson (register operand placement relative to the physical bank layout is a real, silent cost on every architecture examined), not the Kepler-specific bit pattern that exploited it.
  • You assume any of these bit layouts or bank counts hold on the next GPU generation. Both papers are explicit that their findings are architecture-specific and must be rediscovered: "our optimizations are specific to the Volta architecture. Their performance gains will not port to future GPU architectures. Our discovery work needs to be repeated anew on each future architecture."12

Architecture

The mermaid diagram above is the full picture already: registers are grouped into physical banks that gate the operand collector feeding the ALU pipes, and a control word (embedded per-instruction on Volta, shared across a bundle before it) tells the scheduler when to issue the next instruction and which dependency barriers to wait on. Everything below builds on that.

How to use it: read the disassembly, compute the bank, remap

Start from the SASS produced by cuobjdump -sass or nvdisasm (commands and flags in Inline PTX and SASS-level Tuning). Look for the register indices in each source-operand position of the hot instruction, and for a .reuse suffix indicating the compiler already cached that operand.

A minimal, adversarially-validated reference for the bank arithmetic itself, checked against both papers' own worked examples rather than invented from scratch:

# Kepler: register -> bank, as measured on an NVIDIA Tesla K20m.
# Zhang et al., "Understanding the GPU Microarchitecture to Achieve Bare-Metal
# Performance Tuning" (PPoPP'17), Table 3. Hardcoded, not a formula: the paper
# calls the table partial and empirical, so no general rule is claimed beyond
# these measured entries.
KEPLER_BANK_TABLE = {
    0: 0, 2: 0, 8: 0, 10: 0, 16: 0, 18: 0, 24: 0, 26: 0,
    1: 1, 3: 1, 9: 1, 11: 1, 17: 1, 19: 1, 25: 1, 27: 1,
    4: 2, 6: 2, 12: 2, 14: 2, 20: 2, 22: 2, 28: 2, 30: 2,
    5: 3, 7: 3, 13: 3, 15: 3, 21: 3, 23: 3, 29: 3, 31: 3,
}


def volta_bank(register_index: int, num_banks: int = 2) -> int:
    # Jia et al., Ch. 3, Section 3.5: "the bank of a register is the register's index modulo 2."
    return register_index % num_banks


def kepler_conflict(source_regs: list[int]) -> str:
    """Kepler: 32-bit-wide banks serve one operand per bank per cycle, so ANY
    two source registers sharing a bank already serialize (Zhang et al., Table 2:
    2-way conflicts cost 2.33% throughput, 3-way cost 17.17%)."""
    banks = [KEPLER_BANK_TABLE[r] for r in source_regs]
    counts = {b: banks.count(b) for b in set(banks)}
    m = max(counts.values())
    return "3-way" if m >= 3 else ("2-way" if m == 2 else "none")


def volta_conflict(source_regs: list[int]) -> str:
    """Volta: 64-bit-wide banks serve TWO 32-bit operands per bank per cycle, so
    two-of-three sharing a bank is free. A conflict needs all 3 sources in one
    bank (Jia et al., Ch. 1: 'any FFMA instruction that accesses the same bank with all its
    3 source registers is said to have a bank conflict')."""
    banks = [volta_bank(r) for r in source_regs]
    return "3-way" if len(set(banks)) == 1 else "none"


# Kepler ground truth: Zhang et al., Table 2 (FFMA/FMUL throughput by register
# combination), cross-checked against Table 3's bank assignments.
KEPLER_CASES = [
    ([4, 1, 0], "none"),    # FFMA R5,R4,R1,R0  -> Conf 0
    ([2, 1, 0], "2-way"),   # FFMA R5,R2,R1,R0  -> Conf 2-way (R2,R0 both bank 0)
    ([9, 3, 1], "3-way"),   # FFMA R5,R9,R3,R1  -> Conf 3-way (R9,R3,R1 all bank 1)
    ([1, 0], "none"),       # FMUL R4,R1,R0     -> Conf 0
    ([2, 0], "2-way"),      # FMUL R4,R2,R0     -> Conf 2-way
]
for regs, expected in KEPLER_CASES:
    got = kepler_conflict(regs)
    assert got == expected, f"Kepler {regs}: expected {expected}, got {got}"

# Volta ground truth: Jia et al., Ch. 1, Figure 1.1 / Table 1.1. The three
# source operands of each flagged FFMA (an accumulator pattern d = a*b+c reuses
# the destination register as one source).
VOLTA_CONFLICT_CASES = [
    [16, 12, 80],   # "R16, R12 and R80 are in bank 0"
    [25, 13, 81],   # "R13, R81 and R25 are in bank 1"
    [31, 11, 81],   # "R11, R81 and R31 are in bank 1"
    [43, 15, 83],   # "R43, R83 and R43 are in bank 1"
]
for regs in VOLTA_CONFLICT_CASES:
    got = volta_conflict(regs)
    assert got == "3-way", f"Volta {regs}: expected 3-way conflict, got {got}"

# The paper's hand-remapped fix for the first case (Table 1.1, optimized column):
# two of the three sources now share a bank, but Volta's wider port absorbs
# that for free; only all-3-share is a real conflict.
VOLTA_FIXED = [17, 12, 80]
assert volta_conflict(VOLTA_FIXED) == "none", "Volta's 64-bit port should absorb a 2-of-3 share"

print("all register-bank-conflict assertions passed")

Running this prints all register-bank-conflict assertions passed; every case is a literal worked example from one of the two papers, not a synthetic input, so passing it demonstrates the model matches published, measured hardware behavior rather than an invented rule. Adapt KEPLER_BANK_TABLE / volta_bank to your actual target generation's own published or measured mapping before trusting it on anything newer than Volta; neither paper's mapping is asserted to hold beyond the hardware it measured.

How to integrate it

This is a static-analysis step, not a runtime one: run the bank check over the register indices you read out of the disassembly for the loop body that dominates your kernel's time, before deciding whether to touch anything. If the compiler's own allocation is already conflict-free (as the KEPLER_CASES/VOLTA_CONFLICT_CASES "none" outcomes above show is common), stop; there is nothing to fix. If it is not, the two remaining levers, in order of preference, are:

  1. Nudge the compiler. Change register pressure (--maxrregcount, __launch_bounds__, tile-size choice) and re-disassemble; a different register range can shift the allocation out of conflict without touching a single instruction by hand.
  2. Hand-patch the SASS or PTX. Both papers' largest gains came from rewriting the register assignment directly in the disassembled machine code, exactly the kind of last-resort intervention Inline PTX and SASS-level Tuning already frames as technical debt: pinned to one toolchain, one driver, one architecture.

How to run it in production

In production, prefer cuBLAS, cuDNN, CUTLASS, Triton, or torch.compile; this is precisely the tuning they already encode, generated or hand-written by people with lower-level access than a public CUDA programmer has. Reserve hand-tuning at this level for a narrow, dominant, from-scratch kernel where profiling has confirmed real headroom below the vendor library or the roofline ceiling, and treat any hand-patched register layout or control-word tweak as pinned to the exact toolkit and driver version it was validated against, not portable code.

How to maintain it

Re-run the disassembly-and-bank-check after every CUDA toolkit or driver upgrade, and after any source change near the hot loop, even one that looks unrelated. NVCC's register allocator is free to renumber registers on any recompile; it can silently reintroduce a conflict it previously avoided, or remove one a hand patch was built around, with no diagnostic pointing at the regression. Pin the verified register layout behind an automated check (assert the compiled SASS's operand banks still match what was validated, the same way the code above validates against each paper's published examples) rather than a one-time manual pass, and budget for full rediscovery, not just re-verification, on the next GPU generation: neither Zhang et al.'s Kepler table nor Jia et al.'s Volta modulo-2 rule is asserted by its own authors to hold beyond the hardware it measured.

Failure modes

  • Silent conflict, silent regression. A register bank conflict produces correct output at lower throughput; there is no compiler warning and no crash. The only symptom is a throughput number that does not line up with occupancy, coalescing, or ILP counters, which reads as "already at the compute ceiling" when it is not.
  • No dedicated Nsight Compute counter. Unlike shared-memory bank conflicts, which Nsight Compute names directly (see Shared Memory, Bank Conflicts, and Tiling), neither source paper describes a profiler counter for register bank conflicts. The only path to detecting one is the disassembly-plus-bank-arithmetic method above.
  • Treating an architecture-specific mapping as portable. Applying Kepler's 4-bank table or Volta's modulo-2 rule to a different generation without re-measuring. Both papers explicitly disclaim this; carrying either forward unchanged is the single most likely way to waste the effort this page describes.
  • Reintroducing the conflict on recompile. A hand-tuned register allocation is only valid for the exact compiler output it was checked against; an unrelated code change or toolkit bump can shift the allocator's output and quietly reopen the conflict.

References

Related: GPU Execution Model: SMs, Warps, and SIMT · GPU Memory Hierarchy · Inline PTX and SASS-level Tuning · Instruction-Level Parallelism and Warp Stall Analysis · Shared Memory, Bank Conflicts, and Tiling · CUTLASS: Templated GEMM and Kernel Building Blocks · Glossary

  1. Zhang et al., PPoPP'17, Section 3.3 and Table 3: Kepler's register file has 4 banks, 32 bits wide; Observation 1 measures 2-way conflicts costing 2.33% throughput and 3-way conflicts costing 17.17%, in single-issue mode. ↩

  2. Jia et al., arXiv:1804.06826, Chapter 1 and Section 3.5: Volta's register file has 2 banks, 64 bits wide (bank = register index mod 2); a conflict requires all 3 FFMA source registers in one bank. Hand-remapping registers in the paper's worked example raised throughput from 132.05 to 152.43 GFlops/s per SMX (+15.4%). ↩↩

  3. Jia et al., arXiv:1804.06826, Section 2.1 and 2.2: the 6-field, 21-bit control section (reuse flags, wait barrier mask, read/write barrier index, yield flag, stall cycles) documented for Volta, Pascal, and Maxwell; Volta's SM is partitioned into 4 processing blocks with scheduler_id = warp_id mod 4. ↩↩

  4. Jia et al., arXiv:1804.06826, Chapter 2: Volta's 128-bit instruction word carries at least 91 instruction bits, a 2-bit zero header plus the 21-bit control section (23 bits total), and roughly 14 unused bits; pre-Volta architectures instead precede a bundle of instructions with one shared 64-bit control word (7 instructions per bundle on Kepler, 3 on Pascal/Maxwell). ↩

  5. Zhang et al., PPoPP'17, Section 3.3, Observation 2: within Kepler's 8-bit-per-instruction control code, bits 0-3 are a 4-bit stall-cycle count (0x00 = 16-cycle suspend, 0x2n = n-cycle suspend), and bits 4, 5, 7 flag shared-memory, global-memory, and texture-cache dependency barriers respectively. ↩

  6. Zhang et al., PPoPP'17, Section 3.3, Observation 2-3: Kepler's control code also encodes a compiler-scheduled dual-issue mode (0x20 = single-issue, 0x04/0x05 pairing = dual-issue), reverse engineered by sweeping control-code values and measuring FFMA throughput and latency. ↩

  7. Jia et al., arXiv:1804.06826, Section 2.1: "On Pascal and Maxwell, if the combination of this field and the yield flag contain a special combination of bits, the two dispatchers in a processing block can dispatch two consecutive instructions of a warp at the same time (dual issue). On Volta there is only one dispatcher in a processing block, and we do not observe dual issue in the generated code." ↩↩

  8. Jia et al., arXiv:1804.06826, Section 2.2, Table 2.1: single-precision throughput measured while varying which warps issue FFMA sequences across Volta's 4 processing blocks; a single processing block in isolation measures roughly 42 GFLOPS versus roughly 66 GFLOPS when spread across schedulers, indicating at least 128 resident threads are needed to fully use one SM's processing units. ↩

  9. Jia et al., arXiv:1804.06826, Chapter 1, "Why these details matter." ↩

  10. Zhang et al., PPoPP'17, Abstract and Section 5.2.2: incremental optimization breakdown (register-bank-conflict elimination ~10%, LD.128 27-35%, LDG.128 5-12%, dual-issue 84-106%) reaching 3.1 Tflop/s at 88% efficiency on a Tesla K20m versus cuBLAS 7.0's 2.7 Tflop/s at 76%. Table 4 (Section 3.3) separately measures dual-issue FFMA throughput at 186.35 of a 192 theoretical ops/cycle/SM ceiling (97.06%); Section 5.2.3 states FFMA throughput "can achieve 97.67% efficiency" against the same 192-op ceiling, a related but not arithmetically identical figure from a different part of the paper. ↩

  11. Zhang et al., PPoPP'17, Section 6 and Figure 11: the same microarchitectural optimizations applied to a convolution kernel outperform cuDNN 4.0 by 39%, 46%, and 62% on AlexNet, VGG-A, and OverFeat respectively, on a Tesla K20m at batch size 128. ↩

  12. Jia et al., arXiv:1804.06826, Chapter 1: the paper's own stated limitations, that its optimizations require substantial human effort worthwhile mainly for kernels that dominate an application's runtime, and that the findings are Volta-specific and must be rediscovered on each future architecture. ↩↩