simulator.h — CUDA Simulator API#

code/simulator.h is a single-header, dual-target CUDA shim. Include it instead of <cuda_runtime.h> and one kernel source compiles two ways:

under nvcc (__CUDACC__ defined) it is a thin pass-through to the real CUDA runtime — LAUNCH(k, grid, block, ...) becomes k<<<grid,block>>>(...);
under a plain host compiler (MSVC, no nvcc, no GPU) the header is a small CPU simulator of the CUDA execution model, so kernels run locally.

This site documents every symbol the header exposes, grouped exactly like the header’s own section banners.

Contents:

API Reference

Portable spelling#

The only edits a canonical CUDA kernel needs to run on both toolchains:

Canonical CUDA	Portable spelling
`#include <cuda_runtime.h>`	`#include "simulator.h"`
`kernel<<<grid, block>>>(args)`	`LAUNCH(kernel, grid, block, args)`
`kernel<<<grid, block, shmem>>>(args)`	`LAUNCH_DYN(kernel, grid, block, shmem, args)`
`kernel<<<grid, block, 0, stream>>>(args)`	`LAUNCH_STREAM(kernel, grid, block, stream, args)`
`cudaLaunchCooperativeKernel(...)`	`LAUNCH_COOP(kernel, grid, block, args)`
`extern __shared__ T buf[]`	`SHARED_EXTERN(T, buf)`

A kernel named with more than one template argument needs the COMMA escape: LAUNCH(reduce6<int COMMA 256 COMMA true>, g, b, ...).

Quick start#

#include "simulator.h"

__global__ void vectorAdd(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() {
    const int n = 1 << 16;
    std::vector<float> hA(n, 1.0f), hB(n, 2.0f), hC(n);

    float *dA, *dB, *dC;
    cudaMalloc(&dA, n * sizeof(float));
    cudaMalloc(&dB, n * sizeof(float));
    cudaMalloc(&dC, n * sizeof(float));
    cudaMemcpy(dA, hA.data(), n * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(dB, hB.data(), n * sizeof(float), cudaMemcpyHostToDevice);

    const int threads = 256, blocks = (n + threads - 1) / threads;
    LAUNCH(vectorAdd, blocks, threads, dA, dB, dC, n);
    cudaDeviceSynchronize();

    cudaMemcpy(hC.data(), dC, n * sizeof(float), cudaMemcpyDeviceToHost);
    cudaFree(dA); cudaFree(dB); cudaFree(dC);
    return 0;
}

Build it on this PC with no GPU at all:

cl /std:c++20 /EHsc /TP /I code code\test_vector_add.cu && test_vector_add.exe

or on a real GPU with the same source:

nvcc -std=c++20 -arch=sm_75 -o out.exe code\test_vector_add.cu

Debugging the schedule#

Blocks and threads run in a randomised order every launch, which is what turns the simulator into a bug finder: a correct kernel gives the same answer under every schedule.

Knob	Effect
`SIM_SEED=<n>` (env) / simSetSeed(n)	Vary the random block/thread order; sweep it to surface order-dependent bugs.
simSetBlockOrder / simSetThreadOrder	Pin an exact visit order (`OrderThread{3,0,1,2}`, `::sequential()`, `::reverse()`).
simSetStepHook(cb)	Break at each barrier, inspect the segment, reorder the threads, continue.
SIM_CHECK + SIM_LD / SIM_ST	Sanitizer-lite: memcheck, initcheck, racecheck, synccheck.
simSetAtomicPreempt(p)	Preempt at atomics with probability `p` so device critical sections contend.

What is and is not modelled#

The simulator verifies results, not concurrency or timing.

Modelled: a distinct device heap, full 3-D indexing, static and dynamic shared memory, __syncthreads, atomics with yielding, warp primitives under the Volta independent-thread-scheduling model, cooperative launch and this_grid().sync(), a cooperative-groups subset, textures and surfaces, device-side dynamic parallelism, CPU subsets of cuBLAS and Thrust, fp16/bf16, CUDA graphs by stream capture, cuRAND, cuda::barrier, and stream-ordered asynchronous execution.

Not modelled: real timing (every event reads 0 ms), real inter-block parallelism (__threadfence() is a no-op), preemption at arbitrary points, mipmapped/layered textures, explicit CUDA graph construction, the driver API, graphics interop, tensor cores, NCCL/NVML, and cuFFT/cuSPARSE/cuSOLVER/NPP.

API Reference#

See API Reference for the complete symbol reference.