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

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

12 files changed, 1445 insertions, 0 deletions

diff --git a/debug/risc-v-gsp-rm-vm/gpgpu-sim/ADMISSION.md b/debug/risc-v-gsp-rm-vm/gpgpu-sim/ADMISSION.md
new file mode 100644
index 0000000..0cf7585
--- /dev/null
+++ b/debug/risc-v-gsp-rm-vm/gpgpu-sim/ADMISSION.md
@@ -0,0 +1,38 @@
+# Admission records — GreenEye-GPGPU-Sim (DRAFT, pending maintainer filing)
+
+> **Status: drafts.** These five records were prepared with the Phase-0
+> program plan and are reproduced here for review. They are **not yet filed**
+> as public eng-ops pilot admissions: a maintainer must review, name the value
+> axis and file them in the public pilot docs before any of these capabilities
+> is treated as admitted. Nothing below changes a threshold after the fact.
+>
+> Per the project admission rule, each record carries all five fields —
+> scenario (loop), measurable outcome, baseline, threshold with N and
+> environment, and failure signal. A missing field means do-not-build, not
+> "deferred". Thresholds were frozen before the first run of each capability.
+>
+> This overlay is development-only; it never becomes a product or guest
+> component, and no record authorises hardware, module, mdev or guest access.
+
+## Draft records
+
+| Record | Tool | Scenario (loop) | Measurable outcome | Baseline | Threshold (N, env) | Failure signal |
+|---|---|---|---|---|---|---|
+| TOOL-2026-SIM-01 | greeneye-gpgpu-sim build/run | porting/RE loop: run a CUDA kernel off-hardware | functional run completes; perf stats emitted; session dir | "no simulator" → static reading only, or QEMU-only | kernel suite 100% functional; build from clean checkout ≤30 min (N=3, this host) | build not reproducible; run non-deterministic |
+| TOOL-2026-SIM-02 | TU104 model calibration | accuracy loop: answer "does the model match?" | C1–C6 metric table vs published bands | GPGPU-Sim SM75 config (uncalibrated) | latency bands ±15%, bandwidth ±20%, correlation ≥0.9 (N≥5 kernels) | metrics fail 2 consecutive releases or are tuned per-kernel |
+| TOOL-2026-SIM-03 | GSP-in-sim (real firmware) | dynamic RE loop: what does firmware do before hardware exists? | session with `cycle` events + boot-progress milestone | QEMU probe (first fault at `0x4009014`) | reaches ≥1 milestone past probe deterministically (N=3 identical) | non-determinism; unexplained delta on replay |
+| TOOL-2026-SIM-04 | MCU cross-controller model | RE/portability: cross-controller exchange | device-level session + latency table | none (no model) | one exchange reproduced with declared sources; latency within documented band | semantics unprovable in one milestone (kill per F8) |
+| TOOL-2026-SIM-05 | sim debugger backend | debugging loop: hook/watch/Ghidra on simulated target | event-parity session vs QEMU on demo | QEMU workflow (minutes) | parity suite green; workflow ≤50% of QEMU baseline time (N=5) | parity broken or slower than baseline |
+
+## Value-axis mapping (to be confirmed at filing)
+
+| Record | Primary value axes |
+|---|---|
+| TOOL-2026-SIM-01 | development velocity on simulated GSP; analysis throughput |
+| TOOL-2026-SIM-02 | performance analysis; correctness of analysis |
+| TOOL-2026-SIM-03 | faster reverse engineering; new-GSP support time |
+| TOOL-2026-SIM-04 | architectural understanding / SBOM; device support |
+| TOOL-2026-SIM-05 | bug-analysis time; guest-compat readiness |
+
+P4 (fabric/system) maps to a fabric value axis only once a peer model exists;
+it is deliberately not filed here.
diff --git a/debug/risc-v-gsp-rm-vm/gpgpu-sim/BUILD-PORT.md b/debug/risc-v-gsp-rm-vm/gpgpu-sim/BUILD-PORT.md
new file mode 100644
index 0000000..b767f8f
--- /dev/null
+++ b/debug/risc-v-gsp-rm-vm/gpgpu-sim/BUILD-PORT.md
@@ -0,0 +1,222 @@
+# BUILD-PORT.md — GPGPU-Sim v4.0.1 on CUDA 12.x / gcc 13 (WP0.1)
+
+Status: **builds and runs** on this host; `vectorAdd` passes in functional and
+performance modes. This is a development-tool port record, not a support
+claim. Scope: get upstream GPGPU-Sim v4.0.1 to compile and execute a minimal
+CUDA kernel with a modern distro toolchain, and record the minimal patch set.
+
+## 1. Upstream pin
+
+| Item | Value |
+|---|---|
+| Repository | `gpgpu-sim/gpgpu-sim_distribution` (public) |
+| Commit | `90ec3399763d7c8512cfe7dc193473086c38ca38` |
+| Tag | `v4.0.1` |
+| License | BSD-3-Clause (upstream); development copy only |
+
+No upstream fork/history rewrite: reproduce by cloning the pin and applying
+the two patches in `patches/`.
+
+## 2. Toolchain and CUDA used
+
+| Item | Value |
+|---|---|
+| OS | Ubuntu 24.04-class x86-64 |
+| Compiler | gcc/g++ 13.3.0 |
+| make | 4.3 |
+| CUDA toolkit | **12.0 (V12.0.140)**, `/usr/bin/nvcc` (distro package), `CUDA_INSTALL_PATH=/usr` |
+| Extra host package | `xutils-dev` (provides `makedepend`, required by the build) |
+| GPU required | **No** — the simulator is CPU-only; no module load, no `/dev` access |
+
+Notes discovered during the port:
+
+* `/usr/local/cuda-12.9` exists but contains only `gds/` — it has **no
+  `bin/nvcc`**, so it cannot be used as `CUDA_INSTALL_PATH`. CUDA 12.0 from
+  `/usr` is the toolkit actually used.
+* The application must be compiled against the **same CUDA major** as the
+  simulator (`get_app_cuda_version()` asserts major equality): CUDA 12 app +
+  CUDA 12 simulator.
+
+## 3. Reproduction (exact commands)
+
+```bash
+# 1. Host prerequisites
+sudo apt-get install -y build-essential xutils-dev        # makedepend
+#    plus a CUDA toolkit (nvcc) — here /usr (CUDA 12.0)
+
+# 2. Pinned clone
+git clone https://github.com/gpgpu-sim/gpgpu-sim_distribution gpgpu-sim
+cd gpgpu-sim
+git checkout 90ec339        # tag v4.0.1
+
+# 3. Apply the tracked port patches (in order)
+git apply patches/0001-cuda12-toolchain-and-soname.patch
+git apply patches/0002-gcc13-cuda12-source-fixes.patch
+
+# 4. Build (release; default)
+export CUDA_INSTALL_PATH=/usr
+source setup_environment
+make -j4        # use -j2 when a QEMU guest is running on the host
+
+# Artifacts:
+#   lib/gcc-13.3.0/cuda-12000/release/libcudart.so   (+ .so.12 symlink)
+#   build/gcc-13.3.0/cuda-12000/release/             (objects, libs, mcpat)
+```
+
+`setup_environment` prepares `PATH` and `LD_LIBRARY_PATH` so that CUDA
+applications load the simulator's `libcudart.so` instead of the real runtime.
+
+### 3.1 Test application
+
+```cuda
+// vectorAdd.cu
+#include <stdio.h>
+#include <cuda_runtime.h>
+
+__global__ void vectorAdd(const float *A, const float *B, float *C, int n) {
+  int i = blockDim.x * blockIdx.x + threadIdx.x;
+  if (i < n) C[i] = A[i] + B[i];
+}
+
+int main() {
+  int n = 1024;
+  size_t sz = n * sizeof(float);
+  float *hA = (float *)malloc(sz), *hB = (float *)malloc(sz), *hC = (float *)malloc(sz);
+  for (int i = 0; i < n; i++) { hA[i] = 1.0f; hB[i] = 2.0f; hC[i] = 0.0f; }
+  float *dA, *dB, *dC;
+  cudaMalloc((void **)&dA, sz);
+  cudaMalloc((void **)&dB, sz);
+  cudaMalloc((void **)&dC, sz);
+  cudaMemcpy(dA, hA, sz, cudaMemcpyHostToDevice);
+  cudaMemcpy(dB, hB, sz, cudaMemcpyHostToDevice);
+  vectorAdd<<<4, 256>>>(dA, dB, dC, n);
+  cudaDeviceSynchronize();
+  cudaMemcpy(hC, dC, sz, cudaMemcpyDeviceToHost);
+  int ok = 1;
+  for (int i = 0; i < n; i++) if (hC[i] != 3.0f) ok = 0;
+  printf("vectorAdd %s (first=%f)\n", ok ? "PASS" : "FAIL", hC[0]);
+  cudaError_t e = cudaGetLastError();
+  if (e != cudaSuccess) printf("CUDA error: %s\n", cudaGetErrorString(e));
+  return ok ? 0 : 1;
+}
+```
+
+```bash
+nvcc -arch=sm_75 --cudart shared -o vectorAdd vectorAdd.cu
+```
+
+`--cudart shared` makes the binary depend on `libcudart.so.12`; the
+simulator's library is named/interposed accordingly (patches 0001).
+
+### 3.2 Run
+
+```bash
+# run dir: app + a config; the entrypoint hardcodes "-config gpgpusim.config"
+mkdir -p run-functional && cd run-functional
+cp ../vectorAdd .
+cp ../configs/tested-cfgs/SM75_RTX2060/gpgpusim.config .
+cp ../configs/tested-cfgs/SM75_RTX2060/config_turing_islip.icnt .
+sed -i 's/^-gpgpu_ptx_sim_mode 0/-gpgpu_ptx_sim_mode 1/' gpgpusim.config  # functional
+
+export CUDA_INSTALL_PATH=/usr
+source ../setup_environment
+./vectorAdd
+
+# performance mode: use gpgpusim.config unmodified (-gpgpu_ptx_sim_mode 0)
+```
+
+Note: GPGPU-Sim 4.0.1 reads options **only** from the config file
+(`gpgpusim_entrypoint.cc` hardcodes `-config gpgpusim.config`); to switch
+functional/perf, edit that line (or the copied file) rather than passing app
+argv.
+
+## 4. Patch set
+
+Two patches against `90ec339`, exported with `git diff`; both apply cleanly
+to a pristine v4.0.1 tree (`git apply --check` verified).
+
+### 0001-cuda12-toolchain-and-soname.patch
+
+| File | Change | Why |
+|---|---|---|
+| `setup_environment` | CUDA upper bound `11100` → `12999` | version gate rejected CUDA 12.0 before anything else could run |
+| `setup_environment`, `version_detection.mk` | gcc version regex `[0-9]\.[0-9]\.[0-9]` → `[0-9]+\.[0-9]+\.[0-9]+` | single-digit-only pattern failed on gcc 13.3.0 → empty `CC_VERSION` → bogus `gcc-/cuda-12000` build paths |
+| `linux-so-version.txt` | add `libcudart.so.12{}` node | CUDA 12 apps reference symbols as `sym@libcudart.so.12` |
+| `Makefile` | create `libcudart.so.12` / `.12.0` symlinks | CUDA 12 binary `DT_NEEDED` is `libcudart.so.12` |
+
+### 0002-gcc13-cuda12-source-fixes.patch
+
+| File | Change | Why |
+|---|---|---|
+| `src/abstract_hardware_model.h` | `textureReference` placeholder struct when `CUDART_VERSION >= 12000` | CUDA 12 removed the legacy texture-reference type from the runtime headers; the simulator's legacy texture path must still compile (unreachable for CUDA 12 apps) |
+| `src/cuda-sim/cuda-sim.cc` | guard legacy `texref->...` field asserts with `#if CUDART_VERSION < 12000` | fields no longer exist in CUDA 12 |
+| `src/cuda-sim/Makefile` | escape `"` in the generated `ptx_parser_decode.def` (both Linux/macOS rules) | bison comments contain quotes (`/* "end of file" */`) which became illegal C++11 user-defined-literal tokens (`""end`) under gcc 13 |
+| `libcuda/cuda_runtime_api.cc` | `return (unsigned)cudaConfigureCallInternal(...)` in `__cudaPushCallConfiguration` | missing `return` is UB; gcc 13 `-O3` compiled it into a non-returning call/announce loop → stack-overflow SIGSEGV on the first kernel launch |
+
+`libcuda/cuobjdump.l` / `cuobjdump.y` needed **no change**: CUDA 12
+`cuobjdump` output is parsed and the embedded PTX is extracted (see §5).
+
+## 5. Observed results
+
+Kernel: `vectorAdd` (`__global__`, grid (4,1,1), block (256,1,1)), sm_75,
+compiled with CUDA 12.0, config `SM75_RTX2060`.
+
+| Run | Result |
+|---|---|
+| Functional (`-gpgpu_ptx_sim_mode 1`) | `vectorAdd PASS (first=3.000000)`; `Done functional simulation (22528 instructions simulated)` |
+| Performance (`-gpgpu_ptx_sim_mode 0`) | `vectorAdd PASS (first=3.000000)`; `gpu_tot_sim_cycle = 5569`, `gpu_tot_sim_insn = 21504`, `gpu_tot_ipc = 3.8614`, `gpu_silicon_slowdown = 245106x`; L2 and interconnect stats printed |
+
+Clean-build verification: a pristine v4.0.1 worktree + the two patches was
+built from scratch (`make -j4`, exit 0, no errors) and both modes were
+re-run there with the same PASS results.
+
+PTX path: `__cudaRegisterFatBinary` → `cuobjdump` extraction
+(`vectorAdd.1.sm_75.ptx`) → PTX parser → functional/perf execution, all on
+CUDA 12 output without grammar changes.
+
+## 6. Known failures / limitations
+
+* **CUDA 12.0.140 only.** 12.9/13.x untested; `/usr/local/cuda-12.9` has no
+  `nvcc` on this host.
+* **Legacy texture API** (CUDA < 12 API) is a compile placeholder for
+  CUDA ≥ 12; texture execution is untested and unreachable for CUDA 12 apps.
+* **Release build only.** `source setup_environment debug` was not built.
+* **PTXPlus / SASS path** (sm_1x, CUDA ≤ 4.2) is untouched and not applicable
+  to CUDA 12.
+* **OpenCL support not built** (no `libOpenCL` configured).
+* Only `vectorAdd` was validated; no upstream test suite/cuBLAS/DNN kernels.
+* Upstream limitations unchanged: no system-level I/O, no multi-GPU, no
+  device/firmware models (those are later program phases).
+* `makedepend` is required; it is not installed by default on this distro.
+* The host build uses `-O3`; the `__cudaPushCallConfiguration` fix is
+  required because gcc 13 miscompiles the missing-`return` body. If a future
+  compiler changes layout the guard remains correct (explicit return).
+
+## 7. Clean-host reproduction checklist
+
+- [ ] `gcc`/`g++` present (tested 13.3.0); `make` present.
+- [ ] `xutils-dev` installed → `makedepend` on `PATH`.
+- [ ] CUDA toolkit with `nvcc`, `cuobjdump`, `ptxas`; set
+      `CUDA_INSTALL_PATH` to its prefix (tested `/usr`, CUDA 12.0). Verify
+      `$CUDA_INSTALL_PATH/bin/nvcc` exists (a bare `cuda-12.9/` directory may
+      not contain the toolkit).
+- [ ] Clone upstream and `git checkout 90ec339` (tag `v4.0.1`).
+- [ ] `git apply` both patches under `patches/`; confirm `git apply --check`
+      passed on pristine tree.
+- [ ] `export CUDA_INSTALL_PATH=...; source setup_environment; make -jN`
+      exits 0; `lib/gcc-<gccver>/cuda-<ver>/release/libcudart.so` exists and
+      `libcudart.so.12` symlinks to it.
+- [ ] Compile the test app with `nvcc -arch=sm_75 --cudart shared`.
+- [ ] Run functional and perf modes from a run dir containing
+      `gpgpusim.config` (+ copied icnt config); expect `vectorAdd PASS`.
+- [ ] Do not run the app without `setup_environment` (would use the real
+      runtime / physical GPU); the simulator path is CPU-only.
+
+## 8. What is not verified
+
+* No debug build, no other configs (`SM7_QV100`, etc.), no CUDA 12.1–12.9/13.
+* No CUDA regression suite, no multi-kernel or library workloads.
+* No texture, surface, dynamic parallelism or OpenCL execution.
+* No performance/accuracy claims beyond "the timing model runs and prints
+  stats"; the numbers above are one small kernel's simulator output.
+* Patch set was validated on one host/toolchain combination.
diff --git a/debug/risc-v-gsp-rm-vm/gpgpu-sim/README.md b/debug/risc-v-gsp-rm-vm/gpgpu-sim/README.md
new file mode 100644
index 0000000..d9e4ff7
--- /dev/null
+++ b/debug/risc-v-gsp-rm-vm/gpgpu-sim/README.md
@@ -0,0 +1,97 @@
+# GreenEye-GPGPU-Sim overlay (in development)
+
+> **Status: Phase 0 scaffold.** This directory is the tracked overlay that will
+> turn GPGPU-Sim into the execution substrate of the
+> [risc-v-gsp-rm-debugger](../README.md). Today it contains a reproducible
+> build/setup path, a Turing TU104 (RTX 2080 SUPER) **configuration scaffold**,
+> a run wrapper, a config-consistency test, and draft engineering-admission
+> records. **It makes no accuracy, timing, performance or support claim**, and
+> it is **not** calibrated to silicon.
+
+## 1. What this is
+
+GPGPU-Sim is an open-source, data-driven GPGPU performance/functional
+simulator (BSD-3-Clause, University of British Columbia). The debugger program
+needs an off-hardware execution substrate that can host the GPU model, the
+GSP/microcontroller models and the debugger's hooks/watchpoints/session
+recording on one deterministic clock. This overlay is how the project tracks
+its changes to upstream:
+
+* `upstream/` — a pinned clone of upstream GPGPU-Sim (gitignored, never
+  committed; see [UPSTREAM.md](UPSTREAM.md));
+* `patches/` — our overlay patch set, applied in lexical order by `setup.sh`
+  (`0001` CUDA-12 toolchain/soname, `0002` gcc-13/CUDA-12 source fixes; the
+  build record is in [BUILD-PORT.md](BUILD-PORT.md));
+* `configs/` — tracked device configurations. `turing-tu104-2080super/` is the
+  first target (the local TU104 card);
+* `scripts/` — `run_sim.sh` (run a CUDA binary against the simulator) and
+  `config_sanity.py` (config consistency test);
+* `ADMISSION.md` — draft admission records TOOL-2026-SIM-01..05, pending
+  maintainer filing.
+
+**Development-only.** This is analysis tooling. It is never shipped, never
+becomes a product dependency or guest component, and never touches GPU
+hardware, the kernel module, mdev, or a guest. Real-firmware simulation is
+Phase-1 work behind explicit gates; nothing here should be imported by
+product code.
+
+## 2. Relationship to risc-v-gsp-rm-debugger
+
+The debugger already has two execution backends: QEMU (real image probe,
+device-map harvest) and the native reference core under `../sim/`
+(instruction-level, L0/L1). GPGPU-Sim becomes the third backend once the
+Backend-ABI work lands (Phase 0/1.8): the same session format, hooks,
+watchpoints, probes, Ghidra bridge and MCP surface drive all of them. The
+TU104 config scaffold establishes the first GPU-side device profile; the GSP
+and MCU models (Phase 1.5) are planned to run as clock-domain citizens of
+this simulator. Nothing in this directory is on the critical path of the
+existing tool: deleting it leaves the debugger fully functional.
+
+See `../docs/30-GPGPU-SIM-INTEGRATION.md` for the public integration note.
+
+## 3. Quick start
+
+```bash
+cd open-iov/debug/risc-v-gsp-rm-vm/gpgpu-sim
+
+# 1. Clone the pinned upstream, apply patches/, build (see UPSTREAM.md).
+#    --no-build materialises the checkout only. Needs a CUDA toolkit;
+#    CUDA_INSTALL_PATH overrides the default /usr/local/cuda.
+./setup.sh
+
+# 2. Check the TU104 scaffold config is internally consistent.
+python3 scripts/config_sanity.py
+
+# 3. Run a CUDA binary (built with --cudart shared) against the simulator.
+./scripts/run_sim.sh --help
+./scripts/run_sim.sh ~/path/to/vectorAdd
+```
+
+Requirements: git, gcc/g++, make, bison, flex, zlib, and a CUDA toolkit.
+`setup.sh` uses at most 4 compile jobs by default (`JOBS` overrides) to respect
+the host budget.
+
+## 4. What is *not* claimed
+
+* **Not calibrated.** Every latency, timing and bandwidth parameter not tied
+  to a published TU104 number is marked `UNKNOWN:` in the config and is
+  inherited from the upstream SM75 Turing model.
+* **Not a full GPU model yet.** No GSP core, no microcontroller models, no
+  fabric/PCIe/RDMA/NVMe models, no guest co-simulation.
+* **No accuracy claim of any kind** until the Phase-1 calibration report
+  exists (admission TOOL-2026-SIM-02). Config-only modeling without
+  calibration evidence is an explicit anti-goal.
+* **Not a product/guest component** and never a support claim.
+
+## 5. Layout
+
+| Path | Purpose |
+|---|---|
+| `UPSTREAM.md` | pinned upstream commit, provenance, patch policy |
+| `setup.sh` | idempotent clone/checkout + patch + build |
+| `patches/` | overlay patches applied in lexical order |
+| `configs/turing-tu104-2080super/` | TU104/RTX 2080 SUPER scaffold config + Booksim interconnect file |
+| `scripts/run_sim.sh` | LD_LIBRARY_PATH-substitution runner for CUDA binaries |
+| `scripts/config_sanity.py` | stdlib-only config consistency test (also wired into `make test`) |
+| `ADMISSION.md` | draft admission records (pending maintainer filing) |
+| `build/` | build logs and run directories (gitignored) |
diff --git a/debug/risc-v-gsp-rm-vm/gpgpu-sim/UPSTREAM.md b/debug/risc-v-gsp-rm-vm/gpgpu-sim/UPSTREAM.md
new file mode 100644
index 0000000..32e1a09
--- /dev/null
+++ b/debug/risc-v-gsp-rm-vm/gpgpu-sim/UPSTREAM.md
@@ -0,0 +1,71 @@
+# Upstream GPGPU-Sim: pin, provenance and patch policy
+
+> Public-safe reference for the upstream checkout used by this overlay.
+
+## 1. Pinned identity
+
+| Field | Value |
+|---|---|
+| Repository | `https://github.com/gpgpu-sim/gpgpu-sim_distribution` |
+| Commit | `90ec3399763d7c8512cfe7dc193473086c38ca38` |
+| Tag | `v4.0.1` |
+| In-tree version string | `GPGPU-Sim Simulator Version 4.0.0` |
+| License | BSD-3-Clause (Copyright 2009-2011 Tor M. Aamodt and The University of British Columbia) |
+| Checkout location | `gpgpu-sim/upstream/` (gitignored; never committed) |
+
+The pin is reproduced by `setup.sh` (constant `PIN`). A fresh run clones into
+`upstream/`, checks out the pin detached, applies `patches/`, and builds.
+Re-running verifies the checkout is still at the pin and that every patch is
+either applicable or already applied, so the process is idempotent.
+
+## 2. Clone source
+
+`setup.sh` clones from the upstream URL above by default. A local mirror is
+used **only** when `GPGPUSIM_CLONE_SOURCE=<path-or-url>` is set explicitly
+(unset by default); no host path is hardcoded or guessed. The clone source
+never changes the pin: the checkout is always forced to `90ec339`. This is a
+build convenience only; upstream code is never modified in place by hand.
+
+## 3. `origin/dev` (4.2.0): not checked out, tracked as a reference
+
+The repository's `origin/dev` branch is a later line (in-tree version string
+`4.2.0`) that adds, among other things, AccelWattch power modeling, an Ampere
+RTX 3060 (SM86) tested configuration, and Turing L1 fixes. It is **not** the
+pin and is not used by Phase 0. It is recorded here because Phase 1/2 may
+rebase selected changes (architecture config, L1 fixes) into our patch set
+rather than adopting the whole branch. Any such rebase must be an explicit,
+recorded decision with the same reproducibility guarantees (new pin + patch
+set + test evidence), never a silent branch switch.
+
+## 4. How our changes are applied
+
+Design rule: upstream history is never forked. Our work lives entirely in this
+overlay and reaches the upstream tree as **patches**, in this order:
+
+1. `setup.sh` clones/checks out the pin.
+2. `patches/*.patch` are applied in lexical filename order
+   (`0001-*`, `0002-*`, ...) with `git apply`.
+3. For each patch, `git apply --check` decides "apply"; otherwise
+   `git apply --reverse --check` decides "already applied"; if neither holds,
+   `setup.sh` fails loudly rather than guessing.
+4. New source trees that are not clean patches ship as patches too (a patch
+   can add files). There is no separate overlay-copy step in Phase 0.
+
+The first expected patch set is the CUDA-12 build/runtime port (WP0.1), which
+makes `setup_environment`, the PTX path and `cuobjdump` handling work on this
+host's CUDA. Until those patches exist the default build correctly refuses to
+proceed with a clear error.
+
+To reset the checkout: `./setup.sh --reclone` (destroys `upstream/` and clones
+fresh). Nothing outside `gpgpu-sim/upstream/` and `gpgpu-sim/build/` is
+touched by the build; both are gitignored.
+
+## 5. Reproducibility checks
+
+```bash
+git -C gpgpu-sim/upstream rev-parse HEAD          # must equal the pin above
+git -C gpgpu-sim/upstream status --porcelain      # patch-applied delta expected
+sha256sum gpgpu-sim/patches/*.patch               # record with any result
+```
+
+Build logs are written to `gpgpu-sim/build/build.log` and are not committed.
diff --git a/debug/risc-v-gsp-rm-vm/gpgpu-sim/configs/turing-tu104-2080super/config_turing_islip.icnt b/debug/risc-v-gsp-rm-vm/gpgpu-sim/configs/turing-tu104-2080super/config_turing_islip.icnt
new file mode 100644
index 0000000..74a7a30
--- /dev/null
+++ b/debug/risc-v-gsp-rm-vm/gpgpu-sim/configs/turing-tu104-2080super/config_turing_islip.icnt
@@ -0,0 +1,89 @@
+// GreenEye-GPGPU-Sim -- Booksim/ISLIP interconnect scaffold for Turing TU104.
+//
+// Copied from the upstream SM75_RTX2060 `config_turing_islip.icnt` with the
+// topology size adjusted for this config's device scale. Like the upstream
+// Turing config, this file is NOT active by default: `gpgpusim.config` uses
+// `-network_mode 2` (built-in local crossbar). Enable it only together with
+// `-network_mode 1` and `-inter_config_file config_turing_islip.icnt`.
+//
+// Public specs encoded (RTX 2080 SUPER, TU104): 48 SMs, 256-bit GDDR6,
+// 8 GB, ~496 GB/s. This interconnect model is uncalibrated.
+//
+// UNKNOWN: upstream SM75_RTX2060 uses k=52 while this config has 64 modeled
+// nodes (48 SMs + 16 memory modules); the tested configs' k values are not
+// simply SM+memory counts. k=64 is the nearest documented node count for this
+// scaffold and MUST be validated in WP1.3 before any interconnect result is
+// trusted. All other parameters are inherited verbatim from upstream.
+
+//64*1 fly with 32 flits per packet under gpgpusim injection mode
+use_map = 0;
+flit_size = 40;
+
+// currently we do not use this, see subnets below
+network_count = 2;
+
+// Topology
+topology = fly;
+k = 64;
+n = 1;
+
+// Routing
+
+routing_function = dest_tag;
+
+// Flow control
+
+num_vcs     = 1;
+vc_buf_size = 64;
+input_buffer_size = 256;
+ejection_buffer_size = 64;
+boundary_buffer_size = 64;
+
+wait_for_tail_credit = 0;
+
+// Router architecture
+
+vc_allocator = islip; //separable_input_first;
+sw_allocator = islip; //separable_input_first;
+alloc_iters  = 1;
+
+credit_delay   = 0;
+routing_delay  = 0;
+vc_alloc_delay = 1;
+sw_alloc_delay = 1;
+
+input_speedup     = 1;
+output_speedup    = 1;
+internal_speedup  = 2.0;
+
+// Traffic, GPGPU-Sim does not use this
+
+traffic                = uniform;
+packet_size ={{1,2,3,4},{10,20}};
+packet_size_rate={{1,1,1,1},{2,1}};
+
+// Simulation - Don't change
+
+sim_type       = gpgpusim;
+//sim_type = latency;
+injection_rate = 0.1;
+
+subnets = 2;
+
+// Always use read and write no matter following line
+//use_read_write = 1;
+
+
+read_request_subnet = 0;
+read_reply_subnet = 1;
+write_request_subnet = 0;
+write_reply_subnet = 1;
+
+read_request_begin_vc = 0;
+read_request_end_vc = 0;
+write_request_begin_vc = 0;
+write_request_end_vc = 0;
+read_reply_begin_vc = 0;
+read_reply_end_vc = 0;
+write_reply_begin_vc = 0;
+write_reply_end_vc = 0;
diff --git a/debug/risc-v-gsp-rm-vm/gpgpu-sim/configs/turing-tu104-2080super/gpgpusim.config b/debug/risc-v-gsp-rm-vm/gpgpu-sim/configs/turing-tu104-2080super/gpgpusim.config
new file mode 100644
index 0000000..641f8d3
--- /dev/null
+++ b/debug/risc-v-gsp-rm-vm/gpgpu-sim/configs/turing-tu104-2080super/gpgpusim.config
@@ -0,0 +1,233 @@
+# GreenEye-GPGPU-Sim -- Turing TU104 config scaffold: GeForce RTX 2080 SUPER.
+#
+# STATUS: SCAFFOLD, NOT CALIBRATED. This file is derived from the upstream
+# GPGPU-Sim SM75_RTX2060 Turing configuration and adjusts only the
+# device-level parameters that public specifications distinguish between
+# TU106 (RTX 2060) and TU104 (RTX 2080 SUPER). The SM75 microarchitecture
+# parameters (issue widths, latencies, sub-core model, cache geometry) are
+# inherited unchanged. It makes NO accuracy claim; see ../ADMISSION.md
+# (TOOL-2026-SIM-02) and ../README.md.
+#
+# Public specifications encoded below (RTX 2080 SUPER, TU104):
+#   - 48 SMs, 3072 CUDA cores (64 FP32 per SM)            [NVIDIA product specs]
+#   - 8 GB GDDR6, 256-bit bus, 15.5 Gbps effective,
+#     ~496 GB/s (496.1 GB/s nominal)                      [NVIDIA product specs]
+#   - 4 MB L2, 64 ROPs, 192 TMUs                          [TU104 die specs]
+#   - base clock ~1650 MHz, boost clock ~1815 MHz         [NVIDIA product specs]
+#   - compute capability 7.5, PCIe 3.0 x16                [NVIDIA product specs]
+#
+# Machine-checkable spec markers (parsed by ../scripts/config_sanity.py):
+# SPEC: sm_count=48
+# SPEC: memory_type=GDDR6
+# SPEC: vram_gib=8
+# SPEC: bus_width_bits=256
+# SPEC: bandwidth_gbps=496
+# SPEC: l2_mib=4
+# SPEC: rops=64
+# SPEC: tmus=192
+# SPEC: base_mhz=1650
+# SPEC: boost_mhz=1815
+# SPEC: compute_capability=7.5
+#
+# Upstream references for the inherited Turing model:
+# 1- https://www.nvidia.com/content/dam/en-zz/Solutions/design-visualization/technologies/turing-architecture/NVIDIA-Turing-Architecture-Whitepaper.pdf
+# 2- "RTX on--The NVIDIA Turing GPU", IEEE MICRO 2020
+
+# functional simulator specification
+-gpgpu_ptx_instruction_classification 0
+-gpgpu_ptx_sim_mode 0
+-gpgpu_ptx_force_max_capability 75
+
+# Device Limits
+-gpgpu_stack_size_limit 1024
+-gpgpu_heap_size_limit 8388608
+-gpgpu_runtime_sync_depth_limit 2
+-gpgpu_runtime_pending_launch_count_limit 2048
+-gpgpu_kernel_launch_latency 5000
+
+# Compute Capability
+-gpgpu_compute_capability_major 7
+-gpgpu_compute_capability_minor 5
+
+# PTX execution-driven
+-gpgpu_ptx_convert_to_ptxplus 0
+-gpgpu_ptx_save_converted_ptxplus 0
+
+# high level architecture configuration
+# 48 SMs for TU104 (RTX 2080 SUPER). One shader core per cluster.
+-gpgpu_n_clusters 48
+-gpgpu_n_cores_per_cluster 1
+# TU104 has a 256-bit GDDR6 bus. GPGPU-Sim's model counts memory modules;
+# following the upstream Turing convention (SM75_RTX2060 uses 12 modules for
+# a 192-bit bus, i.e. one modeled module per 16-bit channel) we use 16.
+# -gpgpu_n_mem x -gpgpu_n_sub_partition_per_mchannel x 128 KiB = L2 size
+# and buswidth x freq_ratio x DRAM clock x n_mem = modeled bandwidth, both
+# checked by config_sanity.py.
+-gpgpu_n_mem 16
+-gpgpu_n_sub_partition_per_mchannel 2
+
+# Turing TU104 clock domains
+# <Core Clock>:<Interconnect Clock>:<L2 Clock>:<DRAM Clock> (MHz)
+# Base clock 1650 MHz; DRAM command clock = 15.5 Gbps / 4 (GDDR6 QDR).
+-gpgpu_clock_domains 1650.0:1650.0:1650.0:3875.0
+# boost mode (published boost clock 1815 MHz)
+# -gpgpu_clock_domains 1815.0:1815.0:1815.0:3875.0
+
+# shader core pipeline config
+-gpgpu_shader_registers 65536
+-gpgpu_registers_per_block 65536
+-gpgpu_occupancy_sm_number 75
+
+# This implies a maximum of 32 warps/SM
+-gpgpu_shader_core_pipeline 1024:32
+-gpgpu_shader_cta 32
+-gpgpu_simd_model 1
+
+# Pipeline widths and number of FUs
+# ID_OC_SP,ID_OC_DP,ID_OC_INT,ID_OC_SFU,ID_OC_MEM,OC_EX_SP,OC_EX_DP,OC_EX_INT,OC_EX_SFU,OC_EX_MEM,EX_WB,ID_OC_TENSOR_CORE,OC_EX_TENSOR_CORE
+## Turing has 4 SP SIMD units, 4 INT units, 4 SFU units, 8 Tensor core units
+## We need to scale the number of pipeline registers to be equal to the number of SP units
+-gpgpu_pipeline_widths 4,0,4,4,4,4,0,4,4,4,8,4,4
+-gpgpu_num_sp_units 4
+-gpgpu_num_sfu_units 4
+-gpgpu_num_int_units 4
+-gpgpu_tensor_core_avail 1
+-gpgpu_num_tensor_core_units 4
+
+# Instruction latencies and initiation intervals
+# "ADD,MAX,MUL,MAD,DIV"
+# All Div operations are executed on SFU unit
+-ptx_opcode_latency_int 4,13,4,5,145,32
+-ptx_opcode_initiation_int 2,2,2,2,8,4
+-ptx_opcode_latency_fp 4,13,4,5,39
+-ptx_opcode_initiation_fp 2,2,2,2,4
+-ptx_opcode_latency_dp 8,19,8,8,330
+-ptx_opcode_initiation_dp 4,4,4,4,130
+-ptx_opcode_latency_sfu 100
+-ptx_opcode_initiation_sfu 8
+-ptx_opcode_latency_tesnor 64
+-ptx_opcode_initiation_tensor 64
+# UNKNOWN: instruction latencies/initiation intervals are inherited from the
+# upstream SM75_RTX2060 model; they are not measured on TU104. Calibration
+# is Phase-1 work (WP1.1/WP1.4, admission TOOL-2026-SIM-02).
+
+# Turing has four schedulers per core
+-gpgpu_num_sched_per_core 4
+# Greedy then oldest scheduler
+-gpgpu_scheduler gto
+## In Turing, a warp scheduler can issue 1 inst per cycle
+-gpgpu_max_insn_issue_per_warp 1
+-gpgpu_dual_issue_diff_exec_units 1
+
+# shared memory bankconflict detection
+-gpgpu_shmem_num_banks 32
+-gpgpu_shmem_limited_broadcast 0
+-gpgpu_shmem_warp_parts 1
+-gpgpu_coalesce_arch 75
+
+# Turing has sub core model, in which each scheduler has its own register file and EUs
+# i.e. schedulers are isolated
+-gpgpu_sub_core_model 1
+# disable specialized operand collectors and use generic operand collectors instead
+-gpgpu_enable_specialized_operand_collector 0
+-gpgpu_operand_collector_num_units_gen 8
+-gpgpu_operand_collector_num_in_ports_gen 8
+-gpgpu_operand_collector_num_out_ports_gen 8
+# turing has 8 banks dual-port, 4 schedulers, two banks per scheduler
+# we increase #banks to 16 to mitigate the effect of Register File Cache (RFC) which we do not implement in the current version
+-gpgpu_num_reg_banks 16
+-gpgpu_reg_file_port_throughput 2
+
+# <nsets>:<bsize>:<assoc>,<rep>:<wr>:<alloc>:<wr_alloc>:<set_index_fn>,<mshr>:<N>:<merge>,<mq>:**<fifo_entry>
+# ** Optional parameter - Required when mshr_type==Texture Fifo
+-gpgpu_adaptive_cache_config 0
+-gpgpu_l1_banks 4
+-gpgpu_cache:dl1  S:1:128:512,L:L:s:N:L,A:256:8,16:0,32
+-gpgpu_shmem_size 65536
+-gpgpu_shmem_sizeDefault 65536
+-gpgpu_shmem_per_block 65536
+-gpgpu_gmem_skip_L1D 0
+-gpgpu_n_cluster_ejection_buffer_size 32
+-gpgpu_l1_latency 20
+-gpgpu_smem_latency 20
+-gpgpu_flush_l1_cache 1
+# UNKNOWN: L1/shmem latencies are inherited from upstream; not TU104-measured.
+
+# 64 sets, each 128 bytes 16-way for each memory sub partition (128 KB per
+# memory sub partition). With 16 memory modules and 2 sub-partitions each this
+# gives us 4 MB L2 cache, matching the published TU104 L2 size.
+-gpgpu_cache:dl2 S:64:128:16,L:B:m:L:P,A:192:4,32:0,32
+-gpgpu_cache:dl2_texture_only 0
+-gpgpu_dram_partition_queues 64:64:64:64
+-gpgpu_perf_sim_memcpy 1
+-gpgpu_memory_partition_indexing 0
+
+# 128 KB Inst.
+-gpgpu_cache:il1 N:64:128:16,L:R:f:N:L,S:2:48,4
+-gpgpu_inst_fetch_throughput 4
+# 128 KB Tex
+# Note, TEX is deprecated in Volta, It is used for legacy apps only. Use L1D cache instead with .nc modifier or __ldg method
+-gpgpu_tex_cache:l1 N:4:128:256,L:R:m:N:L,T:512:8,128:2
+# 64 KB Const
+-gpgpu_const_cache:l1 N:128:64:8,L:R:f:N:L,S:2:64,4
+-gpgpu_perfect_inst_const_cache 1
+
+# interconnection
+#-network_mode 1
+#-inter_config_file config_turing_islip.icnt
+# use built-in local xbar
+-network_mode 2
+-icnt_in_buffer_limit 512
+-icnt_out_buffer_limit 512
+-icnt_subnets 2
+-icnt_arbiter_algo 1
+-icnt_flit_size 40
+
+# memory partition latency config
+-gpgpu_l2_rop_latency 160
+-dram_latency 100
+# UNKNOWN: L2/ROP and DRAM latencies are inherited from upstream SM75_RTX2060
+# and not calibrated to TU104 GDDR6. WP1.3/WP1.4 work.
+
+# dram model config
+-gpgpu_dram_scheduler 1
+-gpgpu_frfcfs_dram_sched_queue_size 64
+-gpgpu_dram_return_queue_size 192
+
+# Turing has GDDR6
+# http://monitorinsider.com/GDDR6.html
+-gpgpu_n_mem_per_ctrlr 1
+-gpgpu_dram_buswidth 2
+-gpgpu_dram_burst_length 16
+-dram_data_command_freq_ratio 4
+-gpgpu_mem_address_mask 1
+# UNKNOWN: the DRAM address mapping string is inherited from the 12-module
+# RTX 2060 config (4 channel bits). It has not been re-derived for the
+# 16-module TU104 layout; recalibrate in WP1.3 before trusting bandwidth.
+-gpgpu_mem_addr_mapping dramid@8;00000000.00000000.00000000.00000000.0000RRRR.RRRRRRRR.RBBBCCCC.BCCSSSSS
+
+# Use the same GDDR5-derived timing set as upstream, scaled to 3875 MHz
+# (15.5 Gbps / 4). UNKNOWN: true TU104 GDDR6 timings are not calibrated.
+-gpgpu_dram_timing_opt "nbk=16:CCD=4:RRD=10:RCD=20:RAS=50:RP=20:RC=62:
+                        CL=20:WL=8:CDLR=9:WR=20:nbkgrp=4:CCDL=4:RTPL=4"
+
+# select lower bits for bnkgrp to increase bnkgrp parallelism
+-dram_bnk_indexing_policy 0
+-dram_bnkgrp_indexing_policy 1
+
+#-dram_seperate_write_queue_enable 1
+#-dram_write_queue_size 64:56:32
+
+# stat collection
+-gpgpu_memlatency_stat 14
+-gpgpu_runtime_stat 500
+-enable_ptx_file_line_stats 1
+-visualizer_enabled 0
+
+# power model configs, disable until we create a real energy model for Turing
+-power_simulation_enabled 0
+
+# tracing functionality
+#-trace_enabled 1
+#-trace_components WARP_SCHEDULER,SCOREBOARD
+#-trace_sampling_core 0
diff --git a/debug/risc-v-gsp-rm-vm/gpgpu-sim/patches/.gitkeep b/debug/risc-v-gsp-rm-vm/gpgpu-sim/patches/.gitkeep
new file mode 100644
index 0000000..c93f122
--- /dev/null
+++ b/debug/risc-v-gsp-rm-vm/gpgpu-sim/patches/.gitkeep
@@ -0,0 +1,2 @@
+# Overlay patches are applied by ../setup.sh in lexical filename order.
+# Empty is valid in Phase 0; the CUDA 12 build/runtime port (WP0.1) lands here.
diff --git a/debug/risc-v-gsp-rm-vm/gpgpu-sim/patches/0001-cuda12-toolchain-and-soname.patch b/debug/risc-v-gsp-rm-vm/gpgpu-sim/patches/0001-cuda12-toolchain-and-soname.patch
new file mode 100644
index 0000000..af1fb31
--- /dev/null
+++ b/debug/risc-v-gsp-rm-vm/gpgpu-sim/patches/0001-cuda12-toolchain-and-soname.patch
@@ -0,0 +1,56 @@
+diff --git a/Makefile b/Makefile
+index d248211..1ab7c0c 100644
+--- a/Makefile
++++ b/Makefile
+@@ -168,6 +168,8 @@ $(SIM_LIB_DIR)/libcudart.so: makedirs $(LIBS) cudalib
+ 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.10.0 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.10.0; fi
+ 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.10.1 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.10.1; fi
+ 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.11.0 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.11.0; fi
++	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.12 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.12; fi
++	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.12.0 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.12.0; fi
+ 
+ $(SIM_LIB_DIR)/libcudart.dylib: makedirs $(LIBS) cudalib
+ 	g++ -dynamiclib -Wl,-headerpad_max_install_names,-undefined,dynamic_lookup,-compatibility_version,1.1,-current_version,1.1\
+diff --git a/linux-so-version.txt b/linux-so-version.txt
+index 3db07ac..55d3727 100644
+--- a/linux-so-version.txt
++++ b/linux-so-version.txt
+@@ -10,5 +10,7 @@ libcudart.so.10.1{
+ };
+ libcudart.so.11.0{
+ };
++libcudart.so.12{
++};
+ libcuda.so.1{
+ };
+diff --git a/setup_environment b/setup_environment
+index 07d0788..cd054ea 100644
+--- a/setup_environment
++++ b/setup_environment
+@@ -44,11 +44,11 @@ if [ $? = 1 ]; then
+ 	return;
+ fi
+ 
+-CC_VERSION=`gcc --version | head -1 | awk '{for(i=1;i<=NF;i++){ if(match($i,/^[0-9]\.[0-9]\.[0-9]$/))  {print $i; exit 0}}}'`
++CC_VERSION=`gcc --version | head -1 | awk '{for(i=1;i<=NF;i++){ if(match($i,/^[0-9]+\.[0-9]+\.[0-9]+$/))  {print $i; exit 0}}}'`
+ 
+ CUDA_VERSION_STRING=`$CUDA_INSTALL_PATH/bin/nvcc --version | awk '/release/ {print $5;}' | sed 's/,//'`;
+ export CUDA_VERSION_NUMBER=`echo $CUDA_VERSION_STRING | sed 's/\./ /' | awk '{printf("%02u%02u", 10*int($1), 10*$2);}'`
+-if [ $CUDA_VERSION_NUMBER -gt 11100 -o $CUDA_VERSION_NUMBER -lt 2030  ]; then
++if [ $CUDA_VERSION_NUMBER -gt 12999 -o $CUDA_VERSION_NUMBER -lt 2030  ]; then
+ 	echo "ERROR ** GPGPU-Sim version $GPGPUSIM_VERSION_STRING not tested with CUDA version $CUDA_VERSION_STRING (please see README)";
+ 	echo $CUDA_VERSION_NUMBER
+   return
+diff --git a/version_detection.mk b/version_detection.mk
+index ee71a62..693e913 100644
+--- a/version_detection.mk
++++ b/version_detection.mk
+@@ -43,7 +43,7 @@ CUDA_VERSION_STRING:=$(shell $(CUDA_INSTALL_PATH)/bin/nvcc --version | awk '/rel
+ CUDART_VERSION:=$(shell echo $(CUDA_VERSION_STRING) | sed 's/\./ /' | awk '{printf("%02u%02u", 10*int($$1), 10*$$2);}')
+ 
+ # Detect GCC Version 
+-CC_VERSION := $(shell gcc --version | head -1 | awk '{for(i=1;i<=NF;i++){ if(match($$i,/^[0-9]\.[0-9]\.[0-9]$$/))  {print $$i; exit 0 }}}')
++CC_VERSION := $(shell gcc --version | head -1 | awk '{for(i=1;i<=NF;i++){ if(match($$i,/^[0-9]+\.[0-9]+\.[0-9]+$$/))  {print $$i; exit 0 }}}')
+ 
+ # Detect Support for C++11 (C++0x) from GCC Version 
+ GNUC_CPP0X := $(shell gcc --version | perl -ne 'if (/gcc\s+\(.*\)\s+([0-9.]+)/){ if($$1 >= 4.3) {$$n=1} else {$$n=0;} } END { print $$n; }')
diff --git a/debug/risc-v-gsp-rm-vm/gpgpu-sim/patches/0002-gcc13-cuda12-source-fixes.patch b/debug/risc-v-gsp-rm-vm/gpgpu-sim/patches/0002-gcc13-cuda12-source-fixes.patch
new file mode 100644
index 0000000..a9367f8
--- /dev/null
+++ b/debug/risc-v-gsp-rm-vm/gpgpu-sim/patches/0002-gcc13-cuda12-source-fixes.patch
@@ -0,0 +1,77 @@
+diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
+index fd05f55..c6475c1 100644
+--- a/libcuda/cuda_runtime_api.cc
++++ b/libcuda/cuda_runtime_api.cc
+@@ -3595,7 +3595,10 @@ unsigned CUDARTAPI __cudaPushCallConfiguration(dim3 gridDim, dim3 blockDim,
+   if (g_debug_execution >= 3) {
+     announce_call(__my_func__);
+   }
+-  cudaConfigureCallInternal(gridDim, blockDim, sharedMem, stream);
++  // Must return: falling off the end of this non-void function is undefined
++  // behaviour (gcc 13 -O3 miscompiles it into a non-returning loop).
++  return (unsigned)cudaConfigureCallInternal(gridDim, blockDim, sharedMem,
++                                             stream);
+ }
+ 
+ cudaError_t CUDARTAPI __cudaPopCallConfiguration(dim3 *gridDim, dim3 *blockDim,
+diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h
+index 49f3e9f..722dd15 100644
+--- a/src/abstract_hardware_model.h
++++ b/src/abstract_hardware_model.h
+@@ -486,6 +486,18 @@ struct cudaArray {
+ 
+ // Struct that record other attributes in the textureReference declaration
+ // - These attributes are passed thru __cudaRegisterTexture()
++#if defined(CUDART_VERSION) && CUDART_VERSION >= 12000
++// CUDA 12 removed the legacy cudaBindTexture/textureReference runtime API.
++// Provide a placeholder layout for the simulator's legacy texture path so
++// that host-side emulation code still compiles; CUDA 12 apps cannot call the
++// removed API, so this path is unreachable for them.
++struct textureReference {
++  int normalized;
++  int filterMode;
++  int addressMode[3];
++};
++#endif
++
+ struct textureReferenceAttr {
+   const struct textureReference *m_texref;
+   int m_dim;
+diff --git a/src/cuda-sim/Makefile b/src/cuda-sim/Makefile
+index 85d1c8c..1a7454b 100644
+--- a/src/cuda-sim/Makefile
++++ b/src/cuda-sim/Makefile
+@@ -129,9 +129,9 @@ $(OUTPUT_DIR)/instructions.h: instructions.cc
+ 
+ $(OUTPUT_DIR)/ptx_parser_decode.def: $(OUTPUT_DIR)/ptx.tab.c
+ ifeq ($(shell uname),Linux)
+-	cat $(OUTPUT_DIR)/ptx.tab.h | grep "=" | sed 's/^[ ]\+//' | sed 's/[=,]//g' | sed 's/\([_A-Z1-9]\+\)[ ]\+\([0-9]\+\)/\1 \1/' | sed 's/^/DEF(/' | sed 's/ /,"/' | sed 's/$$/")/' > $(OUTPUT_DIR)/ptx_parser_decode.def
++	cat $(OUTPUT_DIR)/ptx.tab.h | grep "=" | sed 's/^[ ]\+//' | sed 's/[=,]//g' | sed 's/"/\\"/g' | sed 's/\([_A-Z1-9]\+\)[ ]\+\([0-9]\+\)/\1 \1/' | sed 's/^/DEF(/' | sed 's/ /,"/' | sed 's/$$/")/' > $(OUTPUT_DIR)/ptx_parser_decode.def
+ else
+-	cat $(OUTPUT_DIR)/ptx.tab.h | grep "=" | sed -E 's/^ +//' | sed 's/[=,]//g' | sed -E 's/([_A-Z1-9]+).*/\1 \1/' | sed 's/^/DEF(/' | sed 's/ /,"/' | sed 's/$$/")/' > $(OUTPUT_DIR)/ptx_parser_decode.def
++	cat $(OUTPUT_DIR)/ptx.tab.h | grep "=" | sed -E 's/^ +//' | sed 's/[=,]//g' | sed 's/"/\\"/g' | sed -E 's/([_A-Z1-9]+).*/\1 \1/' | sed 's/^/DEF(/' | sed 's/ /,"/' | sed 's/$$/")/' > $(OUTPUT_DIR)/ptx_parser_decode.def
+ endif
+ 
+ $(OUTPUT_DIR)/instructions.o: $(OUTPUT_DIR)/instructions.h $(OUTPUT_DIR)/ptx.tab.c
+diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
+index 71f0703..084a889 100644
+--- a/src/cuda-sim/cuda-sim.cc
++++ b/src/cuda-sim/cuda-sim.cc
+@@ -128,7 +128,9 @@ void gpgpu_t::gpgpu_ptx_sim_bindNameToTexture(
+   } else {
+     const struct textureReference *tr = *m_NameToTextureRef[texname].begin();
+     assert(tr != NULL);
++#if CUDART_VERSION < 12000
+     // asserts that all texrefs in set have same fields
++    // (textureReference is opaque from CUDA 12 onward)
+     assert(tr->normalized == texref->normalized &&
+            tr->filterMode == texref->filterMode &&
+            tr->addressMode[0] == texref->addressMode[0] &&
+@@ -139,6 +141,7 @@ void gpgpu_t::gpgpu_ptx_sim_bindNameToTexture(
+            tr->channelDesc.z == texref->channelDesc.z &&
+            tr->channelDesc.w == texref->channelDesc.w &&
+            tr->channelDesc.f == texref->channelDesc.f);
++#endif
+   }
+   m_NameToTextureRef[texname].insert(texref);
+   m_TextureRefToName[texref] = texname;
diff --git a/debug/risc-v-gsp-rm-vm/gpgpu-sim/scripts/config_sanity.py b/debug/risc-v-gsp-rm-vm/gpgpu-sim/scripts/config_sanity.py
new file mode 100755
index 0000000..052695e
--- /dev/null
+++ b/debug/risc-v-gsp-rm-vm/gpgpu-sim/scripts/config_sanity.py
@@ -0,0 +1,278 @@
+#!/usr/bin/env python3
+"""config_sanity.py -- stdlib-only sanity check for the TU104 scaffold config.
+
+Parses gpgpu-sim/configs/turing-tu104-2080super/gpgpusim.config and asserts
+that the key device knobs and the machine-checkable `# SPEC:` markers are
+present and internally consistent:
+
+  * SM count           : n_clusters x n_cores_per_cluster == 48
+  * compute capability : major.minor == 7.5, PTX force max capability == 75
+  * memory type        : GDDR6 (burst length 16, QDR freq ratio 4)
+  * memory capacity    : SPEC vram_gib == 8 (documentation-only; GPGPU-Sim
+                         4.x has no VRAM capacity knob)
+  * L2 size            : sets x bsize x assoc x n_mem x n_sub == 4 MiB
+  * bandwidth          : buswidth x freq_ratio x DRAM clock x n_mem == 496 GB/s
+  * clocks             : base 1650 MHz, DRAM 3875 MHz (15.5 Gbps / 4)
+
+This is a scaffold consistency test, not a calibration: passing says the file
+internally encodes the published TU104 numbers, NOT that the simulator models
+TU104 accurately. Exit code 0 = all checks pass; 1 = at least one mismatch;
+2 = usage/parse error.
+
+Usage:
+    python3 gpgpu-sim/scripts/config_sanity.py [--config PATH] [--quiet]
+"""
+
+from __future__ import annotations
+
+import argparse
+import re
+import sys
+from pathlib import Path
+
+EXPECTED_SPECS = {
+    "sm_count": "48",
+    "memory_type": "GDDR6",
+    "vram_gib": "8",
+    "bus_width_bits": "256",
+    "bandwidth_gbps": "496",
+    "l2_mib": "4",
+    "rops": "64",
+    "tmus": "192",
+    "base_mhz": "1650",
+    "boost_mhz": "1815",
+    "compute_capability": "7.5",
+}
+
+DEFAULT_CONFIG = (
+    Path(__file__).resolve().parents[1]
+    / "configs"
+    / "turing-tu104-2080super"
+    / "gpgpusim.config"
+)
+ICNT_NAME = "config_turing_islip.icnt"
+
+
+def parse_config(path: Path) -> tuple[dict[str, str], dict[str, str]]:
+    """Return (options, spec_markers) parsed from a gpgpusim.config file."""
+    opts: dict[str, str] = {}
+    specs: dict[str, str] = {}
+    for raw in path.read_text(encoding="utf-8").splitlines():
+        spec = re.match(r"^\s*#\s*SPEC:\s*([A-Za-z0-9_]+)\s*=\s*(\S+)\s*$", raw)
+        if spec:
+            specs[spec.group(1)] = spec.group(2)
+        line = raw.split("#", 1)[0].strip()
+        if not line or not line.startswith("-"):
+            continue
+        parts = line.split(None, 1)
+        opts[parts[0]] = parts[1].strip() if len(parts) > 1 else ""
+    return opts, specs
+
+
+class Checker:
+    def __init__(self, quiet: bool = False) -> None:
+        self.quiet = quiet
+        self.passed = 0
+        self.failures: list[str] = []
+
+    def check(self, name: str, condition: bool, detail: str) -> None:
+        if condition:
+            self.passed += 1
+            if not self.quiet:
+                print(f"PASS: {name}: {detail}")
+        else:
+            self.failures.append(f"{name}: {detail}")
+            print(f"FAIL: {name}: {detail}")
+
+    def note(self, text: str) -> None:
+        if not self.quiet:
+            print(f"NOTE: {text}")
+
+
+def parse_int(opts: dict[str, str], key: str) -> int | None:
+    try:
+        return int(opts[key], 0)
+    except (KeyError, ValueError):
+        return None
+
+
+def parse_l2(opts: dict[str, str]) -> tuple[int, int, int] | None:
+    """Parse '<nsets>:<bsize>:<assoc>' from -gpgpu_cache:dl2 S:... ."""
+    m = re.match(r"^S:(\d+):(\d+):(\d+)", opts.get("-gpgpu_cache:dl2", ""))
+    if not m:
+        return None
+    return int(m.group(1)), int(m.group(2)), int(m.group(3))
+
+
+def run_checks(config_path: Path, quiet: bool = False) -> int:
+    if not config_path.is_file():
+        print(f"config_sanity: ERROR: config not found: {config_path}", file=sys.stderr)
+        return 2
+
+    opts, specs = parse_config(config_path)
+    c = Checker(quiet=quiet)
+
+    # 1. spec markers ------------------------------------------------------
+    for key, expected in EXPECTED_SPECS.items():
+        actual = specs.get(key)
+        c.check(
+            f"SPEC {key}",
+            actual == expected,
+            f"expected {expected}, got {actual!r}",
+        )
+
+    # 2. SMs / compute capability -----------------------------------------
+    clusters = parse_int(opts, "-gpgpu_n_clusters")
+    cores = parse_int(opts, "-gpgpu_n_cores_per_cluster")
+    sm_count = None if clusters is None or cores is None else clusters * cores
+    c.check(
+        "SM count",
+        sm_count == int(EXPECTED_SPECS["sm_count"]),
+        f"n_clusters({clusters}) x n_cores_per_cluster({cores}) = {sm_count},"
+        f" expected {EXPECTED_SPECS['sm_count']}",
+    )
+    cc = f"{opts.get('-gpgpu_compute_capability_major', '?')}." \
+         f"{opts.get('-gpgpu_compute_capability_minor', '?')}"
+    c.check(
+        "compute capability",
+        cc == EXPECTED_SPECS["compute_capability"],
+        f"{cc}, expected {EXPECTED_SPECS['compute_capability']}",
+    )
+    c.check(
+        "PTX force max capability",
+        parse_int(opts, "-gpgpu_ptx_force_max_capability") == 75,
+        f"{opts.get('-gpgpu_ptx_force_max_capability', '(missing)')}, expected 75 (sm_75)",
+    )
+
+    # 3. clocks / GDDR6 / memory geometry ---------------------------------
+    clocks = opts.get("-gpgpu_clock_domains", "").split(":")
+    dram_mhz = None
+    if len(clocks) == 4:
+        try:
+            dram_mhz = float(clocks[3])
+        except ValueError:
+            dram_mhz = None
+    c.check(
+        "DRAM clock",
+        dram_mhz == 3875.0,
+        f"{clocks[3] if len(clocks) == 4 else clocks}, expected 3875.0 MHz (15.5 Gbps / 4)",
+    )
+    c.check(
+        "base core clock",
+        len(clocks) == 4 and clocks[0] == "1650.0",
+        f"{clocks[0] if clocks else '(missing)'}, expected 1650.0 MHz",
+    )
+    c.check(
+        "memory type GDDR6",
+        specs.get("memory_type") == "GDDR6"
+        and parse_int(opts, "-gpgpu_dram_burst_length") == 16
+        and parse_int(opts, "-dram_data_command_freq_ratio") == 4,
+        "burst_length="
+        f"{opts.get('-gpgpu_dram_burst_length', '(missing)')}, "
+        "freq_ratio="
+        f"{opts.get('-dram_data_command_freq_ratio', '(missing)')}"
+        " (GDDR6 expected: 16 / 4)",
+    )
+    c.check(
+        "memory chips per controller",
+        parse_int(opts, "-gpgpu_n_mem_per_ctrlr") == 1,
+        f"{opts.get('-gpgpu_n_mem_per_ctrlr', '(missing)')}, expected 1 (upstream Turing convention)",
+    )
+    c.check(
+        "VRAM capacity marker (documentation-only)",
+        specs.get("vram_gib") == "8",
+        "SPEC vram_gib="
+        f"{specs.get('vram_gib')!r}; GPGPU-Sim 4.x has no VRAM capacity knob,"
+        " so 8 GiB is recorded as a spec marker only",
+    )
+
+    # 4. L2 size consistency ----------------------------------------------
+    n_mem = parse_int(opts, "-gpgpu_n_mem")
+    n_sub = parse_int(opts, "-gpgpu_n_sub_partition_per_mchannel")
+    l2 = parse_l2(opts)
+    l2_mib = None
+    if l2 and n_mem and n_sub:
+        sets, bsize, assoc = l2
+        l2_mib = sets * bsize * assoc * n_mem * n_sub / (1024 * 1024)
+    spec_l2 = float(EXPECTED_SPECS["l2_mib"])
+    c.check(
+        "L2 size",
+        l2_mib == spec_l2,
+        f"dl2({l2}) x n_mem({n_mem}) x n_sub({n_sub}) = {l2_mib} MiB,"
+        f" expected {spec_l2} MiB",
+    )
+
+    # 5. bandwidth consistency --------------------------------------------
+    buswidth = parse_int(opts, "-gpgpu_dram_buswidth")
+    ratio = parse_int(opts, "-dram_data_command_freq_ratio")
+    bw_gbps = None
+    if None not in (buswidth, ratio, dram_mhz, n_mem):
+        bw_gbps = buswidth * ratio * dram_mhz * 1e6 * n_mem / 1e9
+    spec_bw = float(EXPECTED_SPECS["bandwidth_gbps"])
+    c.check(
+        "modeled bandwidth",
+        bw_gbps == spec_bw,
+        f"buswidth({buswidth} B) x ratio({ratio}) x DRAM({dram_mhz} MHz)"
+        f" x n_mem({n_mem}) = {bw_gbps} GB/s, expected {spec_bw} GB/s",
+    )
+    bw_bits_ok = False
+    if n_mem and buswidth:
+        bw_bits_ok = (
+            specs.get("bus_width_bits") == "256"
+            and int(EXPECTED_SPECS["bus_width_bits"]) // (8 * n_mem) == buswidth
+        )
+    c.check(
+        "bus width marker",
+        bw_bits_ok,
+        f"SPEC bus_width_bits={specs.get('bus_width_bits')!r}, modeled"
+        f" buswidth={buswidth} B x n_mem={n_mem} (expected {EXPECTED_SPECS['bus_width_bits']} bits"
+        f" = {buswidth} B x {n_mem} modules)",
+    )
+
+    # 6. informational markers not modeled by GPGPU-Sim 4.x ---------------
+    for key, label in (("rops", "ROPs"), ("tmus", "TMUs"), ("boost_mhz", "boost clock")):
+        c.check(
+            f"SPEC {key} present",
+            specs.get(key) == EXPECTED_SPECS[key],
+            f"{label}: {specs.get(key)!r} (expected {EXPECTED_SPECS[key]});"
+            " recorded for documentation, not a simulator knob",
+        )
+
+    # 7. sibling interconnect file (not used with -network_mode 2) -------
+    icnt = config_path.parent / ICNT_NAME
+    c.check(
+        "interconnect config present",
+        icnt.is_file(),
+        f"{icnt}",
+    )
+
+    total = c.passed + len(c.failures)
+    print(
+        f"config_sanity: {c.passed}/{total} checks passed"
+        + (f", {len(c.failures)} failed" if c.failures else "")
+    )
+    if c.failures:
+        for failure in c.failures:
+            print(f"config_sanity: FAILED -> {failure}", file=sys.stderr)
+        return 1
+    print("config_sanity: OK (scaffold consistency only; no accuracy claim)")
+    return 0
+
+
+def main(argv: list[str] | None = None) -> int:
+    parser = argparse.ArgumentParser(
+        description="Sanity-check the GreenEye-GPGPU-Sim TU104 scaffold config."
+    )
+    parser.add_argument(
+        "--config",
+        type=Path,
+        default=DEFAULT_CONFIG,
+        help=f"path to gpgpusim.config (default: {DEFAULT_CONFIG})",
+    )
+    parser.add_argument("--quiet", action="store_true", help="only print failures and summary")
+    args = parser.parse_args(argv)
+    return run_checks(args.config, quiet=args.quiet)
+
+
+if __name__ == "__main__":
+    sys.exit(main())
diff --git a/debug/risc-v-gsp-rm-vm/gpgpu-sim/scripts/run_sim.sh b/debug/risc-v-gsp-rm-vm/gpgpu-sim/scripts/run_sim.sh
new file mode 100755
index 0000000..85dd11e
--- /dev/null
+++ b/debug/risc-v-gsp-rm-vm/gpgpu-sim/scripts/run_sim.sh
@@ -0,0 +1,109 @@
+#!/usr/bin/env bash
+# run_sim.sh -- run a CUDA binary on the GreenEye-GPGPU-Sim TU104 scaffold.
+#
+# Substitutes the CUDA runtime with GPGPU-Sim's libcudart.so via
+# LD_LIBRARY_PATH (the mechanism documented in the upstream README, "Step 3:
+# Run"), points the simulator at the tracked TU104 scaffold config, and
+# executes the binary from a dedicated run directory so simulator output does
+# not pollute the source tree.
+#
+# Development-only tooling: it simulates; it never touches GPU hardware,
+# kernel modules, mdev, or a guest.
+#
+# Usage:
+#   ./run_sim.sh [--workdir DIR] [--build-type release|debug] <cuda-binary> [args...]
+#
+# Prerequisites: ./setup.sh has completed a build, and the binary was linked
+# with `--cudart shared` so it resolves libcudart.so dynamically.
+#
+# Outputs (in the work directory, default gpgpu-sim/build/run/):
+#   gpgpusim.config -> symlink to the tracked TU104 config
+#   config_turing_islip.icnt -> symlink to the tracked interconnect config
+#   gpgpusim_out.txt, gpgpu_inst_stats.txt, ... (upstream simulator outputs)
+set -euo pipefail
+
+HERE="$(cd "$(dirname "${BASH_SOURCE[0]}")/.." && pwd)"
+UPSTREAM="$HERE/upstream"
+DEFAULT_CONFIG_DIR="$HERE/configs/turing-tu104-2080super"
+CONFIG_NAME="gpgpusim.config"
+ICNT_NAME="config_turing_islip.icnt"
+BUILD_TYPE="${GPGPUSIM_BUILD_TYPE:-release}"
+WORKDIR="${GPGPUSIM_RUN_DIR:-$HERE/build/run}"
+CONFIG_DIR="$DEFAULT_CONFIG_DIR"
+
+die() { echo "run_sim.sh: ERROR: $*" >&2; exit 1; }
+
+usage() {
+    cat <<'EOF'
+Usage: run_sim.sh [--workdir DIR] [--build-type release|debug] \
+                  [--config-dir DIR] <cuda-binary> [args...]
+
+  --workdir DIR      directory the binary runs in (default gpgpu-sim/build/run)
+  --build-type TYPE  release (default) or debug
+  --config-dir DIR   config directory to use (default configs/turing-tu104-2080super)
+  -h, --help         this help
+
+The binary must be dynamically linked against libcudart (link with
+"--cudart shared"); LD_LIBRARY_PATH is replaced with the GPGPU-Sim runtime
+so no recompilation for the simulator is needed.
+EOF
+}
+
+while [ $# -gt 0 ]; do
+    case "$1" in
+        --workdir)     WORKDIR="${2:-}"; shift ;;
+        --build-type)  BUILD_TYPE="${2:-}"; shift ;;
+        --config-dir)  CONFIG_DIR="${2:-}"; shift ;;
+        -h|--help)     usage; exit 0 ;;
+        --)            shift; break ;;
+        -*)            usage >&2; exit 2 ;;
+        *)             break ;;
+    esac
+    shift
+done
+
+[ $# -ge 1 ] || { usage >&2; exit 2; }
+[ -n "$WORKDIR" ] || die "--workdir requires a value"
+[ -n "$CONFIG_DIR" ] || die "--config-dir requires a value"
+case "$BUILD_TYPE" in release|debug) ;; *) die "unknown build type '$BUILD_TYPE'" ;; esac
+
+BIN="$1"; shift
+
+case "$BIN" in
+    */*)
+        [ -e "$BIN" ] || die "binary not found: $BIN"
+        BIN="$(cd "$(dirname "$BIN")" && pwd)/$(basename "$BIN")"
+        ;;
+    *)
+        resolved="$(command -v "$BIN" || true)"
+        [ -n "$resolved" ] || die "binary '$BIN' not found in PATH"
+        BIN="$resolved"
+        ;;
+esac
+
+[ -d "$CONFIG_DIR" ] || die "config dir not found: $CONFIG_DIR"
+[ -f "$CONFIG_DIR/$CONFIG_NAME" ] || die "missing $CONFIG_DIR/$CONFIG_NAME"
+
+LIBDIR="$(find "$UPSTREAM/lib" -maxdepth 3 -type d -name "$BUILD_TYPE" 2>/dev/null | sort | head -n 1 || true)"
+[ -n "$LIBDIR" ] || die "no GPGPU-Sim runtime library dir under $UPSTREAM/lib; run ./setup.sh first"
+if ! ls "$LIBDIR"/libcudart.so* >/dev/null 2>&1; then
+    echo "run_sim.sh: WARNING: no libcudart.so in $LIBDIR; the binary may load the system CUDA runtime" >&2
+fi
+
+mkdir -p "$WORKDIR"
+ln -sfn "$CONFIG_DIR/$CONFIG_NAME" "$WORKDIR/$CONFIG_NAME"
+if [ -f "$CONFIG_DIR/$ICNT_NAME" ]; then
+    ln -sfn "$CONFIG_DIR/$ICNT_NAME" "$WORKDIR/$ICNT_NAME"
+fi
+
+export GPGPUSIM_ROOT="$UPSTREAM"
+export LD_LIBRARY_PATH="$LIBDIR${LD_LIBRARY_PATH:+:$LD_LIBRARY_PATH}"
+export CUDA_INSTALL_PATH="${CUDA_INSTALL_PATH:-/usr/local/cuda}"
+
+echo "run_sim.sh: binary      : $BIN"
+echo "run_sim.sh: config      : $CONFIG_DIR/$CONFIG_NAME"
+echo "run_sim.sh: workdir     : $WORKDIR"
+echo "run_sim.sh: runtime libs: $LIBDIR"
+
+cd "$WORKDIR"
+exec "$BIN" "$@"
diff --git a/debug/risc-v-gsp-rm-vm/gpgpu-sim/setup.sh b/debug/risc-v-gsp-rm-vm/gpgpu-sim/setup.sh
new file mode 100755
index 0000000..56e18ad
--- /dev/null
+++ b/debug/risc-v-gsp-rm-vm/gpgpu-sim/setup.sh
@@ -0,0 +1,173 @@
+#!/usr/bin/env bash
+# setup.sh -- materialise and build the pinned GPGPU-Sim upstream checkout.
+#
+# This is the Phase-0 entry point of the GreenEye-GPGPU-Sim overlay:
+#   1. clone upstream GPGPU-Sim at the pinned commit into ./upstream/ (gitignored),
+#   2. apply ./patches/*.patch in lexical order (empty for now is fine),
+#   3. build with the CUDA toolkit selected by CUDA_INSTALL_PATH.
+#
+# Idempotent: re-running is safe (clone is reused, already-applied patches are
+# skipped, the build is incremental). Fails loudly on any ambiguity. Never
+# uses sudo. This is development-only tooling: it never touches GPU hardware,
+# kernel modules, mdev, a guest, or the product tree.
+#
+# Usage:
+#   ./setup.sh [--no-build] [--reclone] [--build-type release|debug]
+#
+# Environment:
+#   CUDA_INSTALL_PATH        CUDA toolkit to build against (default /usr/local/cuda)
+#   JOBS                     compile jobs (default 4; host budget cap)
+#   GPGPUSIM_BUILD_TYPE      release (default) or debug
+#   GPGPUSIM_CLONE_SOURCE    explicit local mirror path or URL (unset by
+#                            default; when unset the upstream URL is used)
+#
+# Exit codes: 0 success; 1 precondition/build failure; 2 usage error.
+set -euo pipefail
+
+HERE="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
+UPSTREAM="$HERE/upstream"
+PATCH_DIR="$HERE/patches"
+BUILD_DIR="$HERE/build"
+PIN="90ec3399763d7c8512cfe7dc193473086c38ca38"
+UPSTREAM_URL="https://github.com/gpgpu-sim/gpgpu-sim_distribution.git"
+
+BUILD=1
+RECLONE=0
+BUILD_TYPE="${GPGPUSIM_BUILD_TYPE:-release}"
+JOBS="${JOBS:-4}"
+
+die() { echo "setup.sh: ERROR: $*" >&2; exit 1; }
+
+usage() {
+    cat <<'EOF'
+Usage: setup.sh [--no-build] [--reclone] [--build-type release|debug]
+
+  --no-build             clone/checkout + apply patches only
+  --reclone              delete ./upstream and clone fresh
+  --build-type TYPE      release (default) or debug
+
+Environment:
+  CUDA_INSTALL_PATH      CUDA toolkit to build against (default /usr/local/cuda)
+  JOBS                   compile jobs (default 4; host budget cap)
+  GPGPUSIM_BUILD_TYPE    release (default) or debug
+  GPGPUSIM_CLONE_SOURCE  explicit local mirror path or URL (unset by default;
+                         the upstream URL is always the default source)
+
+Idempotent; fails loudly; never uses sudo.
+EOF
+}
+
+while [ $# -gt 0 ]; do
+    case "$1" in
+        --no-build)   BUILD=0 ;;
+        --reclone)    RECLONE=1 ;;
+        --build-type) BUILD_TYPE="${2:-}"; shift ;;
+        -h|--help)    usage; exit 0 ;;
+        *)            usage >&2; exit 2 ;;
+    esac
+    shift
+done
+
+case "$BUILD_TYPE" in
+    release|debug) ;;
+    *) die "unknown build type '$BUILD_TYPE' (use release or debug)" ;;
+esac
+
+if [ "$RECLONE" = 1 ] && [ -e "$UPSTREAM" ]; then
+    echo "== --reclone: removing $UPSTREAM"
+    rm -rf "$UPSTREAM"
+fi
+
+# ---------------------------------------------------------------------------
+# 1. upstream checkout at the pinned commit
+# ---------------------------------------------------------------------------
+if [ -d "$UPSTREAM/.git" ]; then
+    echo "== upstream checkout present: $UPSTREAM"
+    head_commit="$(git -C "$UPSTREAM" rev-parse HEAD)"
+    if [ "$head_commit" != "$PIN" ]; then
+        if [ -n "$(git -C "$UPSTREAM" status --porcelain)" ]; then
+            die "upstream is at $head_commit (expected $PIN) and has local changes; --reclone to reset"
+        fi
+        echo "== checking out pinned commit $PIN"
+        git -C "$UPSTREAM" checkout --quiet --detach "$PIN" \
+            || die "pinned commit $PIN not present locally; run with --reclone"
+    else
+        echo "== upstream already at pinned commit ($PIN)"
+    fi
+elif [ -e "$UPSTREAM" ]; then
+    die "$UPSTREAM exists but is not a git checkout; use --reclone"
+else
+    # Default source is the upstream URL.  A local mirror is used only when
+    # GPGPUSIM_CLONE_SOURCE is set explicitly; nothing on this host is guessed.
+    if [ -n "${GPGPUSIM_CLONE_SOURCE:-}" ]; then
+        src="$GPGPUSIM_CLONE_SOURCE"
+        echo "== using explicit clone source from GPGPUSIM_CLONE_SOURCE: $src"
+    else
+        src="$UPSTREAM_URL"
+        echo "== using upstream URL: $UPSTREAM_URL"
+    fi
+    git clone --quiet "$src" "$UPSTREAM" || die "git clone failed from $src"
+    git -C "$UPSTREAM" checkout --quiet --detach "$PIN" \
+        || die "pinned commit $PIN not found in clone from $src"
+fi
+
+# ---------------------------------------------------------------------------
+# 2. overlay patches (lexical order; empty directory is tolerated)
+# ---------------------------------------------------------------------------
+shopt -s nullglob
+patches=("$PATCH_DIR"/*.patch)
+shopt -u nullglob
+
+if [ "${#patches[@]}" -eq 0 ]; then
+    echo "== patches/: empty (Phase 0; the CUDA 12 port patch set lands here)"
+else
+    for p in "${patches[@]}"; do
+        name="$(basename "$p")"
+        if git -C "$UPSTREAM" apply --check "$p" >/dev/null 2>&1; then
+            echo "== applying patch $name"
+            git -C "$UPSTREAM" apply "$p" || die "failed to apply $name"
+        elif git -C "$UPSTREAM" apply --reverse --check "$p" >/dev/null 2>&1; then
+            echo "== patch already applied: $name"
+        else
+            die "patch neither applies cleanly nor is already applied: $name (resolve or --reclone)"
+        fi
+    done
+fi
+
+# ---------------------------------------------------------------------------
+# 3. build (skipped with --no-build)
+# ---------------------------------------------------------------------------
+if [ "$BUILD" = 0 ]; then
+    echo "== --no-build: checkout + patches complete, build skipped"
+    exit 0
+fi
+
+cuda="${CUDA_INSTALL_PATH:-/usr/local/cuda}"
+[ -d "$cuda" ] || die "CUDA_INSTALL_PATH=$cuda does not exist (set the env var to a CUDA toolkit)"
+[ -x "$cuda/bin/nvcc" ] || die "no executable $cuda/bin/nvcc (set CUDA_INSTALL_PATH to a CUDA toolkit)"
+
+mkdir -p "$BUILD_DIR"
+LOG="$BUILD_DIR/build.log"
+echo "== building GPGPU-Sim ($BUILD_TYPE, -j$JOBS, CUDA_INSTALL_PATH=$cuda)"
+echo "== build log: $LOG"
+
+(
+    cd "$UPSTREAM"
+    export CUDA_INSTALL_PATH="$cuda"
+    export PATH="$cuda/bin:$PATH"
+    set +u
+    # shellcheck disable=SC1091
+    source ./setup_environment "$BUILD_TYPE"
+    set -u
+    if [ -z "${GPGPUSIM_CONFIG:-}" ]; then
+        echo "setup.sh: ERROR: setup_environment did not configure a build for CUDA at $cuda." >&2
+        echo "           This host's CUDA may not be ported yet (WP0.1); see UPSTREAM.md." >&2
+        exit 1
+    fi
+    make -j"$JOBS"
+) 2>&1 | tee "$LOG"
+
+echo "== setup complete"
+echo "==   upstream: $UPSTREAM"
+echo "==   runtime libs: $UPSTREAM/lib/<gcc>/<cuda>/$BUILD_TYPE"
+echo "==   next: ./scripts/run_sim.sh <cuda-binary> [args...]"