diff options
Diffstat (limited to 'debug/risc-v-gsp-rm-vm/gpgpu-sim/BUILD-PORT.md')
| -rw-r--r-- | debug/risc-v-gsp-rm-vm/gpgpu-sim/BUILD-PORT.md | 222 |
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. |