PaperlayerView exact revision on arXiv

Source-linked AI summary

Dissecting the NVidia Turing T4 GPU via Microbenchmarking

Zhe Jia, Marco Maggioni, Jeffrey Smith, Daniele Paolo Scarpazza

arXiv:1903.07486v1cs.DC

TL;DR

GPU designers seeking peak performance face rapidly changing architectures and limited disclosure of microarchitectural details, while low-level optimization can demand substantial effort. The paper updates a microbenchmark-based architectural study for Turing and compares it with earlier NVIDIA GPUs. It reports that Turing extends Volta’s instruction set and can express matrix math in fewer instructions, alongside architectural continuity between the two generations.

  • Problem

    Rapid GPU design changes and manufacturers’ reluctance to disclose microarchitectural details complicate efforts by software designers to extract peak performance, while low-level optimization can require substantial effort.

    Source: Low-level details make a difference · p. 5

  • Method

    The authors update their microbenchmark-based architectural discovery study for Turing and compare its revealed details with previous NVIDIA GPU architectures.

    Source: Conclusions · p. 51

  • Results

    Turing shares Volta’s instruction encoding but extends its instruction set, adds uniform registers and more Tensor Core operand types, and lets nvcc express matrix math in fewer instructions than Volta.

    Source: Conclusions · p. 51

  • Takeaways & Limitations

    Turing and Volta show architectural continuity and together represent a significant departure from Kepler and older generations across instruction encoding, memory hierarchy, and processing-unit behavior.

    Source: Conclusions · p. 51

    For int8 and int4, CUTLASS implementations on T4 reach less than 50% of theoretical throughput.

    Source: §4.4 · p. 45

Abstract

from arXiv · show

In 2019, the rapid rate at which GPU manufacturers refresh their designs, coupled with their reluctance to disclose microarchitectural details, is still a hurdle for those software designers who want to extract the highest possible performance. Last year, these very reasons motivated us to dissect the Volta GPU architecture using microbenchmarks. The introduction in August 2018 of Turing, NVidia's latest architecture, pressed us to update our study. In this report, we examine Turing and compare it quantitatively against previous NVidia GPU generations. Specifically, we study the T4 GPU: a low-power board aiming at inference applications. We describe its improvements against its inference-oriented predecessor: the P4 GPU based on the Pascal architecture. Both T4 and P4 GPUs achieve significantly higher frequency-per-Watt figures than their full-size counterparts. We study the performance of the T4's TensorCores, finding a much higher throughput on low-precision operands than on the P4 GPU. We reveal that Turing introduces new instructions that express matrix math more succinctly. We map Turing's instruction space, finding the same encoding as Volta, and additional instructions. We reveal that the Turing TU104 chip has the same memory hierarchy depth as the Volta GV100; cache levels sizes on the TU104 are frequently twice as large as those found on the Pascal GP104. We benchmark each constituent of the T4 memory hierarchy and find substantial overall performance improvements over its P4 predecessor. We studied how clock throttling affects compute-intensive workloads that hit power or thermal limits. Many of our findings are novel, published here for the first time. All of them can guide high-performance software developers get closer to the GPU's peak performance.

Low-level details make a difference

Deep architectural knowledge can improve GPU performance in specific cases, but optimization may require substantial low-level development effort. Whether the gains justify that effort depends on each developer’s circumstances and performance needs.

Low-level details make a difference

  • The section uses a practical example to show that architectural understanding can yield substantial speedups in specific scenarios, at significant development effort.Source: Low-level details make a difference · p. 5
  • Extreme low-level optimization is rare because its effort must be weighed against developers’ circumstances and pressure for performance.It can involve inline PTX, binary patching to obtain specific SASS, and following undocumented instruction encodings without toolchain support.Source: Low-level details make a difference · p. 5
  • Prior Volta work illustrates a possible payoff: patching compiler-emitted code to use the register cache better raised floating-point arithmetic throughput by 15%.Source: Low-level details make a difference · p. 5

6 CHAPTER 1. LOW-LEVEL DETAILS MAKE A DIFFERENCE

The section motivates replacing cuBLAS’s narrower memory accesses for single-precision Saxpy with 128-bit vectorized accesses. It presents this as a way to increase memory-access throughput on Turing.

6 CHAPTER 1. LOW-LEVEL DETAILS MAKE A DIFFERENCE

  • The proposed Saxpy replacement uses 128-bit vectorized memory accesses, while the examined single-precision cuBLAS implementation uses accesses no wider than 64 bits.The analysis identifies the cuBLAS kernel as axpy_kernel_val in CUDA SDK 10.0.Source: §6 · pp. 6–7
  • Turing has half as many load/store units per SM as Volta, and T4 supports half as many threads per SM as V100, motivating wider words per instruction.Source: §6 · p. 7
  • The example processes four array elements per thread at a time and uses vectorized 128-bit loads and stores.Source: §6 · p. 7

8 CHAPTER 1. LOW-LEVEL DETAILS MAKE A DIFFERENCE

The improved Saxpy implementation uses 128-bit global-memory instructions and outperforms cuBLAS except on trivially small arrays. The example illustrates how instruction-set knowledge can expose optimization opportunities.

8 CHAPTER 1. LOW-LEVEL DETAILS MAKE A DIFFERENCE

  • NVCC emits 128-bit global-memory instructions for the improved implementation, as confirmed by inspecting its SASS code.Source: §8 · p. 8
  • The improved Saxpy code is almost twice as fast as cublasSaxpy for large arrays, except that it does not outperform it for arrays smaller than 20 KiB.Source: §8 · p. 8
  • The example presents in-depth Turing instruction-set and architectural performance knowledge as a source of optimization opportunities for software designers.Source: §8 · p. 8

How Turing encodes instructions

The authors examine Turing’s instruction encoding through disassembly and arbitrary code execution. They report that Turing shares Volta’s encoding format, which combines instruction and scheduling-control information in 128 bits.

How Turing encodes instructions

10 CHAPTER 2. HOW TURING ENCODES INSTRUCTIONS

Turing’s instruction control information encodes scheduling and dependency handling alongside each instruction. Its fields support register reuse and track dependencies for variable-latency operations.

10 CHAPTER 2. HOW TURING ENCODES INSTRUCTIONS

  • At least 23 bits encode control information, while the remaining 14 bits appeared unused and were ignored by cuobjdump and the hardware.The reported behavior applies to both Turing and Volta.Source: §10 · p. 10

2.1 Control information

  • Kepler packs a control word with multiple instructions; Turing and Volta associate control information with one instruction in each 128-bit unit.Pre-Volta bundles contain one control word plus three instruction words on Pascal and Maxwell, or seven on Kepler.Source: §2.1 · p. 10
  • Kepler uses seven 8-bit control sections, Pascal and Maxwell use three 21-bit sections, and Turing and Volta use one 21-bit section per instruction.On Turing and Volta, instruction bits precede and follow the control information.Source: §2.1 · p. 10

2.1. CONTROL INFORMATION

  • The control sections share six fields: reuse flags, wait-barrier mask, read- and write-barrier indices, yield flag, and stall cycles.The passages name the fields but do not specify their widths here.Source: §2.1 · p. 11
  • Reuse flags let instructions save up to four source arguments for reuse, reducing register-file accesses and helping reduce register-bank conflicts.Each thread has a two-way associative CAM for each of four conceptual source-register positions.Source: §2.1 · p. 11
  • Write and wait barriers track variable-latency instructions: consumers wait on the barrier associated with the earlier instruction until its results are available.The wait field is a bitmask because an instruction may wait on multiple barriers.Source: §2.1 · p. 11
  • Read dependency barriers protect against write-after-read hazards by keeping registers unchanged while unbuffered instructions write them to memory.The assembler associates such instructions with a barrier.Source: §2.1 · p. 11

12 CHAPTER 2. HOW TURING ENCODES INSTRUCTIONS

On Turing and Volta, warps map to schedulers by index modulo four, and colliding warps share a scheduler. The section also describes instruction scheduling controls, opcode placement, and operand forms.

12 CHAPTER 2. HOW TURING ENCODES INSTRUCTIONS

  • Each variable-latency instruction writing a register is assigned a write barrier, and later instructions consuming that register wait on the barrier.Later instructions wait until the earlier results are available.Source: §12 · p. 12
  • The stall-cycles field tells the scheduler how long to wait before issuing the next instruction, from 0 to 15 cycles.The described special stall/yield-bit combination can enable dual issue on Pascal and Maxwell, but dual issue is not observed in generated Turing or Volta code.Source: §12 · p. 12
  • When the yield flag is cleared, the scheduler prefers another warp, making the next instruction’s register-reuse flags ineffective and costing one warp-switch cycle.When set, the flag instead favors the current warp.Source: §12 · p. 12

2.2 Processing Blocks and Schedulers

  • A Turing SM has four processing blocks, each with a dedicated warp scheduler and dispatch unit; a warp can access only its assigned block’s processing units.Instructions from a warp are allocated to a specific processing block.Source: §2.2 · p. 12, §2.3 · p. 13

2.3. INSTRUCTION WORD FORMAT

  • Warps with the same index modulo 4 map to the same scheduler on Turing and Volta, and their aggregate throughput drops when two active warps collide.The experiment varies two active warp indices and measures aggregate single-precision throughput; the reported mapping is consistent across both architectures.Source: Table 2.1 · p. 12, §2.3 · p. 13
  • Each workload block needs at least 128 threads to fully utilize the processing units on one Turing or Volta SM.The conclusion follows from the observed mapping of warps to schedulers.Source: §2.3 · p. 13

2.3.1 Opcodes

  • Turing and Volta place opcodes in the least significant bits of the first 64-bit instruction word; Turing opcodes span 10 to 13 bits.This differs from Pascal, Maxwell, and Kepler, which place opcodes in the most significant bits.Source: §2.3.1 · p. 13

2.3.2 Operands

  • Turing instruction operands can be typed registers, constant, shared, or global memory addresses, or immediate values; four bits encode predication.The first predication bit is a negation flag, and the remaining three encode a predicate-register index.Source: §2.3.2 · p. 13

Memory hierarchy

This section characterizes Turing’s memory hierarchy by mapping cache organization and measuring cache, constant, shared, and global memory behavior, including address translation through the TLB hierarchy.

3.1.2 Geometry

  • T4 offers twice P4’s L1 data capacity and twice its bandwidth, but experiments detected 7 KiB less than nominal L1 size on both Turing and Volta.With shared memory configured to 64 KiB, cache misses appeared when the scanned array exceeded 25 KiB; the 7-KiB discrepancy was unexplained and not associated with ECC.Source: §3.1 · p. 18, §3.1.2 · p. 20

3.1.1 Latency and bandwidth

  • T4 L1 hit latency was 32 cycles, versus 82 on P4; Kepler’s was 35 cycles, despite K80 running at 875 MHz versus T4’s 1,590 MHz.Source: §3.1.1 · p. 19

3.1.2 Geometry

  • Turing’s L1 line size and load and update granularity match Volta, Pascal, and Maxwell; its replacement policy aims to protect large arrays from sparse-access eviction.When saturated, Turing randomly evicts four consecutive cache lines (128 B), and a second scan can evict more lines from the same set.Source: §3.1.2 · p. 20, §3.2 · p. 21

3.2 Unified L2 cache

3.1.1 Latency and bandwidth

  • 58.83 bytes/cycle/SM on T4 was 3.7× the P4’s 15.7 bytes/cycle/SM for measured L1 data bandwidth.The comparison is meaningful in cycle counts because the cards run at very similar graphics clock frequencies.Source: §3.1.1 · p. 20, Table 3.2 · p. 23

3.2 Unified L2 cache

3.3.2 Size

  • Turing has three instruction-cache levels, named L0, L1, and L2; it has better inverse throughput than predecessors at the second and third levels.The second level is called L1 on Turing and Volta, but L1.5 on Pascal, Maxwell, and Kepler.Source: §3.3.1 · p. 25, §3.3.2 · p. 25

3.3.3 Organization

  • On Turing, L0 instruction cache is private to a scheduler, L1 is private to an SM, and L2 is shared across all SMs.Source: Figure 3.7 · p. 26
  • Turing and Volta have private L0 instruction caches per processing block and private L1 caches per SM; unified L2 is shared across SMs, and they lack older-generation L1.5 instruction caches.On Pascal, Maxwell, and Kepler, L1.5 instruction caches are private to one SM.Source: §3.3.3 · p. 27

3.4 Constant memory hierarchy

  • Turing has three constant-cache levels; its L1.5 constant cache coincides with the L1 instruction cache, and aggressor instruction sequences produced victim miss rates from 0% to 100%.Longer aggressor instruction sequences caused correspondingly higher miss rates in the constant-array scan.Source: §3.4 · pp. 28–29

3.5 Registers

  • Turing and Volta use two dual-ported register banks per processing block; an FFMA stalls when its three source registers map to the same bank, and Turing supports 64 uniform registers.Turing retains 256 regular registers, and the CUDA-kernel register limit is 256 total, including regular and uniform registers.Source: §3.5.1 · p. 30, §3.5.2 · p. 32, §3.5.3 · p. 32

3.6 Shared memory

  • Only V100 has lower shared-memory latency than T4; V100 and P100 provide the highest theoretical and measured shared-memory bandwidth.Measured shared-memory access latency rises with warp conflicts on all examined GPUs except Kepler, whose dual-ported banks tolerate two aliases without penalty.Source: §3.6.1 · p. 34, §3.6.2 · p. 34, Figure 3.12 · p. 34

3.7 Global memory

  • T4 global bandwidth exceeds P4’s, but its actual-to-theoretical bandwidth ratio is lower: 68.8% versus 84.4%.The comparison attributes T4’s higher global bandwidth to GDDR6 memory.Source: §3.7 · p. 36, Figure 3.13 · p. 36

3.8 TLBs

  • Turing’s L1 data cache is virtually indexed and its L2 data cache physically indexed; its L1 TLB covers 32 MiB and its L2 TLB about 8192 MiB.The L1 TLB has 2 MiB page entries; the L2 TLB coverage matches Volta.Source: §3.8 · pp. 36, 38

Instruction latency and throughput

T4’s arithmetic throughput exceeds P4’s for half-precision and int8 matrix multiplication, while the report also finds instruction-latency, Tensor Core, atomic-operation, and throttling differences across GPU generations. Small-form-factor T4 and P4 boards have higher frequency-per-Watt ratings than full-size GPUs but are more prone to clock throttling.

4.1 Native instructions

  • T4 instruction latencies are generally similar to Volta’s, with no apparent latency improvement over Volta; Turing double-precision instruction latencies exceed 40 cycles.The comparison reports Turing’s integer, single-, and half-precision latencies as similar to Volta’s, while Volta and Turing are generally lower-latency than Pascal and older generations.Source: §4.1 · p. 39

4.2. ATOMIC OPERATIONS

  • T4 is not the throughput leader under contention or within a single SM; its best result is across multiple SMs without contention among SMs.Across architectures, Kepler’s shared-memory atomics are 4× to 8× slower than its global-memory atomics because Kepler lacks hardware support; later architectures support low-latency atomics even under contention.Source: Figure 4.1 · p. 41, §4.2 · p. 42, §4.3 · p. 43

4.3 New Tensor Core instructions

  • Turing Tensor Cores add int8, int4, and int1 operands, and compile the same wmma::mma_sync() invocation into four new HMMA instructions instead of Volta’s 16.The Turing instructions use a new .1688 infix.Source: §4.3 · p. 43

4.4. ARITHMETIC PERFORMANCE

  • In half, single, and double precision, cuBLAS outperforms CUTLASS; for int8, cublasLtMatmul throughput is much higher than cublasGemmEx.Except in double precision, the benchmarks do not reach near-peak performance; CUTLASS int8 and int4 implementations achieve less than 50% of theoretical T4 throughput.Source: §4.4 · p. 45
  • T4 delivers higher half-precision and int8 matrix-multiplication throughput than P4, while their single- and double-precision throughput is similar.The report attributes T4’s half-precision and int8 advantage to Tensor Core use; both inference-oriented boards have only two native FP64 cores per SM.Source: §4.4 · p. 45

4.5 Performance throttling

  • T4 and P4 have higher frequency-per-Watt ratings than full-size GPUs, but are more prone to clock throttling.The report attributes the greater susceptibility to smaller heat sinks and manufacturer power limits of 70W on small-form-factor boards versus 250W on full-size boards.Source: §4.5 · p. 46

4.5.2 Thermal throttling

  • On T4, reaching 85°C adds thermal throttling to power-limit throttling and causes a more dramatic clock step-down.Below 85°C, power-limit throttling already reduces the graphics clock as temperature rises.Source: §4.5 · p. 47, §4.5.2 · p. 49

4.5.3 Power-limit throttling across GPU devices

  • Only T4 and P4 throttled in the repeated-kernel comparison; full-height GPUs could not be driven close enough to their power limits to trigger throttling.T4 and P4 initially ran at their highest supported frequencies, but their clocks fell as temperatures rose.Source: §4.5.3 · p. 49

Conclusions

The conclusions characterize Turing as extending Volta while improving selected T4 capabilities over the Pascal P4. They also identify architectural changes and potential software-optimization uses of the findings.

Conclusions

  • The study refreshes the authors’ microbenchmark-based architectural analysis for Turing and compares its details with earlier NVIDIA architectures.Its principal comparison emphasizes the inference-oriented T4 and its Pascal-based predecessor, the P4.Source: Conclusions · p. 51
  • Turing retains Volta’s instruction encoding but extends its instruction set, adds uniform registers, and supports more Tensor Core operand types.The new instructions let nvcc express matrix math in fewer instructions on Turing than on Volta.Source: Conclusions · p. 51
  • Turing’s memory hierarchy resembles Volta’s, with differences in cache sizes, and Turing and Volta show continuity in instruction encoding, memory hierarchy, and processing-unit behavior relative to older generations.The paper describes Turing and Volta together as a significant departure from Kepler and older generations.Source: Conclusions · p. 51
  • Turing’s scheduler-to-cores ratio grew from 1:48 in Kepler to 1:16, alongside a trend in instruction throughput; its L0 instruction cache mitigates penalties from longer instructions.The improved L1 data cache offers lower latency and higher bandwidth; a new replacement policy reduces cache misses without shared memory, and two dual-ported register banks facilitate prevention of bank conflicts compared with Pascal’s four single-ported banks.Source: Conclusions · p. 51
  • T4 delivers higher L1-cache and global-memory bandwidth and higher half-, int8-, and int4-precision matrix-multiplication throughput than P4, while single- and double-precision performance is comparable.The paper attributes comparable single- and double-precision performance to the GPUs having the same number of cores and similar clock frequencies.Source: Conclusions · p. 52
  • Instruction-encoding and memory-hierarchy findings can inform binary-level code optimization, customized SASS assemblers, and working-set choices matched to cache levels.The authors state that these approaches may enable more tightly scheduled code and reduced miss rates, improving overall performance.Source: Conclusions · p. 52

Appendix

The appendix provides opcode information for common instructions, comparing Turing’s encoding with Pascal’s and Volta’s and organizing entries by instruction category.

Appendix

  • The appendix lists opcodes for common instructions in Turing’s encoding alongside Pascal’s and Volta’s for comparison.Its tables use columns labeled Instruction, Pascal, Volta, and Turing.Source: Appendix · pp. 53, 56
  • The listed instruction categories include floating-point instructions and their continuation.The appendix also labels conversion instructions as continued.Source: Appendix · pp. 53, 54, 57
  • The appendix also includes integer, control, and other instruction entries, including continuation sections.The supplied table headings identify instruction columns for Pascal, Volta, and Turing.Source: Appendix · pp. 55, 58, 59

The Authors

The author biographies describe four researchers’ roles, education, and areas of professional interest.

The Authors

  • Zhe Jia is a senior R&D engineer at Citadel whose interests include performance optimization of deep-learning, numerical-modeling, and atmospheric-simulation workloads.He previously worked at Alicloud and Microsoft Beijing and earned physics and meteorology degrees from Peking University.Source: The Authors · p. 65
  • Marco Maggioni is a Citadel senior R&D engineer whose Ph.D. research focused on sparse linear algebra and convex optimization on GPUs.His research established records for sparse matrix-vector multiplication GPU kernels in research and industry.Source: The Authors · p. 65
  • Jeffrey Smith is a senior software engineer at Citadel Securities focused on hardware architectural analysis, latency-focused systems, and parallel software frameworks.He holds a computer science bachelor’s degree from Rose-Hulman Institute of Technology.Source: The Authors · p. 65
  • Daniele Paolo Scarpazza leads Citadel’s High-Performance Computing group and focuses on quantitative performance analysis and optimization across parallel architectures.He previously held research roles at D. E. Shaw Research, IBM, and Pacific Northwest National Laboratory, and is a Gordon Bell prize co-recipient.Source: The Authors · p. 65
Loading 1903.07486v1…