PaperlayerView exact revision on arXiv

Source-linked AI summary

Dissecting the NVIDIA Volta GPU Architecture via Microbenchmarking

Zhe Jia, Marco Maggioni, Benjamin Staiger, Daniele P. Scarpazza

arXiv:1804.06826v1cs.DCcs.PF

TL;DR

The report addresses the difficulty of keeping pace with NVIDIA GPUs when manufacturers disclose few low-level details. It uses latency probing and a representative matrix-multiplication kernel to examine Volta, finding substantial changes in instruction encoding, memory hierarchy, and processing-unit behavior.

  • Problem

    Rapid NVIDIA GPU evolution and limited manufacturer disclosure make microarchitectural understanding difficult for GPU software designers.

  • Method

    The analysis probes instruction latency and uses a representative fixed-size matrix-multiplication kernel to illustrate architectural behavior.

  • Results

    Volta substantially changes instruction encoding, memory hierarchy, and processing-unit behavior, including improved cache behavior and a scheduler-to-core ratio of 1:16 versus 1:48 on Kepler.

  • Takeaways & Limitations

    The findings support binary-level optimization, customized Volta assemblers, cache-matched working sets, lower miss rates, and higher performance.

  • Takeaways & Limitations

    The authors cannot explain a 7-KiB discrepancy and state that the described optimizations may not justify their human effort except for dominant computational kernels.

Abstract

from arXiv · show

Every year, novel NVIDIA GPU designs are introduced. This rapid architectural and technological progression, coupled with a reluctance by manufacturers to disclose low-level details, makes it difficult for even the most proficient GPU software designers to remain up-to-date with the technological advances at a microarchitectural level. To address this dearth of public, microarchitectural-level information on the novel NVIDIA GPUs, independent researchers have resorted to microbenchmarks-based dissection and discovery. This has led to a prolific line of publications that shed light on instruction encoding, and memory hierarchy's geometry and features at each level. Namely, research that describes the performance and behavior of the Kepler, Maxwell and Pascal architectures. In this technical report, we continue this line of research by presenting the microarchitectural details of the NVIDIA Volta architecture, discovered through microbenchmarks and instruction set disassembly. Additionally, we compare quantitatively our Volta findings against its predecessors, Kepler, Maxwell and Pascal.

Summary

The report addresses the lack of public microarchitectural information about rapidly evolving NVIDIA GPUs by extending microbenchmark-based research to Volta and comparing it with earlier architectures.

  • Rapid NVIDIA architectural progression and limited manufacturer disclosure make microarchitectural details difficult for GPU software designers to track.
  • Prior independent studies used microbenchmarks to characterize instruction encoding and memory-hierarchy geometry across Kepler, Maxwell, and Pascal.
  • The report presents NVIDIA Volta’s microarchitectural details using microbenchmarks and instruction-set disassembly.
  • It also quantitatively compares the Volta findings with Kepler, Maxwell, and Pascal.

Why these details matter

The paper argues that deep architectural understanding is needed to reach peak GPU performance, then demonstrates remaining NVCC optimization opportunities on Volta while identifying substantial human-effort costs.

  • Deep architectural understanding is necessary to optimize software for peak performance, with potentially substantial performance left unused otherwise.
  • Plain CUDA and publicly available programming interfaces do not generally provide access to bare-metal peak performance or the optimization level of NVIDIA libraries.
  • Earlier studies found efficiency opportunities in NVCC-generated machine code and established instruction encoding and microarchitectural behavior as necessary for full potential.
  • For Volta, the authors identify NVCC performance opportunities and demonstrate exploiting them with a minimal example.
  • The described optimizations require substantial human effort and may not generally justify their gains outside tight computational kernels dominating application execution.

6 CHAPTER 1. WHY THESE DETAILS MATTER

The chapter presents Volta-specific binary optimization through a matrix-multiplication kernel, targeting register mapping, bank conflicts, and register reuse. It contrasts these optimizations with portability benefits from standard CUDA tooling.

  • Volta-specific optimizations will not port to future GPU architectures, so the discovery work must be repeated for each new architecture.
  • Standard CUDA libraries and NVCC provide portability and efficiency, with gains generally arriving across GPU generations without comparable extra development effort.
  • The motivating kernel multiplies 8-value slices from matrices A and B and accumulates the products into an 8×8 register-mapped matrix tile.
  • NVCC 9.0 produces machine code with removable register-bank conflicts and opportunities to use the register reuse cache more effectively.
  • On Volta, each instruction can access 64 bits from each of two 64-bit register banks per clock, so FFMA instructions reading three same-bank source registers conflict.
  • The binary rewrite adopts a better register mapping that avoids conflicts and leverages register reuse caches, with the mapping shown by register-bank colors.
  • Using 128 threads, the optimized code improves performance over the original machine code, although the supplied passage does not include the resulting values.

8 CHAPTER 1. WHY THESE DETAILS MATTER

Volta performance can still be improved beyond NVCC-generated code by remapping registers to avoid bank conflicts and exploit register reuse caches. With 128 threads, the illustrated optimization raises throughput by 15.4%.

  • 132.05 to 152.43 GFlops/s per SMX (+15.4%) after improving register mapping and register reuse.The optimized code avoids bank conflicts and reuses register caches as much as possible.
  • Volta’s wider register banks allow mapping to avoid conflicts even when row and column input registers share a bank.A third input register from the other bank can eliminate the conflict.
  • Discovering instruction encoding and bank-conflict behavior was necessary to patch the machine code.The report identifies both topics as primary contributions.

Instructions

Volta changes NVIDIA’s instruction representation by combining each instruction and its control information in one 128-bit word. Its encoding uses more instruction bits than earlier architectures and places the opcode in a different location.

  • Volta uses one 128-bit word for each instruction together with its corresponding control information.This differs from earlier designs that used separate 64-bit instruction and control words.
  • nvdisasm displays the 128-bit word as two 64-bit words with instruction and control information distributed across them.The first displayed word only encodes instruction information; the second encodes both instruction and control information.
  • At least 91 bits encode the instruction, while at least 23 bits encode control information.The remaining 14 bits appeared unused in the experiments.

12 CHAPTER 2. INSTRUCTIONS

Volta associates control information directly with each instruction, unlike pre-Volta bundled control words. These fields manage reuse, dependencies, scheduling delays, and warp-yield behavior.

  • Control information: Volta stores one instruction and its associated control information in each 128-bit word.Pre-Volta architectures bundle one control word with multiple instructions.
  • Control information: Pre-Volta bundles place a control word first, followed by 3 instructions on Pascal and Maxwell or 7 on Kepler.The control word affects scheduling within the bundle.
  • Control information: Volta control sections contain 2 leading zeroes and one 21-bit section, with instruction bits before and after the control information.
  • Control information: Reuse flags select four 8-byte source-operand slots whose values can be stored in four reuse caches to mitigate register bank conflicts.The least significant bit controls the first slot, while the most significant bit controls the fourth.
  • Control information: Dependency barriers track variable-latency operations: write barriers mark producers, and wait masks stall consumers until results are available.The six available barriers can be combined because an instruction may wait on multiple barriers.
  • Control information: Read dependency barriers protect registers used by unbuffered writes to memory from later instructions that write those registers.
  • Control information: The 4-bit stall field specifies a delay of 0 to 15 cycles before issuing the next instruction.Pascal and Maxwell can use a special stall/yield combination for dual issue, whereas Volta has one dispatcher and no observed dual issue.
  • Control information: The yield flag controls whether the scheduler favors the current warp or switches to another, with switching costing one extra cycle.

14 CHAPTER 2. INSTRUCTIONS

Volta distributes warps across four processing blocks using scheduler id = warp id%4, and at least 128 threads are needed to fully use the processing units. Its instruction encoding also differs substantially from earlier GPUs.

  • Scheduler: Volta’s SM has four processing blocks, and each warp’s instructions can access only the processing units within its assigned block.
  • Scheduler: scheduler id = warp id%4 maps each warp to a processing block.A two-warp FFMA benchmark tested throughput while varying the active warp indices.
  • Scheduler: At least 128 threads are required to fully use Volta’s processing units.
  • Instruction encoding: Volta uses more instruction-encoding bits than previous architectures.
  • Instruction encoding: Volta places the opcode in the least significant bits of the first 64-bit code-bundle part, unlike Pascal, Maxwell, and Kepler.
  • Instruction encoding: Volta opcodes range from 10 to 13 bits.
  • Instruction encoding: Volta operands include registers, constant/shared/global memory addresses, or immediate values.Predication uses 4 bits: one negation flag and three predicate-register index bits.

Memory hierarchy

The report characterizes Volta V100’s memory hierarchy, caches, TLBs, register banks, and memory performance through microbenchmarking, while comparing findings with earlier NVIDIA architectures. Volta improves several hierarchy and register mechanisms, including L1 latency and bandwidth, replacement behavior, instruction-cache organization, and register-bank design.

  • Volta’s memory-hierarchy study covers cache and TLB geometry, register-file banks and conflicts, and shared/global-memory performance, with comparisons against Pascal, Maxwell, and Kepler.
  • 900 GB/s HBM2 memory accompanies a 6,144 KiB unified L2 cache, and global-memory data is implicitly cached in L1 and L2.
  • 28 cycles is Volta’s average L1 cache-hit latency versus 84 cycles on Pascal, while measured L1 bandwidth reaches 109.1 bytes per cycle per SM versus 31.3 on Pascal.
  • Volta’s replacement policy prioritizes four cache lines from different sets for eviction and better preserves large arrays from sparse-access eviction than commonly used LRU.
  • Volta’s L1 instruction-cache organization introduces compiler-generated 2-cycle stalls, but twice as many schedulers preserve aggregate per-SM L0 instruction-load throughput.
  • Volta instruction caches use private L0 caches per processing block, private L1 caches per SM, and a shared unified L2; the predecessor L1.5 instruction cache is absent.
  • Volta constant caches include two private SMX levels and a shared unified L2, support broadcasting for uniform warp addresses, and serialize divergent addresses.
  • Volta uses two 64-bit register banks, compared with four 32-bit banks previously, while shared-memory latency is lowest among the examined GPUs and bandwidth efficiency reaches 83.3% versus Pascal’s 69.6%.

Instruction latency and throughput

The report measures Volta instruction latency, atomic performance, tensor-core execution, and floating-point throughput against earlier NVIDIA architectures. Volta generally lowers instruction latency and reaches 90.2% of theoretical half-precision throughput, while global-memory atomic latency shows no obvious improvement.

  • Instruction latency: Volta shows a generalized latency improvement over Pascal, including improved shared-memory atomic latency compared with Pascal and Maxwell.The shared-memory atomic improvement is likely related to Volta’s improved shared-memory latency.
  • Instruction latency: Volta most commonly executes integer and single-precision instructions in 4 cycles, while double-precision instructions take 8 cycles and half-precision instructions take 6 cycles.
  • Atomic throughput: Volta shows no obvious improvement in global-memory atomic latency compared with earlier architectures.
  • Atomic throughput: The V100 achieves its best atomic throughput on multiple SMs without contention among SMs, while it does not lead in contention or single-SM scenarios.From Maxwell to Pascal, aggregate atomic throughput increases substantially.
  • Tensor cores: The compiler translates one wmma::mma_sync operation into four sets of four HMMA instructions, with STEP0–STEP3 computing different positions of C.
  • Floating-point performance: 90.2% of theoretical throughput is reached for half-precision matrix multiplication using tensor cores, while single- and double-precision benchmarks approach peak performance.

Volta boards with NVIDIA NVLink support

The report compares V100 PCIe and NVLink boards with Pascal NVLink hardware, separating interconnect effects from Volta microarchitectural behavior. NVLink improves peer-to-peer bandwidth and some performance measures, while cache and atomic latencies remain architecturally consistent across V100 variants.

  • Peer-to-peer communication: V100 NVLink GPUs provide significantly higher peer-to-peer bandwidth than P100 NVLink boards, and NVLink exceeds PCIe bandwidth.
  • Peer-to-peer communication: Peer-to-peer latencies are substantially similar across GPU generations and interconnect types.
  • V100 comparison: The higher V100 NVLink clock, 1,530 MHz versus 1,380 MHz, explains higher floating-point throughput, shared-memory bandwidth, and multi-block atomic throughput than V100 PCIe.
  • V100 comparison: The V100 NVLink GPU achieves the highest atomic bandwidths among the examined GPUs.
  • V100 comparison: Global-memory bandwidth does not differ between PCIe and NVLink V100 versions.
  • V100 comparison: Volta’s shared microarchitecture gives PCIe and NVLink V100 boards the same measured cache, shared-memory, and atomic-operation latencies.
  • V100 comparison: Global-memory latency is 405 versus 391 clock cycles but 26.5 versus 28.3 ns for NVLink and PCIe, respectively.

Conclusions

Volta differs substantially from earlier architectures in instruction encoding, memory hierarchy, and processing-unit behavior. Its scheduler, caches, and register banks improve instruction throughput, cache performance, and conflict handling, while the findings support binary-level and cache-aware optimization.

  • Conclusions: Volta introduces substantial changes to instruction encoding, memory hierarchy, and processing-unit behavior compared with previous architectures.
  • Conclusions: The scheduler-to-core ratio grows from 1:48 on Kepler to 1:16 on Volta, facilitating higher instruction throughput.
  • Conclusions: The new L0 instruction cache mitigates the penalty associated with Volta’s larger instruction size.
  • Conclusions: Volta’s improved L1 cache lowers latency, raises bandwidth, and reduces miss rates without shared-memory use; redesigned register banks mitigate bank conflicts.
  • Conclusions: The findings enable software designers to optimize Volta code at the binary level and construct customized assemblers for tighter scheduling and higher performance.
  • Conclusions: Disclosed memory-hierarchy information supports selecting working sets that match cache levels, reducing miss rates and improving performance.

Appendix

The appendix provides instruction opcodes encoded for Volta and, for comparison, Pascal.

  • Appendix: The appendix lists opcodes for common instructions in Volta’s instruction encoding and provides Pascal encodings for comparison.
  • Appendix: The appendix includes a section covering floating-point instructions.
Loading 1804.06826v1…