PantheonSim / Docs

Docs

PantheonSim, the GPU simulator for CUDA and ROCm, is a functional simulator: it answers every call a CUDA or ROCm 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 and the HIP runtime build without a GPU toolkit; the CUDA libraries, and so vgpu run for CUDA programs, build when a CUDA toolkit is installed. ROCm is only needed for its compiler, hipcc, to build the programs you run.

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 run from the machine code (SASS) the binary carries for the simulated card’s architecture, and from PTX, which nvcc embeds by default, otherwise. 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.

Run a ROCm program

Build the HIP program with ROCm’s own hipcc, for the card’s target, and run it on a simulated Instinct. The simulator is the HIP runtime the program loads (libamdhip64, for ROCm 6 or 7):

$ hipcc --offload-arch=gfx942 vectoradd.hip -o vectoradd
$ ./build/vgpu run --gpu amd/mi300x ./vectoradd
AMD Instinct MI300X (gfx942), 304 CUs, wavefront 64
ok: c[7] = 21

Build for the card’s own target, as on hardware: gfx942 for the MI300X and MI325X, gfx950 for the MI350X. A program built for another target is refused by name, the way the real runtime refuses it.

The AMD side is newer than the NVIDIA one. The HIP runtime covers devices, memory, streams, events, graphs and peer copies; the kernels run as 64-lane wavefronts over the CDNA instructions clang emits for the programs tested so far. Anything else is refused by name rather than guessed at. In the repository’s current source, AMD’s own libraries such as rocBLAS run on it too. The playground offers the three Instinct cards and the Radeon RX 6800 and RX 6700 XT with ROCm 7.1 and hipcc. More in running ROCm and HIP without an AMD GPU.

Simulate a machine

vgpu shell gives you a shell on the machine you describe. nvidia-smi on an NVIDIA machine, rocm-smi and amd-smi on an AMD one, and lspci, dmesg and uname on both report it, and so does /proc/driver/nvidia/version where the shell can create a mount namespace (the playground cannot). nvcc or hipcc, 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
./build/vgpu shell -y --gpu amd/mi300x --count 8 --rocm 7.1.0

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

CUDA and ROCm 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.

ROCm is simpler: the HIP runtime answers to both of ROCm’s names for it, libamdhip64.so.6 and .so.7, so programs built by hipcc from ROCm 6.4, 7.0, 7.1 or 7.2 load it as they are. vgpu shell --rocm sets the version the machine reports (6.2.0 unless you give one).

nvidia-smi, NVML and rocm-smi

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.

On an AMD machine, rocm-smi (the concise table, -a, --showproductname, --showmeminfo vram, --json, --csv and the rest), amd-smi and rocm_agent_enumerator answer from the profile the same way, and rocprofv3, AMD’s profiler, runs unmodified and counts each kernel’s waves and instructions exactly.

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 runs unmodified under vgpu run and vgpu shell; elsewhere it needs a one-line hook (nvidia/tools/vgpu_triton.py in the repository). PyTorch’s official CUDA 13 build runs on the simulated NVIDIA GPUs when the simulator is built from its current source and started with vgpu run --preload (guide), but the playground does not support it yet and says so when it is imported; VGPU_PY_NOTICE=0 turns that off. CuPy does not work: it links the CUDA runtime statically.

Find bugs

Every error in a CUDA program 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. The scheduler orders only the PTX engine, so set VGPU_SASS=0 with it:

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

ROCm programs get the memory checks: an out-of-bounds or use-after-free access in a kernel is caught and reported. --race, --strict and --counters are CUDA-only for now; for AMD counters, run the program under rocprofv3.

Test on every GPU

vgpu test --matrix runs one CUDA binary on every measured NVIDIA profile (fourteen) and compares the outputs to the first. --gpus a,b,c picks profiles, --all adds the unmeasured ones and --show-diff prints the first output that differs. 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
...
14 profiles: 13 identical to nvidia/a10, 0 different

Environment

ROCm programs read VGPU_GPU, VGPU_DEVICE_COUNT, VGPU_THREADS and VGPU_QUIET; the rest apply to CUDA programs only, for now.

Limits

• CUDA kernels run from SASS when the binary carries code for the card’s architecture, and from PTX otherwise. A kernel the simulator cannot run fails with an error that names the instruction.
• ROCm kernels must be built for the simulated card’s own target (gfx942 for the MI300X and MI325X, gfx950 for the MI350X); a code object for another target is refused by name.
• CUDA programs must link the CUDA runtime shared (-cudart shared). This is what stops CuPy, which links it statically.
• Performance is not modelled, and passing here does not replace a run on physical GPUs.
• On AMD machines, the HIP runtime and the CDNA instructions the tested programs use are implemented; anything else is refused by name. AMD’s own libraries, such as rocBLAS, run on the simulator in the repository’s current source.

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