PantheonSim / Docs

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 or third-party library is needed. The core builds without a CUDA toolkit; the CUDA libraries, and so vgpu run, build when one is installed.

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 as 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, --max-steps, --quiet, --preload, --print-env. vgpu run --help describes each. --quiet silences error messages as well as the banner, so leave it off while debugging.

Simulate a machine

vgpu shell gives you a shell on the machine you describe. nvidia-smi, lspci, dmesg and uname report it, and so does /proc/driver/nvidia/version where the shell can create a mount namespace (the playground cannot). 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

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

CUDA versions

The simulator’s CUDA libraries are built against one CUDA toolkit and carry its version, so a program has to be compiled with that same major version. A real driver is more forgiving, because programs ship their own runtime. The playground offers two, each with the driver NVIDIA pairs with it:

On your own machine, the build uses whichever toolkit it finds. vgpu run says so plainly when a program was built against a different major version than the simulator carries.

nvidia-smi and NVML

The table, -L, -q, --query-gpu with --format=csv, --query-compute-apps, -q -x, topo -m 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 inside vgpu shell, or while a simulated program is running: NVML reads the live simulator, and on its own reports the driver as not loaded. Queries the simulator has no data for return NVML_ERROR_NOT_SUPPORTED, which those tools show as N/A.

topo -m reports every pair of GPUs as PHB: simulated devices reach each other through the host, and no profile carries NVLink data measured from a card, so none is claimed.

Python

Numba runs unmodified: it compiles Python to PTX itself and talks to the driver directly. It is preinstalled in the playground.

from numba import cuda
import numpy as np

@cuda.jit
def add(a, b, out):
    i = cuda.grid(1)
    if i < out.size:
        out[i] = a[i] + b[i]

Triton works with a one-line hook (tools/vgpu_triton.py in the repository). PyTorch and CuPy do not: PyTorch’s bundled CUDA runtime checks the driver’s identity and refuses a simulated one, and CuPy links the runtime statically. The playground says so when either is imported; VGPU_PY_NOTICE=0 turns that off.

Find bugs

Every error returns the documented CUDA code and prints what happened, where. The messages come from VirtualGPU, the simulator’s engine:

[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 '_Z6vecAddPfPKfi', PTX line 36
  lane 16
  instruction: ld.global.f32 %f1,[%rd6]
  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 interleaves warps as often as the model permits, the order most likely to expose one:

[vgpu] cudaLaunchKernel: VirtualGPU error [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, so which one wins is not something the program decided. Hardware usually hides this because warps advance together; it is a real race either way
  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

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.