diff options
Diffstat (limited to 'debug/risc-v-gsp-rm-vm/docs')
16 files changed, 3023 insertions, 15 deletions
diff --git a/debug/risc-v-gsp-rm-vm/docs/00-OVERVIEW.md b/debug/risc-v-gsp-rm-vm/docs/00-OVERVIEW.md index 0df8d08..50b4428 100644 --- a/debug/risc-v-gsp-rm-vm/docs/00-OVERVIEW.md +++ b/debug/risc-v-gsp-rm-vm/docs/00-OVERVIEW.md @@ -32,7 +32,7 @@ developer Frida-like instrumentation over its memory and execution: It is deliberately not a GDB wrapper: a pure-Python RSP client makes the whole stack scriptable, and QEMU is just the first execution backend. -## Status (v0.1, 2026-10-02) +## Status (v0.1 + 2026-10-03 simulator wave) | Capability | State | |---|---| @@ -44,24 +44,27 @@ whole stack scriptable, and QEMU is just the first execution backend. | Plugin loader + example plugin | working | | Native RV64IMAC simulator (M1, `--backend native`) | working; QEMU-parity differential test | | Native exec/access trace (M1.5, R4): `--trace`, analyzer coverage | working; trace blob + determinism test | -| Test suite | 194 tests green (live QEMU + native differential included) | +| Test suite | full suite green on the 2026-10-03 review host (298 tests incl. live QEMU + native differential; count varies with local prerequisites) | | Ghidra bridge | scripts + docs | -| Real GSP `.fwimage` probe (first-fault report) | working (honest: it does not boot without a TU104 model — see 06) | +| Real GSP `.fwimage` probe (first-fault report) | working (first fault + lenient continuation) | +| GSP device model + versioned profiles (S1) | in development: `sim/devices/`, `sim/profiles/`, `sim/gsp_boot.py`; B0–B2 on the real 570.144 image with 0 unknown MMIO (36) | +| Modeled cycle accounting | increment 0: `sim/cycles.py`; modeled defaults, not silicon timing (32) | +| Simulator-vs-QEMU A/B + new-version triage | working: frozen thresholds, verdict `above` on the S1 path per docs/37 §4.3; `harness/triage.py` (37, 40) | +| GPGPU-Sim overlay | Phase 0 scaffold + CUDA-12 build port record; Phase G gated (30) | | TCG full-memory tracing plugin | source + optional build (verified live) | | Automation harness (mutation/portability/version diff) + value demos | working (18/19/21/25) | | Host RM oracle harness (`rmharness`) | wrappers + capture working; device tables depend on driver permissions (17) | | First-touch device map + lenient native MMIO | working (24) | | MCP server (`plugins/mcp_server.py`) | working (20); stdlib-only, read-only by default | | Porting automation + trial packaging (28/29) | scaffolds + tests present; unadmitted pilots (dry-run default) | -| Native M2+ device model | roadmap only (08, 16) | The in-development notices were removed on 2026-10-02 after the tool audit ([27-TOOL-AUDIT](27-TOOL-AUDIT.md)), the 194-test green suite, and the end-to-end s01–s10 value demonstrations ([25-VALUE-VALIDATION](25-VALUE-VALIDATION.md)). Open items stay listed as -open: docs 28/29 are unadmitted pilots, the device model and custom simulator -M2+ remain roadmap, and nothing here is a stability guarantee beyond the -audited v0.1 surface. +open: docs 28/29 are unadmitted pilots, the device model is partial (S1, +B0–B2; docs/36) and custom-simulator M2+ remains roadmap, and nothing here is +a stability guarantee beyond the audited v0.1 surface. ## Quick start @@ -113,6 +116,17 @@ Interactively: | [27-TOOL-AUDIT](27-TOOL-AUDIT.md) | tool audit: purpose, tests, usage, status | | [28-PORTING-AUTOMATION](28-PORTING-AUTOMATION.md) | driver-version sweep and new-device scaffolds (pilot) | | [29-PACKAGING-AND-TRIALING](29-PACKAGING-AND-TRIALING.md) | trial simulation packages (pilot) | +| [30-GPGPU-SIM-INTEGRATION](30-GPGPU-SIM-INTEGRATION.md) | GPGPU-Sim overlay scaffold and build port (Phase G) | +| [31-TURING-TU104-PARAMETERS](31-TURING-TU104-PARAMETERS.md) | Turing TU104 parameter dossier | +| [32-GSP-CYCLE-MODEL](32-GSP-CYCLE-MODEL.md) | modeled cycle accounting (not silicon timing) | +| [33-AMPERE-PARAMETERS](33-AMPERE-PARAMETERS.md) | Ampere parameter dossier | +| [34-BLACKWELL-SOURCES](34-BLACKWELL-SOURCES.md) | Blackwell source classification | +| [35-GREENEYE-SIMULATOR-DESIGN](35-GREENEYE-SIMULATOR-DESIGN.md) | standalone GSP simulator design and gates | +| [36-GSP-DEVICE-MODEL](36-GSP-DEVICE-MODEL.md) | S1 device model, profiles, boot ladder | +| [37-SIM-QEMU-LINE](37-SIM-QEMU-LINE.md) | simulator-vs-QEMU A/B and measured verdict | +| [38-GSP-EM-API](38-GSP-EM-API.md) | GSP-emulator/system-simulator TLV boundary (design) | +| [39-RE-WORKFLOW-REVIEW](39-RE-WORKFLOW-REVIEW.md) | reverse-engineering workflow review | +| [40-NEW-VERSION-TRIAGE](40-NEW-VERSION-TRIAGE.md) | one-command new-version triage | | [RISC-V-GSP-RM-DEBUGGER-ROADMAP](../RISC-V-GSP-RM-DEBUGGER-ROADMAP.md) | long-term program H1–H7 | ## Where this fits in the Open-IOV workflow diff --git a/debug/risc-v-gsp-rm-vm/docs/06-ACCURACY.md b/debug/risc-v-gsp-rm-vm/docs/06-ACCURACY.md index ebc9226..a9a91aa 100644 --- a/debug/risc-v-gsp-rm-vm/docs/06-ACCURACY.md +++ b/debug/risc-v-gsp-rm-vm/docs/06-ACCURACY.md @@ -123,14 +123,17 @@ Before trusting a risc-v-gsp-rm-debugger session: ## 8. Known accuracy gaps (living list) -1. No TU104 device model: no boot beyond the first unmapped access. +1. Partial TU104 device model (S1, docs/36): the real 570.144 image is now + served past the first unmapped access to B2 (relocation copy, 0 unknown + MMIO) with a recorded boot shim; B3/B4 (CSB poll, message loop) are not + reached, and the model is uncalibrated and partial. 2. No time/concurrency model; multi-core and partition-switch behaviour is - approximated. + approximated. Cycles are modeled defaults, not measured (docs/32). 3. PMU/SEC2/ACR/Booter behaviour is message-level only. 4. Re-hosted uproc code is compiled for `qemu virt`, so compiler/layout differences with the shipped image must be accounted for when comparing control flow. 5. Interrupt delivery and MSI semantics differ; error-path M5 is the weakest - metric until the device model exists (08). + metric until interrupt delivery is modelled (docs/36 §6). Each gap maps to a roadmap item; do not silently ignore them in reports. diff --git a/debug/risc-v-gsp-rm-vm/docs/20-MCP-SERVER.md b/debug/risc-v-gsp-rm-vm/docs/20-MCP-SERVER.md index 46e28a4..0e5fede 100644 --- a/debug/risc-v-gsp-rm-vm/docs/20-MCP-SERVER.md +++ b/debug/risc-v-gsp-rm-vm/docs/20-MCP-SERVER.md @@ -127,7 +127,7 @@ the audit trail). | `initialize` | returns `protocolVersion` (echoes `2024-11-05`, `2025-03-26` or `2025-06-18`; otherwise the server default), `serverInfo`, `capabilities.tools` | | `notifications/initialized` | accepted silently (notifications never get a response) | | `ping` | `{}` | -| `tools/list` | the 24-tool catalogue with JSON input schemas | +| `tools/list` | the 31-tool catalogue with JSON input schemas | | `tools/call` | `{"name": ..., "arguments": {...}}` → `{"content":[{"type":"text","text":"<json>"}]}`, or `isError: true` for tool errors | JSON-RPC errors: `-32700` parse error, `-32600` invalid request, `-32601` @@ -144,6 +144,8 @@ them deterministically. | Tool | Arguments | Class | Notes | |---|---|---|---| | `session_info` | — | read | backend, PC, symbols, recording, breakpoints/watches, enabled capabilities | +| `backends_list` | — | read | QEMU/native availability and features; `active` backend; `launchable` (standalone only) | +| `session_launch` | `backend=native\|qemu`, `kernel`, `record=true`, `name`, `trace=false` | session | protocol-level backend selection; replaces the served session; **standalone mode only**; `trace` is native-only; allowed without `--allow-write` (it replaces the local served session and writes only its own recording; it never touches target memory) | | `symbols_resolve` | `target` | read | symbol, `symbol+0xNN` or hex address | | `bp_set` | `target`, `kind=sw\|hw` | observe | hits recorded as `bp_hit` | | `bp_clear` | `target` | observe | removes every breakpoint at that address | @@ -157,8 +159,13 @@ them deterministically. | `snapshot` | `addr`, `size` (max 16 MiB), `label=mcp` | read | zlib + sha256 into the session | | `call_function` | `target`, `args=[]`, `timeout=5.0` | **write** | requires `--allow-write`; Frida-like remote call | | `inject_service` | `service`, `slot=6`, `load_addr`, `cmd=1`, `arg0`, `arg1`, `path=remote\|firmware\|both`, `settle=0.25` | **inject** | requires `--allow-inject`; service ELF built with `make sdk` | +| `service_list` | — | read | services loaded into this session (injection registry) + the demo handler slots | | `trace_start` | `exec=true`, `start`, `end`, `max_events` | trace | native backend + recording session only | | `trace_stop` | — | trace | flushes `blobs/trace.jsonl`, returns counts | +| `sim_cycles` | — | read | native only: `insns`, modeled `cycles`, cycle-model table, fidelity label | +| `sim_device_state` | `limit=16` | read | native only: unknown-MMIO policy/log, device-map size, optional `sim/devices` model state | +| `sim_profile` | — | read | native only: firmware, bias, stop/profile metadata, trace config, model identity, session meta | +| `sim_regions` | `limit=16` | read | native only: RAM/UART windows, loaded firmware segments, device-map entries, observed unmodelled MMIO pages | | `analyze_current` | `max_rows=40` | read | `render_summary` for the active recording | | `compare_sessions` | `a`, `b`, `symbols=true`, `max_rows=40` | read | A/B recorded sessions | | `export_hits` | `session`, `out` | read | Ghidra `greeneye_hits.tsv` | @@ -231,7 +238,10 @@ recorded, and the session can be compared like any other run. 3. **Read-only by default.** `--allow-write` enables `write_mem` and `call_function`; `--allow-inject` enables `inject_service`. The agent's client config decides which manifest class it gets; a read/observe agent is - configured without those flags. + configured without those flags. Session-class tools (`session_info`, + `session_launch`) stay available without the flags: launching selects or + replaces the local served session (standalone mode) and never mutates + target memory. 4. **Sessions are the audit trail.** Writes, calls, injections and snapshots land in `events.jsonl`; run with `--record` (the trace tools require it). A finding without a session is not a finding. @@ -259,14 +269,16 @@ The tool list is stable regardless of what is present. ## 9. Testing ```bash -python3 -m unittest tests.test_mcp -v # 34 tests, no external deps +python3 -m unittest tests.test_mcp -v # 55 tests, no external deps ``` The suite drives `initialize`/`tools/list`/`tools/call` over in-memory streams (including unknown method/tool, parse errors, write/inject gating, no-stdout-pollution and structured `not_available`), plus a standalone subprocess test on the native backend that skips cleanly when the native -simulator is absent. +simulator is absent. Native-backend tools have additional end-to-end tests on +the bundled demo firmware (skipped when it is not built): `step`/`run` budgets, +`sim_cycles` accounting, `sim_regions` metadata and native backend detection. ## 10. Roadmap diff --git a/debug/risc-v-gsp-rm-vm/docs/26-ARCHITECTURE.md b/debug/risc-v-gsp-rm-vm/docs/26-ARCHITECTURE.md index fce7a96..0aefabd 100644 --- a/debug/risc-v-gsp-rm-vm/docs/26-ARCHITECTURE.md +++ b/debug/risc-v-gsp-rm-vm/docs/26-ARCHITECTURE.md @@ -254,7 +254,7 @@ flowchart LR All results are one JSON text block; absent modules return a structured `not_available` error; stdout carries protocol frames only (20 §4, §8). Full -argument schemas for all 24 tools are in 20 §5. +argument schemas for all 31 tools are in 20 §5. ## 9. Automation harness diff --git a/debug/risc-v-gsp-rm-vm/docs/30-GPGPU-SIM-INTEGRATION.md b/debug/risc-v-gsp-rm-vm/docs/30-GPGPU-SIM-INTEGRATION.md new file mode 100644 index 0000000..74830f2 --- /dev/null +++ b/debug/risc-v-gsp-rm-vm/docs/30-GPGPU-SIM-INTEGRATION.md @@ -0,0 +1,104 @@ +# 30 — GPGPU-Sim integration (in development) + +> **Status: in development (Phase 0 scaffold).** This page describes a +> development-only overlay that is being prepared to give the debugger an +> off-hardware GPU execution substrate. It makes **no accuracy, timing, +> performance or support claim**, is **not calibrated to silicon**, and is +> not a product or guest component. Nothing here authorises GPU hardware, +> kernel-module, mdev or guest access. Thresholds for the capabilities below +> exist only as unfiled drafts in +> [`../gpgpu-sim/ADMISSION.md`](../gpgpu-sim/ADMISSION.md). + +## 1. Purpose + +QEMU (with the `vgsp` host module) and the native reference core under `sim/` +cover the debugger's instruction-level needs today. The next step is an +execution substrate that can host, on one deterministic clock, both a +GPU microarchitecture model and the GSP/microcontroller models the program +needs, while reusing the debugger's session format, hooks/watchpoints, Ghidra +bridge and MCP surface. + +GPGPU-Sim is the chosen base: it is open source (BSD-3-Clause, University of +British Columbia), data-driven, and already models Turing SM75 in its tested +configurations. The overlay in +[`../gpgpu-sim/`](../gpgpu-sim/README.md) tracks the project's changes to +upstream without forking its history: a pinned clone plus an ordered patch +set. + +## 2. Status (Phase 0) + +What exists today: + +| Item | State | +|---|---| +| Overlay layout + pinned upstream policy | done — [`../gpgpu-sim/UPSTREAM.md`](../gpgpu-sim/UPSTREAM.md) (pin `90ec339`, tag v4.0.1) | +| Idempotent setup/build script | done — `../gpgpu-sim/setup.sh` (clone/checkout, ordered patches, CUDA-`CUDA_INSTALL_PATH` build; `--no-build` supported) | +| Turing TU104 (RTX 2080 SUPER) config scaffold | done — `../gpgpu-sim/configs/turing-tu104-2080super/` | +| Config consistency test | done — `../gpgpu-sim/scripts/config_sanity.py`, wired into `make test` | +| Run wrapper for CUDA binaries | done — `../gpgpu-sim/scripts/run_sim.sh` (LD_LIBRARY_PATH substitution per upstream) | +| Admission records | drafts only — [`../gpgpu-sim/ADMISSION.md`](../gpgpu-sim/ADMISSION.md), pending maintainer filing | +| CUDA 12.x build/runtime port | **done (WP0.1)** — patches `0001`/`0002`; `../gpgpu-sim/BUILD-PORT.md` records the pinned CUDA 12.0 / gcc 13 build and `vectorAdd` functional+perf green | +| Calibrated TU104 model | **not done** (Phase 1) | +| GSP core + Turing MCU models | **not done** (Phase 1.5, behind device-map and model gates) | + +## 3. What the TU104 scaffold encodes, and how + +The config is derived from the upstream SM75_RTX2060 Turing config, with only +the device-level parameters public specifications distinguish. Values carried +from public sources: + +* 48 SMs (3072 CUDA cores at 64 FP32/SM) → `-gpgpu_n_clusters 48`; +* 8 GB GDDR6, 256-bit, 15.5 Gbps effective → recorded as machine-checkable + `# SPEC:` markers plus the DRAM model (burst 16, QDR ratio 4, DRAM command + clock 3875 MHz); **memory capacity itself is not modelled by GPGPU-Sim 4.x** + and is documented, not asserted; +* ~496 GB/s → the modeled module count/geometry is chosen so + `buswidth × ratio × DRAM clock × modules` equals 496 GB/s, and + `config_sanity.py` recomputes it; +* 4 MB L2 → cache geometry × modules × sub-partitions equals 4 MiB, also + recomputed by the test; +* 64 ROPs / 192 TMUs → recorded as `# SPEC:` markers (GPGPU-Sim 4.x has no + count knobs for these); +* base ~1650 MHz / boost ~1815 MHz → base clock active, boost mode commented; +* compute capability 7.5, PCIe 3.0 x16 → CC asserted; PCIe is not modelled. + +Anything not tied to a published number keeps an explicit `# UNKNOWN:` +marker (latencies, L1/L2/DRAM timing, address mapping, Booksim topology +size). See the config header for the full list. + +## 4. How to build and run + +```bash +cd open-iov/debug/risc-v-gsp-rm-vm/gpgpu-sim + +./setup.sh # clone pin + apply patches + build +python3 scripts/config_sanity.py +./scripts/run_sim.sh --help +./scripts/run_sim.sh ~/path/to/vectorAdd +``` + +`setup.sh` is idempotent, never uses `sudo`, and caps compile jobs at 4 by +default (`JOBS` overrides). The CUDA toolkit is selected with +`CUDA_INSTALL_PATH` (default `/usr/local/cuda`). On hosts whose CUDA is not +yet ported, the build step stops with a clear error; `--no-build` still +materialises the checkout and patches. + +## 5. What is NOT claimed + +* No accuracy, cycle, latency, bandwidth or correlation claim: the config is + a documented scaffold, and config-only modeling is treated as a failure + mode, not a result. +* No calibrated model, no GSP/MCU/fabric/guest simulation yet; the real GSP + image is not loaded by this overlay. +* No support or product claim of any kind; development-only tooling, never + shipped, never in a guest. +* No hardware access from any script in this overlay. + +## 6. Admission drafts + +The five draft records (build/run, TU104 calibration, GSP-in-sim, MCU +cross-controller, sim debugger backend) with scenario, measurable outcome, +baseline, threshold and failure signal live in +[`../gpgpu-sim/ADMISSION.md`](../gpgpu-sim/ADMISSION.md). They are drafts and +require maintainer filing under the project admission rule; until then these +capabilities are pilots at most. 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. diff --git a/debug/risc-v-gsp-rm-vm/docs/32-GSP-CYCLE-MODEL.md b/debug/risc-v-gsp-rm-vm/docs/32-GSP-CYCLE-MODEL.md new file mode 100644 index 0000000..018e07c --- /dev/null +++ b/debug/risc-v-gsp-rm-vm/docs/32-GSP-CYCLE-MODEL.md @@ -0,0 +1,187 @@ +# 32 — GSP cycle model (increment 0, in development) + +> **Status: in development (Phase 1.5, work package WP1.5.1 increment 0).** +> This page describes a table-driven **modeled** cycle count added to the +> native reference core (`sim/`) and a bounded, deterministic "cycle-0" +> profile of the real 570.144 GSP firmware. It makes **no silicon timing, +> performance, latency or accuracy claim**: the numbers are modeled defaults +> under an explicit table, not measurements. The simulator's fidelity +> statement is unchanged — L0 (instruction semantics) and L1 (memory +> behaviour) are modelled; **L2 (time/concurrency) is not** (`06-ACCURACY.md`). +> Development-only: nothing here authorises GPU hardware, kernel modules, +> mdev or guest access. + +## 1. Purpose + +The native simulator reports virtual time as an instruction count (`insns`). +That is enough for functional debugging but not for comparing instruction +mixes, for calibration-table work, or for a first, honest cycle-shaped +artifact from a real firmware run. Increment 0 adds: + +* `sim/cycles.py` — an isolated, per-instruction-class latency table with a + deterministic cost function; +* a `cycles` counter on `sim.machine.Machine`, advanced by that table + alongside the existing instruction count; +* `sim/gsp_profile.py` — a bounded, read-only run of the real + `firmware/57X-GSP/570.144/gsp-570.144.fwimage.bin` that records a normal + `greeneye-session/1` directory with the modeled cycle totals; +* `tests/test_sim_cycles.py`. + +Everything is stdlib-only Python 3.10+, offline, and read-only with respect to +the firmware image. + +## 2. The model + +`sim/cycles.CycleModel` is a plain Python table. All defaults are **modeled** +from public, generic in-order RISC-V pipeline assumptions — one cycle for +plain ALU work, conservative multi-cycle M-extension costs, a fixed extra +latency for a completed memory access — and are deliberately *not* derived +from the GSP core, from silicon or from vendor data. + +| Class | Default | Notes | +|---|---|---| +| ALU / shift / `lui` / `auipc` / op-imm | 1 | add, sub, logical, compare | +| branch | 1 | taken/not-taken penalty deliberately not modelled | +| `jal` / `jalr` | 1 | | +| `mul` / `mulh` / `mulhsu` / `mulhu` | 3 | conservative | +| `div` / `divu` / `rem` / `remu` | 20 | conservative, data-independent | +| load / store | 1 | plus a completed-access latency (below) | +| AMO (`lr`/`sc`/`amo*`) | 2 | plus one completed-access latency | +| CSR | 2 | | +| trap (`ecall`/`ebreak`/`mret`) | 3 | | +| `fence` / `fence.i` | 1 | no ordering model (single hart) | +| illegal | 1 | execution stops immediately | + +Memory latency is a **parameter**, not a cache model: + +| Parameter | Default | Meaning | +|---|---|---| +| `memory_latency` | 3 | extra cycles when a RAM access completes | +| `mmio_latency` | 12 | extra cycles when an unmodelled-MMIO access completes | + +The cost of one instruction is `instruction_cost(kind, sub)`; a completed +load/store/AMO adds `memory_cost(is_mmio)` on top. Cost depends only on the +decoded instruction class (`sim/decode.py` kind/sub) and the access region: +no host clock, no entropy, no caches, no hazards. `CycleModel.describe()` +emits the full table for session metadata. + +## 3. Machine integration + +`sim.machine.Machine` keeps its previous behaviour exactly; cycles are +additive: + +```python +from sim.machine import Machine + +m = Machine() # DEFAULT_MODEL active +m.run(10_000) +m.insns # virtual time (unchanged semantics) +m.cycles # modeled cycles, advanced by the table +m.cycle_model.name # "rv64imac-modeled/1" + +off = Machine(cycle_model=False) # cycles stays 0; disabling path for A/B +slow = Machine(cycle_model=CycleModel(memory_latency=9)) +``` + +`greeneye.sim_backend.launch_native(..., cycle_model=...)` threads the choice +through the Session API (`None` = default table, `False` = off, or an +instance). The counter increments for the instruction the machine stops on +under a watchpoint/breakpoint, mirroring how `insns` counts it; the memory +latency is only added when the access actually completes. Determinism is by +construction: two identical runs produce identical instruction and cycle +totals (tested). + +Measured overhead on this host (CPython 3.12, 200k instructions of the real +image, best of three): ~0.073 s with accounting disabled vs ~0.082 s enabled +(~12%); disabled runs are byte-for-byte the pre-increment behaviour. + +## 4. The cycle-0 GSP profile + +```bash +# bounded, deterministic, read-only; writes a session under dumps/ +python3 -m sim.gsp_profile --instructions 200000 + +# explicit session directory and machine-readable result +python3 -m sim.gsp_profile --record dumps/gsp-cycle0 --json + +# start from a harvested first-touch map and keep what the run learned +python3 -m sim.gsp_profile --device-map dumps/<probe>/device-map.json --learn --learn-out /tmp/map.json + +# image absent (e.g. LFS not fetched): exits 2 with a skip reason, no session +python3 -m sim.gsp_profile --image /nonexistent.bin +``` + +The profile mirrors the established device-map/probe path: the ELF's PT_LOAD +plan is relocated by `bias=0x80000000` (`greeneye/elf.py`), RAM is sized to +cover the image as `greeneye/device_map.py` computes it, and all unmodelled +accesses are served leniently (zero or the mapped value) through a +`greeneye-device-map/1`. The instruction budget is exact, not wall-clock: +`Machine.run(budget)` is called directly, so the run is deterministic. + +Session metadata (`session.json`) and a `gsp_cycle_profile` event carry: +`firmware_sha256`, `bias`, `instructions_executed`, `cycles`, +`cycle_model`, `cycle_model_table`, `cycle_fidelity`, `stop_reason`/`stop_pc` +(both relocated and firmware-space), `mmio_unknown` and the first unknown +MMIO access. + +**Observed output** (this tree, 2026-10-03; sha256 +`2d1bb5c6…c3ceea`, budget 200 000): + +``` +instructions: 200000 +cycles (modeled): 200012 +stop: budget pc=0x84009070 (fw 0x4009070) +unknown MMIO accesses: 1 +first unknown MMIO: pc=0x84009014 addr=0x2000000000000a00 size=4 rw=r +``` + +This is a bounded lenient first-touch profile, **not a boot**: it reproduces +the documented QEMU probe first-touch address and continues under the modeled +lenient policy for the budget. The stop PC is a budget sample inside the +firmware's early path, not a milestone. + +## 5. What is not claimed + +* No silicon timing, latency, throughput, frequency or pipeline claim; the + table is a modeled default and says so in code, metadata and reports + (`sim/cycles.py::NOTE`, `CYCLE_FIDELITY`). +* No cache/memory-hierarchy model: `memory_latency`/`mmio_latency` are fixed + parameters, not simulated caches or DRAM. +* No device model: L3 remains absent (`06-ACCURACY.md`); the profile runs + under lenient MMIO, so its control flow after the first unknown access is + conditioned on that policy, not on device behaviour. +* No accuracy metric may be computed from this increment alone. Calibration + (`RISC-V-GSP-RM-DEBUGGER-ROADMAP.md`, `16-CUSTOM-SIMULATOR.md`) still + requires sourced latency tables and, for any product-facing claim, + supervised physical confirmation. + +## 6. Tests + +`tests/test_sim_cycles.py` (21 tests, all offline): + +* table sanity (every decoded kind has a positive cost; ALU/branch/jump unit + cost; M-extension ordering; memory latency is a parameter; labeling); +* counter arithmetic on hand-assembled programs (RAM vs MMIO latency); +* disabling the model leaves `cycles` at 0 and changes no execution state; +* determinism double-runs at machine and profile level; +* the bundled demo executes identically with and without the model; +* real-image bounded profile: first touch matches the documented probe + address, session metadata carries the cycle totals, two runs are identical; +* absent image skips cleanly. + +## 7. Next steps + +1. Device model (WP1.5.2): falcon/SEC2 windows, CSB/TRANSCOM, WPR2/FWSEC, + PTIMER, IRQ routing — the run then stops modelling MMIO leniently. +2. Boot shims with recorded assumptions (WP1.5.3) so the profile can report a + real milestone instead of a budget sample. +3. Sourced instruction/access latency tables and a documented calibration + metric (`06-ACCURACY.md`, `16-CUSTOM-SIMULATOR.md`) before any + timing-shaped statement; per-MCU cycle domains come with + WP1.5.4/WP1.5.5. +4. Additive `cycle` event class in the session schema (roadmap V2, "formats + and versioning") once the event vocabulary is reviewed. + +The increment is deliberately a table and a counter: every later step can +replace defaults without touching the execution loop (the model is passed in) +or the session format (metadata is additive). 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). diff --git a/debug/risc-v-gsp-rm-vm/docs/34-BLACKWELL-SOURCES.md b/debug/risc-v-gsp-rm-vm/docs/34-BLACKWELL-SOURCES.md new file mode 100644 index 0000000..f0e28df --- /dev/null +++ b/debug/risc-v-gsp-rm-vm/docs/34-BLACKWELL-SOURCES.md @@ -0,0 +1,190 @@ +# 34 — Blackwell sources and known-public parameters (planned / uncertain) + +> **Status: planned, uncertain, do not implement.** Phase 3 (U10/U17) is +> gated behind the Phase 2 exit and a legal/source review +> (`handoff-simulator.md` §5, Phase 3). This document is a bibliography and +> a confidence-labelled parameter register only. It is **not** a model spec +> and must not be used to seed runnable configuration until Phase 3 entry. +> +> **Rules applied here:** +> * every number has a source and a confidence level; +> * values that are not public are written as **open questions** — never +> invented, interpolated or copied from a sibling architecture; +> * first-party NVIDIA publications outrank aggregators; where only an +> aggregator exists the row is labelled accordingly; +> +> **Confidence legend:** **A** = first-party NVIDIA publication (whitepaper, +> newsroom, datasheet, product page); **B** = first-party claim corroborated +> by at least one independent secondary source; **C** = secondary source only +> (aggregator/press, single-source or unverifiable against NVIDIA); +> **U** = unknown / not found — an explicit open question. + +## 1. Bibliography + +| # | Source | Type / URL | +|---|---|---| +| [RBW] | NVIDIA, *NVIDIA RTX Blackwell GPU Architecture* whitepaper (Jan 2025, v1.x) | First-party whitepaper, consumer GB20x. https://images.nvidia.com/aem-dam/Solutions/geforce/blackwell/nvidia-rtx-blackwell-gpu-architecture.pdf | +| [B200-N] | NVIDIA Newsroom, *NVIDIA Blackwell Platform Arrives to Power a New Era of Computing* (GTC, Mar 2024) | First-party press release, B100/B200/GB200. https://nvidianews.nvidia.com/news/nvidia-blackwell-platform-arrives-to-power-a-new-era-of-computing | +| [GB200] | NVIDIA, *GB200 NVL72* product page | First-party product/spec page. https://www.nvidia.com/en-us/data-center/gb200-nvl72/ | +| [GB300] | NVIDIA, *GB300 NVL72* product page | First-party product/spec page. https://www.nvidia.com/en-us/data-center/gb300-nvl72/ | +| [SPARK] | NVIDIA, *NVIDIA DGX Spark* product page (Specifications section) | First-party product page, GB10. https://www.nvidia.com/en-us/products/workstations/dgx-spark/ | +| [SPARK-DS] | NVIDIA, *DGX Spark* datasheet (GTC 2025 spring) | First-party datasheet, GB10. https://nvdam.widen.net/s/tlzm8smqjx/workstation-datasheet-dgx-spark-gtc25-spring-nvidia-us-3716899-web | +| [CUDA-CC] | NVIDIA, *CUDA C Programming Guide* — compute capability 10.x / 12.x | First-party ISA/feature statement. https://docs.nvidia.com/cuda/cuda-c-programming-guide/ | +| [OCP-MX] | Open Compute Project, *OCP Microscaling Formats (MX) Specification v1.0* (2024) | Industry spec defining MXFP4/MXFP6. https://www.opencompute.org/documents/ocp-microscaling-formats-mx-v1-0-spec-final-pdf | +| [MX-PAPER] | Rouhani et al., *Microscaling Data Formats for Deep Learning*, arXiv:2310.10537 (2023) | Published format analysis. | +| [WIKI-BW] | Wikipedia, *Blackwell (microarchitecture)* (retrieved 2026-10-03), citing TechPowerUp GPU database and AnandTech | Secondary aggregation; provenance for per-die core counts. https://en.wikipedia.org/wiki/Blackwell_(microarchitecture) | + +## 2. Datacenter Blackwell (GB100 / B100 / B200 / GB200) + +| Parameter | Value | Confidence | Source | Notes | +|---|---|---|---|---| +| Die | GB100, TSMC 4NP, reticle-limit size | A | [B200-N], [WIKI-BW] | | +| Transistors per die | 104 B | C | [WIKI-BW] (citing AnandTech) | NVIDIA does not headline a per-die count. | +| Package (B100/B200) | two GB100 dies, 208 B transistors total | A | [B200-N] | | +| Die-to-die link | 10 TB/s (NV-HBI) | A | [B200-N] | B200 acts as one unified GPU with cache coherency. | +| NVLink 5 per GPU | 1.8 TB/s bidirectional | A | [B200-N] | NVL72 supports up to 576 GPUs. | +| NVLink-C2C (GB200 → Grace CPU) | 900 GB/s | A | [B200-N] | "ultra-low-power NVLink chip-to-chip interconnect". | +| NVLink switch total (NVL72) | 130 TB/s | A | [GB300] (identical fabric class) | Cited on the GB300 page; GB200 NVL72 is the same 72-GPU fabric. | +| CUDA cores per die | 18,432 | C | [WIKI-BW] (TechPowerUp) | Not published by NVIDIA. | +| Tensor cores per die | 576 | C | [WIKI-BW] (TechPowerUp) | 4/die? No — per-die count not first-party; treat as unconfirmed. | +| L2 per die | 60 MB | C | [WIKI-BW] (TechPowerUp) | | +| Memory interface | 8,192-bit HBM3e per package | C | [WIKI-BW] | | +| B200 memory | 192 GB HBM3e, ~8 TB/s | B/C | widely reported; **not captured from a first-party page in this survey** | open question §5.1 | +| B100 memory | 192 GB HBM3e (variant differences open) | C | [WIKI-BW] | B100/B200 memory split varies by SKU; verify. | +| Compute capability | 10.x | A | [CUDA-CC] | | +| Low-precision formats | MXFP6, MXFP4 (OCP microscaling), FP8, INT8/INT4, binary | A | [B200-N] (Transformer Engine), [OCP-MX], [MX-PAPER] | 5th-gen tensor core with micro-tensor scaling. | +| FP4 performance | 20 PFLOPS (GB200 dual-GPU, excluding 2× sparsity claim) | A | [GB200] (as cited by [WIKI-BW]) | Per-GPU and per-die numbers need the technical brief; see §5.1. | +| Products | B100/B200 SXM, GB200 Grace-Blackwell superchip | A | [B200-N] | | + +## 3. Consumer Blackwell (GB202 / GB203 / GB205, RTX 50 series) + +All rows in this section are **[A]** from the RTX Blackwell whitepaper +unless noted; this is the **best-documented Blackwell part** and the natural +first Phase-3 target on the consumer side. + +| Parameter | GB202 full | RTX 5090 | Source | +|---|---|---|---| +| GPCs / TPCs / SMs | 12 / 96 / 192 | 11 / 85 / 170 | [RBW] §GB202, Table 3 | +| CUDA cores / SM | 128 | 128 | [RBW] | +| CUDA cores / GPU | 24,576 | 21,760 | [RBW] | +| Tensor cores / SM / GPU | 4 (5th gen) / 768 | 4 (5th gen) / 680 | [RBW] Table 3 | +| RT cores | 192 (4th gen) | 170 (4th gen) | [RBW] | +| Register file / SM | 256 KB | 256 KB | [RBW] §SM | +| L1/shared per SM | 128 KB | 128 KB | [RBW] §SM | +| L2 | 128 MB | 96 MB | [RBW] §GB202 | +| Memory | 512-bit GDDR7 (16 × 32-bit MCs) | 32 GB GDDR7, 28 Gbps, 1,792 GB/s | [RBW] Table 3 | +| Boost clock | — | 2,407 MHz | [RBW] Table 3 | +| FP32 (non-tensor) | — | 104.8 TFLOPS | [RBW] Table 3 | +| FP16/BF16/INT32 (non-tensor) | — | 104.8 / 104.8 / 104.8 TFLOPS | [RBW] Table 3 (INT32 fully unified per SM) | +| FP4 tensor | — | 1,676 / 3,352 TFLOPS (non-sparse / sparse) | [RBW] Table 3 | +| FP8 tensor (FP16 acc) | — | 838 / 1,676 TFLOPS | [RBW] | +| FP16 tensor (FP16 acc) | — | 419 / 838 TFLOPS | [RBW] | +| FP16 tensor (FP32 acc) / BF16 | — | 209.5 / 419 TFLOPS | [RBW] | +| TF32 tensor | — | 104.8 / 209.5 TFLOPS | [RBW] | +| INT8 tensor | — | 838 / 1,676 TOPS | [RBW] | +| FP64 | — | 1/64 of FP32 rate, 384 FP64 cores (2/SM) on the full die | [RBW] §GB202 note | `[derived]` ≈ 1.64 TFLOPS for the full chip. | +| Transistors / die size | 92.2 B / 750 mm² | 92.2 B / 750 mm² | [RBW] Table 3 | +| TGP | — | 575 W | [RBW] Table 3 | +| SM changes | INT32 fully unified with FP32 (2× many integer ops; unified cores are either FP32 or INT32 per cycle); one 4th-gen RT core + four 5th-gen tensor cores per SM | | [RBW] §SM | + +## 4. GB10 (Grace Blackwell superchip; DGX Spark class) and GB300 + +### 4.1 GB10 + +| Parameter | Value | Confidence | Source | +|---|---|---|---| +| CPU | 20-core Arm: 10 × Cortex-X925 + 10 × Cortex-A725 | A | [SPARK] | +| GPU | Blackwell architecture, 5th-gen tensor cores, 4th-gen RT cores | A | [SPARK] | +| Tensor performance | up to 1 PFLOP FP4 (theoretical, sparsity) | A | [SPARK], [SPARK-DS] | +| System memory | 64 GB (OEM) or 128 GB LPDDR5X, coherent unified | A | [SPARK], [SPARK-DS] | +| Memory interface / bandwidth | 256-bit / 273 GB/s | A | [SPARK], [SPARK-DS] | +| NVLink-C2C | "5× the bandwidth of PCIe Gen 5" (exact GB/s not stated) | A (ratio only) | [SPARK-DS] | +| Networking | ConnectX-7 200 Gb/s; up to 4 systems connected | A | [SPARK] | +| TDP | 140 W (GB10 chip, incl. CPU+GPU) | A | [SPARK] | +| SMs / CUDA cores | 48 SMs / 6,144 CUDA cores | C | [WIKI-BW] (TechPowerUp) | +| Tensor / RT cores | 384 / 48 | C | [WIKI-BW] | +| L1 / L2 | 6 MB / 50 MB | C | [WIKI-BW] | +| Compute capability | 12.x (consumer Blackwell family; confirm GB10 SKU) | U | — | open; see §5.2 | + +### 4.2 GB300 / Blackwell Ultra + +| Parameter | Value | Confidence | Source | +|---|---|---|---| +| Product | GB300 NVL72: 72 Blackwell Ultra GPUs + 36 Grace CPUs | A | [GB300] | +| FP4 dense uplift | 1.5× dense FP4 FLOPS vs Blackwell (GB200-class) | A | [GB300] | +| Attention performance | 2× higher attention vs Blackwell | A | [GB300] | +| NVLink fabric | 130 TB/s per rack | A | [GB300] | +| GPU memory (rack) | 20 TB HBM3e class ("Fast Memory 37 TB" total incl. CPU) | A | [GB300] table | +| HBM bandwidth (rack) | up to 576 TB/s aggregate | A | [GB300] table | +| CPU memory | 17 TB LPDDR5X, 14 TB/s | A | [GB300] | +| CPU cores | 2,592 Arm Neoverse V2 cores | A | [GB300] | +| Rack compute | FP4 1,440 / 1,080 PFLOPS (sparse / without sparsity); FP8-FP6 720 PFLOPS; INT8 24 POPS; FP16/BF16 360 PFLOPS; TF32 180 PFLOPS; FP32 6 PFLOPS; FP64 100 TFLOPS | A | [GB300] table | +| Per-GPU derived | 288 GB HBM3e, ~8 TB/s HBM, 20 PFLOPS FP4 sparse / 15 PFLOPS dense | B | `[derived]` from [GB300]: 20 TB ÷ 72 ≈ 278–288 GB (nominal 288 GB), 576 TB/s ÷ 72 = 8 TB/s, 1,440 ÷ 72 = 20 PFLOPS | mark derived; NVIDIA has separately quoted 288 GB per Blackwell Ultra GPU. | +| GPU clocks, SM counts, cache sizes | — | U | — | open; see §5.3 | + +## 5. Open questions (explicit; never fill by invention) + +### 5.1 GB100 / B200 +1. Per-die SM count, CUDA/tensor core counts and per-SM composition are not + in a first-party public source found in this survey ([WIKI-BW] only). +2. HBM3e capacity/bandwidth per SKU (B100 vs B200, 192 GB vs other) needs + the NVIDIA Blackwell technical brief or product brief. +3. L2 organization/coherence across the two dies beyond "cache coherent" + ([B200-N]) is unpublished. +4. NV-HBI latency/bandwidth characteristics beyond the 10 TB/s peak. +5. No per-instruction latency table exists for Blackwell comparable to + [AMP]/[T4]. A microbenchmark campaign (own or literature) is required + before any C1 claim. + +### 5.2 GB10 +6. Exact NVLink-C2C bandwidth in GB/s (only the "5× PCIe Gen 5" ratio is + first-party). +7. Whether GB10's compute capability is 12.0 or 12.x for the relevant SKU. +8. SM/core/cache counts are secondary-source only. +9. LPDDR5X timing and real achievable bandwidth (273 GB/s is the + first-party theoretical interface rate). + +### 5.3 GB300 / Blackwell Ultra +10. Per-GPU clocks, SM counts, cache sizes and tensor shape/issue details. +11. Whether per-GPU HBM is exactly 288 GB on every SKU (the rack arithmetic + is approximate). +12. NVLink-C2C generation/bandwidth in GB300 (assume same class as GB200 + only as a hypothesis, not a model input). +13. FP4/MXFP4 numerical semantics for simulation (block scaling, exponent + sharing) beyond the OCP spec. + +### 5.4 Cross-cutting +14. A public **Hot Chips** presentation with GB100/GB202 microarchitecture + details was **not found** in this survey pass. If one exists (or is + released later), add it here with the exact deck title/URL — this is the + single highest-value missing source. +15. Blackwell memory-latency, sectoring, cp.async-successor semantics and + SM issue details are unpublished; Phase 3 must either measure locally + (supervised, hardware-permitting) or ship explicit uncertainty bands. +16. Blackwell datacenter firmware/boot material is out of scope here and + remains gated by the RM 610 no-fly zone (docs/23) — no 610 artifact may + be fetched or analysed for Phase 3 without the approval token. + +## 6. What a Phase-3 smoke parameter set may and may not contain + +* **May** use §3 rows (all first-party) as the consumer starting point, with + per-row uncertainty labels stored in the simulator session metadata. +* **May** use §2/§4 first-party *aggregate* claims as sanity bands. +* **May not** use any [C] row as a calibration target; [C] rows are leads to + verify, not values to fit. +* **May not** interpolate GB100 from GB202 or from Hopper; die topology, + clock domains and cache hierarchy are different designs. +* Any physical Blackwell capture is supervised-only and never an agent task + (RUNBOOK P4/P7); any firmware/BSP input obeys docs/23. +* Exit evidence for Phase 3 is defined in WP3.1–WP3.3: source dossier (this + document), an explicit uncertainty report, and a device profile — not an + accuracy claim. + +## 7. Bibliography maintenance + +Add to §1 whenever any of these appear: NVIDIA Blackwell technical brief +(datacenter), NVIDIA GB10/C2C technical brief, Hot Chips/ISSCC decks, +peer-reviewed Blackwell microbenchmarking papers, official CUDA PTX ISA +additions for `sm_100`/`sm_120` (e.g. tcgen05/5th-gen tensor instructions), +and OCP MX follow-ups. Keep [C] rows attached to their aggregator and +re-verify before promotion. diff --git a/debug/risc-v-gsp-rm-vm/docs/35-GREENEYE-SIMULATOR-DESIGN.md b/debug/risc-v-gsp-rm-vm/docs/35-GREENEYE-SIMULATOR-DESIGN.md new file mode 100644 index 0000000..83954d4 --- /dev/null +++ b/debug/risc-v-gsp-rm-vm/docs/35-GREENEYE-SIMULATOR-DESIGN.md @@ -0,0 +1,291 @@ +# 35 — GreenEye simulator design (single design document) + +> **Status:** in development (2026-10-03). Phase S (standalone GSP simulator) +> is the only implementation scope until its pause gate; Phase G (GPGPU-Sim +> board model) is designed and gated, not started. +> +> **Audience:** agents and maintainers of risc-v-gsp-rm-debugger. +> + +--- + +## 0. What this document is + +One consolidated design for the simulator program requested for +risc-v-gsp-rm-debugger. It covers, in order of execution: + +1. **Phase S — standalone GSP simulator** (now): execution of the real + firmware under the modelled, evidence-labelled device environment (L0/L1; + silicon accuracy is **not** claimed), full tool/workflow integration, and + the above-the-QEMU-line contract on the frozen checks (docs/37 §3). + **The program pauses at Phase S exit** before any GPGPU-Sim work. +2. **Phase G — GPGPU-Sim full-board model** (later, gated): TU104 board + devices (GMMU, FBIF, FALCON→GMMU DMA, vidmem heap, VBIOS), GSP↔core + interactions, PCIe/VFIO bridge for a vGSP VF, QEMU-gdb time sync, MCP + extension, multi-GPU/fabric backlog. +3. **Plan B** — the narrowly scoped reference plan for "a simulator simply + better than the QEMU baseline for the open-iov workflow" + (`../Handoff-Simulator-Plan-B.md`, private), with an explicit Plan A/Plan B + comparison and decision procedure. + +Owner directives are numbered U1–U45 in the private handoff (§1 and §13); +this document restates the operative ones and never contradicts them. + +## 1. Definitions + +### 1.1 The QEMU line (the bar to beat) + +The **QEMU baseline** is the v0.1/v0.2 tool behaviour measured on the real +GSP image: RSP transport, hooks/breakpoints/watchpoints, memory/struct/ +snapshot tooling, sessions, symbols/analyze/compare, probes/plugins, Ghidra +symbol + control bridges, TCG trace, injection SDK, `Session.call`/ +`run_until`; and on the real `gsp-570.144.fwimage.bin`: relocation-free flat +load, first exception at firmware vaddr `0x4009014` (load access fault +reading `0x2000000000000A00`), then `mtvec=0` fetch loop. Planned QEMU +features (device model H2.2, boot shims H2.3, deterministic multi-core H2.4, +RM service injection H1.3, trace→session folding R4, debugger parity H1.5) +are **part of the bar** even though QEMU has not implemented them. + +**"Above the QEMU line"** means, at minimum: + +* executes the real 570.144 image deterministically to at least one recorded + milestone **past** the QEMU first fault (ladder in §3.3), with the device + model and shims documented; +* every QEMU feature available against the native backend, plus the planned + QEMU features above; +* every run produces `greeneye-session/1` evidence with simulator build hash, + firmware hash, profile/hash, shim assumptions and fidelity labels; +* run-to-run determinism proven; cycle counts emitted and labelled *modeled*; +* a real-workload demonstration whose eng-ops baseline/threshold were frozen + before the run. + +### 1.2 Fidelity labels + +`L0` instruction semantics, `L1` memory behaviour, `L2` time/concurrency, +`L3` device environment (extended from docs/06). The standalone GSP +simulator targets **L0–L2 for the GSP core and its device model**, with L2 +explicitly labelled *modeled* (table/source-derived) until calibrated. The +GPGPU-Sim board model adds board-level L3. No silicon-accuracy claim without +calibration evidence and a declared uncertainty band. + +### 1.3 Accuracy metrics + +* `M1–M6` — the existing metric contract (inventory, replies, ordering, state, + error paths, memory side effects; docs/06). +* `C1–C6` — cycle calibration (instruction latency agreement, memory latency, + bandwidth, kernel correlation, determinism, clock-domain ratio); bands are + frozen in the private handoff §9.2 and must be recorded before first use. + +## 2. The workflow this exists to serve + +The simulator must move one of the project's OODA loops faster with evidence +(eng-ops §3): **dynamic RE** (what does the firmware do?), **portability** +(new firmware/device triage), **adversarial** (unexpected behaviour), and +**correctness** (does our emulation match?). Time-to-value is the currency: +every capability below names the baseline it replaces and the number that +decides keep/redesign/retire. + +## 3. Phase S — standalone GSP simulator (NOW) + +### 3.1 Architecture + +``` +greeneye control plane (Python; unchanged public API) + Session · hooks · bp/watch · probes · plugins · recorder · analyze/compare + Ghidra bridge · MCP · injection SDK + │ Backend ABI (typed; QEMU and native are peers) + ▼ +native backend (sim/) + machine.py RV64IMAC core, deterministic virtual time, cycle model + devices/ NEW: GSP device model (falcon/SEC2 windows, CSB/TRANSCOM, + WPR2/FWSEC state, PTIMER, IRQ routing, DMA/IOMMU, PMU stub) + profiles/ NEW: versioned device profiles (JSON) + boot shims + trace.py exec/access/cycle events → session blob + moments (S2) snapshot/restore + deterministic replay +``` + +*The QEMU backend stays untouched and remains the differential reference.* +Backend selection is explicit (`--backend qemu|native`); no shared mutable +state; the session format evolves additively (`greeneye-session/1` + new +event classes `cycle`, `device`, `mcu`). + +### 3.2 Work packages + +| WP | Name | Deliverable | Agent | +|---|---|---|---| +| S0 | baseline (done) | native core, session parity, M1.5 trace, modeled cycles, real-image bounded run | — | +| S1 | GSP device model + boot ladder | `sim/devices/` + `sim/profiles/` + milestone sessions + tests | S1 | +| S2 | debugger parity | catchpoints, core/thread selection, memory-region metadata, moments/replay, parity matrix | S2 | +| S3 | MCP + scripting + injection | MCP tools against native backend; injection against the simulated GSP; docs | S3 | +| S4 | real-workload validation | boot-ladder A/B vs QEMU probe; scenario harness; value demo + report | S4 | +| S5 | stability/packaging | long-run soak, error taxonomy, session size budgets, quickstart, release notes | successor | + +### 3.3 Boot ladder (S1 acceptance) + +Each rung is a recorded session with a session hash; each rung must be +reproducible twice identically and have its device requirements listed. + +| Rung | Milestone | Current status | +|---|---|---| +| B0 | flat image loads; first MMIO observed at `0x2000000000000A00` | done (native, lenient) | +| B1 | device model serves the whole boot path: 0 unknown MMIO, 0 unhandled in windows; deterministic 2M-instruction run stopping at the boot-halt/report loop (`0x4010928`) | done (S1 device model; sessions `dumps/gsp-boot-b1-run1/2`) | +| B2 | boot stage progress: vector table/TCM/MPU/trap handling exercised with the shim; passes the bootargs consumer and reaches the relocation copy, then stops on the known DMEM-runspace unknown (`fault @ 0x400EA0C`; docs/36 §4) | done (shim) | +| B3 | first CSB/TRANSCOM mailbox poll with correct read-back semantics | target | +| B4 | message loop reached (`appMain`/equivalent); waits on host message deterministically | target | +| B5 | one scripted host-side RM message served end-to-end with a recorded reply | target (gated by S3 injection) | + +The ladder is deliberately versioned-data driven: a new firmware version is a +new profile + shim record, not a code fork. + +### 3.4 What is explicitly not in Phase S + +GPGPU-Sim changes, board devices beyond the GSP complex, guest/PCIe models, +VBIOS, VF bridge, Ampere/Blackwell, fabric — all Phase G or backlog (§10). + +### 3.5 Pause gate + +Phase S ends when: S1–S4 exits met; suite green incl. new tests; a real +workload demo exists with a frozen baseline and measured time-to-value; +`docs/30–40` updated; adversarial review (§9) signed off. **Pause and report +before any Phase G work.** + +## 4. Phase G — GPGPU-Sim board model (later, gated) + +This is the previously requested GPGPU-Sim program, now explicitly sequenced +after Phase S. It is *designed* now so Phase S artifacts (device map, +profiles, session/event classes, metrics, MCP surface) fit it without rework. + +### 4.1 Board device inventory (TU104 family) + +Expanded from owner requirements (GMMU, FBIF, FALCON→GMMU DMA, vidmem heap) +and the porting dossier. + +| Device | Model target | Evidence class | +|---|---|---| +| GMMU (+ PTEs, page tables, ATS/IOMMU seam) | page-table walk, TLB, fault reporting, VA→PA translation for firmware descriptors | register/field maps + formats; A/B vs probe | +| FBIF (FB interface, compression, ECC) | FB access path, alignment/partition hashing, vidmem heap model | register maps; documented assumptions | +| FALCON→GMMU DMA line | descriptor-driven DMA with cycle/latency accounting | boot/mainline firmware usage; source traces | +| vidmem heap stack | allocator map, reserved regions, WPR2 carve-outs | boot state + shims | +| GSP core + CSB/TRANSCOM | message queues, doorbells, interrupts | Phase S artifacts | +| PMU / SEC2 / booter / ACR/SOE | interfaces, mailboxes, state machines, latencies | Phase S MCU models | +| VBIOS/PROM | image layout, init tables, device declaration for guests | public parsing + profile | +| PCIe endpoint | config space, BARs/MMIO windows, MSI-X, TLPs, D2D | descriptor + QEMU/GDB observations | +| VF bridge | vGSP VF traffic → simulator BAR/MMIO instead of real GPU | design-only until human approval | + +### 4.2 Light-touch VF bridge (design) + +* Userspace gpgpu-sim hosts the board model; a future (separately approved) + kernel module presents a PCIe/VFIO device FD. +* Protocol: shared-memory queue + eventfd doorbells; simulator services + MMIO/BAR reads/writes, MSI-X injection and DMA windows as messages; no + changes to `vgsp` primitives or data paths. +* **Hard rule:** no `vgsp` kernel-module edits without explicit human + approval. Phase S/G implement everything up to that boundary; the module is + design + tests on stubs only. +* A conformance harness replays recorded QEMU/VFIO traffic against the + simulator model and compares replies (accuracy loop), so the bridge is + proven before any kernel work is proposed. + +### 4.3 QEMU-gdb time sync (design) + +Guest QEMU and the simulator exchange a clock-sync protocol (periodic +time-base beacons + event ordering barriers). Deterministic mode: simulator +virtual time is the master; QEMU is granted step/sync credits. Never a +security boundary; no confidentiality claims. + +### 4.4 MCP extension (design) + +The existing 31-tool MCP surface grows simulator-native tools: run/step with +budget, cycle queries, device-state dumps, MCU state, boot-ladder report, +board topology, DMA/TLP traces. Read/observe default; writes/step gated by +flags; every call attributable to a session. + +## 5. Feature parity matrix + +| Capability | QEMU (current) | QEMU (planned) | Native GSP (now) | Native GSP target | +|---|---|---|---|---| +| RSP-free transports/launch | yes | — | launch/attach | yes | +| Hooks/watchpoints/breakpoints | yes | catchpoints/core select | hooks/watch/bp | full parity (S2) | +| Memory/struct/snapshots | yes | region metadata | memory/struct | snapshots/moments (S2) | +| Sessions/analyze/compare | yes | trace folding | yes + cycles | trace folded, coverage (S2) | +| Probe/record/plugins | yes | — | backend peer | yes | +| Ghidra symbol/control bridges | yes | watch/heat-map/nav | protocol-compatible | full (S2/S3) | +| Injection SDK | demo | RM-core interception H1.3 | demo backend | services vs GSP model (S3) | +| Device model / real boot | no (faults) | H2.2/H2.3/H2.4 planned | ladder B0–B2 (S1, device model + shim) | B4/B5 (S1/S3) | +| Determinism/replay | wall-clock | H2.4 budgets | instruction-count | event replay (S2) | +| MCP | yes (QEMU) | — | native sim tools + backend selection (20) | wider GSP-model tools (S3) | +| Tiering | — | — | modeled cycles | calibrated bands (S4) | + +## 6. Boundaries and invariants + +1. All changes live under `debug/risc-v-gsp-rm-vm/` (private handoff copies at + workspace root excepted). Never edit project directories outside it. +2. **No `vgsp` kernel-module edits without explicit human approval** (U33). +3. No physical GPU, mdev, guest, module loads, `/dev` or reboots from agents; + RM 610 is a no-fly zone (docs/23); no unattended cycles (RUNBOOK P1–P9). +4. The simulator is a study environment, never a product dependency, never in + a guest, no security claims (RUNBOOK P5/P6). +5. Upstream GPGPU-Sim stays a pinned read-only checkout; our changes are + tracked patches + overlay (Phase G), never a silent fork. +6. The QEMU path must remain green while native features land (differential + suite is the regression net). + +## 7. eng-ops admissions + +`gpgpu-sim/ADMISSION.md` holds draft records TOOL-2026-SIM-01..05 (build/run, +TU104 calibration, GSP-in-sim, MCU model, debugger backend) with the five +required fields; Phase S adds one admission record per S work package (S1 +device model, S2 parity, S3 MCP/injection, S4 validation). Thresholds are +frozen before the first run and never moved silently; every result carries a +session directory and hashes. + +## 8. Real-workload proof (S4) + +The "real workload" for Phase S is the open-iov workflow on the real image: +boot the 570.144 GSP to the highest ladder rung, expose device/MCU state +through the standard tooling (CLI/Python/MCP/Ghidra), serve one scripted +message, and record sessions. Baselines: (a) QEMU `probe` first fault; +(b) static analysis time estimates from the project handoff. The demo is +accepted only if it beats the frozen threshold; otherwise the failure signal +triggers redesign per eng-ops §4. + +## 9. Adversarial review + +| Risk | Failure signal | Mitigation | +|---|---|---| +| "Accuracy" theatre (lenient MMIO masquerades as a model) | rung claimed with lenient policy | each rung requires explicit device semantics; lenient runs labelled | +| Device model guesses without evidence | unexplained delta in A/B | every register backed by source or observed firmware use; record assumption | +| Parity breaks the QEMU path | differential suite red | QEMU untouched; feature-flag additions | +| Scope creep into Phase G | files outside `sim/`+docs | gate checks; private handoff scope contract | +| Non-determinism | replay delta non-empty | virtual-time only; double-run tests | +| False "above QEMU" claim | missing A/B session | acceptance requires session pair + hashes | +| Kernel-module temptation | any `open-iov/kernel/` diff | hard rule U33; reviews | +| Time-to-value negative | measured baseline worse | per-WP admission thresholds | + +## 10. Plan B comparison plan + +`Handoff-Simulator-Plan-B.md` (private) defines the narrow fallback: the +smallest simulator that is simply better than QEMU for the open-iov workflow. +At Phase S exit (and earlier if a work package fails twice), compare Plan A +output vs the Plan B design on: rung reached, feature coverage, determinism, +cycle evidence, integration effort, maintenance cost, and measured +time-to-value. Decision: **keep Plan A**, **apply Plan B lessons in place**, +or **start the Plan B implementation** (with a migration note and re-pilot). +The comparison and decision are recorded in the private handoff. + +## 11. Status and wave log + +| Wave | Date | Result | +|---|---|---| +| Recon + master handoff | 2026-10-03 | `handoff-simulator.md`, local sources register | +| W1 GPGPU-Sim build | 2026-10-03 | CUDA 12 port patches; `vectorAdd` functional+perf green | +| W1 TU104 config/overlay | 2026-10-03 | config scaffold + sanity 26/26 + admissions draft | +| W1 GSP cycle-0 | 2026-10-03 | modeled cycles; real image 200k instr deterministic; 298 tests green (2026-10-03 suite run; count varies with prerequisites) | +| W1 dossier | 2026-10-03 | Turing/Ampere/Blackwell parameter docs + cross-generation survey | +| W2 Phase S (S1 + S4) | 2026-10-03 | `sim/devices` + `sim/profiles` device model: B0–B2 on the real image (0 unknown MMIO), 2-run determinism; A/B harness discovers the profile and records verdict `above` on the frozen checks (docs/36, docs/37 §4.3, docs/40) | +| W2 Phase S | launched | S2 parity, S3 MCP/injection, S4 validation | + +Backlog (recorded, not started): Ampere/Blackwell models; fabric (PCIe, +NVLink/C2C, RDMA/NIC/switch/cable/transceiver, NVMe TLPs, bounce buffers, +DRAM tiers, rendezvous hashing); 8-GPU chassis; H100 SuperPOD RA; QEMU-gdb +time sync; VBIOS; VF bridge module; Verilog/FPGA; multi-version sweeps. diff --git a/debug/risc-v-gsp-rm-vm/docs/36-GSP-DEVICE-MODEL.md b/debug/risc-v-gsp-rm-vm/docs/36-GSP-DEVICE-MODEL.md new file mode 100644 index 0000000..9516fd5 --- /dev/null +++ b/debug/risc-v-gsp-rm-vm/docs/36-GSP-DEVICE-MODEL.md @@ -0,0 +1,175 @@ +# 36 — GSP device model (in development) + +> **Status: in development (2026-10-03, S1 work package).** This page +> documents the `sim/devices/` device model, the versioned +> `sim/profiles/` documents and the real-firmware boot-ladder results. It is +> the companion to `35-GREENEYE-SIMULATOR-DESIGN.md` §3 (Phase S) and +> `06-ACCURACY.md`. Development-only: nothing here authorises GPU hardware, +> kernel modules, mdev, guest access or reboots. +> +> **See also:** `37-SIM-QEMU-LINE.md` (the B0–B2 device model on the frozen +> simulator-vs-QEMU checks) and `40-NEW-VERSION-TRIAGE.md` (one-command triage +> of a new image through this ladder). + +## 1. What this is + +The device model gives the native RV64 core (`sim/machine.py`) real MMIO +semantics for the windows the real GSP firmware touches while booting, so a +bounded run advances **past the QEMU first fault** with evidence instead of +lenient zero-reads. Modeled devices take precedence over the lenient +device-map/unknown fallback; an access inside a modeled window that the +device does not implement is still reported as unknown (never silently +zeroed). Every modeled register carries an **evidence kind**: + +* **header** — a hardware register/field map (offset, width, access type, + field positions, reset/alignment); +* **firmware use** — the disassembled firmware code that reads/writes it; +* **observed access** — the address/width/direction/PC/value sequence seen in + a real run; +* **shim** — reset-time state the real SEC2 booter/host would provide; always + an explicit, session-recorded assumption, never a device claim. + +If a window cannot be proven, none of it is modeled and it is listed as an +unknown (§6). + +## 2. Architecture + +``` +sim/profiles/<name>.json versioned profile: windows, chip id, load layout, + │ DMA aperture windows, boot-shim values + ▼ +sim/devices/__init__.py registry_from_profile(profile, machine) + │ + ├─ base.py DeviceRegistry: range+size decode, access log, + │ JSON state/snapshot, peek() debugger surface + ├─ pmc.py PMC/master block (NV_PMC_BOOT_42 chip ID) + ├─ falcon.py falcon block: IRQ, MAILBOX0/1, DEBUGINFO, DMA, + │ FBIF control/transcfg + ├─ priscv.py RISC-V-local block: IRQ mask, trace, error regs + ├─ timer.py PTIMER (modeled virtual time; opt-in) + ├─ csb.py CSB/TRANSCOM message window seam (B3 landing) + ├─ fwsec.py WPR2 address pair (FWSEC state unmodeled) + ├─ pmu.py PMU stub: observes, never fakes + ├─ dma.py aperture → RAM translation seam (FBGPA/SYSGPA) + └─ irq.py falcon + RISC-V-local IRQ aggregation + ▼ +sim/machine.py Machine(devices=..., translator=..., reset_regs=...) + ▼ +sim/gsp_boot.py bounded real-firmware run + session + rung +``` + +`Machine` consults devices after RAM/UART and before the device map/lenient +fallback; the aperture translator (`dma.py`) is consulted first so +FBGPA/SYSGPA memory accesses resolve to RAM. Overlapping device windows +resolve last-registered-first, and a window that does not model an access +defers to the next matching window (safe fallback), so a narrow overlay can +serve some offsets without hiding the wider device. `Machine.device_state()` +and the `greeneye-sim-devices/1` snapshot expose all device state as +serialisable data (snapshot/restore support is additive and not yet wired to +moments). + +## 3. Modeled inventory + +Offsets are relative to each block base. `RW` semantics are storage unless +noted; every unlisted offset is unhandled and logged. + +| Block | Window(s) (PRIV unless noted) | Modeled registers | Evidence | +|---|---|---|---| +| PMC | `+0x0` (size 0x1000) | `BOOT_42` @0xA00: ARCHITECTURE 28:24, IMPLEMENTATION 23:20, MAJOR 19:16, MINOR 15:12, MINOR_EXT 11:8, CHIP_ID 28:20 (TU104 = 0x164), read-only | header + firmware use + observed (first MMIO) | +| Falcon | `+0x110000`, alias `0x300000000` | `IRQSSET` 0x000 (`irq_status \|=`), `IRQSCLR` 0x004 (`&= ~`), `IRQSTAT` 0x008 (read), `MAILBOX0/1` 0x040/0x044, `OS` 0x080, `DEBUGINFO` 0x094, `CPUCTL` 0x100, `BOOTVEC` 0x104, `DMACTL` 0x10c, `DMATRFBASE` 0x110, `DMATRFMOFFS` 0x114, `DMATRFCMD` 0x118 (transfer on write), `DMATRFFBOFFS` 0x11c, `DMATRFBASE1` 0x128, `FBIF_TRANSCFG(0..7)` 0x600, `FBIF_CTL` 0x624 | header + firmware use + observed | +| RISC-V local | `+0x111000`, alias `0x300001000` | `IRQMSET` 0x2ac, `IRQMCLR` 0x2b0, `IRQMASK` 0x2b4 (read), `IRQDEST` 0x2b8, `IRQTYPE` 0x2bc, `TRACECTL` 0x344, `TRACE_RDIDX` 0x348, `TRACE_WTIDX` 0x34c, `PRIV_ERR_STAT` 0x360, `PRIV_ERR_INFO` 0x364 (read), `PRIV_ERR_ADDR` 0x368 (read), `HUB_ERR_STAT` 0x36c | header + observed | +| PTIMER | `+0x9000` | `TIME_0` aperture 0x9400 / block `+0x400` (NSEC 31:5), `TIME_1` aperture 0x9410 / block `+0x410` (NSEC 28:0); value = instruction count × `ns_per_insn` (**modeled**, not calibrated) | header + **modeled** time | +| WPR2 | `+0x1FA000` | `WPR2_ADDR_LO` 0x1FA824, `WPR2_ADDR_HI` 0x1FA828 (4 KiB alignment; HI points at the last 128 KiB block) | header + firmware use | +| CSB/TRANSCOM | profile-declared | none by default; a profile may declare individual registers with width/value and `rw`/`ro`/`wo`/`w1c`/`set`/`clear` semantics | none yet (B3 landing point) | +| PMU | profile-declared | none; the stub records accesses and returns unhandled | none yet | +| DMA seam | FBGPA/SYSGPA (profile) | GPU-physical → RAM windows for falcon DMA and aperture loads/stores; unmapped transfers are refused and recorded | address-map header + firmware use | + +Falcon DMA follows the header field semantics: `size = 1 << (SIZE+2)`, +`IMEM`/`DMEM` aperture select, `WRITE` direction, address +`((BASE1<<32 | BASE) << 8) + FBOFFS`, `IDLE` set on success and `ERROR` on a +refused transfer. Which FB/SYS aperture a DMA address targets (the +`FBIF_TRANSCFG.TARGET` decode) is **not proven yet** and is listed in §6. + +### 3.1 Boot shim (declared assumptions) + +The firmware's `_start` consumes boot arguments the real SEC2 booter/host +writes; the simulator reconstructs them from the documented structure layout +and records each as an assumption in `session.json.boot_shim.assumptions`: + +* `NV_RISCV_BOOT_PARAMS` at a RAM address (bootType `NVRM`, size 120, + version 1, core-dump buffer, one `WPRM` libos init arg); +* a `GspFwWprMeta` (magic/revision/verified and a self-consistent FB layout); +* raw-image mirror at the DMEM aperture base so the app's own relocation copy + is self-consistent (it is not: see §6); +* `a0`–`a5` register arguments (bootargs PA/VA, ELF PA/size, load base, WPR + id). + +The DEBUGINFO **host-consume** behaviour is also a shim: the firmware writes a +boot-stage marker and polls bits 22:20 until they clear, and no self-clearing +path exists in the firmware, so the host side is modelled as consuming the +stage marker on read. The device never changes a read value except for this +documented shim, and the shim is profile-enabled and session-recorded. + +## 4. Boot ladder results (observed, 2026-10-03) + +| Rung | Claim | Evidence | +|---|---|---| +| **B0** | flat image loads; first MMIO is the unmodelled PMC `BOOT_42` read at `0x2000_0000_0000_0A00` | `tests/test_sim_devices.py` (B0 classification) | +| **B1** | device model serves the whole boot path: **0 unknown MMIO, 0 unhandled in windows**, deterministic 2M-instruction run stopping at the boot-halt/report loop (`0x4010928`) | `run_boot(use_shim=False)`; sessions `dumps/gsp-boot-b1-run1`, `dumps/gsp-boot-b1-run2`; test `test_no_shim_b1_device_model_serves_boot` | +| **B2** | with the shim, the firmware passes the chip gate, MPU/vector setup, the bootargs consumer (`0x4007A30`) and reaches its **relocation copy** (`0x4007B7C`), then stops at a known unknown (DMEM runspace) | sessions `dumps/gsp-boot-b2-run1`, `dumps/gsp-boot-b2-run2`; test `test_shim_reaches_relocation_b2` | + +Reproduction (read-only; no hardware): + +```bash +cd open-iov/debug/risc-v-gsp-rm-vm +python3 -m sim.gsp_boot --instructions 2000000 --record dumps/gsp-boot-b2-run1 +python3 -m unittest tests.test_sim_devices -v +``` + +Two-run determinism is asserted by `test_two_runs_identical` and was checked +at the recorded B2 sessions: identical instruction count (14 682), stop +reason/PC (`fault @ 0x400EA0C`), device-state SHA-256 +(`3fa264b9…`) and register file. The B1 no-shim run is likewise identical +across runs (2 000 000 instructions, budget stop at `0x4010928`, device-state +SHA-256 `88eb07b7…`). + +## 5. Session evidence + +Each `run_boot` invocation writes a `greeneye-session/1` directory with: + +* `session.json` — firmware SHA-256, profile name, load layout, rung, + milestones, stop reason/PC, unknown-MMIO counts, device-state SHA-256, + device-access counts and the full boot-shim assumption list; +* `blobs/device-state.json` — the `greeneye-sim-devices/1` snapshot of every + device (registers, IRQ state, DMA windows/counters, shim flags); +* `events.jsonl` — the normal session event stream (`gsp_boot` note). + +## 6. Unknowns (explicitly not modeled / not verified) + +1. **DMEM runspace content.** The app's relocation copy reads the DMEM + aperture (`0x1_0000_0000+`) and writes it to the `0x400E000` runtime + region; the bootloader-provided content is not reconstructed, so the copy + writes zeros and the run stops at `0x400EA0C` (illegal instruction). This + is the B2 ceiling today. +2. **CSB/TRANSCOM register semantics.** No message-window poll has been + observed yet, so no register is modeled. B3 waits for the first observed + poll; the `csb.py` seam and profile register declarations are in place. +3. **PMU.** No PMU register access observed; stub only. +4. **FWSEC state machine.** Only the WPR2 address pair is modeled; FWSEC + execution is not. +5. **DMA aperture selection.** `FBIF_TRANSCFG.TARGET`/`MEM_TYPE` decode is + not proven; only the profile's DMA windows decide what an address means. +6. **Interrupt delivery.** Falcon/RISC-V IRQ state is aggregated and + serialisable, but mstatus/mie/mip/mtvec trap entry and interrupt injection + are not modelled. +7. **PTIMER calibration.** Values are modeled from the instruction count with + a placeholder `ns_per_insn`; no firmware PTIMER read has been observed and + no timing claim is made. +8. **MPU translation.** The firmware programs its MPU CSRs (observed), but + the machine does not enforce translations; the libos private mapping is + served by explicit device aliases and the FBGPA/SYSGPA seam, not by an MPU + model. + +Nothing in this page is a silicon-accuracy claim; `06-ACCURACY.md` fidelity +levels apply (L0/L1 modelled, L2 modelled-not-calibrated, L3 partial and +listed above). diff --git a/debug/risc-v-gsp-rm-vm/docs/37-SIM-QEMU-LINE.md b/debug/risc-v-gsp-rm-vm/docs/37-SIM-QEMU-LINE.md new file mode 100644 index 0000000..42c76cc --- /dev/null +++ b/debug/risc-v-gsp-rm-vm/docs/37-SIM-QEMU-LINE.md @@ -0,0 +1,249 @@ +# 37 — Simulator vs the QEMU line: evaluation harness + real RE proof + +> **Status: in development (2026-10-03).** This document is the S4 (value +> validation) record for the standalone GSP simulator workstream. It defines +> the evaluation harness, freezes the thresholds used, reports the first +> measured A/B run on the real TU104 570.144 image, and records the U51 +> real-reverse-engineering demonstration. Nothing here is a product claim: +> simulator evidence is L0/L1 (docs/06). +> +> **Update (2026-10-03, later):** S1 (device model + profiles + boot shim) +> landed and the harness was wired to it; §4.3 records the re-measured A/B +> where the native side is the S1 `sim.gsp_boot` path (verdict `above` on the +> frozen checks). §4.1/§4.2 remain the pre-S1 measured record, and §7's +> "S1 not landed" limit applies only to those sections. +> +> **Artifacts:** `harness/sim_qemu_line.py` (A/B harness), +> `tests/test_sim_qemu_line.py` (shape + determinism + QEMU regression), +> `harness/re_demo.py` + `harness/re-demo-report.md` (real RE proof). +> **Line definition source:** docs/35-GREENEYE-SIMULATOR-DESIGN.md §1.1 and +> the boot ladder in §3.3. + +--- + +## 1. The line being evaluated + +The QEMU line (docs/35 §1.1) is the v0.1/v0.2 tool behaviour on the real +image: hooks/watchpoints, memory/struct/snapshot tooling, sessions, symbols +and analyze/compare, probes/plugins, Ghidra bridges, trace, injection SDK and +`Session.call`/`run_until`, plus the documented first fault at firmware vaddr +`0x4009014` (4-byte read of `0x2000000000000A00`). + +The boot ladder (docs/35 §3.3) is `pre-B0` (image loads, first MMIO observed +*as a fault*) < `B0` (first MMIO observed; lenient) < `B1` (first fault past +the QEMU point, explained by device semantics) < `B2` (boot-stage progress) < +`B3` (first mailbox poll) < `B4` (message loop) < `B5` (one scripted message +served). + +**"Above the QEMU line"** requires *all* of the frozen checks in §3, +including the docs/35 §1.1 requirement that the milestone be reached "with +the device model and shims documented". Lenient MMIO is an observation +mechanism, not a device model, so it cannot satisfy the contract by itself. + +## 2. Method + +Everything is stdlib-only, bounded (`<= 2,000,000` native instructions, +`<= 10 s` QEMU), and touches no hardware, GPU, mdev, module, guest or network. + +```bash +cd open-iov/debug/risc-v-gsp-rm-vm +python3 harness/sim_qemu_line.py # A/B harness -> harness/qemu_line/ +python3 -m unittest tests.test_sim_qemu_line -v +python3 harness/re_demo.py # U51 proof -> harness/re-demo-report.md +``` + +* **QEMU side** — `harness.common.open_flat_session` (the same relocation and + `guest_errors,unimp,int` logging as `greeneye probe`), bounded run, first + exception parsed from the bounded head of the log (trap loops flood it), + first-touch device map harvested, session recorded with the log blob. +* **Native side** — a discovered S1 profile runs `sim.gsp_boot` (device model + + recorded boot shim) twice and reports the observed rung; without one, the + lenient `sim.gsp_profile.run_profile` baseline (deterministic RV64IMAC, + modeled cycles) runs twice with `greeneye.device_map` for MMIO. S1's + profile is discovered defensively (`sim.devices` Python entry points, then + `sim/profiles/*.json` matched by firmware sha / `sim/devices/*.json`); when + absent the run is the lenient `B0` baseline and the report says so. +* **Feature matrix** — exercised on the bundled demo firmware + (`firmware/uproc_demo`): hooks, breakpoints, watchpoints, memory, struct + decode, sessions, analyze/compare, exec trace, injection SDK. MCP is + recorded `not_run` here to keep this harness independent of the MCP surface + (the native MCP tools are tested in `tests/test_mcp.py`; docs/20). + +## 3. Frozen thresholds (2026-10-03, before the first measurement) + +Stored in `harness/sim_qemu_line.py:FROZEN_THRESHOLDS` and +`harness/re_demo.py:FROZEN_THRESHOLDS`; they are never moved silently. + +| check | threshold | why | +|---|---|---| +| `qemu_first_fault_reproduced` | QEMU reproduces fw `0x4009014` / `0x2000000000000A00` | otherwise the two sides are not comparable | +| `native_above_qemu` | native rung >= QEMU rung + 1 | the line's rung rule | +| `identical_runs` | 2 identical native runs (key fields) | determinism is a docs/35 §1.1 requirement | +| `native_pair_wall_s` | <= 20 s for the two native runs | keeps the harness usable | +| `device_model_documented` | a discovered S1 device map/profile is used | docs/35 §1.1 "with the device model and shims documented" | + +`re_demo.py` additionally freezes: emulator wall <= 120 s, 5 answers present, +2 identical lenient runs, milestone past the QEMU point; the manual static +baseline is a documented **floor estimate** of 600 s (method note in the +script: ~90 instructions across three sections read by hand and correlated +with the QEMU fault report; not measured). + +## 4. Measured results (2026-10-03) + +### 4.1 A/B harness on the real image + +Firmware `gsp-570.144.fwimage.bin`, sha256 `2d1bb5c60b208cfd…`; S1 absent at +run time (`sim/devices/` and `sim/profiles/` not present), so the native side +ran lenient. + +| side | ran | rung | instructions | cycles (modeled) | stop | wall s | evidence | +|---|---|---|---:|---:|---|---:|---| +| QEMU probe | yes | `pre-B0` | – | – | `fault_load` fw `0x4009014` -> `0x2000000000000A00` | 0.95 | `harness/qemu_line/sessions/qemu-probe` | +| native run 1 | yes | `B0` | 200,000 | 200,012 | `budget@0x4009070` | 0.21 | `harness/qemu_line/sessions/native-1` | +| native run 2 | yes | `B0` | 200,000 | 200,012 | `budget@0x4009070` | 0.22 | `harness/qemu_line/sessions/native-2` | + +Wall figures are from one measured run (host-load dependent); re-run the +command for current numbers. Session hashes for the record are in +`harness/qemu_line/sim-qemu-line.json`. + +* first-touch match with the documented reference: **true** (same address, + same 4-byte read, same fw PC); +* identical across the two runs: **true**; session compare ran (breakpoint, + watchpoint and written-page deltas); +* verdict: **`above_lenient_only`** — rung progress exists (`B0` is one past + `pre-B0`) but the device-model check fails, so `above_qemu_line = false`. + This is the honest underclaim: the native run observes and continues past + the QEMU fault, but lenient MMIO is not device semantics. +* native feature matrix (demo firmware): `hooks`, `breakpoints`, `watch`, + `memory`, `struct`, `sessions`, `analyze_compare`, `trace`, `injection` all + **pass**; `mcp` **not_run** (S3-owned). +* session hashes for every side/session are in + `harness/qemu_line/sim-qemu-line.json` and rendered in `.md`. + +Expected delta when S1 landed (realised; see §4.3): `device_model_used = true`, +rung `B2` because the model serves the QEMU first touch and the shim reaches +the relocation stage, and `verdict = above` because the run stays deterministic +and within the wall bound. S1's profile is picked up automatically from +`sim/profiles/*.json` when its `firmware.sha256` matches (or pass +`--device-map FILE` for an observational map). + +### 4.2 Real RE proof (U51) + +Scenario: *what happens on the real GSP image between entry and the first +fault/milestone, and what does the state tell us about boot?* + +| side | method | wall s | answers | +|---|---|---:|---| +| baseline | QEMU probe (measured, 0.5 s bound) | 0.99 | first fault only | +| baseline | manual static reading (floor estimate) | 600 | entry->fault path by hand | +| baseline | **total** | **600.99** | full question | +| emulator | 3 recorded native sessions (`harness/re_demo.py`) | **0.13** | full question + evidence | + +Frozen threshold: emulator wall <= 120 s, 5/5 answers present, two identical +lenient runs, milestone past the QEMU point. Result: **PASS** (all checks +true). Session hashes and the JSON answers are in `harness/re_demo/result.json` +and `harness/re-demo-report.md`. + +Answer (evidence-backed; simulator L0/L1, no silicon claim): + +1. **First touch:** a 4-byte read of PA `0x2000000000000A00` at fw PC + `0x4009014`; it is exactly the QEMU-documented first fault. +2. **Path:** the entry state is all-zero (no bootloader handoff is modelled); + `.bootstrap` materialises the device PA with five arithmetic instructions, + reads it, then runs a chip-ID comparison chain — 25 executed instructions + from entry to the halt (exec trace). +3. **Field:** the read value's bits `[28:20]` are extracted (`srli 20`, + `andi 0x1ff`) and compared against `354..360` (Turing IDs) and `368` + (Ampere ID). +4. **Dispatch:** `354..360` jumps to `.tu10x_text` at `0x4007000`; `368` + jumps to `.ga100_text` at `0x4005000`; any other value (including the + lenient zero) falls into the fail-closed `j .` loop at `0x4009070`. +5. **Next blocker (B1 candidate):** with a Turing ID served from a one-entry + device map, `.tu10x_text` programs CSRs `0x7c8–0x7cc` (region + descriptors), enables `mstatus` FP/SIMD state, builds the **absolute** + link-time address `0x40070f8` and jumps to it. Because the loader + relocates segments but not absolute immediates, that jump lands outside + RAM and faults (first fault past the QEMU point, 74 instructions). The + next rung therefore needs a chip-ID device model **plus** load-at-link + address (or alias/translation of absolute immediates); the modeled ID + value itself is an assumption (the QEMU baseline faults before reading + it), while the field position and the compared IDs are firmware evidence. + +### 4.3 Re-measurement after S1 landed (2026-10-03, later) + +The harness now discovers the S1 `greeneye-sim-profile/1` document +(`sim/profiles/*.json`, matched by `firmware.sha256`) and runs the native side +through `sim.gsp_boot` (device model + recorded boot shim) instead of lenient +MMIO. Frozen thresholds unchanged. Measured run on the same image +(sha256 `2d1bb5c60b208cfd…`): + +| side | ran | rung | instructions | cycles (modeled) | stop | wall s | evidence | +|---|---|---|---:|---:|---|---:|---| +| QEMU probe | yes | `pre-B0` | – | – | `fault_load` fw `0x4009014` -> `0x2000000000000A00` | 0.98 | `harness/qemu_line/sessions/qemu-probe` | +| native run 1 | yes | `B2` | 14,682 | 27,294 | `fault@0x400EA0C` (DMEM runspace unknown) | 0.16 | `harness/qemu_line/sessions/native-1` | +| native run 2 | yes | `B2` | 14,682 | 27,294 | `fault@0x400EA0C` | 0.17 | `harness/qemu_line/sessions/native-2` | + +* `device_model_used = true`, `device_model_documented = true`, + `identical_runs = true` (device-state sha256 + `3fa264b9f81081e8…` on both runs), `native_pair_wall_s = 0.33`, + `native_above_qemu = true`; **verdict `above`** on all five frozen checks. +* milestones observed: `bootargs_consumer`, `relocation_copy`; + `unknown MMIO = 0`, `device_unhandled = 0`. +* What this does **not** claim: B3/B4 are not reached (the run stops at the + documented DMEM-runspace unknown); the shim values are declared assumptions; + cycles are modeled (docs/32); no timing/device-environment claim. +* Command: `python3 harness/sim_qemu_line.py` (regenerates + `harness/qemu_line/sim-qemu-line.json|.md`). Wall figures are host-load + dependent. + +## 5. Failure signals (when to disable, per U48/U56) + +* Any frozen check false **with QEMU available and the firmware present**: + keep `above_qemu_line = false`, do not claim the line. +* `device_model_documented` still false at the Phase S pause review: the S1 + work package has not met its acceptance; the harness keeps reporting + `above_lenient_only`. +* A feature matrix entry `fail`: disable that feature flag and record it as a + failed experiment; an entry silently stuck at `not_run` blocks the S4 exit. +* `re_demo.py` pass false: U51 is not satisfied; fix the tool chain before + any Phase G design work. +* Thresholds change only with a dated entry in this document; a moved + threshold without a measurement is a review failure. + +## 6. Re-running / reproducing + +```bash +cd open-iov/debug/risc-v-gsp-rm-vm +python3 harness/sim_qemu_line.py # full A/B + feature matrix +python3 harness/sim_qemu_line.py --no-qemu # native only, no baseline claim +python3 -m unittest tests.test_sim_qemu_line -v # shape/determinism/QEMU gate +python3 harness/re_demo.py # real RE proof +``` + +Without the 570.144 image the commands skip cleanly (exit 2 for the CLIs, +skipped tests), because the tool tree is used on machines without the +firmware. The QEMU side also needs `qemu-system-riscv64`; without it the +report is emitted with `unverified_qemu_absent`. + +## 7. Honest limits + +* **S1 partial** — the device model reaches B2 (relocation copy) with the + recorded boot shim; B3/B4 (CSB poll, message loop) are not reached and the + DMEM-runspace unknown is the documented ceiling (docs/36 §6). The §4.1 + lenient record predates S1; §4.3 is the post-S1 measurement. +* **L2/L3 partial** — cycles are modeled (`rv64imac-modeled/1`, docs/32), not + measured; the device model is partial and carries per-register evidence + (docs/36). No timing or full device-environment claim is made. +* **MCP native surface not exercised here** — the A/B feature matrix records + it `not_run`; the native MCP tools have their own suite (docs/20 §9). +* **Modeled chip-ID value is an assumption** — only the field extraction and + the compared ID set come from the image; in the S1 profile the `0x164` + value is declared by the profile, and the boot shim values are declared + assumptions recorded in the session. +* **Relocation artifact** — the B1-candidate fault at `0x40070f8` is a + property of the bias-relocated harness, not of silicon; it is labelled as + such and not counted as a device milestone. (The S1 path loads at the link + addresses and does not use that relocation.) +* **Manual baseline is an estimate** — labelled a floor with its method; the + measured parts of the baseline (QEMU probe) are reproducible with the + harness. diff --git a/debug/risc-v-gsp-rm-vm/docs/38-GSP-EM-API.md b/debug/risc-v-gsp-rm-vm/docs/38-GSP-EM-API.md new file mode 100644 index 0000000..311e8e1 --- /dev/null +++ b/debug/risc-v-gsp-rm-vm/docs/38-GSP-EM-API.md @@ -0,0 +1,139 @@ +# 38 — GSP-Em API: TLV boundary between the GSP emulator and a full-system simulator + +> **Status:** design (2026-10-03). Implementation gated by the standalone +> GSP simulator's validation gate (handoff-simulator.md §13.5/§14). The +> bridge build flag `GREENEYE_GSP_EM_BRIDGE` is **OFF by default** until then. +> +> **Purpose:** define one stable, documented, TLV-framed protocol so the +> standalone GSP emulator can be driven by (and drive) any full-system +> simulator — GPGPU-Sim first, any API-compliant alternative later — without +> either side depending on the other's internals. +> +> **See also:** `35-GREENEYE-SIMULATOR-DESIGN.md` §3.4/§4 (what is and is not +> in Phase S vs Phase G) and `36-GSP-DEVICE-MODEL.md` (the device model behind +> this API). + +## 1. Roles and boundary + +``` + Standalone GSP emulator (this tool) Full-system simulator (GPGPU-Sim or alternative) + ┌──────────────────────────────┐ GSP-Em API ┌────────────────────────────────────────┐ + │ real firmware · device model │◄──(TLV/UDS)───►│ SMs · caches · DRAM · fabric · host │ + │ MCUs · cycles · sessions │ │ consumes/implements the same API │ + └──────────────────────────────┘ └────────────────────────────────────────┘ +``` + +* The GSP emulator is the **authority** for GSP/device/MMIO semantics. +* The system simulator is the **authority** for SM/memory/fabric timing. +* Neither side parses the other's logs: all interaction is API messages. +* Same API is used for A/B replay of captured QEMU/VFIO traffic. + +## 2. Transport and framing + +* Transport: UNIX `SOCK_SEQPACKET` (ordered, message-preserving) at a path + the caller supplies; shared-memory doorbell mode considered later, same + framing. Local only (never a network). +* Every message is one frame: 16-byte header + payload. Little-endian. + +``` +header: magic 'G''S''E''1' (4B) | version u16 | flags u16 | seq u32 | type u16 | rsvd u16 | payload_len u32 +payload: repeated TLV fields +TLV: tag u16 | len u16 | value[len] (pad to 4B boundary with zeroes) +``` + +* Unknown TLVs MUST be ignored (forward compatibility). Unknown message types + MUST be answered with `E_UNIMPLEMENTED` unless `NO_REPLY` is set. +* Version: major in header version; additive minor via `CAPS`. Breaking + changes require a new major and a shim. + +## 3. Message types + +| Type | Name | Direction | Payload essentials | +|---|---|---|---| +| 0x01 | HELLO | both | PROTO_MIN, PROTO_MAX, BUILD_HASH, FW_HASH, PROFILE_ID | +| 0x02 | HELLO_ACK | both | selected version, session ID | +| 0x03 | CAPS | both | CAP_MASK (bitmask), limits (max frame, max DMA), mode flags | +| 0x10 | MMIO_READ | system→gsp | ADDR, SIZE, DEADLINE_CYCLE; reply: DATA or error | +| 0x11 | MMIO_WRITE | system→gsp | ADDR, SIZE, DATA; reply: STATUS | +| 0x12 | IRQ_RAISE | gsp→system | IRQ_VECTOR, DOMAIN, CYCLE | +| 0x13 | DMA_READ | both | GSP_ADDR/PA, SYSTEM_ADDR, SIZE, DEADLINE_CYCLE | +| 0x14 | DMA_WRITE | both | as above | +| 0x15 | CLOCK_TICK | system→gsp | DOMAIN, CYCLE_DELTA; reply: CYCLE_NOW | +| 0x16 | CLOCK_SYNC | both | EPOCH, GRANULARITY, MODE (free-run / credited) | +| 0x17 | RESET | system→gsp | RESET_KIND (cold/warm/domain) | +| 0x18 | STATE_SNAPSHOT | both | REGION selector; reply: STATE blob TLV (opaque, hashed) | +| 0x19 | EVENT_SUBSCRIBE | system→gsp | MASK (mmio, dma, irq, mcu, trap, cycle) | +| 0x1A | EVENT | gsp→system | EVENT_KIND + fields + CYCLE | +| 0x7F | ERROR | both | STATUS, TEXT, ORIGINAL_TYPE, SEQ | +| 0x7E | BYE | both | — | + +Common TLV tags: `ADDR u64`, `SIZE u32`, `DATA bytes`, `VALUE u64`, +`STATUS u32`, `ERROR_CODE u32`, `TEXT utf8`, `CYCLE u64`, `DOMAIN u16`, +`IRQ_VECTOR u32`, `CAP_MASK u64`, `PROFILE utf8`, `FW_HASH bytes32`, +`BUILD_HASH bytes32`, `SESSION_ID bytes16`. + +## 4. Semantics + +* **Ordering:** per-domain FIFO; MMIO/DMA replies carry the `seq` of their + request. The system simulator may issue in-order per domain only; the GSP + emulator may reorder internally but must preserve observable order. +* **Time:** in `credited` mode the system clock advances only by `CLOCK_TICK` + grants; the GSP emulator returns `CYCLE_NOW` so both sides share one + timeline. In `free-run` mode the emulator advances by its own modeled + cycles and reports them in events. Wall clock never enters a decision. +* **State:** snapshots are opaque blobs with a hash; format is the emulator's + `greeneye-session/1` state section. A system simulator may store them but + not parse them. +* **Determinism:** same inputs + same grants ⇒ byte-identical transcripts. + Every conformance run must demonstrate this twice. +* **Security:** study environment only; no confidentiality claim; no message + may be used to defeat a hardware boundary (RUNBOOK P5). + +## 5. Error model + +`OK=0`, `E_UNIMPLEMENTED=1`, `E_RANGE=2`, `E_SIZE=3`, `E_STATE=4`, +`E_BUSY=5`, `E_PROTO=6`, `E_TIMEOUT=7`, `E_DENIED=8`. An error never +terminates the session; `BYE` is graceful. Timeouts are caller-budgeted. + +## 6. GPGPU-Sim bridge (first system simulator) + +* Build-time: `GREENEYE_GSP_EM_BRIDGE=1` compiles the adapter into the + overlay (`gpgpu-sim/overlay/`). Default OFF. +* Runtime: `--gsp-em-endpoint <path>`; absent endpoint ⇒ GPGPU-Sim runs + standalone exactly as upstream (light-touch guarantee). +* Adapter maps: BAR/MMIO trap → `MMIO_READ/WRITE`; GSP doorbell/IRQ → + `IRQ_RAISE`; firmware DMA descriptors → `DMA_*`; simulator clock domains → + `CLOCK_TICK/SYNC`. +* The adapter is additive: no upstream data structure or algorithm is + modified; issues go through the existing device-model seam only. + +## 7. API-compliant backup (replacing GPGPU-Sim) + +The same API must allow a different system simulator (or a purpose-built +service simulator) to take GPGPU-Sim's place: + +1. **Interface contract:** implement the *system side* of §2–§5 (client for + MMIO/DMA/IRQ, source for CLOCK_TICK, event consumer). +2. **Conformance suite:** `gsp-em-conformance` replays golden transcripts + (hello/caps, MMIO R/W, IRQ, DMA, clock, snapshot, error paths) and must be + byte-identical twice. Any implementation passing it is interchangeable. +3. **Replacement criteria:** passes conformance; sustains required + cycles/s and frame rates; preserves session hashes for the same stimulus; + no changes to the GSP emulator. +4. **Build flag:** the system side selects implementation by + `GREENEYE_SYSTEM_SIM=gpgpu-sim|api-compliant-ref` (design; default + `gpgpu-sim`). + +## 8. Validation plan for this spec + +* Unit: frame codec round-trips, unknown-TLV tolerance, error mapping. +* Integration: emulator ↔ stub system client; credited-clock determinism; + snapshot hash stability across runs. +* Conformance harness is the acceptance gate for both the GPGPU-Sim bridge + and any backup implementation. + +## 9. Non-goals + +No remote/network use; no guest access; no hardware writes; no product +dependency. This document specifies the boundary only; implementation is +gated by the standalone simulator's value proof (handoff §14.5). diff --git a/debug/risc-v-gsp-rm-vm/docs/39-RE-WORKFLOW-REVIEW.md b/debug/risc-v-gsp-rm-vm/docs/39-RE-WORKFLOW-REVIEW.md new file mode 100644 index 0000000..2b94746 --- /dev/null +++ b/debug/risc-v-gsp-rm-vm/docs/39-RE-WORKFLOW-REVIEW.md @@ -0,0 +1,678 @@ +# 39 — Reverse-engineering workflow review: tools, real usage, time sinks, stale assumptions and pragmatic upgrades + +> **Status:** review (2026-10-03). Written by a read-only analysis pass over the +> tool tree, the public state page, the private project session logs (cited by +> PART/ADDENDUM and date only; no private paths appear here) and the +> engineering-ops framework. It is a documented review, not a product change; +> nothing in it authorises hardware, module, guest, mdev or RM 610 work. +> +> **Scope of the review:** every reverse-engineering tool/harness/debugger this +> project has built or used, how each has actually been used, where wall-clock +> time is lost, which tool-design assumptions are now stale, how the +> documentation can be made cheaper to maintain, and the smallest set of +> upgrades that would move the dominant cost first. +> +> **Method for reading it:** claims are labelled **[E]** (evidence: a file, +> line, session or recorded event) or **[I]** (inference: the analyst's +> judgement). Section references use `docs/NN` for the tool's documentation set +> and `open-iov/...` for tracked project files. Private project session records +> are cited as `[log PART nn/Add. nn, date]`; they live outside the tracked +> tree and are not named here. + +--- + +## §1 Method and sources + +### 1.1 Question + +The review answers one question for the risc-v-gsp-rm-debugger program: **what +actually reduces the time from "a new GSP/driver release or a live product +wall" to "a supported `vgsp`"** — and where is that time being lost today? + +It is deliberately not a feature wish-list. Every recommendation in §8 carries +the engineering-ops admission fields (scenario, measurable outcome, baseline, +threshold, failure signal) in the shape required by +`open-iov/docs/eng-ops.md` §2. + +### 1.2 Sources read + +| Class | Source | What it provided | +|---|---|---| +| Public live state | `open-iov/STATE-CURRENT.md` (snapshot 2026-10-02 22:12) | milestone ladder, frontier walls, physical-backend record, tool status paragraph | +| Tool documentation | `debug/risc-v-gsp-rm-vm/docs/00`–`38`, README, roadmap | design, accuracy contract, audit, value demos, porting/CI, new simulator work (30–38) | +| Code and harnesses | `greeneye/`, `sim/`, `harness/` (`*.py`, `value_demos/`, `qemu_line/`, `re_demo/`), `plugins/`, `tests/`, `qemu/`, `ghidra/`, `rmharness/`, `gpgpu-sim/` overlay, `sdk/` | what exists and what its tests exercise | +| Recorded evidence | 458 session directories under `dumps/` (counted 2026-10-03), `harness/qemu_line/sim-qemu-line.json`, `harness/re_demo/result.json`, value-demo artifact JSON | what has actually been run, not just built | +| Project ops | `open-iov/docs/eng-ops.md`, `future-eng-ops.md`, `RUNBOOK.md`, `docs/05-operations/03`, `04`, `docs/07-developers/02` | admission discipline, unused-capability catalogue, scenario runbooks | +| Private project logs | the session-part logs and the tool-program ledger (2026-09-24 → 2026-10-03), the progress-analysis log (2026-09-29 → 2026-10-02) | real walls, real costs, real tool usage (cited by PART/date only) | +| Project tooling | Ghidra corpora + headless helper scripts, CUDA porting tools, guest diagnostics, host RM oracle, physical claim harness | the incumbent methods the tool must beat | + +### 1.3 Method and limits + +* **Behaviour framing.** The review describes components and behaviour; it + names no source locations, vendor file paths or filenames. +* **Evidence over recollection.** Cost claims are anchored to recorded numbers + (docs/25 §3–§6, `STATE-CURRENT.md`, PART/date) rather than estimates, except + where explicitly marked as an estimate. +* **No new runs were performed.** No hardware, GPU, module, mdev, guest or + reboot was touched; no test suite was re-run. Test-count claims are therefore + quoted from their source with the source named, and the count discrepancy + found in §5 is reported, not resolved. +* **Sampling.** The 458 session directories were classified by name and + timestamp; individual sessions were sampled, not all read. +* **Fidelity rules apply.** All simulator evidence is L0/L1 unless a source + says otherwise (`docs/06`); every "blocked" real-firmware projection in this + review is marked as such. + +--- + +## §2 Complete RE/tooling inventory + +### 2.1 The tool tree (`open-iov/debug/risc-v-gsp-rm-vm/`) + +Status vocabulary follows `docs/27` §0: **verified useful** (exercised and/or +wired into tests/demos/docs), **in development** (owner-directed work landed +but unfinished; docs 30–38 and the untracked `sim/devices`/`gpgpu-sim` trees), +**pilot** (unadmitted, per `future-eng-ops` F2–F4), **built, not used yet** +(no consumer). + +| # | Tool | Location | Purpose | Who uses it | Evidence of use | Status | +|---|---|---|---|---|---|---| +| 1 | Launcher + CLI (`info/demo/run/attach/launch/probe/record/analyze/compare/symbols/plugins/inject/ghidra`) | `risc-v-gsp-rm-debugger.sh`, `greeneye/cli.py` | entry surface; run/inspect real or demo firmware | agents, owner | audit smokes (`docs/27` §1); sessions in `dumps/` (`demo`, `probe-*`, `inject`) | verified useful | +| 2 | Session/scripting API (`launch`, `attach`, `Session`, hooks, watchpoints, `call`, `run_until`, memory/structs/snapshots) | `greeneye/session.py`, `memory.py`, `structs.py`, `record.py`, `analyze.py` | Frida-like core; sessions as evidence | scripts, tests, value demos | `docs/27` §2; s01–s10 | verified useful | +| 3 | Pure-Python RSP client (QEMU backend) | `greeneye/rsp.py`, `target.py` | transport; no GDB dependency | default backend | demo/probe live-verified (`docs/01` §3) | verified useful | +| 4 | Native RV64IMAC backend (M1) | `sim/decode.py`, `sim/machine.py`, `greeneye/sim_backend.py` | deterministic in-process backend; QEMU parity | tests, demos, new sim work | differential test (`docs/16` §1.1); `sim_diff.py` `PARITY: OK` | verified useful (reference core) | +| 5 | Native exec/access trace (M1.5/R4) | `sim/trace.py` | full instruction/access trace as session blob | record/analyze, value demos | `docs/16` §1.1a; trace tests | verified useful | +| 6 | Modeled cycle counter | `sim/cycles.py` | table-driven modeled cycles (explicitly not silicon timing) | 21 offline tests | `docs/32`; `results.json` of `test_sim_cycles` | in development (increment 0) | +| 7 | GSP device model + profiles + bounded boot runner | `sim/devices/`, `sim/profiles/`, `sim/gsp_boot.py`, `sim/gsp_profile.py` | real 570.144 image past the QEMU first fault (B0–B2) | S1 workstream | `dumps/gsp-boot-b1-run1/2`, `b2-run1/2`; 2-run determinism asserted (`docs/36` §4) | in development (S1) | +| 8 | Declarative probe specs | `greeneye/probes.py`, `examples/probes/` | JSON hooks/conditions without scripts | audit smoke; documented pattern | `docs/27` §3 (3 probes fired) | verified useful | +| 9 | Plugin loader + example plugin | `greeneye/plugins.py`, `plugins/example_plugin.py` | extension point | audit smoke | `plugins` CLI lists all three | verified useful | +| 10 | MCP server | `plugins/mcp_server.py` | agent/CI machine surface (JSON-RPC over stdio/loopback) | no live agent consumer found | 55 tests (`docs/20` §9); tool catalogue 31 (`docs/20` §5) | implemented; **unused in real workflow** (F1 `partially used`) | +| 11 | Injection SDK + `inject` | `sdk/`, `greeneye/inject.py` | position-independent RV64 service load/call, direct + mailbox | tests, s06 | `docs/11`/`12`; audit smoke `0xbeae` both paths | verified useful | +| 12 | Ghidra bridge (symbol export, hit annotation, control) | `ghidra/*.java`, `greeneye/ghidra.py`, `gdbridge.py` | static↔dynamic address join; run-until from decompiler | Ghidra headless (program ledger wave 2) | 25-command bridge tested (`docs/27` §2); Java verified in the tool ledger | verified useful | +| 13 | TCG full-trace plugin | `qemu/greeneye_trace.c` | complete exec/access JSONL without stops | audit live run | 8,012,129 events in ~1 s (`docs/27` §0) | verified useful | +| 14 | Mutation harness | `harness/mutation_harness.py` | bounded deterministic one-field mutation; outcome classes | tests, s08 | seed replay byte-identical (`docs/18` §2) | verified useful | +| 15 | Portability matrix | `harness/portability_matrix.py` | probe+scenario+compare across targets | tests | self-comparison all-zero deltas | verified useful | +| 16 | Release triage pipeline | `harness/version_diff.py`, `harness/ci_release.py` | old/new image diff + hashed bundle; RM610 gate | tests; CI docs | self-diff zero; gate refusals tested (`docs/21`) | verified useful | +| 17 | Control discovery + reachability | `harness/control_discovery.py`, `harness/reach.py` | unknown-call sweep; "did it run / what args" | demo sweep only | 189 probes ≈ 36 s (`docs/19` §3.4); `--surface rm` exits clean | `--surface demo` verified; real-RM surface **built, not used yet** (F5) | +| 18 | Simulator-vs-QEMU A/B harness | `harness/sim_qemu_line.py` | frozen-threshold evaluation on the real image | S4 workstream | verdict `above_lenient_only`, `device_model_used=false` (`harness/qemu_line/sim-qemu-line.json`) at review time; **same-day post-review**: S1 profile discovery wired, verdict `above`, native B2 (docs/37 §4.3) | in development (S4) | +| 19 | Real-RE proof harness | `harness/re_demo.py` | U51 demonstration: entry→first-fault question | S4 workstream | `harness/re_demo/result.json`; baseline floor estimate 600 s vs 0.13 s emulator | in development | +| 20 | Driver-version sweep | `harness/port_driver_versions.py` | 540→current per-version evidence | none in a real sweep yet | 23 tests; dry-run default; versioned inputs absent locally | pilot (unadmitted F2) | +| 21 | New-device scaffold | `harness/port_new_device.py` | read-only device profile capture + pairing | none yet | tests; no profile captured in a workflow | pilot (unadmitted F3) | +| 22 | Trial packaging | `harness/package_sim.py` | deterministic sim+injection bundle for review | none yet | 6 tests; `--pack-image` pre-injected demo | pilot (unadmitted F4) | +| 23 | Value demos s01–s10 | `harness/value_demos/` | measured analog time-to-value vs recorded real baselines | tests; cited in docs 25/27/eng-ops | `artifacts/results.json` (last local run: 3 demos, s01/s04/s05, 1.72 s) | implemented; artifact currently partial (§5) | +| 24 | Host RM oracle harness | `rmharness/`, wrappers for `tools/rmtool`, `rmtop`, `rmq` | read-only host RM identity/memory/engine captures into sessions | audit capture smoke | 3 tests outside CI; read-only allowlist (`docs/17`) | working with caveats (device tables permission-dependent) | +| 25 | GPGPU-Sim overlay | `gpgpu-sim/` | Phase G full-board substrate (gated) | W1 build agents | CUDA 12 build + `vectorAdd` functional/perf (program ledger §13.3); `BUILD-PORT.md` | in development, deferred behind Phase S | +| 26 | Test suite + CI | `tests/` (22 files), `.github/workflows/risc-v-gsp-rm-debugger.yml` | regression net | CI, every change | 155 at audit; 194 claimed at flag removal; 298 `def test_` statically countable and 298 collected in the 2026-10-03 R2 run (live QEMU present) | verified useful; count claim drifted (§5) | +| 27 | Documentation set | `docs/00`–`38`, README, roadmap | design/usage/proof/portability | agents | link checker 61 files OK at wave 8 | mixed: 00–29 audited; 30–38 unindexed elsewhere | + +### 2.2 Project tooling outside the tool tree (incumbent methods) + +| Tool | Location (public ref) | Purpose | Evidence of use | Relationship to this tool | +|---|---|---|---|---| +| Ghidra corpus + headless helper scripts (`DumpFuncs`, `Xrefs`, `Disasm`, `ScanBranches`, `MemCheck`) | `open-iov/tools/ghidra/` | decompile/disassemble/xref the firmware | `[log PART 92/Add.25, 2026-10-02]` decode of the GR commit assert; `[log PART 91/Add.9]` GPFIFO fix | static half of the loop; the bridge joins it to runtime | +| CUDA/userspace porting tools (`cudatool.py`, `diff_corpus.py`, `annotate_corpus.py`, symbol exports) | `open-iov/userspace/cuda/tools/` | userspace version porting | `docs/22` §2; F2 in production use for 570.144 | separate axis (CUDA userspace); the tool handles the firmware axis | +| Guest diagnostics (`tools/guest-diag/`: ioctl tracer, kprobe/tracker scripts, pagemap tools, VM restart scripts) | `open-iov/tools/guest-diag/` | what the stock driver requested/where it blocks | `STATE-CURRENT.md` (guest-wall fixes); `[log PARTs 29–64]` | complements the firmware question; the tool does not replace it | +| Kernel-module debug planes (default-off params, `/sys/kernel/debug/vgsp/*`, `gsp-phys-state`) | `open-iov/kernel/`, `open-iov/docs/eng-ops.md` §10 | live emulation/physical state | every recorded physical cycle | supplies the real-A/B reference for M1–M6 comparisons | +| Supervised physical claim harness | `tools/guest-diag/phys_claim_cycle.sh` (per `STATE-CURRENT.md`) | safe GPU ownership/boot cycles | 3/3 green pre-lock cycles; all later cycles | the expensive oracle the simulator exists to reduce | +| x86 QEMU/guest harness (`vgpu_vm.py`, `arcd`, guest agent) | `open-iov/scripts/` | drive the stock guest | smoke runs (`ALL OK` 64/64) | produces the payload/failure stimulus | +| Evidence archives (per-cycle logs, rings, decodes) | local evidence archive (untracked) | manual evidence packs | cited per milestone in `STATE-CURRENT.md` | the manual method that sessions/analysis should replace | +| Hourly check-back loop | `checkback/hourly.sh` (per progress analysis §1aq) | commits/pushes `STATE-CURRENT.md` monthly-minutes | running since 2026-09-30 | process automation, not an RE tool | + +### 2.3 Actual usage versus build inventory + +**[E]** Of 458 session directories in `dumps/` (2026-10-03): 413 are +`unittest-*` runner sessions, 5 are `opencode-*` agent sessions, 5 are +`gsp-boot-b*` device-model runs, and the rest are demo/probe/inject/record +sessions from self-verification (list sampled). No session directory in the +tree is tagged with a live product-wall question (e.g. a GR-promote or +UVM-push question). The only real-image sessions are `probe-*` first-fault +runs and the `gsp-boot` B1/B2 ladder runs. + +**[I]** The tool has been *used* overwhelmingly to validate itself and to +characterise the real image's first fault; it has not yet been used as the +daily instrument on a live product wall. That is consistent with the tool's own +honesty (`docs/25` §2.1 marks every real case `blocked_on_device_model`) and +with the private logs: the one place the tool is named as a next route for a +live wall — the GR engine-context promote/commit wall — records it as route 3, +unused so far `[log PART 92/Add.25, 2026-10-02; STATE-CURRENT.md:122-125]`. + +--- + +## §3 Usage history: recurring RE tasks, tools, costs, bottlenecks + +The project's real work decomposes into eight recurring task classes. Each row +gives the recorded cost and the bottleneck that made the cost. + +### 3.1 Task class A — guest-initiation walls (userspace driving the emulated vGPU) + +**Typical form:** the stock guest driver/client blocks or returns an opaque +error; find the first unsatisfied expectation. + +**Recorded instances:** `cuCtxCreate` rc=999 and the 20 per-channel +completion counters (five days, `[log PARTs 29–64, 2026-09-24…29]`); UVM +tracker/semaphore waits; the kernel UVM map wait `[log PART 22]`; clock-domain +probe decode (`STATE-CURRENT.md:520-558`). + +**Tools used:** guest diagnostics (strace on ioctls, guest kprobes, pagemap, +counter/tracker samplers), kernel-module debug planes (channel dump, map +add/probe), QMP `gva2gpa`, serial console, and the inline CUDA driver-API smoke +test. **Recorded cost:** ~5 days for the counter campaign; multiple +days-scale walls before that. **Bottleneck [E]:** guest-side observability — +ptrace was broken in the guest (`[log PART 22]`), so every question cost a +custom sampler; the host's VA→PA map was rebuilt/rescanned under live traffic +and the decisive page (counter page, pool page) had no translation until a raw +live dump proved it (`docs/25` §5 s02). + +### 3.2 Task class B — RM control/RPC gaps and opaque server failures + +**Typical form:** the server faults or reports a generic status; the first +error is invisible. + +**Recorded instances:** the 0x59 RM-init chase (cycles 2–13; **~5 h 10 m**, +10+ supervised boots, reply was a generic `0x59`; `docs/25` §3 s01 and §5 s01); +the missing pre-boot messages (fn72/fn73) found only after the NOCAT storm +collapsed 851→5 records `[log PART 91/Adds.2–6]`; the RM log-ring id bug that +hid the first error (`STATE-CURRENT.md:439-457`). + +**Tools used:** raw host-memory reads of message buffers, the server's own +journal ring (once mapped), static source review, supervised cycles; the +Ghidra corpus + headless scripts for the commit path `[log PART 92/Add.25]`. +**Bottleneck [E]:** one experiment per supervised boot and an invisible first +error; static review closed host-side leads but the decisive field was only +visible in a raw live page. + +### 3.3 Task class C — GR engine/channel bring-up and execution + +**Typical form:** allocate channel, commit context, submit work, get a +fault or a soft `INVALID_STATE`. + +**Recorded instances:** the GR ring layout and 20-channel autopin +(`[log PART 30 addenda]`); the first GPFIFO submit fault +(`docs/25` s04/s09: cycles 18–25, **~1 h 41 m**, 8 supervised boots); the +engine-context promote/commit campaign of 2026-10-02 (modes 4–22, **a full +day**, 730+ lines of firmware RE, three wedge-reboot cycles +`[log PART 92/Adds.13–25]`, and the campaign continued into +`STATE-CURRENT.md:98-131`). + +**Tools used:** the module's own ladder modes and debugfs, the decoded status +queue, Ghidra headless (commit assert `0x52a8b04` classified as a *secondary* +failure), host reads. **Bottleneck [E]:** every hypothesis cost one full +boot+cycle (~5–6 min) and the failing path is server-internal; the one +hypothesis with 730+ lines of static RE produced an accepted root cause +(externally-owned VAS) but the next gate still wedged and required a reboot. + +### 3.4 Task class D — physical boot chain and backend bring-up + +**Typical form:** FWSEC → booter → RISC-V → RM init → clean release on the +real GPU. + +**Recorded instances:** FWSEC-only carve proved non-recoverable (fuse-locked), +requiring an approved reboot and retiring the one-way stage +(`STATE-CURRENT.md:466-519`); the first real boot and first GSP-RM message +(`STATE-CURRENT.md:645-705`); cycles 4–13 to init-done +(`STATE-CURRENT.md:72`, `:495-519`); every cycle ran under the supervised +harness. + +**Tools used:** the module-owned claim/boot path, the power-domain D0 harness, +debugfs state (`gsp-phys-state`), pstore inspection, the fail-safe. **Cost:** +10+ supervised cycles for one question; each failure originally needed a +reboot, until the server's own shutdown supplied the missing handoff bit and +made clean releases possible (`STATE-CURRENT.md:672-705`). +**Bottleneck [E]:** the release path: fused parts cannot be host-recovered, so +any wedge converts a 6-minute experiment into a reboot that kills every other +agent on the host. + +### 3.5 Task class E — guest work forwarding and hardware execution (P5/P6) + +**Typical form:** guest-originated pushes must execute on the physical engine +with correct data and exactly-once releases. + +**Recorded instances:** P6a guest-page DMA (`STATE-CURRENT.md:242-270`), P6b +waves 1/1.5/2, P6c relay, all on 2026-10-02; the P6a→P6c span was +**~02:20–08:30 (~6 h)** with one supervised cycle per rung on average +(`docs/25` §3 s06). Two false walls: the guest-RAM file opened without the +large-file flag (root-caused in-run) and the stale shared-memory image that +poisoned cross-run state and cost **7 cycles** +`[log PART 93/Add.29, 2026-10-02]`. + +**Tools used:** the module's forwarding parameters/counters, a guest pagemap +probe, offline reconstruction scripts against the shared guest-RAM image +`[log PART 93/Add.30]`, and the CUDA smoke as the acceptance gate. +**Bottleneck [E]:** cross-run state contamination and the absence of a +running, symbol-aware firmware view; the decisive 12.7 KB push was decoded +*offline from a dead run's memory image*, not from a live session. + +### 3.6 Task class F — firmware portability/version questions + +**Typical form:** does firmware version X behave differently from anchor Y? + +**Recorded instances:** 570.144 vs 570.211 pairing disproven at one full boot +each (`docs/25` s07; `STATE-CURRENT.md:485-490`); CUDA/userspace porting via +the corpus tools (`docs/22` §2, F2). **Tools used:** manual probe + boot +comparison; `version_diff.py` exists but has no admitted real sweep; `run_all` +s07 demonstrates the shape on the analog. **Bottleneck [E]:** each version +question needs hardware or a hand diff; symbols/corpus transfer warnings +(~1% wrapper re-resolution) are a known trap (`future-eng-ops.md` F2). + +### 3.7 Task class G — evidence assembly and status upkeep + +**Typical form:** turn a cycle into a reviewable claim. + +**Recorded instances:** every milestone's evidence pack is assembled by hand +from per-cycle directories (multi-MB logs/rings/decodes) plus the session log; +the value demo s10 packages the same contract in 4.92 s on the analog, and its +"real case" is explicitly device-model-blocked (`docs/25` §3 s10). Status +upkeep is continuous: an hourly check-back loop commits `STATE-CURRENT.md`, +and the session logs are append-only — the main log alone is ~15k lines +(measured 2026-10-03) with the newest entries prepended. +**Bottleneck [E]:** there is no machine-readable cross-cycle evidence index; +every claim is re-narrated in prose, and prose drifts (see §5). + +### 3.8 Task class H — self-verification and audit + +**Typical form:** prove the tool works so its claims can be trusted. + +**Recorded instances:** two adversarial passes (AR1/AR2/AR3), the 67-surface +audit, flag removal, value demos, 194-test claims; the current wave is +device-model B1/B2 with 2-run determinism (`docs/36` §4) and the A/B harness +verdict `above_lenient_only` (`docs/37` §4.1). **Cost:** substantial but +front-loaded; it is the reason the honest `blocked_on_device_model` labels +exist. **Bottleneck [I]:** self-verification consumes the sessions that the +live workflow should be producing, and there is no separate "live-question" +session tag (§2.3). + +--- + +## §4 Time-sink analysis: top 10 recurring costs, ranked + +Ranking is by recorded project minutes lost × frequency × recoverability. +Each item names the evidence and the tool-side lever (if any). + +| # | Time sink | Recorded evidence | Why it costs | Tool-side lever | +|---|---|---|---|---| +| 1 | **Supervised hardware cycle as the only oracle** | ~5–6 min/cycle (`docs/25` §6); 10+ cycles ≈ 5 h 10 m for the 0x59 question; 8 cycles ≈ 1 h 41 m for the ladder; 10 NOCAT cycles; P6a–c ≈ 6 h | every hypothesis = one full boot + fail-safe discipline; cross-question batching is manual | finish device model (S1) so hypotheses resolve off-hardware | +| 2 | **Server wedge → reboot loop** | 22 reboot-request lines in the main log since 2026-10-01 (counted); 8 host-freeze mentions; 3 wedge-reboots in one afternoon `[log PART 92/Adds.13–24]`; a user-paused reboot at 10:25 2026-10-02 | fused-part WPR2 cannot be cleared; a reboot kills all other agents' sessions and breaks continuity | none in-tool today; a recoverable wedge or snapshot-resume would collapse this | +| 3 | **Guest-side observability gaps** | ~5 days for the 20-counter wall; ptrace broken in guest; custom samplers per question `[log PARTs 29–64]` | host emulation state and guest expectations are only visible through ad-hoc probes | pair each guest wall with a firmware-side session instead of a new sampler | +| 4 | **Opaque server-internal failures** | 0x59 invisible until the log ring was mapped (id bug); promote `0x40`/CTXNOTVALID; 730+ lines of static RE for one root cause | first error hidden behind generic statuses and unmapped rings | hooks/watchpoints on the rehosted failure path; session compare | +| 5 | **Work loss / tree hygiene** | PART 44/45 fixes lost from the tree, PART 47 re-root-cause `[log PART 47]`; modes 4–22 uncommitted at the campaign's end; progress analysis notes repeated "highest-loss-risk" moments | long uncommitted experiments die with a reboot/host reset | commit-per-mode discipline; sessions preserve the evidence even when code is lost | +| 6 | **Documentation sprawl and contradiction** | main log ~15k lines + progress analysis ~4.8k + tool ledger ~1k + state page ~0.9k; ≥8 contradictory live status claims found (§5) | every agent re-derives state; wrong assumptions send work down dead ends | single generated status/index; archive superseded logs | +| 7 | **Cross-run state contamination** | stale shared guest-RAM image cost 7 cycles `[log PART 93/Add.29]`; stale-PTE data-check residual (`STATE-CURRENT.md:195-205`); evidence dirs reused | runs are not hermetic; a phantom wall looks like a firmware bug | pre-run state hash + hermetic temp state; session records inputs | +| 8 | **Manual evidence assembly/comparison** | per-cycle multi-MB dirs; s10 shows the same contract in 4.92 s on the analog but is device-model-blocked; M-tables hand-filled | every claim re-narrated; comparisons by eye | one cycle → session → `compare` pipeline (rmharness pattern, extended read-only) | +| 9 | **Long-run guest stalls and timeouts** | agent/SSH timeouts from a trace reader flooding the ring; UVM spins; one no-watcher guest kernel crash (`STATE-CURRENT.md:577-582`) | half-finished runs and VM restarts | timeout budgets in the harness; no-watcher runs marked unsafe | +| 10 | **Planning churn** | owner directives U24→U31–U45→U46–U58 within 24 h (program ledger §13–§14), superseding §5 phases twice | plans and docs rewritten; in-flight work re-scoped | one live plan with an append-only directive ledger, not a rewritten one | + +**[I]** Rank 1 is the root cause; ranks 3, 4, 6, 8 are amplifiers. If only one +thing improves, it must be the number of supervised cycles per question — by +resolving hypotheses off-hardware (rank 1) or by extracting more per cycle +(rank 8). + +--- + +## §5 Design-assumption audit: stale or contradicted claims + +The tool's design and status text contains assumptions that its own newer work +has invalidated. Each row names the document section, the stale claim, the +contradicting evidence, and the correction. These are documentation defects, +not fidelity claims: the underlying work is honest; the text lagged it. + +| # | Stale claim (exact location) | Contradicting evidence | Correction | +|---|---|---|---| +| 1 | "**No TU104 device model**: no boot beyond the first unmapped access" — `docs/06` §8 gap 1; "does not boot the real GSP image to its message loop yet (**no TU104 device model**)" — agent brief, field-guide status text (2026-10-02); "Native M2+ device model — roadmap only" — `docs/00` status table line 56 | `sim/devices/`, `sim/profiles/`, `sim/gsp_boot.py` exist (untracked, 2026-10-03); B1: "**0 unknown MMIO, 0 unhandled in windows**" and B2 reaches the relocation copy; sessions `gsp-boot-b1/b2` (`docs/36` §3–§4) | say: "a partial GSP device model exists (B1/B2); the image does not yet reach the message loop; B3/B4 are the targets" | +| 2 | "L2 — **not modelled** (do not claim)" — `docs/06` §2; "It is not: a time, interrupt or device model (L2/L3 absent)" — `docs/26` §1; "L3 absent" — `docs/16` §3.1 table | `docs/32` models per-instruction cycles (labelled modeled, not calibrated); `docs/35` §1.2 states the standalone simulator "targets L0–L2 … with L2 explicitly labelled *modeled*"; `docs/36` models device windows, IRQ aggregation and PTIMER (partial L3) | the canonical accuracy doc should add a "modelled, uncalibrated" state between L2 absent and L2 calibrated, and all status text should use it | +| 3 | Documentation indexes stop at **docs 29**: README map lines 75–109, `docs/00` map, agent field-guide index/table | docs 30–38 exist and are untracked/in development (git status 2026-10-03): GPGPU-Sim integration, Turing/Ampere/Blackwell dossiers, cycle model, simulator design, device model, sim-QEMU line, GSP-Em API | add 30–38 to every map with an "in development" tag; the new docs are currently discoverable only by directory listing | +| 4 | "**24-tool** JSON-RPC surface" — agent brief, field-guide MCP rows; F1 "24 tools" — `future-eng-ops` F1 | `docs/20` §5 lists **31** tools and §9 says 55 tests (the doc itself is modified in the working tree) | pick one number from the code (`tools/list`) and generate it into all status text | +| 5 | Test-count claims: **194** (`docs/00` line 47; agent brief 194-test references; `eng-ops` register; project doc §6), **155** at audit (`docs/27` §0), **215** in the simulator ledger §13.3; 272 `def test_` statically countable 2026-10-03 | three different live claims plus a fourth countable value; no source states which suites were skipped in which environment | one generated badge from a recorded `make test` output; claims must name the environment (QEMU present?) | +| 6 | "Not implemented: device model (R1), custom simulator (R2), **MCP plugin (R3)** — all roadmap" — `open-iov/docs/05-operations/03` §5 line 97 | `docs/00` line 54: MCP "working (20)"; device model in development (`docs/36`); MCP 54 tests | project page: MCP implemented; device model in development; only custom simulator M2+ remains roadmap | +| 7 | Roadmap gates: "Items here are **not** v0.1. Do not start them ad hoc; each has a gate" and R1 "starts only when a portability task or an accuracy measurement requires it" — `docs/08` lines 5–8, 31–37; `docs/16` lines 2, 12–15 "nothing here exists yet" | owner-directed S1/B1/B2 device-model work and the Phase S program started 2026-10-03 (private program ledger §13–§14; `docs/35`, `docs/36`) | mark R1/R3/R4 started/superseded with the directive citation; keep the gate text for M2+ | +| 8 | R4 trace "**roadmap** … available once the QEMU plugin headers are vendored" — `docs/08` R4; MCP R3 framed as original design sketch even while "implemented" | audit: trace plugin verified live (8M events); native M1.5 trace implemented (`docs/16` §1.1a) | R4: implemented in both forms; session folding remains the only roadmap piece | +| 9 | `STATE-CURRENT.md` self-contradiction: line 160 "(2026-10-02 08:36, **in development**)" vs lines 841–844 "in-development notices were removed 2026-10-02 after … 194-test suite" | same file, same day, two statuses | the dated frontier bullets should be labelled historical and superseded inline | +| 10 | Value-demo numbers: "`run_all` aggregate (**~18.4 s total**)" — `docs/25` §3; "value demos 10/10 in **22.5 s**" — tool ledger wave-8 final; local `artifacts/results.json` records only **3 demos, 1.72 s** | the artifact is gitignored and currently a partial `--only s01,s04,s05` run; doc says "do not quote a stale number without re-running it" (`docs/25` §10) | regenerate before quoting; make `run_all.py` record the selection list and date into the doc via a generator | +| 11 | Private tool ledger header declares v0.1 acceptance frozen and "MCP is roadmap-only / explicitly not v0.1 … no custom simulator", while waves 2–8 and the simulator program superseded it | the ledger's own change log (waves 2–8) and `docs/35` §13 | treat the ledger as historical; point readers at README + docs 00/27 instead | +| 12 | Design assumption: "QEMU is just the first execution backend" and "the whole stack scriptable" — `docs/00`; "the only production backend is QemuTarget + RspClient" — `docs/16` §1 | M1 native backend is a peer with a differential test; device model only exists natively, not in QEMU | state that native is the only backend with device semantics today; QEMU remains the only one able to host the real image meaningfully? (Neither boots it; QEMU gives the first fault) | +| 13 | Design assumption in `docs/25` §2.1: projections become real "once H2.2 boots the image" | the program replaced H2.2-first sequencing with Phase S (standalone simulator) and Plan B; H2.2 is a legacy-handoff concept not mapped to the public docs 30–38 | map legacy H-numbers to docs 30–38 in one migration table | + +**[I]** The pattern is consistent: docs 00–29 describe the audited v0.1 +surface correctly, but the 2026-10-03 simulator wave (docs 30–38, untracked +`sim/devices`, `gpgpu-sim/`, new harnesses) has not been propagated into the +status claims. The cheapest durable fix is a generated status block plus a +claim scanner (§6, §8-U3). + +--- + +## §6 Documentation efficiency + +### 6.1 What is duplicated + +1. **Documentation maps (5 copies).** README §Documentation map (`README.md` + lines 75–109), `docs/00` §Document map, the agent field-guide 22-row + discoverability index, the project user guide §10, and the project master + log's "when to use it" table. All five enumerate the same files; only some + include docs 28/29, none include 30–38. **[E]** +2. **Status blocks (≈20 copies).** Nearly every doc opens with the same + "Status: implemented, documented and tested (audit docs/27 …; value demos + docs/25 …)" line. A status change requires ~20 edits; the test-count drift + (§5 #5) proves the mechanism fails. +3. **The "it is / it is not" framing** appears verbatim in `docs/26` §1, the + project user guide §1 and (paraphrased) the project page §2. +4. **Gating text** for RM 610 / no-hardware appears in README, docs 06/11/12/ + 14/15/19/20/21/22/23/24/25/28/29/35/36/38 and the project docs. Honest, but + repeated in full rather than linked. + +### 6.2 What is stale + +The 13 rows of §5. The highest-impact three: device model existence (#1), +L2/L3 fidelity wording (#2), and maps/status stopping at 29 (#3). + +### 6.3 What is unlinked + +| Item | Where it should be linked | Why | +|---|---|---| +| docs 30–38 | README map, `docs/00` map, agent field-guide index, project page §5 | currently invisible to a cold reader | +| `SIMULATOR-WORK-BOUNDARY.txt` | README layout + docs/35 §6 | the scope contract for the current wave | +| `gpgpu-sim/` overlay (`ADMISSION.md`, `BUILD-PORT.md`) | docs/30 and README | the only place the pinned methodology is recorded | +| `harness/qemu_line/`, `harness/re_demo*` artifacts | docs/37 (it references them) | evidence exists but is outside any rendered report someone will find | + +### 6.4 What is too verbose + +* **The main session log** (~15k lines, newest-first addenda) — excellent + evidence, poor orientation. It answers "what happened" but not "what is + true now". +* **The agent field guide** — a 22-row index plus a 12-row matrix plus F–K + examples, duplicating docs/00 and the user guide. +* **The private simulator handoff** — §4/§5 (the original plan) are superseded + by §13/§14 in the same file; readers must internalise two plans and a + supersession order. **[E]** +* **The historical tool-program ledger** — a wave-by-wave journal whose + header claims to be "the source of truth for what is being built"; now + history. **[E]** + +### 6.5 Concrete trim/re-link plan (cheap, ordered) + +| Step | Action | Files touched | Measurable effect | +|---|---|---|---| +| T1 | Generate `docs/STATUS.md` from the test runner output + `git describe` + harness JSON (`qemu_line`, `re_demo`, `results.json`); replace the ≈20 status blocks with one link and a date | tool docs (mechanical) | status edits: 20 → 1; no stale count survives a test run | +| T2 | Add a claim scanner to the link checker: flag numeric claims (`\d+-test`, `\d+-tool`, `L2 not modelled`, `roadmap only`) whose source value differs from the generated status | `scripts/check-docs-links.py` | contradictions detectable in CI | +| T3 | Add docs 30–38 (and the boundary file) to README, docs/00 map, agent field-guide index, project page; tag each "in development (Phase S)" | 4 files | cold-start discovery | +| T4 | Mark superseded sections in the private simulator handoff (§5 superseded by §13; §4 phases by §14) and the tool ledger (historical header); keep them as evidence | 2 private files | halves the reading needed to find the live plan | +| T5 | Collapse the four doc maps into README as the single map; other maps link to it | README, docs/00, user guide, agent brief | one map to maintain | +| T6 | Add a retention rule for `dumps/`: keep named/live-question sessions, prune `unittest-*` older than N days in CI or a script | `.gitignore` already excludes; add script | 107 MB today, >90 % runner noise; keeps real evidence findable | +| T7 | Replace hand-assembled M-tables with a generator over session pairs (the `compare` output + M1–M6 skeleton already exist) | harness | evidence pack becomes machine-produced | +| T8 | One "where to start" flow: README → user guide → STATUS; remove the duplicate tables | maps | orientation time down; this review's §5 defects become non-recurring | + +--- + +## §7 First-principles re-derivation: what actually minimises time + +For each dominant problem in §4, the minimum-time approach, including methods +not currently used. + +### 7.1 Problem: hypotheses are evaluated at 5–6 minutes per experiment + +**Current method:** each candidate fix/hypothesis is compiled into the module +and paid for with a supervised hardware cycle; failures can wedge the server +and cost a reboot. + +**First principles:** the marginal cost of an experiment should be the cost of +*evaluating the hypothesis*, not the cost of *rebooting the device*. +Two independent moves: + +1. **Move evaluation off-hardware against real firmware.** The only currently + viable route is the S1 device model + rehosted execution (docs/35 §3.3). + Until B4 exists, the tool cannot answer a live firmware question at all — + every session in `dumps/` is either self-verification or a first-fault + probe. +2. **Until then, batch per cycle.** A single boot can be instrumented to + answer many questions (the module already exposes debugfs state, counters + and ladders). The recorded campaigns did one hypothesis per cycle even when + the per-cycle data would have covered several (`docs/25` §5 s01: ten boots + for ten candidate fixes). A "one boot, N recorded observations" harness is + cheaper than N boots and needs no new fidelity. + +**Not currently used:** *surgical rehost* — compiling and running only the +failing function region with recorded inputs, instead of booting the whole +image. The initial failure is localized (a commit helper, a parser), and +rehosting a function does not require the full device model — only its +observed input/output state. This is the shortest path from "730 lines of +static RE" to "watched it decide". + +### 7.2 Problem: the server wedge makes failures catastrophic + +**Current method:** fail-safe holds the device; recovery is a fuse-limited +teardown or an approved host reboot. + +**First principles:** a failure should stop the experiment, not the host. The +tool-side analogue is *deterministic replay from a snapshot* (moments, F7): +experiment on a state image, never on the only copy. That is gated on M2/M3 +today, but a weaker version is available now: record the full stimulus and +state hash so a wedge cannot lose the evidence and the next cycle can +reproduce the exact pre-failure state by construction (the session already +carries snapshots; the missing piece is the capture path for the physical +cycle). + +### 7.3 Problem: opacity (first error invisible) + +**Current method:** map the journal ring, read raw host memory, static review. + +**First principles:** the first error is a data question. Watchpoints on the +reply/state object plus before/after diff answer it directly — the same shape +as value demos s01/s05, which are measured in ~1 s on the analog. The tool +has these primitives; what it lacks is a running firmware to attach them to. +Again: the device model is the unblocker, and *surgical rehost* can provide it +locally for one failure path. + +### 7.4 Problem: documentation and status re-derivation + +**First principles:** state should be produced by the system that changes it, +not retyped by hand. Test results, session manifests, harness JSON and git +describe already exist; a generated status page costs one small script and +removes an entire class of stale claims (§5). The session format is already +the evidence contract; the status page should consume it. + +### 7.5 Problem: phantom walls from state contamination + +**First principles:** make runs hermetic and make inputs visible. A pre-run +state manifest (image hashes, shared-RAM file identity/hash, mdev/module +versions, boot id) recorded in the session would have made the stale +shared-memory image immediately obvious instead of seven cycles later +`[log PART 93/Add.29]`. This is cheap, local and applies to both the physical +and simulator paths. + +### 7.6 What not to re-derive + +Do not rebuild what already works: QEMU execution, the session format, hooks/ +watchpoints, the Ghidra join, the injection SDK. The review found no +evidence that reimplementing these would pay; the evidence points entirely at +evaluation cost and evidence plumbing. + +--- + +## §8 Ranked pragmatic upgrade list + +Ranking rule: expected reduction of §4 rank-1/2/8 costs per engineering week, +with the eng-ops admission fields (`open-iov/docs/eng-ops.md` §2). Thresholds +are stated for the record and must be frozen before the first measurement run. + +### U1 (top) — Land the S1 device model to B4 (message loop) on the real image + +* **Scenario (dynamic-RE loop):** "what does the real 570.144 GSP firmware do + between entry and its message loop, deterministically, without hardware?" +* **Measurable outcome:** boot-ladder rung reached, session hash, device state + hash, unknown-MMIO count, first host message; `sim_qemu_line` verdict field. +* **Baseline (recorded):** QEMU `probe` first fault at fw `0x4009014` + (`docs/06` §3) plus static reading (floor estimate 600 s in `re_demo.py`); + every real-firmware value demo marked `blocked_on_device_model`. +* **Threshold:** reach B3 (first CSB/TRANSCOM poll with read-back) and B4 + (message loop) twice with byte-identical key fields; one scripted host + message served end-to-end with a recorded reply (B5); `device_model_used = + true` and `above_qemu_line = true` in `harness/sim_qemu_line.py`; N=3 runs, + one host. +* **Failure signal:** two consecutive image/profile versions require + unabsorbed code changes (profile seam failed), or replay key fields differ + after three attempts (scheduler/determinism failed); or lenient MMIO is + still load-bearing in a rung claimed as modelled. +* **Why first:** every measured value projection in `docs/25` and every + blocked question in §4 is gated on this one capability. + +### U2 — Prove the tool on the live wall: surgical rehost of the GR +engine-context promote/commit path (and the UVM large-push parser regression) + +* **Scenario (dynamic-RE):** the project's current product wall is a + server-internal commit failure; answer its open hypotheses (the V6 + soft-`INVALID_STATE` site; the V5 commit assert's *first* failure) from + rehosted execution with hooks/watchpoints, not from another 730-line static + read or another reboot. +* **Measurable outcome:** questions resolved per session; supervised cycles + and reboot requests avoided; recorded session pair per hypothesis. +* **Baseline (recorded):** modes 4–22 across 2026-10-02 (a full day), three + wedge-reboot cycles, one accepted root cause after 730+ lines of static RE + `[log PART 92/Adds.13–25; STATE-CURRENT.md:98-131]`. +* **Threshold:** ≥2 open hypotheses answered with A/B sessions and **zero** + hardware cycles; the answer agrees with at least one live/session artefact; + total effort ≤2 engineering days on the rehost path. N=1 path, 2 sessions. +* **Failure signal:** after two bounded attempts the failing code cannot be + driven (required device semantics absent) or rehosted outcomes disagree with + live evidence; then stop, record the gap, and fall back to static+physical. +* **Why:** it converts the largest single-day cost on record into a + device-free artifact and produces the first non-probe real-firmware sessions, + which is exactly the missing usage in §2.3. + +### U3 — One-command physical-cycle evidence capture → session → `compare` + +* **Scenario (correctness/evidence loop):** every supervised cycle leaves a + machine-readable session and an automatic delta against the previous cycle. +* **Measurable outcome:** human minutes per cycle evidence pack; M1–M6 + skeleton rows auto-filled; unexpected-delta list. +* **Baseline (recorded):** manual multi-MB per-cycle directories, rings + decoded by hand, M-tables filled in prose; s10 takes 4.92 s on the analog + (`docs/25` §6). +* **Threshold:** one cycle's read-only captures (module state, counters, + ring/journal excerpts, pstore status) become a session + `compare` report in + ≤2 wall-clock minutes, N=3 cycles; a field change between cycles is + attributed or listed as unexplained. +* **Failure signal:** the capture misses the decisive field twice (tool does + not see what the failure needs) or deltas are non-reproducible; retire to a + checklist and keep sessions for the simulator only. +* **Why:** it attacks rank-1/rank-8 cost directly without waiting for the + device model, is read-only by construction, and reuses `rmharness`'s proven + recording pattern (`docs/17`). + +### Runners-up (below the top 3, kept for the next review) + +* Get one agent to actually use the MCP server end-to-end on a release-triage + rehearsal (F1 threshold) or retire the claim of agent-driven monitoring. +* Extend value demos to the rehosted path (F6: ≥6/10 on rehosted code or + stop calling them value validation). +* Session retention/prune + a session index (`dumps` is >90 % test noise). +* Fold `port_driver_versions` into `ci_release` rather than maintaining a + third orchestrator (`future-eng-ops` F2 adjudication). +* Calibrate cycles (C1–C6) only when a supervised capture window is approved; + until then keep the modeled label and claim nothing. + +### Do-not-build list (with the trigger that would reopen each) + +| Do not build | Why | Reopen trigger | +|---|---|---| +| GPGPU-Sim board model (GMMU/FBIF/DMA/vidmem/VBIOS), Phase G | Phase S is the owner-locked scope; designed-only | Phase S pause + a funded board question | +| Ampere/Blackwell architecture models | parameter dossiers exist; no consumer | a device port actually scheduled | +| Verilog/FPGA soft processor (F12/F13) | multi-quarter, no funded timing claim | a funded cycle-accuracy claim | +| SIMD/kernel execution on device (F11) | unbounded research; hardware executes kernels today | funded thesis/cycle-accuracy use | +| Moments/time-travel before B4 | deferred by the pragmatic scope lock; convenience, not support latency | message loop reached + a real replay need | +| TCG-trace session folding (R4 remainder) | trace already exists; folding is convenience | trace consumers ask for coverage reports | +| Ghidra heat-maps/bidirectional nav (B2–B4) | bridges work; polish | a real navigation complaint | +| MCP port-chain automation (F1) until a release rehearsal runs | no consumer has exercised the chain | first agent-driven release run | +| Packaging/trialing (F4) as a workflow | no reviewer demand yet | a reviewer asks for a replayable bundle | +| A second multi-version orchestrator | duplicates `ci_release` | F2 threshold unmet **without** a bundle precondition | +| RM 610 anything | hard human gate (`docs/23`) | explicit named-human approval | +| SBOM/observation inventory (F15) | legal/review gates; no consumer | SBOM review asks for per-image rows | +| New MCP transports / network exposure | loopback/stdio is the safe surface | never (security rule) | + +--- + +## §9 Frida-like capability gap analysis + +### 9.1 What exists + +**[E]** hooks (enter/leave on RV64 ABI, `Session.hook`), breakpoints (sw/hw), +watchpoints (write/read/access with clear/step/re-arm to match QEMU stop +semantics), memory read/write/hexdump/search/snapshot/diff, schema struct +decode, recorded sessions, `run_until` (run to cursor), `call` (remote +function call with scratch stack and ABI save/restore), injection of +position-independent services (direct and firmware-mailbox paths), exec/access +trace (native and TCG), plugins, MCP, and Ghidra in/out bridges. Source: +`docs/00` §Status, `docs/01` §4, `docs/11`–`13`, `docs/27`. + +### 9.2 What is missing, and whether it matters day-to-day + +| Frida-class capability | In the tool? | Day-to-day matter | Verdict | +|---|---|---|---| +| Backtrace / stack unwind on a hit | no (only `ra`-based leave hooks) | high: "who called this handler" is a daily RE question | **build later** — an `ra`/frame-pointer walk is bounded; pairs with `run_until` | +| Semantic call tracing (Stalker-like call graph) | partial: full exec trace, no call graph/summary | medium-high: call ordering at scale | **extend trace** with a call/extent fold when B4 exists | +| Catchpoints (trap/exception class selection) | partial: traps stop as plain stops; catchpoints planned (S2) | high for error-path RE (fault handling, asserts) | **build now (thin)** — approved in the pragmatic scope table | +| Hart/thread selection | no (single hart) | low today; needed for multi-core libos | defer with H2.4 | +| Interceptor *replace* (patch a function body, not just insert) | partial: SDK services + debugger writes; no atomic replace | medium: fault injection/return-value forcing | keep injection SDK; add a documented return-patch helper only if a scenario needs it | +| Module/symbol enumeration at runtime | n/a (stripped bare-metal; symbol DB from static exports) | n/a | not applicable | +| DWARF type introspection | no (manual schemas) | medium: manual schema per structure is a cost | acceptable; schemas are versioned data by design | +| Memory allocation tracking | no | low for firmware servers | do not build | +| Remote/network attach, ASLR/module-load handling, Java/ObjC bridges, debugger-side allocation hooks | no (deliberately) | none | not applicable; keep out | +| Allocation/protection changes (mprotect-like) | no; MPU is not enforced (`docs/36` §6.8) | medium for fault experiments | defer; documented gap | + +### 9.3 The real gap + +**[I]** The Frida-like feature list is not the constraint. The constraint is +that none of it can yet run against the real firmware to the point where a +live question exists (B4). Every missing feature above is worth less than the +device model. The Frida-like work that *is* worth doing before B4 is small and +already approved: crash/trap catchpoints, hart/core metadata, and memory +region metadata (S2 "thin"); a backtrace would be the next candidate because +it converts the most common live question into one call. + +**[I]** Conversely, several Frida-ish ideas have no firmware use case: +network attach, module enumeration, allocation tracking, OS-level process +injection, Java/ObjC bridges. They should stay out of the roadmap's active +list to keep the feature-count pressure off the critical path. + +--- + +## §10 What to scrap or disable + +Nothing is deleted by this review; these are recorded recommendations for the +next audit (`docs/27`) and the `future-eng-ops` §4 prune/demote review. + +| Item | Action | Evidence / rule | +|---|---|---| +| `Session.symbol_at` (no caller) | **prune at next audit** unless a caller appears | already the first prune candidate (`future-eng-ops` §4) | +| "Value validation" framing of s01–s10 | **demote** to "tooling regression suite" until ≥6/10 class transfer on rehosted code | F6 kill criterion: fewer than four classes transfer | +| `--surface rm` control discovery | keep code, **remove from advertised next steps** until H1.3/H2.2 exposes the table | F5: `built, not used yet`; exits clean | +| `harness/port_driver_versions.py` as a third orchestrator | **fold into `ci_release.py`** unless the F2 threshold is met without a bundle precondition | F2 adjudication note | +| `harness/port_new_device.py` / `package_sim.py` without a consumer | keep as pilots, **do not advertise as admitted**; expire per eng-ops pilot rule | F3/F4 unadmitted | +| Trace session folding / heat-maps / bidirectional nav | remove from "next" lists; keep issues recorded | pragmatic scope lock (§14.2 of the private program ledger) | +| MCP tools that are QEMU-only on the native backend | keep the structured `not_available`, but stop counting them toward "native parity" | `docs/20` §8; parity matrix in `docs/35` §5 | +| `dumps/` growth | add retention/prune for runner sessions; never track | 458 dirs / 107 MB, >90 % `unittest-*` (measured 2026-10-03) | +| GPGPU-Sim overlay further work | freeze until Phase S pause | owner scope lock; overlay build only | +| Historical tool ledgers | mark historical; stop treating as live plan | §5 #11; §6.4 | +| Test-count claims | replace with generated value; delete hand-written numbers | §5 #5 | + +**Disable-now candidates (if a hard choice is needed):** stop advertising +`--surface rm`, port/packaging pilots and trace folding as available +workflows, and stop quoting hand-maintained test/tool counts. Each currently +consumes review attention without a consumer. + +--- + +## §11 Review conclusion (one paragraph) + +**[I]** The tooling is real, honest about fidelity, and already proven on its +own terms, but the evidence shows it has not yet paid for itself on the +project's dominant cost: the supervised hardware cycle. The single +highest-value move is to finish the device model to the message loop so the +hundreds of blocked questions become sessions; the second is to prove the tool +on the current product wall by rehosting the failing path; the third is to +capture every physical cycle as a session so evidence stops being hand-built. +Everything else — more features, more docs, more architectures — should wait +until those three produce recorded, measurable time savings. + +--- + +*Change note: new document; scope: review only; benefit: consolidates the +2026-10-03 tool/simulator/wall evidence into actionable, gated upgrades; +no product, hardware or guest action authorised.* diff --git a/debug/risc-v-gsp-rm-vm/docs/40-NEW-VERSION-TRIAGE.md b/debug/risc-v-gsp-rm-vm/docs/40-NEW-VERSION-TRIAGE.md new file mode 100644 index 0000000..613e966 --- /dev/null +++ b/debug/risc-v-gsp-rm-vm/docs/40-NEW-VERSION-TRIAGE.md @@ -0,0 +1,296 @@ +# 40 — New-version triage (one command) + +> **Status: in development (2026-10-03).** Implemented, tested and measured on +> the TU104 570.144 image; the measured value proof is section 5. Interfaces +> and report fields may still change. Simulator-only: never touches hardware, +> GPU, mdev, a kernel module or a guest; nothing is downloaded. +> +> **Audience:** agents and maintainers of risc-v-gsp-rm-debugger. +> +> This document is public-safe: it describes components, behaviour and +> measured numbers only. + +--- + +## 1. Why this exists (U46) + +The program has one goal: cut the time from a new NVIDIA RM driver / GSP +firmware release to supported in `vgsp`. Today a new image is triaged by hand +by reading docs/07 P1 and invoking four different tools in sequence (probe, +native boot, A/B, version diff), then collecting session paths and writing a +summary. + +`harness/triage.py` turns that into **one bounded command** that reuses the +existing tools - it does not reimplement any of them: + +| Step | Reuses | Answers | +|---|---|---| +| `identity` | `greeneye.elf`, `sha256` | what is the image, which version, which baseline | +| `probe` | `harness/sim_qemu_line.py:run_qemu_baseline` (the v0.1 probe path) | where does it first fault / first touch under QEMU | +| `native_boot` | `python3 -m sim.gsp_boot` (S1 device model + boot ladder) | how far does the native ladder get (B0–B4) | +| `ab_eval` | `harness/sim_qemu_line.py:run_harness` | QEMU vs native A/B, determinism, frozen thresholds | +| `version_diff` | `harness/version_diff.py` (only with `--baseline`) | image/symbol/corpus/probe/behavior delta, RM 610 gate respected | +| `features` | none (introspection) | hooks/watch/memory/session API, native sim, TCG plugin, MCP, demo | +| `support_readiness` | derived | blockers and ordered next actions for `vgsp` support work | + +Every step is wrapped: a failing step is recorded as `status: error`, the +remaining steps still run, and the report is still written. Every session a +step used is listed with its sha256 hashes. + +--- + +## 2. Usage + +Run from the tool root (`open-iov/debug/risc-v-gsp-rm-vm`): + +```bash +# Plan only (no targets opened, no --out created) +python3 harness/triage.py \ + --firmware ../../firmware/57X-GSP/570.144/gsp-570.144.fwimage.bin \ + --version 570.144 --out dumps/triage-570.144 --dry-run + +# Single new image: identity + probe + bootstrap ladder + A/B + features +python3 harness/triage.py \ + --firmware <new.fwimage.bin> --version <label> --out dumps/triage-<label> + +# With a baseline: adds the full version diff (docs/21) +python3 harness/triage.py \ + --firmware <new.fwimage.bin> --version <new-label> \ + --baseline <old.fwimage.bin> --baseline-label <old-label> \ + --out dumps/triage-<old>-<new> + +# One command via the launcher or make +./risc-v-gsp-rm-debugger.sh triage --firmware <new> --out dumps/triage-new +make triage FW=<new> OUT=dumps/triage-new TRIAGE_ARGS="--baseline <old>" +``` + +Flags: `--firmware` (required), `--baseline`, `--version`, `--baseline-label`, +`--out` (required unless `--dry-run`), `--budget` (native instructions per +attempt, default 200000, hard cap 2000000), `--qemu-duration` (default 0.5s), +`--step-timeout` (per-step wall-clock net, default 300s), `--release`, +`--approval-token`, `--dry-run`. + +Safety: `--dry-run` opens no target and creates no `--out`. The RM 610 gate +runs first in every mode (docs/23) and refuses with exit `3`; labels checked +are `--release`, `--version`, `--baseline-label` and both file basenames. +Exit codes: `0` report written (skipped steps included), `1` at least one step +errored, `2` tool/input error (or firmware absent), `3` RM 610 refusal. + +Output tree under `--out`: + +```text +<out>/ + triage-report.json stable machine-readable report + triage-report.md rendered report + sessions/qemu-probe/ QEMU probe session + sessions/native-boot/ S1 boot-ladder session + ab-eval/ docs/37 A/B report + sim-qemu-line.json|.md + sessions/ + version-diff/ only with --baseline (docs/21 tree) +``` + +--- + +## 3. Report schema (`greeneye-new-version-triage/1`) + +```json +{ + "format": "greeneye-new-version-triage/1", + "tool": "harness/triage.py", + "generated_at": "2026-10-03T04:44:24Z", + "rm610_gate": {"status": "passed", "labels": ["570.144"]}, + "input": { + "firmware": {"path": "...", "sha256": "...", "size": 28528288, + "version": "570.144W", "version_source": "gsp_version.txt sidecar"}, + "baseline": null, "baseline_supplied": false, + "budget": 200000, "qemu_duration": 0.5, "triage_command": "..." + }, + "thresholds": {"frozen_on": "2026-10-03", "manual_baseline_s": 1800.0, + "triage_wall_s_max": 600.0}, + "steps": [ + {"name": "identity", "status": "ok", "wall_s": 0.099, "firmware": {...}, "elf": {...}}, + {"name": "probe", "status": "ok", "rung": "pre-B0", "first_fault": {...}, + "stopped_pc": 0, "sessions": [...]}, + {"name": "native_boot", "status": "ok", "rung": "B2", + "instructions_executed": 14682, "stop": {...}}, + {"name": "ab_eval", "status": "ok", "native_rung": "B2", + "native_identical": true, "above_qemu_line": true, "verdict": "above"}, + {"name": "version_diff", "status": "skipped", "reason": "no --baseline supplied"}, + {"name": "features", "status": "ok", "checks": {...}, "unavailable": []} + ], + "sessions": [{"step": "probe", "name": "qemu-probe", "dir": "sessions/qemu-probe", + "present": true, "hashes": {"session.json": "...", "events.jsonl": "...", "summary.md": "..."}}], + "support_readiness": {"verdict": "sim-observable", "baseline_compared": false, + "observable_rung": "B2", "blockers": [], "next_actions": ["..."]}, + "threshold_checks": {"triage_wall_s": 2.781, "triage_wall_s_max": 600.0, + "wall_within_max": true, "manual_baseline_s": 1800.0, + "speedup_x": 647.2, "beats_manual_baseline": true, + "required_steps_missing": []}, + "wall_s": 2.781, + "artifacts": {"report_json": "<out>/triage-report.json", "report_md": "<out>/triage-report.md"} +} +``` + +The `wall_s`/`threshold_checks` values above are the **measured 2026-10-03 +sample** and are host-load dependent, not fixed constants. + +Key fields: + +* `steps[].status` is `ok | skipped | error`; `skipped` carries a reason and + never fails the run. +* `support_readiness.verdict` is `blocked` (identity/step errors), + `first-touch-only` (no runtime observation yet) or `sim-observable` (native + ladder reached `B2`+, the point where the firmware runs past boot and can be + instrumented). +* `support_readiness.next_actions` is the ordered work list for `vgsp` support + (device-model gaps, symbol DBs, A/B above-the-line work, missing inputs). +* `sessions[].hashes` is the docs/07 evidence-checklist hash list; attach it + with the recorded sessions. + +--- + +## 4. Worked example — 570.144 (measured 2026-10-03) + +Commands (real, read-only; `/usr/bin/time` around the wrapper): + +```bash +./risc-v-gsp-rm-debugger.sh triage \ + --firmware ../../firmware/57X-GSP/570.144/gsp-570.144.fwimage.bin \ + --version 570.144 --out dumps/triage-570.144 +``` + +Measured output (trimmed): + +```text +[triage] identity: status=ok +[triage] probe: status=ok reason=ok +[triage] native_boot: status=ok +[triage] ab_eval: status=ok +[triage] version_diff: status=skipped reason=no --baseline supplied +[triage] features: status=ok +[triage] readiness: sim-observable (baseline_compared=False) +[triage] next: native ladder is at B2 (runtime-observable): inspect the milestone/stop with hooks ... +[triage] next: native A/B is not yet above the QEMU line (verdict=above_lenient_only): device model ... +[triage] next: re-run with --baseline <previous.fwimage> for the symbol/behavior delta ... +[triage] wall: 2.781s vs manual baseline 1800.0s (speedup 647.2x) +``` + +The `2.781s`/`647.2x` line is the **measured 2026-10-03 sample** (host-load +dependent, not a fixed constant). + +What the report recorded: + +| Step | Measured | Meaning | +|---|---|---| +| identity | `570.144W` via `gsp_version.txt`, sha256 + 28,528,288 bytes | image paired and hashed before any run | +| probe | QEMU first fault `pc_fw 0x4009014`, load access at `0x2000000000000A00` | matches the docs/37 QEMU reference first touch | +| native_boot | rung **B2**, 14,682 instructions, stop fault `pc_fw 0x400EA0C`, 0 unknown MMIO | S1 ladder is two rungs past the QEMU fault; boot shim consumed, relocation milestone hit | +| ab_eval | native `B2` (14,682 instructions, 0 unknown MMIO), identical across the repeat, `above_qemu_line: true`, verdict `above` | post-S1 re-run (2026-10-03, later): the discovered `sim/profiles` model + shim now satisfy the docs/37 line contract; the pre-S1 run was `B0`/`above_lenient_only` | +| features | greeneye API, native backend, boot ladder, A/B, version_diff, TCG plugin, MCP, demo, QEMU all present | capability surface for the follow-up work | + +Post-S1 re-run (2026-10-03, later; same command, after S1 and the harness +wiring landed): `ab_eval` reports native rung **B2**, all five frozen +thresholds true and `above_qemu_line: true`; the "native A/B is not yet above +the QEMU line" next-action line no longer appears. Wall ≈2.7 s +(host-load dependent). The pre-S1 trace above is kept as the recorded +before/after evidence. + +Baseline variant (same image both sides, an unchanged-version rehearsal with +the docs/21 path): + +```bash +./risc-v-gsp-rm-debugger.sh triage \ + --firmware ../../firmware/57X-GSP/570.144/gsp-570.144.fwimage.bin \ + --baseline ../../firmware/57X-GSP/570.144/gsp-570.144.fwimage.bin \ + --version 570.144 --baseline-label 570.144 --out dumps/triage-570.144-baseline +``` + +Measured (2026-10-03, host-load dependent): **11.388s** internal wall +(11.64s including interpreter startup); +`version_diff` = 8.526s, sections `image/probe/behavior: ok`, +`symbols/corpus: skipped` with reasons; image self-diff `same: true`; 7 session +directories hashed (probe, native boot, 2 A/B native, A/B QEMU, version-diff +probe + behavior). Readiness stayed `sim-observable`, `baseline_compared: +true`, with next actions for the missing symbol/corpus inputs. + +Frozen thresholds (recorded in `harness/triage.py` on 2026-10-03 **before** +the runs, per handoff §14.5): manual baseline `1800s`, triage wall max +`600s`. Result: `wall_within_max: true`, `beats_manual_baseline: true`, +speedup `647x` (single image) / `158x` (with baseline) — both wall times are +**measured 2026-10-03 samples (host-load dependent)**, not fixed constants. The +manual baseline is +a conservative understatement of the steps replaced (read docs/07 P1, choose +budgets, run probe + boot ladder + A/B, run version diff, collect session +hashes and write the summary for one image). + +--- + +## 5. Value proof and its limits + +* The frozen contract is in the report itself (`thresholds`, + `threshold_checks`); the same numbers are restated in the module docstring + and here. Never move them silently. +* The measured numbers above are from this checkout on 2026-10-03; reproduce + with the exact commands in section 4. +* **Conservative baseline:** 1800s is an *estimate of manual effort*, not a + timed human study. The speedup numbers are reported against it, and the + report flags them as such. +* **What the run proves:** one bounded command produces the identity, probe, + boot-ladder, A/B and capability evidence with hashed sessions in seconds - + the mechanical part of docs/07 P1. +* **What it does not prove:** that the remaining human work (interpreting the + deltas, extending the device model, the accuracy loop, the `vgsp` change) is + automated; it is not, and the next-actions list is where it stays explicit. + +--- + +## 6. Release-to-support mapping + +| Triage output | Next workflow (docs) | +|---|---| +| identity + hash + version | docs/07 P1 step 1, docs/22 pairing checks, docs/28 sweep input | +| probe first fault / first touch | docs/07 P1 step 2, docs/24 device map harvest | +| native boot rung + stop | docs/35 §3.3 ladder; the `vgsp` support scope starts at `B2`+ | +| A/B verdict | docs/37 line contract (U50) | +| version_diff sections + symbol/corpus skips | docs/21, docs/14 §6 evidence pack | +| readiness next actions | docs/07 P1.7 acceptance, the actual `vgsp` support work items | +| RM 610 gate | docs/23 (never bypassed) | + +--- + +## 7. Not verified (explicit) + +* **No real 570.211+ new-version pair exists locally**; the cross-version + delta path is exercised only by the self-diff rehearsal in section 4. +* **Symbols and corpus are skipped for raw `.fwimage` targets** (no `.symtab`, + no Ghidra exports supplied): the handler-inventory delta is not measured + until those inputs are supplied. +* **The S4 A/B verdict is `above` on the frozen checks after S1 landed** + (B2, documented device model + shim, deterministic, within the wall + bound); B3/B4 are not reached and triage makes no "beats QEMU" claim beyond + the documented rung contract (`above_lenient_only` was the pre-S1 result). +* **Fidelity:** L0/L1 execution facts only; no L2 timing and no full L3 + device-environment claims (docs/06, docs/36). Cycles are modeled; the boot + shim values are declared assumptions. +* **No physical backend, `/dev`, module, mdev, guest or network access** was + used or is performed by the tool. + +--- + +## 8. Tests + +```bash +python3 -m unittest tests.test_triage -v # 14 tests +make test # full suite +``` + +`tests/test_triage.py` covers the version-label extraction chain, the report +shape and artifact writes, session hashing, partial-failure handling +(injected step failure still yields a report), missing-firmware behaviour +(reported, not raised), budget bounds, RM 610 refusal (exit `3`), dry-run +side-effect freedom, and a tiny-budget integration run against the real +570.144 image that skips cleanly when the image or QEMU is absent. + +<!-- change (2026-10-03): new document recording the one-command new-version +triage workflow (U46), its measured 570.144 value proof and its limits; +scope: docs/40 only; benefit: the release-to-support loop has one documented +entry point and one attachable evidence bundle. --> diff --git a/debug/risc-v-gsp-rm-vm/docs/41-STABILITY-AND-TIME-TO-VALUE.md b/debug/risc-v-gsp-rm-vm/docs/41-STABILITY-AND-TIME-TO-VALUE.md new file mode 100644 index 0000000..77e1298 --- /dev/null +++ b/debug/risc-v-gsp-rm-vm/docs/41-STABILITY-AND-TIME-TO-VALUE.md @@ -0,0 +1,122 @@ +# 41 — Phased practical improvements to risc-v-gsp-rm-debugger: stability and time-to-value report + +> **Date:** 2026-10-03. **Scope:** the simulator/debugger workstream inside +> `debug/risc-v-gsp-rm-vm/` only. **Verdict:** the minimum bar for marking the +> *tool and standalone GSP simulator* as **stable (conditional)** is passed: +> 298 tests green twice, deterministic real-firmware runs, an independently +> reproduced A/B verdict "above the QEMU line", and reproducible value +> measurements. Conditions are listed in §7 — chiefly that the device model +> currently reaches boot rung B2 (not the message loop) and cycle numbers are +> modelled, not silicon-calibrated. +> +> Local notes: `../../../../handoff-simulator.md` (directive ledger, top work), +> `../LOCAL-REVIEW-NOTES.md` (local provenance). Public text describes +> components and behaviour only. + +## 1. What we were doing before (the QEMU baseline) + +* The only execution engine was `qemu-system-riscv64` via the gdbstub. On the + real `gsp-570.144.fwimage.bin` it loaded flat, executed a handful of + instructions, faulted at firmware vaddr `0x4009014` reading + `0x2000000000000A00`, then looped on `mtvec=0`. +* Questions about GSP boot progress could only be answered by static reading + of decompiled code; every dynamic question stalled at the first fault. +* Tool runs produced sessions, but with no device semantics, no cycle context, + and no way to prove "the emulator knows more than the probe". +* New-firmware triage was a manual sequence of probe + version diff + reading. + +## 2. What we are doing now (the upgrade, phased) + +| Phase | Improvement | Status | +|---|---|---| +| P1 | **GSP device model** (`sim/devices/`, `sim/profiles/`) with PMC/FALCON/PRISCV/timer/CSB/FWSEC/PMU/DMA/IRQ windows, versioned profile, boot shims | B2: 14,682 instr, fault `0x400EA0C`, 0 unknown MMIO, two-run identical | +| P2 | **`sim.gsp_boot` boot-ladder runner** with rung classification and serialisable device state | B0–B2 (B3/B4 targets) | +| P3 | **One-command release triage** (`./risc-v-gsp-rm-debugger.sh triage`) orchestrating identity/QEMU probe/native boot/A-B eval/version diff/readiness | 570.144 in ~2.7 s; baseline pair ~11.4 s | +| P4 | **QEMU-line A/B evaluation harness** (`harness/sim_qemu_line.py`, frozen checks) | verdict `above`, all threshold checks computed from live runs | +| P5 | **Real RE demonstration** (`harness/re_demo.py`) answering a concrete bootstrap question with simulator tools | PASS 5/5, 0.13 s vs 601 s baseline floor | +| P6 | **MCP native-backend tools** (launch, cycles, device state, profile, regions, budgets) | 55 MCP tests green; read-only default preserved | +| P7 | **Documentation and workflow review** (`docs/39`), docs 30–41, README/Makefile integration, link checks | 143 tool files link-checked, 0 broken | + +## 3. Proof of value (before → after) + +| Scenario | Before | Now | Evidence | +|---|---|---|---| +| What happens in GSP boot before the first wall? | QEMU: fault at `0x4009014`, stop | Device model serves the path; deterministic B2 fault `0x400EA0C`; unknown accesses 0 | `python3 -m sim.gsp_boot`; sessions `dumps/gsp-boot-b2-run1/2` | +| Triage a new firmware image | manual probe + reading, ~1800 s estimated | one command, **2.68 s**, JSON+MD report with hashes | `./risc-v-gsp-rm-debugger.sh triage --firmware … --version 570.144` | +| Is the emulator better than QEMU? | no mechanism | validators run both, compute rung/feature deltas; verdict `above` | `python3 harness/sim_qemu_line.py` | +| Answer a bootstrap RE question | static reading only | 5/5 recorded answers in 0.13 s | `python3 harness/re_demo.py` | +| Scripted/agent operation against the simulator | QEMU-only MCP | native backend tools tested | `tests/test_mcp.py` | + +Baselines are explicitly labelled estimates where they are not measured human +studies (frozen thresholds: 600 s RE-time floor, 1800 s triage baseline); all +wall times are labelled host-load dependent. Every number above has a +one-line command in the report/doc that produced it. + +## 4. Stability evidence + +* `python3 -m unittest discover -s tests -t .` → **298 tests, OK**, twice + (rounded identity across runs); MCP 55, triage 14, device model 34 new + checks included. +* Determinism: B1/B2 device-state SHA identical across runs; triage sessions + hashed; A/B harness repeats native runs and compares. +* Back-compat: `Session`, `RspClient`, `SimClient` and CLI keep prior + behaviour; additions are keyword/default forms; QEMU path untouched + (toggle `--backend qemu|native`). +* Boundary compliance: no writes outside the tool dir; no hardware, module, + guest, mdev, reboot or RM 610 action; MCP read/observe default unchanged. +* Fidelity honesty: cycles are labelled *modeled*; L2 is not claimed + silicon-accurate; L3 device semantics are declared assumptions in-session; + "above QEMU" is computed against the frozen rung contract, not asserted. + +## 5. Adversarial review record + +* **Round 1** (review + fix): 3 P1s found and fixed — device model not wired + into the A/B harness/triage, dead PTIMER window offsets, MCP device-state + returning a non-JSON object. P2 backlog listed. +* **Round 2** (independent, harsher): falsification runs confirmed all + headline numbers; 2 P1s found and fixed (docs/35 B1 overclaim; personal + host path in `gpgpu-sim/setup.sh`), 8 P2s fixed (timer opt-in, FWSEC state + derivation, DMA refusal counting, PMC chip-ID duplication, overlapping + window fallback, stale doc counts, `session_launch` documentation, docs/40 + measurement labels). One latent item recorded: WPR2 window offsets versus + profile base (firmware never touches it in B0–B2). +* Verdict after fixes: conditional stability as stated in §7; no unverified + "beats silicon" or "boots RM" statement anywhere. + +## 6. Top work for the next session (time-to-value first) + +1. **Push the device model to B3/B4** (CSB/TRANSCOM mailbox semantics, + DMEM runspace source, PMU/FWSEC stubs) — this is the gate to serving a + real RM message and the largest remaining time-to-value lever. +2. **Surgical rehost of the live GR engine-context promote/commit wall** + (docs/39 top-3): convert the manual, multi-cycle physical campaign into + device-free hook/watch/compare sessions. +3. **Physical-cycle capture → session → compare** helper (read-only), to cut + the dominant recorded cost of manual evidence assembly. +4. Everything else remains in the private handoff: GPGPU-Sim board model, + VF bridge (`GREENEYE_GSP_EM_BRIDGE`, default OFF), VBIOS, fabric, + Ampere/Blackwell, and the "not useful yet" list. + +## 7. Conditions on the stability claim + +1. Device model ceiling is B2; B3/B4 pending profile semantics. +2. Cycle counts are modelled defaults, not calibrated bands. +3. Manual-baseline speedups use labelled estimates, not timed human studies. +4. Cross-version triage was exercised as a self-diff rehearsal (no second + local firmware version); RM 610 remains human-gated. +5. The stability claim covers the tool and its standalone GSP simulator, not + any GPGPU-Sim integration or silicon equivalence. + +## 8. Reproduce everything + +```bash +cd open-iov/debug/risc-v-gsp-rm-vm +python3 -m unittest discover -s tests -t . # 298 tests +python3 -m sim.gsp_boot # B2 +python3 -m sim.gsp_boot --no-shim # B1 budget stop +python3 harness/sim_qemu_line.py # verdict "above" +python3 harness/re_demo.py # PASS 5/5 +./risc-v-gsp-rm-debugger.sh triage \ + --firmware ../../../firmware/57X-GSP/570.144/gsp-570.144.fwimage.bin \ + --version 570.144 --out /tmp/triage-570.144 +``` |