cuda4cpu

**使用常规 C++ 编译器在 CPU 上运行 CUDA 和 HIP 程序，无需修改。**无需 GPU、nvcc 或 CUDA 工具包：GCC 或 Clang 可编译你的 .cu 文件，内核会作为原生代码在每个核心上运行。

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

• **在没有 GPU 的地方测试 CUDA 代码：**在 GitHub 托管运行器、笔记本电脑或课堂上进行 CI 测试。
• 让 AI Agent 测试它们编写的内核：cuda4cpu-mcp 是面向 Claude Code、Cursor 或任何 MCP 客户端的 MCP 服务，可编译、运行并检查 CUDA 和 HIP 代码，并报告错误结果、缺失的屏障和内存错误所在的源代码行。
• **使用你已有的工具调试内核：**gdb 可在单个 CUDA 线程中暂停，AddressSanitizer、UBSan 和 Valgrind 可指出内核的源代码行。
• **捕获 GPU 会隐藏的错误：**缺少 __syncthreads()，以及块中的线程到达不同屏障。
• **真实代码也能运行：**NVIDIA 的 cuda-samples、Rodinia 和 AMD 的 rocm-examples 中，有 154 个程序的 130 个无需修改即可构建并通过测试，CI 会检查它们是否持续通过。
• **你今天编写的 CUDA 代码：**共享内存、warp 洗牌与投票、协作组（包括 grid.sync()）、原子操作、半精度、bfloat16 和 8 位浮点数、流、事件，以及支持流捕获的 CUDA 图。HIP 代码也能编译。

目录

• 快速开始
• 看看它如何捕获错误
• 从 AI Agent 使用（MCP）
• 在项目中使用
• 支持情况
• 工作原理
• 对比
• 常见问题
• 路线图与贡献

快速开始

你需要在 x86-64 上使用 Linux（或 WSL 2），并安装 GCC 14 或更高版本以及 Python 3，例如 Ubuntu 24.04（sudo apt-get install g++-14 python3）。将最新版本安装到 ~/.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

然后使用 cuda4cpu-c++ 编译任意 .cu 文件，它接受与 C++ 编译器相同的选项；也可以使用 cuda4cpu-nvcc，它接受 nvcc 的选项：

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

使用 Docker 时，无需安装其他内容：镜像中已包含 cuda4cpu、GCC 和 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

从源代码构建还需要 CMake 3.25+ 和 Ninja，并会构建和运行测试：

git clone https://github.com/javier-cabezas/cuda4cpu.git
cd cuda4cpu
cmake --workflow --preset release
cmake --install build/release --prefix /usr/local

其他安装方式和构建选项，请参阅 docs/usage.md。

看它捕获 Bug

缺少 __syncthreads()。 在 GPU 上，这个内核可能长期通过测试：只有每个 warp 的第一个线程会读取另一个 warp 写入的内容。cuda4cpu 会按索引顺序运行一个线程块中的线程，直到遇到屏障，这也会掩盖这个 Bug。改为按逆序或随机顺序运行，就能发现它（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

越界访问，由 AddressSanitizer 报告，并指出内核中的对应行（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)

该数组有 1100 个浮点数，共 4400 字节：cuda4cpu 会将分配大小向上取整到 256 字节，并对其余部分进行毒化，因此会报告末尾之后的第一个元素，并标记为 use-after-poison。

分歧屏障。 CUDA 要求线程块中的所有线程都到达同一个 __syncthreads()，cuda4cpu 会报告线程未全部到达的线程块：

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.

在 CUDA 线程中使用 gdb。 每个 CUDA 线程都是拥有独立堆栈的纤程，因此断点、单步执行和回溯都能像在任何程序中一样正常工作。cuda4cpu 的 gdb 脚本还会添加内置变量：

(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。

通过 AI Agent（MCP）使用

没有 GPU 的机器上的编码 Agent 可以编写 CUDA，但无法运行它，因此也无法判断内核是否有效。cuda4cpu-mcp 为它提供了验证方法：这是一个 MCP 服务，可使用 cuda4cpu 编译、运行和测试 CUDA 与 HIP 代码，并为它运行上述检查。将它注册到你的 Agent 中，例如 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

其他客户端也使用相同的命令，例如在 Cursor 的 mcp.json 或 Claude Desktop 的配置中设置：{"mcpServers": {"cuda4cpu": {"command": "cuda4cpu-mcp"}}}。该服务通过 stdio 使用 MCP 协议，并在 MCP Registry 中列为 io.github.javier-cabezas/cuda4cpu。在 Docker 中，它只能看到你挂载的目录（上面的 -v "$PWD:$PWD"）：以内联方式提供的源代码则不需要挂载。

工具	Agent 获取的内容
compile	nvcc 或 hipcc 在源代码对应行给出的诊断，以及 cuda4cpu 尚未实现的代码所用功能
run	退出代码、输出内容和 cuda4cpu 的警告
check	检查结论，以及标出源代码行的发现：内存错误和未定义行为（AddressSanitizer 和 UBSan）、随线程块内线程顺序变化的结果（缺少 __syncthreads() 或 __syncwarp()、竞争）、分歧屏障、方向错误的复制操作、挂起
test_kernel	无需编写程序即可测试单个内核：Agent 提供启动配置和参数（自行提供或生成），以及预期值或 C++ 参考函数；结果会列出不匹配的元素和 check 的发现
capabilities	cuda4cpu 已实现和未实现的功能，让 Agent 能区分代码中的 Bug 与缺失的功能

结果为 JSON，摘要在前。对于上文提到的缺少 __syncthreads() 的情况，其程序在默认顺序下可以通过，check 会返回：

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.

对于通过 blockDim.y 而不是 blockDim.x 计算索引的 SAXPY，使用逐元素循环的参考函数调用 test_kernel，会返回：

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

check 未发现问题，表示在它运行的线程顺序和检查中没有发现异常，并不代表程序不存在竞争。程序会以你的权限运行，各自在临时目录和独立进程组中执行，并会在超时后被终止：该服务不会对其进行沙箱隔离。 docs/debugging.md 介绍了每种工具及其参数。

在项目中使用

CMake。 获取 cuda4cpu，并将 .cu 文件添加到目标中 (examples/cmake-project)：

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

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

安装 cuda4cpu 后，find_package(cuda4cpu REQUIRED) 会完成相同的操作。 HIP 源文件通过 cuda4cpu_add_hip_sources 添加。

任意构建系统。 cuda4cpu-c++ 是一个编译器驱动程序：像使用 C++ 编译器一样使用它，并将 .cu 文件作为输入之一。cuda4cpu-hipcc 为 HIP 项目执行相同的操作，并移除 hipcc 中面向 AMD GPU 的选项，因此为 hipcc 编写的构建脚本无需修改即可使用：

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。 托管运行器没有 GPU，但可以使用 cuda4cpu 运行 CUDA 测试。该 action 会安装一个发行版，将驱动程序放入 PATH，并将 CMake 包放在 find_package(cuda4cpu) 能找到的位置：

jobs:
  test:
    runs-on: ubuntu-24.04
    env:
      CXX: g++-14
    steps:
      - uses: actions/checkout@v4
      - uses: javier-cabezas/[email protected]
      - 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

其输出包括 toolkit 目录，可供为 nvcc 编写的 Makefile 使用（make CUDA_PATH=...），以及供 CMake 的 CUDA 和 HIP 语言使用的 toolchain-file。通过 FetchContent 获取 cuda4cpu 的项目无需使用 action。

支持情况

• 内核语言： __global__、__device__、__shared__（静态和动态）、__constant__、__syncthreads()、所有 warp shuffle、vote 和 match 操作、__activemask() 和 __syncwarp()、原子操作、cuda_fp16.h、cuda_bf16.h、cuda_fp8.h、cuComplex.h、数学和整数内建函数、舍入模式内建函数、clock64()。
• 协作组： 线程块、最多 1024 个线程的 tile、合并组、分区、reduce 和 scan，以及协作启动中的 grid.sync()。
• 运行时 API： 设备、内存（包括跨步、3D、统一、固定和按流顺序分配）、符号、流、事件、占用率和 CUDA 图：显式图支持除外部信号量以外的所有节点类型、条件节点、流捕获和 cudaGraphExecUpdate。
• HIP，与 HIP 的 NVIDIA 平台一致：使用 HIP 名称的运行时 API、hipLaunchKernelGGL、HIP_DYNAMIC_SHARED、不带掩码的 warp 函数，以及 HIP 自有的扩展，例如指针属性、虚拟内存管理和上下文。

完整列表：CUDA 功能、HIP、限制。兼容性测试套件展示了实际程序仍缺少哪些功能：主要是纹理和动态并行。

工作原理

1. 小型重写器会将 C++ 无法表达的两种 CUDA 语法——kernel<<<grid, block>>>(args) 和 extern __shared__ 数组——转换为库调用。其他内容都是带宏的 C++：__global__、threadIdx 等由头文件定义。行号保持不变，因此诊断信息会指向你的文件。
2. 线程块通过 OpenMP 在 CPU 核心上并行运行。
3. 线程块中的 CUDA 线程依次运行。 不含屏障的内核就是一个遍历线程的普通循环，并将内核内联其中：向量相加的速度与使用 OpenMP 编写的相同循环一样快。
4. 遇到屏障或 warp 操作时，每个 CUDA 线程都会切换到下一个线程：CUDA 线程是 fiber，各自拥有堆栈，并通过十几条指令完成上下文切换。单核上，线程在 __syncthreads() 处停留约 6.5 ns，其中包括调度时间。
5. **在循环中，一个线程块的线程每隔几次迭代轮流运行，**因此它们会像在 GPU 上一样同时访问相邻地址。每个启动点都会测量切换频率。网格跨步复制的运行速度提升了 5.7 倍。

docs/how-it-works.md 逐一解释各个部分，docs/usage.md 则展示了重写器的处理方式。

对比

其他项目可在 CPU 上运行 CUDA 或 HIP 代码：

• HIP-CPU 是 AMD 的仅头文件库，可在 CPU 上运行 HIP 代码。
• chipStar 使用基于 LLVM 的工具链将 CUDA 和 HIP 编译为 SPIR-V，并在 OpenCL 或 Level Zero 设备上运行，包括通过 PoCL 等 OpenCL CPU 运行时在 CPU 上运行。
• CuPBoP 是一个研究框架，使用 Clang 将 CUDA 编译为 LLVM IR， 再将其转换为适用于 CPU 的代码。

cuda4cpu 无需特殊工具链：内核是普通的 C++ 函数，由你的 GCC 或 Clang 编译。因此，所有 CPU 工具都能像处理程序其他部分一样处理这些内核。代价是，每个 CUDA 线程遇到屏障时都要进行一次纤程切换；通过围绕屏障重构内核来避免这一点的编译器不会有此问题：每隔几条指令就进行同步的内核是 cuda4cpu 最慢的情况。

常见问题

它是模拟器吗？ 不是。内核会编译为适用于 CPU 的原生代码，使用编译器的优化，并在所有核心上运行。不会解释执行 PTX。

我可以保留现有的构建系统吗？ 通常可以。为 CUDA 编写的 CMake 项目（LANGUAGES CUDA、find_package(CUDAToolkit)）可以使用 -DCMAKE_TOOLCHAIN_FILE=<prefix>/lib/cuda4cpu/toolkit/cuda4cpu.cmake 构建（需要 CMake 4.1 或更高版本），NVIDIA 自带 CMakeLists.txt 的示例项目也可以。cuda4cpu-nvcc 接受 nvcc 的选项，而 <prefix>/lib/cuda4cpu/toolkit 的目录结构与 CUDA 工具包相同，因此使用 make CUDA_PATH=<prefix>/lib/cuda4cpu/toolkit 即可构建调用 $(CUDA_PATH)/bin/nvcc 的 Makefile。参见在项目中使用 cuda4cpu。

需要安装 CUDA 吗？ 不需要。cuda4cpu 提供自己的 cuda_runtime.h、cooperative_groups.h、cuda_fp16.h 等文件。

我需要修改代码吗？ 通常不需要。使用纹理、动态并行、内联 PTX、驱动 API 或 CUDA 库（cuBLAS、 Thrust、CUB 等）的程序目前尚无法构建。参见限制。

我的内核会很快吗？ 可能比你预期的快，但它运行在 CPU 上：应期待达到并行化良好的 CPU 程序的性能，而不是 GPU 的性能。不含屏障的内核表现最佳。

支持 macOS 或 Windows 吗？ 在 Windows 上，可通过 WSL 2 使用。除此之外，cuda4cpu 需要运行 glibc 的 Linux 系统和 64 位 CPU，并且已在 x86-64 上测试。

路线图与贡献

ROADMAP.md 列出了计划内容：纹理、动态并行、剩余的 API 缺口以及更快的屏障。提交附有失败程序的 Bug 报告是最有帮助的贡献。请参阅 CONTRIBUTING.md，了解如何构建、测试和提交更改。

如果 cuda4cpu 对你有帮助，请点个星标，让更多人发现它。

cuda4cpu 根据 Apache License, Version 2.0 授权。