# 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.