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

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

1 files changed, 222 insertions, 0 deletions

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.