1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
275
276
277
278
279
280
281
282
283
284
285
286
287
288
|
# 33 — Ampere parameter dossier (Phase 2 model inputs)
> **Status:** research dossier for the GreenEye-GPGPU-Sim Phase 2 model
> (WP2.1, `handoff-simulator.md` §5). Public-source only. Rows marked
> `[scaled]` are derived by an explicit arithmetic rule; rows marked
> `[inferred]` are reasoning, not measurement. Nothing here is a
> silicon-accuracy claim. Phase 2 is gated (U24): this document is written
> now, implemented later.
>
> **Scope:** A100 / GA100 (the microbenchmark paper's target), GA102
> (RTX 3090), and GA106 (RTX 3060, the SM86 consumer case named in the
> handoff). Host/device modeling only; not GSP/MCU.
## 1. Sources and citation labels
| Label | Source | What it provides |
|---|---|---|
| **[AMP]** | H. Abdelkhalik, Y. Arafa, N. Santhi, A.-H. Badawy, *Demystifying the Nvidia Ampere Architecture through Microbenchmarking and Instruction-level Analysis*, arXiv:2208.11174v1 (2022) — measurements on an NVIDIA "AI100" (A100) Tesla GPU | Per-instruction CPI (dependent + independent), memory-unit latencies, tensor-core (WMMA) latencies/throughput and PTX→SASS mapping, PTX/SASS tables. |
| **[A100-WP]** | NVIDIA, *NVIDIA A100 Tensor Core GPU Architecture* whitepaper, https://www.nvidia.com/content/dam/en-zz/Solutions/Data-Center/nvidia-ampere-architecture-whitepaper.pdf | GA100/A100 architecture: SM count, L2, HBM2, MIG, cp.async, async barriers, sparsity, NVLink 3, per-precision peaks. |
| **[A100-DS]** | NVIDIA, *NVIDIA A100 Tensor Core GPU* datasheet (Jun 2021), https://www.nvidia.com/content/dam/en-zz/Solutions/Data-Center/a100/pdf/nvidia-a100-datasheet-us-nvidia-1758950-r4-web.pdf | Form-factor-resolved A100 table: memory, bandwidth, TDP, MIG, NVLink/PCIe, precision peaks. |
| **[GA102-WP]** | NVIDIA, *NVIDIA Ampere GA102 GPU Architecture* whitepaper, v1/v2.x, https://www.nvidia.com/content/PDF/nvidia-ampere-ga-102-gpu-architecture-whitepaper-v2.1.pdf | GA102/GA10x SM (128 FP32/SM), RTX 3090 table, GDDR6X, 28.3 B transistors, 628.4 mm². |
| **[RTX30-W]** | Wikipedia *GeForce RTX 30 series* product tables (retrieved 2026-10-03), citing NVIDIA product pages, TechPowerUp GPU database and AnandTech reviews: https://en.wikipedia.org/wiki/GeForce_RTX_30_series | GA106/RTX 3060 board numbers, GA102 board clocks/bandwidth cross-checks. |
| **[T4]** | arXiv:1903.07486 (Turing) | Turing comparison points. |
## 2. Method and caveats
1. **[AMP] measured one A100 board** (SXM-class, 108 SMs, HBM2). Its
*instruction* and *memory-latency* numbers are architecture-level for
GA100; its board-level bandwidth/power numbers are not transferable to
GA10x.
2. GA100 (datacenter) and GA10x (consumer) **SM designs differ**: GA100 has
64 FP32 + 64 INT32 per SM; GA10x has 128 FP32-capable cores per SM with
INT32 sharing. See §4.
3. `cp.async` and asynchronous barriers are **compute-capability 8.0 (GA100)
features**; SM86 (GA10x consumer) does not expose them ([A100-WP] §CUDA
Asynchronous Copy; CUDA programming guide compute-capability notes).
Do not model cp.async on the RTX 3060 profile.
4. Dependent-issue latencies dominate a statically scheduled core; [AMP]
distinguishes dependent vs independent CPI (§6.2). Use both.
5. Board clocks, memory rates and TDP are board properties (§3); cache
geometry and cycle latencies are architecture properties.
## 3. Chip and board parameters
### 3.1 GA100 / A100
| Parameter | Value | Source | Notes |
|---|---|---|---|
| Full GA100 | 128 SMs, 8 GPCs, 8 TPCs/GPC, 6 HBM2 stacks, 12 512-bit MCs | [A100-WP] §GA100 | |
| A100 (product) | 108 SMs, 7 GPCs, 5 HBM2 stacks, 10 512-bit MCs | [A100-WP] §GA100 | 432 3rd-gen tensor cores. |
| Boost clock | 1,410 MHz | [A100-WP] A100 vs V100 table | |
| L1/shared per SM | 192 KB (vs V100 128 KB) | [A100-WP] §overview | 1.5× aggregate capacity per SM. |
| L2 | 40 MB, partitioned crossbar, 2.3× V100 L2 read bandwidth | [A100-WP] §40 GB HBM2 and 40 MB L2 | ~7× V100 capacity. |
| HBM2 (40 GB SXM/PCIe) | 1,555 GB/s | [A100-DS] | 1,215 MHz DDR, 5 stacks × 8 dies [A100-WP]. |
| HBM2e (80 GB SXM / PCIe) | 2,039 / 1,935 GB/s | [A100-DS] | |
| TDP | 400 W SXM4; 250 W PCIe (300 W 80 GB PCIe) | [A100-DS] | |
| PF32 | 19.5 TFLOPS | [A100-DS] | |
| FP64 / FP64 TC | 9.7 / 19.5 TFLOPS | [A100-DS] | |
| FP16 / BF16 (non-tensor) | 78 / 39 TFLOPS | [A100-WP] Table 1 | |
| TF32 TC | 156 TFLOPS (312 sparse) | [A100-DS] | |
| FP16 TC / BF16 TC | 312 TFLOPS (624 sparse) | [A100-DS] | |
| INT8 TC / INT4 TC | 624 TOPS (1,248 sparse) / 1,248 TOPS (2,496 sparse) | [A100-DS] | |
| NVLink 3 | 12 links, 25 GB/s per direction per link, 600 GB/s aggregate | [A100-WP] §Third-Generation NVLink, [A100-DS] | 50 Gbit/s per signal pair. |
| PCIe | Gen4 ×16, ~31.5 GB/s per direction (64 GB/s bidirectional per datasheet) | [A100-WP], [A100-DS] | SR-IOV supported. |
| MIG | Up to 7 instances; isolated crossbar ports, L2 banks, memory controllers, DRAM buses | [A100-WP] §MIG, [A100-DS] | 40 GB: 7 × 5 GB profiles; 80 GB: 7 × 10 GB. |
### 3.2 GA102 / RTX 3090
| Parameter | Value | Source | Notes |
|---|---|---|---|
| Full GA102 | 7 GPCs, 42 TPCs, 84 SMs, 10752 CUDA cores, 336 tensor cores, 84 RT cores, 28.3 B transistors, 628.4 mm² | [GA102-WP] §GA102 Key Features | |
| RTX 3090 | 82 SMs, 10496 CUDA cores, 328 tensor cores (4/SM, 3rd gen), 82 RT cores, 24 GB GDDR6X, 384-bit | [GA102-WP] Table 7, [RTX30-W] | |
| Boost clock (3090 FE) | 1,695 MHz | [GA102-WP] Table 7 | Base 1,395 MHz [RTX30-W]. |
| FP32 (non-tensor) | 35.6 TFLOPS | [GA102-WP] Table 7 | 128 FP32/SM. |
| FP16/BF16 (non-tensor) | 35.6 / 35.6 TFLOPS | [GA102-WP] | |
| INT32 | 17.8 TOPS | [GA102-WP] | IMAD-based. |
| FP16 TC (FP16 acc) | 142 (284 sparse) TFLOPS | [GA102-WP] Table 7 | |
| FP16 TC (FP32 acc) | 71 (142 sparse) TFLOPS | [GA102-WP] | |
| TF32 TC | 35.6 (71 sparse) TFLOPS | [GA102-WP] | |
| INT8 / INT4 TC | 284 (568 sparse) / 568 (1136 sparse) | [GA102-WP] | |
| L1/shared per SM | 128 KB (10496 KB per 82-SM GPU) | [GA102-WP] Table 7 | |
| L2 | 6 MB | [GA102-WP] Table 7 | |
| Memory | GDDR6X 19.5 Gbps, 936 GB/s | [GA102-WP] Table 7 | 1.5× Titan RTX. |
| TGP | 350 W | [GA102-WP] | |
| NVLink 3 | 4 × x4 links, ~112.5 GB/s bidirectional | [GA102-WP] §NVLink | Consumer 3090-class only. |
### 3.3 GA106 / RTX 3060 (SM86)
| Parameter | Value | Source | Notes |
|---|---|---|---|
| CUDA cores / SMs / tensor cores | 3,584 / 28 / 112 (4/SM) | [RTX30-W] GA106-300 row | SM count 28. |
| Clocks (base/boost) | 1,320 / 1,777 MHz | [RTX30-W] | |
| FP32 (boost) | 12.74 TFLOPS | [RTX30-W] (9.462 × 1.777 GHz computed by the table) | `[scaled]` consistent with 28 SM × 128 × 2 × 1.777 GHz. |
| Memory | 12 GB GDDR6, 192-bit, 15 Gbps, 360 GB/s | [RTX30-W] | |
| L2 | 3 MB | [RTX30-W] (GA106 row) | |
| TGP | 170 W | [RTX30-W] | |
| Compute capability | 8.6 | NVIDIA CUDA programming guide (as cited by [RTX30-W]) | No cp.async (8.0-only). |
## 4. Ampere SM architecture
| Feature | GA100 (A100) | GA10x (GA102/GA106) | Source |
|---|---|---|---|
| FP32 cores / SM | 64 | 128 | [A100-WP] Fig. 7, [GA102-WP] §GA102 Key Features |
| INT32 cores / SM | 64 | (128 FP32 cores double as INT32 for many ops) | [A100-WP], [GA102-WP] |
| Tensor cores / SM | 4 (3rd gen) | 4 (3rd gen) | [A100-WP], [GA102-WP] Table 7 |
| Partitioning | 4 processing blocks | 4 processing blocks | [AMP] §V-C ("each SM … 4 tensor cores") |
| L1/shared per SM | 192 KB unified | 128 KB unified | [A100-WP], [GA102-WP] Table 7 |
| New in generation | `cp.async`, async barriers (`mbarrier`), L2 residency control, compute data compression, 2:4 structured sparsity, full-rate FP64 tensor cores | same tensor/L2 features; SM doubles FP32 per SM | [A100-WP], [GA102-WP] |
| L2 residency control | yes: persistent window (`accessPolicyWindow`), set-aside carveout | yes on GA10x | [A100-WP] §L2 residency |
| Compression | up to 4× DRAM/L2 bandwidth, up to 2× L2 capacity | not documented for GA10x | [A100-WP] §overview |
**Data types on the 3rd-gen tensor core:** FP16, BF16, TF32 inputs
(FP16/FP32 accumulators), plus INT8/INT4/binary for inference; sparsity
(2:4) doubles effective throughput [A100-WP] §Third-Generation Tensor Core.
## 5. Memory latencies (A100, cycles)
From [AMP] Table IV (dependent-issue pointer chase; global-memory path
bypasses caches with `.cv`):
| Memory type | CPI (cycles) | Source |
|---|---|---|
| Global memory | 290 | [AMP] Table IV, §V-B |
| L2 cache | 200 | [AMP] |
| L1 cache | 33 | [AMP] |
| Shared memory load / store | 23 / 19 | [AMP] |
Comparison points quoted by [AMP] for Turing: global 434, L2 188, L1 32
(the same values as [T4]'s measured 32/188/296/616 global-path figures,
with 434 corresponding to the DRAM-latency band). Use [T4] Table 3.1 for
the Turing shared-memory row (19 cycles no-conflict); [AMP]'s 23/19 is the
Ampere measurement.
## 6. Instruction performance (A100, cycles)
### 6.1 Full PTX→SASS latency table [AMP] Table V
`cycles` = dependent-issue CPI measured by [AMP] on A100.
| PTX | SASS | cycles | PTX | SASS | cycles |
|---|---|---|---|---|---|
| add.u16 | UIADD3 | 2 | mul.wide.u16 | LOP3.LUT+IMAD | 4 |
| addc.u32 | IADD3.X | 2 | mul.wide.u32 | IMAD | 4 |
| add.u32 | IADD | 2 | mul.lo.u16 | LOP3.LUT+IMAD | 4 |
| add.u64 | UIADD3.x+UIADD3 | 4 | mul.lo.u32 | IMAD | 2 |
| add.s64 | UIADD3.x+UIADD3 | 4 | mul.lo.u64 | IMAD | 2 |
| add.f16 | HADD | 2 | mul24.lo.u32 | PRMT+IMAD | 3 |
| add.f32 | FADD | 2 | mul24.hi.u32 | UPRMT+USHF.R.U32.HI+IMAD.U32+PRMT | 9 |
| add.f64 | DADD | 4 | mul.rn.f16 | HMUL2 | 2 |
| mad.lo.u16 | LOP3.LUT+IMAD | 4 | mul.rn.f32 | FMUL | 2 |
| mad.lo.u32 | FFMA | 2 | mul.rn.f64 | DMUL | 4 |
| mad.lo.u64 | IMAD | 2 | fma.rn.f16 | HFMA2 | 2 |
| mad24.lo.u32 | SGXT.U32+IMAD | 4 | fma.rn.f32 | FFMA | 2 |
| mad24.hi.u32 | USHF.R.U32.HI+UIMAD.WIDE.U32+2*UPRMT+IADD3 | 11 | fma.rn.f64 | DFMA | 4 |
| mad.rn.f32 | FFMA | 2 | abs.s16 | PRMT+IABS+PRMT | 4 |
| mad.rn.f64 | DFMA | 4 | abs.s32 | IABS | 2 |
| sad.u16/s16 | (2*LOP3)+ULOP3+VABSDIFF | 6 | abs.s64 | UISETP.LT.AND+UIADD3.X+UIADD3+2*USEL | 11 |
| sad.u32/s32 | VABSDIFF+IMAD | 3 | abs.f16 | PRMT | 1 |
| sad.u64/s64 | UISETP.GE.U32.AND+UIADD+IADD | 10 | abs.ftz.f32 | FADD.FTZ | 2 |
| rem/div.u16/s16 | multiple instrs | 290 | abs.f64 | DADD or (DADD+UMOV) | 4 |
| rem/div.s32/u32 | multiple instrs | 66 | neg.s16 | UIADD3+UPRMT | 5 |
| rem/div.u64/s64 | multiple instrs | 420 | neg.s32 | IADD3 | 2 |
| div.rn.f32 | multiple instrs | 525 | neg.s64 | IMAD.MOV.U32+HFMA2.MMA+MOV+UIADD3 | 10 |
| div.rn.f64 | multiple instrs | 426 | neg.f32 | FADD or IMAD.MOV.U32 | 2 |
| sqrt.rn.f32 | incl. MUFU.RSQ | 190–235 | neg.f64 | DADD+(UMOV) | 4 |
| sqrt.approx.f32 | incl. MUFU.SQRT | 2–18 | rcp.rn.f32 | incl. MUFU.RCP | 198 |
| sqrt.rn.f64 | incl. MUFU.RSQ64 | 260–340 | rcp.approx.f32 | incl. MUFU.RCP | 23 |
| rsqrt.approx.f32 | incl. MUFU.RSQ | 2–18 | rcp.rn.f64 | incl. MUFU.RCP64H | 244 |
| rsqrt.approx.f64 | MUFU.RSQ64H | 8–11 | ex2.approx.f32 | FSTEP+FMUL+MUFU.EX2+FMUL | 14 |
| sin.approx.f32 | FMUL+MUFU.SIN | 8 | ex2.approx.f32 | FSETP.GEU.AND+2*FMUL+MUFU.EX2 | 18 |
| cos.approx.f32 | FMUL.RZ+MUFU.COS | 8 | ex2.approx.f16 | MUFU.EX2.F16 | 6 |
| lg2.approx.f32 | FSETP.GEU.AND+FMUL+MUFU.LG2+FADD | 18 | tanh.approx.f32 / .f16 | MUFU.TANH / MUFU.TANH.F16 | 6 |
| popc.b32 | POPC | 6 | and.b16 | LOP3.LUT or 1.5*LOP3.LUT | 2 |
| popc.b64 | 2*UPOPC+UIADD3 | 7 | and.b32 | LOP3.LUT | 2 |
| clz.b32 | FLO.U32+IADD | 7 | and.b64 | ULOP3.LUT | 2–3 |
| clz.b64 | UISETP.NE.U32.AND+USEL+UFLO.U32+2*UIADD3 | 13 | or / xor (b16/b32/b64) | LOP3 family | 2–4 |
| bfind.u32 | FLO.U32 | 6 | not.b16 / not.b32 | LOP3.LUT | 2 |
| bfind.u64 | FLO.U32+ISETP.NE.U32.AND+IADD3+BRA | 164 | not.b64 | 2*ULOP3.LUT | 4 |
| bfind.s32 | FLO | 6 | lop3.b32 | IMAD.MOV.U32+LOP3.LUT | 4 |
| bfind.s64 | multiple instrs | 195 | cnot.b16/b32/b64 | ULOP3…/UISETP…/multiple | 5/4/11 |
| testp.normal.f32 | IMAD.MOV.U32+2*ISETP.GE.U32.AND | 0 or 6 | bfe.s32/.u32 | 3*PRMT+2*IMAD.MOV+SHF.R.U32.HI+SGXT | 11 |
| testp.subnor.f32 | ISETP.LT.U32.AND | 0 or 6 | bfe.u64 | UMOV+USHF.L.U32+(UIADD3+ULOP3.LUT) | 5 |
| testp.normal.f64 | 2*UISETP.LE.U32.AND+2*UISETP.GE.U32.AND | 13 | bfe.s64 | multiple instrs | 14 |
| testp.subnor.f64 | UISETP.LT.U32.AND+2*UISETP.GE.U32.AND.EX | 8 | bfi.b32 | 3*PRMT+2*IMAD.MOV+SHF.L.U32+BMSK+LOP3.LUT | 11 |
| min.u16 | ULOP3.LUT+UISETP.LT.U32.AND+USEL | 8 | bfi.b64 | UMOV+USHF.L.U32+(UIADD3+ULOP3.LUT) | 5 |
| min.u32 | IMNMX.U32 | 2 | brev.b32 | BREV+SGXT.U32 | 2 |
| min.u64 | UISETP.LT.U32.AND+2*USEL | 8 | brev.b64 | 2*UBREV+MOV | 6 |
| min.s16 | PRMT+IMNMX | 4 | copysign.f32 | 2*LOP3.LUT or 1.5*LOP3.LUT | 4 |
| min.s32 | IMNMX | 2 | copysign.f64 | 2*ULOP3.LUT+IMAD.U32+*MOV | 6 |
| min.s64 | UISETP.LT.U32.AND+UISETP.LT.AND.EX+2*USEL | 8 | cvt.rzi.s32.f32 | F2I.TRUNC.NTZ | 6 |
| min.f16 | HMNMX2+PRMT | 4 | setp.ne.s32 | ISETP.NE.AND | 10 |
| min.f32 | FMNMX | 2 | mov.u32 %clock | CS2R.32 | 2 |
| min.f64 | DSETP.MIN.AND+IMAD.MOV.U32+UMOV+FSEL | 10 | dp4a.u32.u32 | IMAD.MOV.U32+IDP.4A.U8.U8 | 135–170 |
| bar.warp.sync | NOP | changes | dp2a.lo.u32.u32 | IMAD.MOV.U32+IDP.2A.LO.U16.U8 | 135–170 |
| fns.b32 | multiple instrs | 79 | | | |
### 6.2 Dependence and issue effects
[AMP] Table II (dependent vs independent CPI):
| Instruction | dependent | independent |
|---|---|---|
| add.f16 | 3 | 2 |
| add.u32 | 4 | 2 |
| add.f64 | 5 | 4 |
| mul.lo.u32 | 3 | 2 |
| mad.rn.f32 | 4 | 2 |
[AMP] §V-A notes: 2 `add` + 2 `mad` complete in ~4 cycles total, i.e. ~1
cycle each, because MAD issues on the FP pipeline and ADD on the integer
pipeline (they overlap). `mad` always runs on the FP pipeline even for
integers (mapped to FFMA). Signed/unsigned pairs map identically except
`bfind`, `min`, `max`. Input initialization can change PTX→SASS mapping
(e.g. `neg.f32` → `FADD` vs `IMAD.MOV.u32`).
## 7. Tensor core (WMMA) latencies and throughput, A100 [AMP] Table III
| PTX shapes | Input | Acc | Cycles | Measured theoretical | PTX instruction | SASS |
|---|---|---|---|---|---|---|
| m16n16k16 / m8n32k16 / m32n8k16 | f16 | f16 | 16 | 311–312 GB/s | `wmma.mma.sync...m16n16k16.f16.f16` | 2× HMMA.16816.F16, 8 cyc each |
| same | f16 | f32 | 16 | 310–312 GB/s | `...f16.f32` | 2× HMMA.16816.F32, 8 cyc each |
| same | bf16 | f32 | 16 | 310–312 GB/s | `...f32.bf16.bf16.f32` | 2× HMMA.16816.F32.BF16, 8 cyc each |
| m16n16k8 | tf32 | f32 | 16 | 132–156 GB/s | `...f32.tf32.tf32.f32` | 4× HMMA.1684.F32.TF32, 4 cyc each |
| m8n8k4 | f64 | f64 | 16 | 19–19.5 GB/s | `...m8n8k4.f64...rn` | 1× DMMA.884, 16 cyc |
| m16n16k16 / m32n8k16 / m8n32k16 | u8 | u32 | 8 | 594–624 GB/s | `...s32.u8.u8.s32` | 2× IMMA.16816.U8.U8, 4 cyc each |
| m8n8k32 | u4 | u32 | 4 | 1229–1248 GB/s | `...s32.u4.u4.s32` (row.col) | 1× IMMA.8832.U4.U4, 4 cyc |
Additional findings from [AMP] §V-C:
* Each Ampere SM holds 4 tensor cores; the physical TC implementation can
perform 8×4×8 [AMP, citing Sun et al. arXiv:2206.02874].
* PTX and SASS shapes differ; the SASS instruction count depends on how
many SASS-shaped ops tile the PTX shape (hence 2 vs 4 SASS per WMMA).
* For the same data type, **different PTX shapes do not change** the
measured latency on Ampere (unlike Turing, where shape matters [AMP,
citing Yan et al.]).
* FP16/BF16 WMMA requires a `MOVM.16.MT88` matrix-load/transpose
instruction; its count depends on row/column layout.
* Volta had only `HMMA.884`; Turing added `HMMA.1688`; Ampere adds
`DMMA.884` (FP64), `IMMA.16816` (U8), `IMMA.8832` (U4) [AMP] §V-C.
## 8. Generation changes the model must express (not just retune)
| Change | Turing (Phase 1) | Ampere (Phase 2) | Source |
|---|---|---|---|
| FP32/SM | 64 | 64 (GA100) / 128 (GA10x) | [A100-WP], [GA102-WP] |
| Concurrent FP32+INT32 | separate 16+16 per block | GA10x doubles FP32 and shares; GA100 keeps separate INT32 | [TWP], [GA102-WP] |
| Global memory latency | 434 cycles | 290 cycles | [AMP] |
| L2 latency | 188 | 200 | [T4], [AMP] |
| L1 latency | 32 | 33 | [T4], [AMP] |
| L2 size | 4 MB (TU104) | 40 MB (A100) / 6 MB (GA102) / 3 MB (GA106) | [TWP], [A100-WP], [GA102-WP], [RTX30-W] |
| cp.async | absent | present on sm_80 only | [A100-WP] |
| Async barrier | absent | present | [A100-WP] |
| Sparsity | absent | 2:4 structured, 2× tensor rate | [A100-WP] |
| Tensor dtypes | FP16, INT8/4/1 | + BF16, TF32, FP64; INT4; binary | [TWP], [A100-WP] |
| L2 residency / compression | no / no | yes / yes (A100) | [A100-WP] |
| MIG | no | A100 up to 7 instances | [A100-WP] |
## 9. Open / uncertain items (do not invent)
1. GA100 **L1 sectoring/line behaviour** and L2 request coalescing are not
published; [AMP] gives only aggregate latencies.
2. cp.async **issue/completion timing** and mbarrier latency are not
published.
3. GA10x (SM86/SM89) L1/L2 microarchitecture differs from GA100; [AMP]
does not measure GA10x. Treat GA10x memory latencies as Turing-band
until a GA10x source exists; mark the profile's uncertainty.
4. MIG slice-level bandwidth/latency QoS curves are not published.
5. GDDR6X timing set (PAM4) and the 936 GB/s achievable ratio are not
public; model as JEDEC-class + calibrated efficiency.
6. A100 80 GB HBM2e (2,039 GB/s) changes memory timing but not
instruction/cache cycles; keep memory parameters separate.
7. Clock-throttling envelopes for GA102/GA106 boards are not modelled from
published curves.
8. The A100 microbenchmark table does not include every PTX op (e.g.
`cp.async` variants, `mbarrier` ops); these need direct measurement or
Triton/PTX docs before any cycle claim.
## 10. Validation hooks
* C1: §6.1 table (≥85% of entries within ±1 cycle).
* C2: §5 latencies (≥80% within ±15%).
* C3: A100 HBM2 1,555 GB/s theoretical vs achievable; GDDR6X 936 GB/s
theoretical vs achievable.
* C4: kernel correlation using triangular GEMM sizes from [AMP] (WMMA
throughputs) plus a reduction/streaming suite.
* Phase 2 entry additionally requires rebasing or superseding the
upstream `origin/dev` 4.2.0 RTX 3060 SM86 config (WP2.1/WP2.2).
|
