# 31 — Turing TU104 parameter dossier (Phase 1 model inputs)

> **Status:** research dossier for the GreenEye-GPGPU-Sim Phase 1 model
> (WP1.1, `handoff-simulator.md` §5). Public-source only; every row carries a
> citation. Values that are **scaled** from a measured board to a different
> board or clock are marked `[scaled]`; values that are **inferred** rather
> than directly measured/published are marked `[inferred]`. Nothing in this
> document is a silicon-accuracy claim — measure before claiming (docs/06).
>
> **Scope:** modelling inputs for the full TU104 (RTX 2080 Super class,
> 48 SM) and the T4 board variant (40 SM) measured by the microbenchmarking
> literature. Not in scope: GSP/MCU parameters (docs/32+, Phase 1.5).

## 1. Sources and citation labels

| Label | Source | What it provides |
|---|---|---|
| **[T4]** | Z. Jia, M. Maggioni, J. Smith, D. P. Scarpazza, *Dissecting the NVidia Turing T4 GPU via Microbenchmarking*, arXiv:1903.07486 (2019), 65 pp. | Microbenchmarked TU104 (T4 40-SM board) memory hierarchy geometry/latency/bandwidth, instruction latencies, atomics, tensor-core throughput, SM scheduling/encoding facts. |
| **[TWP]** | NVIDIA, *NVIDIA Turing GPU Architecture* whitepaper, WP-09183-001_v01 (2018), https://www.nvidia.com/content/dam/en-zz/Solutions/design-visualization/technologies/turing-architecture/NVIDIA-Turing-Architecture-Whitepaper.pdf | Per-SM composition, TU102/TU104/TU106 product tables, clock/bandwidth/L2/register-file, NVLink 2, tensor/RT core counts. |
| **[AMP]** | H. Abdelkhalik, Y. Arafa, N. Santhi, A.-H. Badawy, *Demystifying the Nvidia Ampere Architecture through Microbenchmarking and Instruction-level Analysis*, arXiv:2208.11174 (2022) | Ampere A100 numbers, used here only for **cross-generation comparisons** to Turing (global-memory latency 434 cycles attributed to Turing via the paper's reference [13]). |
| **[2080S]** | NVIDIA RTX 20-series product table as published on Wikipedia *GeForce 20 series* (retrieved 2026-10-03), citing NVIDIA product pages and TechPowerUp GPU database: https://en.wikipedia.org/wiki/GeForce_RTX_20_series | Board-level RTX 2080 Super (TU104-450) clocks, memory rate, bandwidth, power; TU104 die data. |
| **[VOLTA]** | Z. Jia et al., *Dissecting the NVIDIA Volta GPU Architecture via Microbenchmarking*, arXiv:1804.06826 (2018) | Volta comparison points quoted inside [T4]. |

> **Provenance rule.** This document cites public material and recorded
> observations; register-level details are cross-checked against the tracked
> field catalogue.

## 2. Method and scaling caveats

1. **[T4] is a 40-SM, 16 GiB GDDR6 (10 Gbps) TU104 variant running at
   1,590 MHz** [T4 §3, Table 3.1]. The Phase-1 target is the full
   **TU104-450 / RTX 2080 Super: 48 SM, 8 GiB GDDR6 (15.5 Gbps), 1650/1815
   MHz** [2080S].
2. Cache geometry **per SM** and latencies **in cycles** are
   microarchitecture properties and apply to all TU104 variants; per-chip
   aggregate bandwidth and absolute GFLOPS/GB/s **must be scaled** by SM
   count and clock as noted. All scaled rows are marked `[scaled]`.
3. [T4] measured "actual" bandwidths under its own microbenchmarks (pointer
   chase / streaming loops). These are *achievable* numbers on that setup;
   they are not peak specifications. Keep both rows distinct in the model.
4. Instruction-latency measurements are dependent-issue latencies on a
   statically scheduled core; independent-issue throughput may differ
   (see §7).
5. The GDDR6 DRAM **timing parameters** (tCL/tRCD/tRP/tRAS/refresh) are not
   published for TU104; the model must use a JEDEC-class GDDR6 timing set
   and treat it as a calibrated unknown (see §10).

## 3. Chip and device parameters

### 3.1 TU104 die / board

| Parameter | Value | Source | Notes |
|---|---|---|---|
| GPU chip | TU104 | [TWP], [2080S] | Full die supports 48 SMs (Quadro RTX 5000, RTX 2080 Super); 46-SM (RTX 2080/2070S) and 40-SM (Tesla T4) also ship. |
| Transistors | 13.6 B | [TWP] Table 4, [2080S] | |
| Die size | 545 mm² | [TWP] Table 4, [2080S] | |
| Process | TSMC 12 nm FFN | [TWP] | |
| GPC / TPC / SM (full) | 6 GPCs / 24 TPCs / 48 SMs | [TWP] (Quadro RTX 5000 column) | 2 SMs per TPC. |
| T4 board | 5 GPCs / 20 TPCs / 40 SMs | [TWP] Table 5, [T4] Table 3.1 | The measured board. |
| Max graphics clock | 1,590 MHz (T4) | [T4] Table 3.1, [TWP] Table 5 | |
| Base/boost clock (2080 Super) | 1,650 / 1,815 MHz | [2080S] | Board-dependent; FE is 1,650/1,815, reference differs. |
| Base/boost clock (Quadro RTX 5000) | 1,620 / 1,815 MHz | [TWP] Table 4 | |
| Clock throttle behavior | compute workloads can throttle below app clock | [T4] §4.5 | Model needs a clock envelope, not a fixed clock. |

### 3.2 SM (per-SM parameters; applies to all TU104 variants)

| Parameter | Value | Source | Notes |
|---|---|---|---|
| FP32 CUDA cores / SM | 64 | [TWP] §TU102, Table 1/4 | 4 processing blocks × 16 FP32. |
| INT32 cores / SM | 64 | [TWP] §Turing SM | 4 blocks × 16 INT32; concurrent execution with FP32 on a separate datapath. |
| Tensor cores / SM | 8 (2nd gen) | [TWP] Table 4 | 2 per processing block. |
| RT cores / SM | 1 | [TWP] §TU102 | |
| Texture units / SM | 4 | [TWP] §TU102 | |
| Register file / SM | 256 KB (64 Ki × 32-bit) | [TWP] Table 4, [T4] Table 3.1 | 16,384 32-bit registers per processing block [T4] §3.5.1 = 64 Ki per SM. |
| Register banks | 2 banks, 64-bit wide | [T4] Table 3.1 | Dual-ported; conflicts only when all 3 FFMA sources are in the same bank [T4] §3.5.1. |
| Warp schedulers / dispatcher | 4 blocks each 1 scheduler + 1 dispatch unit | [TWP] §Turing SM, [T4] §2.2 | No dual issue observed [T4] §2.1. |
| Warp→scheduler mapping | `scheduler = warp_id % 4` | [T4] §2.2, Table 2.1 | Same rule as Volta. |
| Threads / SM (max) | 1,024 (32 warps) | [T4] Table 3.1 | This is **half** of Volta's 2,048 [T4 §1]. |
| Independent FP32+INT32 issue | yes | [TWP] §Turing SM | The key Turing SM addition. |
| L1 data / shared memory (unified) | 96 KB configurable | [TWP] §TU102 | Whitepaper figure for the TU10x SM. |
| Shared memory (max user configuration) | 64 KB (T4 measured config) | [T4] Table 3.1, §3.1.2 | Whitepaper: 32 or 64 KB L1 data usage [T4 §3.1.2]; model as 0–64 KB shared partitioning of a 96 KB unified array. |
| L0 instruction cache | ~16 KB | [T4] Table 3.1 | |
| L1 instruction cache | ~46 KB | [T4] Table 3.1 | |

## 4. Memory hierarchy (cycle latencies and geometry)

All cycle numbers are measured on the T4 at 1,590 MHz; latencies in cycles are
treated as architecture-level (see caveat §2.2). "Hit latency" rows are
dependent-issue latencies.

### 4.1 Latency table

| Access | Latency (cycles) | Source | Notes |
|---|---|---|---|
| L1 data cache hit | 32 | [T4] Table 3.1, Fig. 3.5 | Volta 28, Kepler 35. |
| L1 data – same line secondary access | 32 | [T4] Fig. 3.5 | |
| L2 hit (L1 miss) | ~188 | [T4] Table 3.1, Fig. 3.5 | Volta ~193. |
| L2 miss, TLB hit (DRAM) | 296 | [T4] Fig. 3.5 | |
| L2 miss + TLB miss (first touch) | 616 | [T4] Fig. 3.5 | First access in the chase. |
| Global memory (Ampere paper cross-check) | 434 (Turing) vs 290 (Ampere A100) | [AMP] §V-B, Table IV, citing [13] | Independent corroboration of the Turing DRAM latency band. |
| L1 constant broadcast | ~26 | [T4] Table 3.1 | |
| L1.5 constant broadcast | 92 | [T4] Table 3.1 | ~46 KiB L1.5 const cache. |
| L2 constant broadcast | ~215 | [T4] Table 3.1 | |
| Shared memory no-conflict | 19 | [T4] Table 3.1 | Volta 19. |
| Atomic shared (no contention) | 8 | [T4] Table 4.2 | see §7.3. |
| Atomic global (no contention) | 76 | [T4] Table 4.2 | T4 is slower than V100 (36). |

### 4.2 Geometry and bandwidth

| Parameter | Value | Source | Notes |
|---|---|---|---|
| L1 data size (nominal) | 32 or 64 KiB | [T4] Table 3.1, §3.1.2 | |
| L1 detectable size | "7 KiB short of nominal"; misses observed once working set > 25 KiB with 64 KiB shared configured | [T4] §3.1.2, Table 3.3 | Unexplained replacement-policy effect; model as effective capacity < nominal (calibrate). |
| L1 line size | 32 B | [T4] Table 3.1 | load granularity 32 B; update granularity 128 B. |
| L1 load bandwidth (measured) | 58.83 B/cycle/SM | [T4] §3.1.1, Table 3.2 | P4 15.7; V100 108.3. `[scaled]` per-chip ×SM×clock. |
| L1 load bandwidth (theoretical) | LSU count (16) × bytes/inst | [T4] §3.1.1 | 16 LSU per SM on Turing vs 32 on Volta [T4 §1]. |
| L2 size | 4,096 KiB | [TWP] Table 4, [T4] Table 3.1 | 16-way set-associative, 64 B line [T4] §3.2. |
| L2 latency | ~188 cycles | [T4] Table 3.1 | |
| L2 load throughput (measured) | 1,270 GB/s (T4) | [T4] Table 3.4 | V100 2,155 GB/s; P4 979 GB/s. `[scaled]` for 48 SM/clock. |
| L2 read sectors per request | not published | — | open (see §10). |
| Constant caches | L1 const 2 KiB, 8 sets, 4-way, 64 B line; L1.5 ~46 KiB, 256 B line | [T4] Table 3.1 | |
| L1 TLB coverage | 32 MiB, 2 MiB pages | [T4] Table 3.1 | |
| L2 TLB coverage | ~8 GiB, 32 MiB pages | [T4] Table 3.1 | |
| Shared memory per SM | 32 or 64 KiB | [T4] Table 3.1 | 32 banks/SM, 4 B bank width. |
| Shared memory per chip (T4) | 1,280 or 2,560 KiB | [T4] Table 3.1 | `[scaled]`: ×48/40 for full TU104. |
| Shared throughput (T4 theoretical / actual) | 4,070 / 3,662 GiB/s | [T4] Table 3.1 | Chip-level, as published. |
| LSU count / SM | 16 | [T4] Table 3.1 | |
| DRAM type (T4) | GDDR6, 5,001 MHz, 256-bit | [T4] Table 3.1 | |
| DRAM bandwidth (T4 theoretical / actual) | 320 / 220 GB/s | [T4] Table 3.1 | Actual = 68.8% of theory. |
| DRAM bandwidth (2080 Super) | 496 GB/s @ 15.5 Gbps, 256-bit | [2080S] | Theoretical; model the achievable/theoretical ratio as a calibrated parameter. |
| Frame buffer (T4 / 2080 Super) | 16,384 MiB / 8,192 MiB | [T4] Table 3.1, [2080S] | Size does not affect timing beyond capacity. |
| NVLink 2 (TU104, where enabled) | 2 x8 links, 25 GB/s per direction each → 100 GB/s aggregate bidirectional | [TWP] §Second-Generation NVLink | Consumer RTX 20 cards do not expose it; model seam only (Phase 4). |

## 5. Compute throughput (board-level, as published)

| Metric | T4 (40 SM, 1,590 MHz) | Source | 2080 Super scaling `[scaled]` |
|---|---|---|---|
| Peak FP32 | 8.1 TFLOPS | [TWP] Table 5 | ×(48/40)×(1815/1590) ≈ **11.1 TFLOPS** — matches [2080S] 10.138/11.151 (base/boost), confirming the scaling rule. |
| Peak FP16 (non-tensor) | 16.2 TFLOPS | [TWP] Table 5 | ~22.2 TFLOPS scaled. |
| FP64 | 253 GFLOPS (1/32 of FP32; 2 FP64 units/SM) | [T4] Table 4.3, [TWP] note | 1/32 rate is architectural [TWP §TU102 note]. |
| FP32 (cuBLAS measured) | 7,174 GFLOPS | [T4] Table 4.3 | Achieved not peak. |
| FP16 (cuBLAS measured) | 41,616 GFLOPS | [T4] Table 4.3 | Tensor path. |
| INT8 (measured) | 74,934 GOPS | [T4] Table 4.3 | |
| INT4 (measured) | 114,384 GOPS | [T4] Table 4.3 | |
| INT1 (measured) | 552,230 GOPS | [T4] Table 4.3 | |
| FP16 tensor (with FP16 acc), published peak | 65 TFLOPS (T4) | [TWP] Table 5 | |
| FP16 tensor (with FP32 acc), published peak | 65 TFLOPS (T4) | [TWP] Table 5 | |
| INT8 tensor peak | 130 TOPS (T4) | [TWP] Table 5 | |
| INT4 tensor peak | 260 TOPS (T4) | [TWP] Table 5 | |
| RTX 2080 (not Super) tensor peaks | 80.5/84.8 FP16-acc; 40.3/42.4 FP32-acc; 161.1/169.6 INT8; 322.2/339.1 INT4 (ref/FE) | [TWP] Table 4 | Use RTX 2080 column only as a cross-check; 2080 Super is 48 SM like Quadro RTX 5000 (89.2 TFLOPS FP16-acc at FE clocks). |

## 6. Tensor core instruction behavior

| Item | Value | Source | Notes |
|---|---|---|---|
| Unit | 2nd-gen Tensor Core, 8/SM | [TWP] | Supports FP16 inputs (FP16/FP32 acc) plus INT8/INT4/INT1 inference modes. |
| `wmma.mma.sync` m16n16k16 (FP16) | 4 × `HMMA.1688.F32` on Turing (vs 16 × `HMMA.884` on Volta) | [T4] §4.3 | Instruction count per warp-level primitive. Model issue/interconnect at SASS level. |
| SASS shapes in use | `HMMA.1688`, `HMMA.884` | [AMP] §V-C (Turing comparison) | |
| Sparsity | not present on Turing | [AMP] §I vs Ampere | Ampere adds 2:4 sparsity [A100 whitepaper]; model as Turing-absent. |
| Tensor instruction latency | not tabulated per-instruction in [T4] | — | open; calibrate from SASS issue model (see docs/06 metric C1). |

## 7. Instruction latency and throughput

### 7.1 Native instruction latency (dependent issue), cycles

From [T4] Table 4.1:

| Instruction class | Turing | Volta | Pascal |
|---|---|---|---|
| IADD3, SHF, LOP3, SEL, MOV, FADD, FFMA, FMUL, ISETP, FSET, FSETP | **4** | 4 | 6 |
| IMAD, FMNMX, DSET, DSETP | **5** | 5 | 6 / 12 |
| HADD2, HMUL2, HFMA2 | **6** | 6 | 6 |
| POPC, FLO, BREV, MUFU | **~15** | ~10–14 | ~14 |
| DADD, DMUL | **~48** | 8 | 8 |
| DFMA, DSET, DSETP | **~54** | 8 | 8 |
| IMUL, IMAD (Pascal only) | — | — | ~86 (emulated) |

Notes: Turing reduces integer/SP latency vs Pascal, but **not** vs Volta;
FP64 latency regresses sharply (48–54 cycles) because only 2 FP64 units/SM
exist and are optimized for correctness [T4] §4.1, [TWP] note.

### 7.2 Independent-issue throughput

[T4] §4.1 states FP32 throughput is limited by latency and register-bank
conflicts, not by a hard issue limit; long sequences of identical FFMA show
2-cycle stalls on some GPUs [T4] §4.1 setup note. The model should use the
4-cycle dependent FFMA latency and the measured B/cycle numbers rather than
assuming full dual-issue.

### 7.3 Atomic latency (cycles), [T4] Table 4.2

| Contention | Shared T4 | Shared V100 | Global T4 | Global V100 |
|---|---|---|---|---|
| none | 8 | 6 | 76 | 36 |
| 2 threads | 10 | 7 | 72 | 31 |
| 4 threads | 14 | 11 | 73 | 32 |
| 8 threads | 22 | 18 | 81 | 41 |
| 16 threads | 37 | 24 | 97 | 58 |
| 32 threads | 69 | 66 | 116 | 76 |

### 7.4 Instruction encoding facts relevant to the model

* Turing and Volta instructions are **128-bit** (opcode + 64-bit control
  section) [T4] §2.1; opcode in the least-significant bits of the first
  64-bit word, 10–13 bits long [T4] §2.3.1.
* Control fields: 4-bit reuse flags, 6-bit wait-barrier mask, read/write
  barrier indices, yield flag, 4-bit stall cycles [T4] §2.1.
* 6 dependency barriers for variable-latency instructions [T4] §2.1.
* Reuse CAM: 2-way associative per source-operand slot, up to 4 operands
  [T4] §2.1.
* A warp switched to another warp costs one extra cycle [T4] §2.1 (yield).

## 8. Differences T4 (40 SM) vs full TU104 (48 SM) — scaling summary

| Quantity | Rule | Worked example | Notes |
|---|---|---|---|
| SM-count quantities (CUDA/tensor/RT/L1/shared totals) | ×48/40 = 1.2 | 2,560 → 3,072 CUDA cores | [TWP] confirms 3,072 for Quadro RTX 5000. |
| Compute rate | ×1.2 × (clock ratio) | FP32 8.1 → 11.1 TFLOPS at 1,815 MHz | Matches [2080S] published 11.151. |
| Shared/L1 bandwidth per chip | ×1.2 × clock ratio | 4,070 → ~5,575 GiB/s at 1,815 MHz (≈6.0 TB/s) | Treat as bound; calibrate. |
| L2 size, latency | unchanged | 4 MB, 188 cycles | Architecture-level. |
| DRAM bandwidth | board-specific | T4 320; 2080 Super 496 GB/s | Do **not** scale T4 DRAM numbers. |
| Memory latencies (cycles) | unchanged | 32/188/296/616 | Architecture-level. |

## 9. What the model can safely assert now vs what needs calibration

* **Assert (public evidence):** SM composition, cache geometry, cycle
  latencies, instruction latency table, atomic table, tensor instruction
  decomposition, scheduler mapping, encoding control fields.
* **Calibrate (unknowns/board-dependent):** DRAM timings and effective
  bandwidth ratio, L1 effective capacity (7 KiB discrepancy), replacement
  policy behavior, crossbar/NoC topology and per-hop latencies, L2
  sectoring/partition behavior, throttling envelope, tensor-instruction
  latency.

## 10. Open / uncertain items (do not invent)

1. TU104 GDDR6 timing set (tCL/tRCD/tRP/tRAS, refresh, bank groups) — not public.
2. L2 sector/queue/partition microarchitecture and per-request serialization.
3. L1 replacement policy (the 7 KiB effect in §4.2 is unexplained by [T4]).
4. Exact crossbar (NoC) latency from SM to L2/FB partitions.
5. Tensor-core per-instruction latency/throughput as a function of HMMA shape
   beyond the wmma-level decomposition.
6. Max boost behavior under power/thermal limits for the specific
   2080 Super Max-Q board (the local card variant throttles at 80 W, well
   below the desktop 250 W part [2080S]); the model must not assume desktop
   clocks for a Max-Q profile.
7. NVLink 2 availability on TU104 consumer boards (present on the die, not
   exposed by RTX 20 cards) — model as a Phase-4 seam only.

## 11. Validation hooks (docs/06 C-metrics)

* C1 instruction-latency agreement: use §7.1 table (target ≥85% within
  ±1 cycle).
* C2 memory-latency agreement: use §4.1 table (target ≥80% within ±15%).
* C3 bandwidth agreement: use §4.2/§5 achieved numbers (target ≥80% within
  ±20%).
* C4 kernel correlation: [T4] reports cuBLAS/CUTLASS matrix-multiply
  throughputs (Table 4.3) that can seed a smoke suite.
* Any physical capture on the local TU104 board is supervised-only
  (RUNBOOK P4/P7) and is never an unattended-agent requirement.