vgsp: host CUDA bridge + sema release recovery; greeneye simulator wave

- kernel: gsp_cuda_bridge (QMD launch capture, debugfs job/out/done, param+output resolution cached from inline writes), sema drain fixes (never hold/starve a channel), GR ring autopin + counter autoreg, per-page push reads, counter-derive gate; module builds clean
- tools/cuda-bridge: CUDA driver-API daemon executing mail/vadd kernels on the physical GPU, CUBB2/CBD2 v2 protocol, layout check + self-tests
- tools/guest-diag + tools/cuda-ab: scanner/pin fixes, victory push-only mode
- debug/risc-v-gsp-rm-vm: Phase S simulator wave (GSP device model, boot ladder, triage/re-demo harnesses, docs 30-41), GPGPU-sim scaffolding
- STATE-CURRENT/docs updates

1 files changed, 288 insertions, 0 deletions

diff --git a/debug/risc-v-gsp-rm-vm/docs/33-AMPERE-PARAMETERS.md b/debug/risc-v-gsp-rm-vm/docs/33-AMPERE-PARAMETERS.md
new file mode 100644
index 0000000..929db3d
--- /dev/null
+++ b/debug/risc-v-gsp-rm-vm/docs/33-AMPERE-PARAMETERS.md
@@ -0,0 +1,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).