Docs

PantheonSim is a functional GPU simulator: it answers every call a CUDA program makes and runs its kernels on the CPU. It models behaviour, never performance. Everything here is also in the repository, with more depth.

Getting started

You need Linux, CMake 3.20 or newer and a C++20 compiler (gcc 13 or clang 17). No GPU, CUDA toolkit or third-party library is needed to build it.

git clone https://github.com/pantheongpu/pantheonsim
cd pantheonsim
./scripts/build.sh
./scripts/test.sh
./build/vgpu list-gpus

Or skip installing and use the playground.

Run a CUDA program

Build the program from its unmodified source with the real CUDA toolkit, linking the runtime shared, then run it on a simulated GPU:

nvcc -cudart shared my_app.cu -o my_app
./build/vgpu run --gpu nvidia/h200 ./my_app

vgpu run puts the simulator’s CUDA libraries in front of the real ones and runs the program in place, so its exit code and signals are its own. It also checks the binary first for the two things that otherwise fail silently: a CUDA library version this build does not carry, and an RPATH that would load the real libraries ahead of it (--preload gets past that one).

Kernels have to carry PTX, which nvcc embeds by default. Building for a card’s architecture works the way it does on hardware — -arch=compute_75 runs on a T4 and everything newer — and nvidia-smi --query-gpu=compute_cap reports it.

Options: --gpu, --count, --vram-mb, --threads, --race, --strict, --counters, --trace, --print-env. vgpu run --help describes each.

Simulate a machine

vgpu shell gives you a shell on the machine you describe. nvidia-smi, lspci, dmesg, uname and /proc/driver/nvidia/version all report it, and nvcc and any program you build run against it.

./build/vgpu shell -y --gpu nvidia/h100 --count 8 --os ubuntu:24.04 \
  --cuda 12.6 --driver 560.35.05 --hostname dgx-h100

Without -y it asks for each value. Inside, nvcc -arch=native resolves to the simulated card.

nvidia-smi and NVML

The table, -L, -q, --query-gpu with --format=csv, --query-compute-apps and -i by index, UUID or bus id behave as they do on a real driver, down to the column widths scripts cut on. Memory and utilization are the simulator’s real state; power, temperature and clocks are a model driven by it, and vgpu smi --explain says which is which.

NVML is served as libnvidia-ml.so.1, so pynvml, gpustat and nvitop work. Queries the simulator has no data for return NVML_ERROR_NOT_SUPPORTED, which those tools show as N/A.

Find bugs

Every error returns the documented CUDA code and prints what happened, where:

[vgpu] cuLaunchKernel: VirtualGPU error [out-of-bounds]: device memory read at
0x7fff00003140 is 0 bytes past the end of the 64-byte allocation at 0x7fff00003100
  in kernel 'vecAdd', PTX line 30
  lane 16
  instruction: ld.global.f32 %f1,[%rd8]
  GPU profile: nvidia/h100

--race checks the rule a block promises: two warps may touch the same shared word without a barrier only if both are reading. VGPU_SCHEDULER=adversarial runs the warp order most likely to expose one:

[vgpu] data race: read-write race on shared memory at byte offset 0: warp 2 and warp 0 both reach it with no bar.sync between them
  in kernel '_Z5tallyPi', PTX line 38
  instruction: st.shared.u32[_ZZ5tallyPiE5local], %r4
  GPU profile: nvidia/h100

--strict also refuses integer division by zero and storing a register nothing has written. It is off by default because correct compiler output trips both.

--counters prints exact instruction, memory, sector and bank-conflict counts for each launch. They are counted rather than sampled, so they do not move between runs. There are no timings: inventing them would be worse than not having them.

Test on every GPU

vgpu test --matrix runs one binary on every measured profile and compares the outputs to the first. Build it for the oldest architecture you support, or the older profiles will refuse its kernels:

$ nvcc -cudart shared -arch=compute_75 vectoradd.cu -o vectoradd
$ vgpu test --matrix ./vectoradd
profile                    exit  output
nvidia/a10                    0  baseline
nvidia/a100                   0  identical
nvidia/h100                   0  identical
nvidia/rtx3060                0  identical
nvidia/a100-sxm4-40gb         0  identical
...
11 profiles: 10 identical to nvidia/a10, 0 different

Environment

Variable	Meaning	Default
`VGPU_GPU`	The GPU profile to simulate	nvidia/h100
`VGPU_DEVICE_COUNT`	How many identical GPUs	1
`VGPU_VRAM_MB`	Memory per GPU, overriding the profile	the profile
`VGPU_THREADS`	Host threads that run blocks; 1 is strictly serial	every core
`VGPU_RACE`	1 reports unordered shared-memory access between warps; 2 also reports stores that change nothing	off
`VGPU_SCHEDULER`	Warp order: `adversarial` searches for the order that breaks a missing barrier	fixed
`VGPU_STRICT`	1 refuses integer division by zero and storing a register nothing has written	off
`VGPU_COUNTERS`	1 prints exact per-launch counters	off
`VGPU_MAX_STEPS`	Raises the runaway-kernel step budget; 0 removes it	built in
`VGPU_QUIET`	1 silences diagnostics on stderr	off
`VGPU_TRACE`	1 logs every runtime and driver entry point	off

Limits

Kernels must carry PTX. A binary built for machine code alone is refused with a precise error.
Programs must link the CUDA runtime shared (-cudart shared). This is what stops CuPy, which links it statically, and PyTorch’s bundled runtime refuses the simulated driver.
Performance is not modelled, and passing here does not replace a run on physical GPUs.
AMD profiles exist for discovery only; AMD execution is not implemented.

Found something that behaves differently from a real card? Open an issue with the program and the profile.