1 files changed, 260 insertions, 0 deletions
diff --git a/debug/risc-v-gsp-rm-vm/docs/31-TURING-TU104-PARAMETERS.md b/debug/risc-v-gsp-rm-vm/docs/31-TURING-TU104-PARAMETERS.md
new file mode 100644
index 0000000..a68fc32
--- /dev/null
+++ b/debug/risc-v-gsp-rm-vm/docs/31-TURING-TU104-PARAMETERS.md
@@ -0,0 +1,260 @@
+# 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.