cuda4cpu

Run CUDA and HIP programs on your CPU, unmodified, with your regular C++ compiler. No GPU, no nvcc, no CUDA toolkit: GCC or Clang compile your .cu files, and the kernels run as native code on every core.

$ cuda4cpu-c++ -O2 vecadd.cu -o vecadd

$ ./vecadd c[42] = 126, computed on cuda4cpu: AMD Ryzen 7 8845HS w/ Radeon 780M Graphics

• Test CUDA code where there's no GPU: in CI on GitHub's hosted runners, on a laptop, or in a classroom.
• Let AI agents test the kernels they write: cuda4cpu-mcp, an MCP server for Claude Code, Cursor or any MCP client, compiles, runs and checks CUDA and HIP code, and reports wrong results, missing barriers and memory errors at their source lines.
• Debug kernels with the tools you already use: gdb stops in a single CUDA thread, and AddressSanitizer, UBSan and Valgrind point at the kernel's source line. A kernel that crashes is reported with its CUDA thread and block.
• Catch the bugs a GPU hides: missing __syncthreads(), threads of a block that reach different barriers, reads of memory never written, and allocations never freed.
• Real code works: 132 of 159 programs from NVIDIA's cuda-samples, Rodinia and AMD's rocm-examples build unmodified and pass, and CI checks that they keep passing.
• The CUDA you write today: shared memory, warp shuffles and votes, cooperative groups (including grid.sync()), atomics, half precision, bfloat16 and 8-bit floats, streams, events, and CUDA graphs with stream capture. HIP code compiles too.

Contents

• Quick start
• See it catch bugs
• Use it from an AI agent (MCP)
• Use it in your project
• What's supported
• How it works
• How it compares
• FAQ
• Roadmap and contributing

Quick start

You need Linux (or WSL 2) on x86-64, with GCC 14 or later and Python 3, as on Ubuntu 24.04 (sudo apt-get install g++-14 python3). Install the latest release in ~/.local:

mkdir -p ~/.local

curl -fsSL https://github.com/javier-cabezas/cuda4cpu/releases/latest/download/cuda4cpu-linux-x86_64.tar.gz | tar -xz -C ~/.local --strip-components=1

Then compile any .cu file with cuda4cpu-c++, which takes the same options as your C++ compiler, or with cuda4cpu-nvcc, which takes nvcc's:

cuda4cpu-c++ -O2 vecadd.cu -o vecadd

./vecadd

With Docker, there's nothing else to install: the image has cuda4cpu, GCC and CMake.

docker run --rm -u "$(id -u):$(id -g)" -v "$PWD:/src" -w /src ghcr.io/javier-cabezas/cuda4cpu \

cuda4cpu-nvcc -O2 vecadd.cu -o vecadd

From source, which needs CMake 3.25+ and Ninja too, and builds and runs the tests:

git clone https://github.com/javier-cabezas/cuda4cpu.git

cd cuda4cpu cmake --workflow --preset release cmake --install build/release --prefix /usr/local

See docs/usage.md for the other ways to install it and the build options.

See it catch bugs

A missing `__syncthreads()`. On a GPU, this kernel can pass its tests for a long time: only the first thread of each warp reads what another warp wrote. cuda4cpu runs the threads of a block in index order between barriers, which hides the bug too. Run them in reverse or in random order, and it shows (examples/bugs/missing_barrier.cu):

__global__ void pair_sum(const int *in, int *out)

{ shared int tile[256]; int t = threadIdx.x; int i = blockIdx.x * blockDim.x + t; tile[t] = in[i]; // Bug: a __syncthreads() is missing here out[i] = tile[t] + (t > 0 ? tile[t - 1] : 0); }

$ ./missing_barrier

PASSED: 0 wrong results $ CUDA4CPU_SCHEDULE=reverse ./missing_barrier FAILED: 65279 wrong results $ CUDA4CPU_SCHEDULE=random ./missing_barrier cuda4cpu: running threads in random order, seed 1277581447 (CUDA4CPU_SCHEDULE=random:1277581447 reproduces it) FAILED: 32593 wrong results

An out-of-bounds access, reported by AddressSanitizer with the kernel's line (examples/bugs/out_of_bounds.cu):

$ cuda4cpu-c++ -g -fsanitize=address out_of_bounds.cu -o out_of_bounds

$ ./out_of_bounds ==8457==ERROR: AddressSanitizer: use-after-poison on address 0x7c81561e2630 ... READ of size 4 at 0x7c81561e2630 thread T4 #0 0x65457dfe1ad5 in scale(float*, int) .../out_of_bounds.cu:27 ... 0x7c81561e2630 is located 4400 bytes inside of 4608-byte region [0x7c81561e1500,0x7c81561e2700) ... cuda4cpu: error: CUDA thread (76, 0, 0) of block (4, 0, 0) made the error that the sanitizer reports above, in the kernel launched at out_of_bounds.cu:37.

The array has 1100 floats, 4400 bytes: cuda4cpu rounds allocations up to 256 bytes and poisons the rest, so the first element past the end is reported, as use-after-poison. Without the sanitizer, the crash is reported with its CUDA thread and block too.

Divergent barriers. CUDA requires every thread of a block to reach the same __syncthreads(), and cuda4cpu reports blocks whose threads don't:

cuda4cpu: warning: threads of block (0, 0, 0) reached different __syncthreads() calls,

at kernel.cu:12 and kernel.cu:15. CUDA requires every thread of a block to reach the same one.

gdb in a CUDA thread. Each CUDA thread is a fiber with its own stack, so breakpoints, stepping and backtraces work as in any program. cuda4cpu's gdb script adds the built-in variables:

(gdb) break kernel.cu:42 if $threadIdx("x") == 6 && $blockIdx("x") == 2

(gdb) run (gdb) cuda4cpu thread (6, 0, 0) of block (2, 0, 0), lane 6; blocks of (64, 1, 1) threads, grid of (8, 1, 1) blocks

docs/debugging.md has the details.

Use it from an AI agent (MCP)

A coding agent on a machine without a GPU can write CUDA, but it can't run it, so it can't tell whether a kernel works. cuda4cpu-mcp gives it a way to find out: an MCP server that compiles, runs and tests CUDA and HIP code with cuda4cpu, and runs the checks above for it. Register it with your agent, for example Claude Code:

claude mcp add cuda4cpu -- cuda4cpu-mcp    # installed from a release, or build/release/tools/cuda4cpu-mcp

claude mcp add cuda4cpu -- docker run -i --rm -v "$PWD:$PWD" ghcr.io/javier-cabezas/cuda4cpu

Other clients take the same command, for example in Cursor's mcp.json or Claude Desktop's configuration: {"mcpServers": {"cuda4cpu": {"command": "cuda4cpu-mcp"}}}. The server speaks MCP over stdio, and is listed in the MCP Registry as io.github.javier-cabezas/cuda4cpu. In Docker, it sees only the directories you mount (-v "$PWD:$PWD" above): sources given inline need none.

Results are JSON, with a summary first. For the missing __syncthreads() above, whose program passes in the default order, check answers:

problems found

[depends on the thread order] running the threads of each block in reverse order (reverse) ends with exit code 1 instead of exit code 0; 8 of 8 other orders change the result a thread reads data that another thread of its block writes without a syncthreads() between them (or, within a warp, a syncwarp()), or threads race on the same data.

And for a SAXPY that computes its index with blockDim.y instead of blockDim.x, test_kernel with a reference that loops over the elements answers:

failed

argument 3: 998 of 1000 elements differ (largest error 1998) [1] 5, expected 3 [2] 13, expected 5 ...

A clean check means that nothing showed up in the orders and checks it ran, not that a program has no race. Programs run with your permissions, each in a temporary directory and its own process group, killed at a time-out: the server doesn't sandbox them. docs/debugging.md describes every tool and its arguments.

Use it in your project

CMake. Fetch cuda4cpu and add your .cu files to a target (examples/cmake-project):

include(FetchContent)

FetchContent_Declare(cuda4cpu GIT_REPOSITORY https://github.com/javier-cabezas/cuda4cpu.git GIT_TAG v0.5.0) FetchContent_MakeAvailable(cuda4cpu)

add_executable(my_app main.cpp) cuda4cpu_add_cuda_sources(my_app kernels.cu solver.cu)

With cuda4cpu installed, find_package(cuda4cpu REQUIRED) does the same. HIP sources go through cuda4cpu_add_hip_sources.

Any build system. cuda4cpu-c++ is a compiler driver: use it like your C++ compiler, with .cu files among the inputs. cuda4cpu-hipcc does the same for HIP projects, and drops hipcc's options for AMD GPUs, so build scripts written for hipcc work unchanged:

cuda4cpu-c++ -O2 -o my_app main.cpp kernels.cu

cuda4cpu-hipcc -O2 --offload-arch=gfx90a -o my_app main.cpp kernels.cpp

GitHub Actions. Hosted runners have no GPU, but they can run your CUDA tests with cuda4cpu. The action installs a release, with the drivers on the PATH and the CMake package where find_package(cuda4cpu) finds it:

jobs:

test: runs-on: ubuntu-24.04 env: CXX: g++-14 steps:

• uses: actions/checkout@v4
• uses: javier-cabezas/cuda4cpu@v0.5.0
• run: cmake -S . -B build -DCMAKE_BUILD_TYPE=Release
• run: cmake --build build
• run: ctest --test-dir build --output-on-failure

Again with the threads of each block in another order, which exposes

missing __syncthreads()

• run: CUDA4CPU_SCHEDULE=reverse ctest --test-dir build --output-on-failure

Its outputs give the toolkit directory, for Makefiles written for nvcc (make CUDA_PATH=...), and the toolchain-file for CMake's CUDA and HIP languages. A project that fetches cuda4cpu with FetchContent needs no action.

What's supported

• Kernel language: __global__, __device__, __shared__ (static and dynamic), __constant__, __syncthreads(), every warp shuffle, vote and match, __activemask() and __syncwarp(), atomics, cuda_fp16.h, cuda_bf16.h, cuda_fp8.h, cuComplex.h, math and integer intrinsics, rounding-mode intrinsics, clock64().
• Cooperative groups: thread blocks, tiles of up to 1024 threads,