PantheonSim / Run ROCm and HIP without an AMD GPU

How to run ROCm and HIP programs without an AMD GPU

PantheonSim, the GPU simulator for CUDA and ROCm, runs HIP programs built by ROCm's own hipcc on a computer with no AMD GPU. The simulator is the HIP runtime the program loads, so the binary runs unmodified and its kernels execute on the CPU. This page covers the build rule you must follow, the commands to run a program, and where the AMD side is still younger than the NVIDIA one.

What happens when you run it

You compile with the real hipcc from a ROCm installation. The program then asks for libamdhip64, and the simulator answers under both of ROCm's names for it (.so.6 and .so.7). Programs built by ROCm 6.4, 7.0, 7.1 and 7.2 run. The simulator decodes the AMD machine code in the binary and executes it on the host cores, in 64-lane wavefronts on the CDNA 2 to CDNA 4 Instinct cards and 32-lane ones on the Radeon (RDNA) cards. rocm-smi, amd-smi and rocm_agent_enumerator report the simulated machine, and AMD's profiler rocprofv3 runs unmodified.

Like the NVIDIA side, it is functional: it checks what your code does, never how fast it does it. There is no timing model, so timings from a simulated run mean nothing.

Try it in the browser

The playground offers the Instinct MI300X, MI325X and MI350X with ROCm 7.1 and hipcc, in a browser terminal, with no install and no sign-up. To run it on your own machine you need Linux, CMake 3.20 or newer, a C++20 compiler, and ROCm's hipcc to build the programs you run. No AMD driver and no GPU are needed.

git clone https://github.com/pantheongpu/pantheonsim
cd pantheonsim
make doctor
make build
build/vgpu list-gpus

The gfx target rule

On hardware, a HIP binary carries code for specific GPU architectures, named by gfx targets. The simulator keeps that rule: build for the simulated card's own gfx target, or the program is refused the way the real runtime refuses it. A program that carries code for the wrong GPU stops with hipErrorNoBinaryForGpu. The target comes from the card you choose:

To see the target of any profile, ask the simulator: build/vgpu info --gpu amd/mi300x prints it on the gfx target line. The GPU index lists the target for every profile, and the source repository's AMD notes cover the cards that are newer still.

Build and run a HIP program

Save this as devices.hip. It prints what the runtime reports about device 0, then adds two vectors on the GPU and checks the result:

#include <hip/hip_runtime.h>
#include <cstdio>

__global__ void add(const float* a, const float* b, float* c, int n) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n) c[i] = a[i] + b[i];
}

int main() {
  hipDeviceProp_t prop;
  if (hipGetDeviceProperties(&prop, 0) != hipSuccess) { printf("no HIP device\n"); return 1; }
  printf("%s (%s), %d CUs, wavefront %d\n", prop.name, prop.gcnArchName,
         prop.multiProcessorCount, prop.warpSize);

  const int n = 1024;
  float ha[n], hb[n], hc[n];
  for (int i = 0; i < n; i++) { ha[i] = (float)i; hb[i] = 2.0f * i; }
  float *a, *b, *c;
  hipMalloc(&a, sizeof ha); hipMalloc(&b, sizeof hb); hipMalloc(&c, sizeof hc);
  hipMemcpy(a, ha, sizeof ha, hipMemcpyHostToDevice);
  hipMemcpy(b, hb, sizeof hb, hipMemcpyHostToDevice);
  add<<<(n + 255) / 256, 256>>>(a, b, c, n);
  hipMemcpy(hc, c, sizeof hc, hipMemcpyDeviceToHost);
  for (int i = 0; i < n; i++)
    if (hc[i] != 3.0f * i) { printf("FAIL at %d\n", i); return 1; }
  printf("PASS\n");
  return 0;
}

hipcc --offload-arch=gfx942 devices.hip -o devices
build/vgpu run --gpu amd/mi300x ./devices

The first line it prints names the simulated card and its target, and the last is PASS. A program that hipcc built also runs directly when the simulator's libraries are first on the library path, which is what the GitHub Action arranges; under vgpu run that is done for you. Pick another card with --gpu and another count with --count, building for that card's target each time.

For a whole simulated machine, vgpu shell gives you a shell where rocm-smi and amd-smi answer for the GPUs you asked for:

build/vgpu shell -y --gpu amd/mi300x --count 8 --rocm 7.1.0
rocm-smi
rocm_agent_enumerator
hipcc --offload-arch=gfx942 devices.hip -o devices && ./devices

The HSA_OVERRIDE_GFX_VERSION note for gfx1031

Many ROCm libraries, PyTorch's ROCm wheels among them, ship code for gfx1030 and none for gfx1031, the Radeon RX 6700 XT. On a real RX 6700 XT, people set HSA_OVERRIDE_GFX_VERSION=10.3.0 so the runtime reports the card as gfx1030 and that code loads. The simulated amd/rx6700xt honours the same variable: with it set, the card reports itself as gfx1030 and the gfx1030 code loads; without it, that code is refused. The simulator ignores any other value.

HSA_OVERRIDE_GFX_VERSION=10.3.0 build/vgpu run --gpu amd/rx6700xt python my_script.py

Libraries and Python

AMD's own libraries run on the simulator unmodified in the repository's current source, starting with rocBLAS (and hipBLASLt, which it links); AMD's own rocBLAS quick tests pass. PyTorch built for ROCm runs on the simulated AMD GPUs too, and so does RCCL across several of them. The details, and the instructions that are implemented for each architecture, are in the AMD notes in the repository. The playground ships hipcc for building and running HIP programs.

Limits

• Functional, not timed. There is no timing or cache model. Passing here does not replace a run on an Instinct before a release.
• The AMD side is newer. What is not implemented is refused by name rather than guessed at. Several profiles, including the MI300X and MI350X, are built from public documents and are marked as not yet measured against a physical card; the GPU index says which.
• Checks differ from the NVIDIA side. Out-of-bounds and use-after-free accesses in a kernel are caught. For counters, run the program under rocprofv3, and for stepping through a kernel use vgpu debug.
• Some reported fields are N/A. VBIOS, serial numbers, firmware versions and energy counters have no value on a simulated card, so rocm-smi and amd-smi print N/A instead of inventing one.

Questions

Next steps

Run the same idea for CUDA in how to run CUDA programs without a GPU, put HIP tests in GitHub Actions (name an amd/ profile and the action installs hipcc for you), read the docs, or look at the simulated rocm-smi and amd-smi. To see how this differs from translation layers, see the comparison.