cuda4cpu

CUDA および HIP プログラムを変更せずに、手元の C++ コンパイラーで CPU 上で実行できます。 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、ノート PC、教室などで実行できます。
• AI エージェントに、作成したカーネルをテストさせられます： Claude Code、Cursor、その他の MCP クライアント向けの MCP サーバーである cuda4cpu-mcp は、CUDA および HIP コードをコンパイル、実行、検証し、誤った結果、バリアの不足、メモリエラーを該当するソース行とともに報告します。
• 使い慣れたツールでカーネルをデバッグできます： gdb は単一の CUDA スレッドで停止し、AddressSanitizer、UBSan、Valgrind はカーネルのソース行を指し示します。
• GPU では見つかりにくいバグを検出できます： __syncthreads() の不足や、ブロック内のスレッドが異なるバリアに到達する問題を検出します。
• 実際のコードも動作します： NVIDIA の cuda-samples、Rodinia、AMD の rocm-examples に含まれる 154 個のプログラムのうち 130 個が、変更なしでビルドされ、テストに合格します。CI では、引き続き合格することも確認します。
• 現在の CUDA で使える機能： 共有メモリ、ワープのシャッフルと投票、協調グループ（grid.sync() を含む）、アトミック操作、半精度、bfloat16 および 8 ビット浮動小数点、ストリーム、イベント、ストリームキャプチャを使った CUDA グラフに対応しています。HIP コードもコンパイルできます。

目次

• クイックスタート
• バグ検出の例
• AI エージェントから使用する（MCP）
• プロジェクトで使用する
• サポート対象
• 仕組み
• 比較
• FAQ
• ロードマップとコントリビューション

クイックスタート

Ubuntu 24.04 のように、Linux（または WSL 2）を搭載した x86-64 環境と、GCC 14 以降および Python 3 が必要です（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

次に、C++ コンパイラーと同じオプションを受け取る cuda4cpu-c++、または nvcc のオプションを受け取る cuda4cpu-nvcc を使って、任意の .cu ファイルをコンパイルします。

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 をご覧ください。

バグを検出する様子を見る

__syncthreads() の欠落。 GPU では、このカーネルは長い間テストに合格することがあります。別のワープが書き込んだ値を読み取るのが、各ワープの最初のスレッドだけだからです。cuda4cpu はバリア間でブロックのスレッドをインデックス順に実行するため、同じようにバグが隠れます。逆順またはランダムな順序で実行すると、バグが表面化します（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 個の float があり、サイズは 4400 バイトです。cuda4cpu は割り当てサイズを 256 バイト単位で切り上げ、残りを poison するため、末尾を越えた最初の要素が 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 エージェントから使用する（MCP）

GPU のないマシン上でも、コーディングエージェントは CUDA を記述できますが、実行できないため、カーネルが動作するかどうかを確認できません。cuda4cpu-mcp を使うと確認できます。cuda4cpu で CUDA および HIP コードをコンパイル、実行、テストし、上記のチェックも実行する MCP サーバーです。エージェントに登録してください。たとえば 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 として登録されています（https://registry.modelcontextprotocol.io）。Docker では、マウントしたディレクトリ（上記の -v "$PWD:$PWD"）だけが見えます。インラインで指定したソースにはマウントは不要です。

ツール	エージェントが得られる内容
compile	ソース内の行番号とともに表示される nvcc または hipcc の診断、および cuda4cpu が実装していない使用機能
run	終了コード、出力、cuda4cpu の警告
check	判定と、ソース内の行番号を含む検出内容：メモリエラーと未定義動作（AddressSanitizer と UBSan）、ブロックのスレッド順序によって変わる結果（__syncthreads() または __syncwarp() の欠落、競合状態）、バリアの分岐、方向が誤ったコピー、ハング
test_kernel	プログラムを記述せずに、単一のカーネルをテストします。エージェントは起動設定と引数（指定または生成）に加え、期待値または C++ 参照関数を渡し、異なる要素と check の検出内容を受け取ります
capabilities	cuda4cpu が実装している機能と実装していない機能。これにより、エージェントはコードのバグと機能不足を見分けられます

結果は 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.

また、SAXPY のインデックス計算に blockDim.x ではなく blockDim.y を使った場合、要素をループする参照関数を使った 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 プロジェクトでも同様に使用でき、AMD GPU 向けの hipcc のオプションを取り除くため、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 テストを実行できます。このアクションはリリースをインストールし、ドライバーを 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

出力には、nvcc 用に書かれた Makefile で使用する toolkit ディレクトリ（make CUDA_PATH=...）と、CMake の CUDA および HIP 言語用の toolchain-file が含まれます。FetchContent で cuda4cpu を取得するプロジェクトでは、アクションは不要です。

サポート対象

• カーネル言語： __global__、__device__、__shared__（静的および動的）、__constant__、__syncthreads()、すべての warp シャッフル、vote および match、__activemask() と __syncwarp()、アトミック操作、cuda_fp16.h、cuda_bf16.h、cuda_fp8.h、cuComplex.h、数学および整数の組み込み関数、丸めモードの組み込み関数、clock64()。
• 協調グループ： スレッドブロック、最大 1024 スレッドのタイル、合体グループ、パーティション、reduce と scan、および協調起動での grid.sync()。
• ランタイム API： デバイス、メモリ（ピッチ付き、3D、マネージド、ピン留め、およびストリーム順序付きの割り当てを含む）、シンボル、ストリーム、イベント、占有率、CUDA グラフ：外部セマフォ以外のすべてのノード型を含む明示的なグラフ、条件付きノード、ストリームキャプチャ、cudaGraphExecUpdate。
• HIP： HIP の NVIDIA プラットフォームと同様に、HIP の名前によるランタイム API、hipLaunchKernelGGL、HIP_DYNAMIC_SHARED、マスクなしの warp 関数、ポインター属性、仮想メモリ管理、コンテキストなど、HIP 独自の追加機能をサポートします。

全機能の一覧：CUDA の機能、HIP、制限事項。互換性テストスイートでは、実際のプログラムで未対応の機能（主にテクスチャと動的並列処理）を確認できます。

仕組み

1. 小規模なリライターが、C++ では表現できない CUDA 構文の 2 つの部分、kernel<<<grid, block>>>(args) と extern __shared__ 配列を、ライブラリー呼び出しに変換します。それ以外はマクロを使った C++ です。__global__、threadIdx などはヘッダーで定義されています。行番号は変わらないため、診断メッセージは元のファイルを指します。
2. スレッドブロックは OpenMP を通じて CPU のコア上で並列に実行されます。
3. ブロック内の CUDA スレッドは 1 つずつ実行されます。 バリアーのないカーネルは、カーネルがインライン化された、スレッドを巡回する単純なループです。ベクトル加算は、OpenMP で書いた同じループと同じ速度で実行されます。
4. バリアーまたは warp 操作では、各 CUDA スレッドが次のスレッドに切り替わります。 CUDA スレッドはそれぞれ独自のスタックを持つファイバーで、コンテキストの切り替えには十数個の命令を要します。1 つのコアで __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 以降）。独自の CMakeLists.txt を使う NVIDIA のサンプルも同様です。cuda4cpu-nvcc は nvcc のオプションを受け付けます。また、<prefix>/lib/cuda4cpu/toolkit は CUDA ツールキットと同じ構成になっているため、$(CUDA_PATH)/bin/nvcc を呼び出す Makefile は、make CUDA_PATH=<prefix>/lib/cuda4cpu/toolkit でビルドできます。プロジェクトでの cuda4cpu の使用を参照してください。

CUDA のインストールが必要ですか？ いいえ。cuda4cpu は独自の cuda_runtime.h、cooperative_groups.h、cuda_fp16.h などを提供します。

コードを変更する必要がありますか？ 通常は必要ありません。テクスチャ、動的並列処理、インライン PTX、ドライバー API、または CUDA のライブラリ（cuBLAS、Thrust、CUB など）を使用するプログラムは、まだビルドできません。制限事項を参照してください。

カーネルは高速ですか？ 予想より高速かもしれませんが、あくまで CPU です。GPU の性能ではなく、並列化が適切な CPU プログラムと同程度の性能を想定してください。バリアを含まないカーネルが最も良い性能を発揮します。

macOS や Windows で使えますか？ Windows では WSL 2 上で動作します。それ以外の場合、cuda4cpu には glibc を搭載した 64 ビット CPU の Linux が必要で、x86-64 でテストされています。

ロードマップと貢献

ROADMAP.md には、テクスチャ、動的並列処理、残る API の不足部分、より高速なバリアなど、今後の予定が記載されています。失敗するプログラムを添えたバグ報告が、最も役立つ貢献です。ビルド、テスト、変更の送信方法については、CONTRIBUTING.mdを参照してください。

cuda4cpu が役に立った場合は、スターを付けると、ほかの人が見つけやすくなります。

cuda4cpu はApache License, Version 2.0に基づいてライセンスされています。