1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
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.
|
