バーラCUDA ― AMD GPU 向けに設計されたオープンソースの CUDA コンパイラー

BarraCUDA

AMD GPU を対象にしたオープンソースの CUDA コンパイラで、今後さらに多くのアーキテクチャを追加予定です。C99 で 15,000 行程度のコードから構成され、LLVM に依存していません。

.cu

.hsaco

これは NVIDIA のウォールガーデンを眺め、「そんなに難しいものなの？」と考えた結果です。答えは「実際にはかなり難しい」ですが、私はそれでも挑戦しました。
※ 現在の Tenstorrent 実装を試したい場合は、そのブランチをクローンしてください。

何ができるか

CUDA C ソースコード（

nvcc

.cu

┌───────────────────────────────────────┐
│             BarraCUDA Pipeline        │
├───────────────────────────────────────┤
│ Source (.cu)                          │
│   ↓                                   │
│ Preprocessor → #include, #define, …   │
│   ↓                                   │
│ Lexer → Tokens                        │
│   ↓                                   │
│ Parser (Recursive Descent) → AST      │
│   ↓                                   │
│ Semantic Analysis → Type checking, …  │
│   ↓                                   │
│ BIR (BarraCUDA IR) → SSA form, …     │
│   ↓                                   │
│ mem2reg → Promotes allocas to SSA regs│
│   ↓                                   │
│ Instruction Selection → AMDGPU ops    │
│   ↓                                   │
│ Register Allocation → VGPR/SGPR       │
│   ↓                                   │
│ Binary Encoding → GFX11 instruction  │
│   ↓                                   │
│ ELF Emission → .hsaco                 │
│   ↓                                   │
│ Your kernel runs on your silicon      │
└───────────────────────────────────────┘

すべてのエンコーディングは

llvm-objdump

ビルド方法

# C99 コードです。gcc でビルドできます。
make

CMake や autoconf は不要です。ビルド手順は 47 ステップではなく、簡潔です。もし失敗したら、GCC が壊れている可能性があります。

必要条件

• C99 コンパイラ（gcc, clang 等）
• （任意）挑戦心

LLVM は 不要 です。BarraCUDA は独自に命令をエンコードします。

使用方法

# AMD GPU バイナリへコンパイル
./barracuda --amdgpu-bin kernel.cu -o kernel.hsaco

# IR をダンプ（デバッグや好奇心のため）
./barracuda --ir kernel.cu

# ただ AST を解析して表示
./barracuda --ast kernel.cu

# セマンティック分析を実行
./barracuda --sema kernel.cu

対応機能

以下の CUDA 機能が GFX11 マシンコードとして正しく動作します。

__global__

__device__

__host__

__shared__

threadIdx

blockIdx

__syncthreads()

s_barrier

atomicAdd

atomicSub

__shfl_sync

__ballot_sync

&&

__half

__launch_bounds__

cooperative_groups::this_thread_block()

コンパイラ機能

• 完全な C 前処理器：

  #include

  ,

  #define/#undef

  , 関数型マクロ, 条件付きディレクティブ,

  #pragma

  ,

  #error

• エラー回復（複数エラーを報告しても停止しない）
• IR ダンプ時のソース位置追跡
• 構造体の値渡し

例

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

$ ./barracuda --amdgpu-bin vector_add.cu -o vector_add.hsaco
wrote vector_add.hsaco (528 bytes code, 1 kernels)

LLVM は不要です。

アーキテクチャ

lexer.c

preproc.c

parser.c

sema.c

bir.c + bir_lower.c

bir_mem2reg.c

bir_print.c

amdgpu_isel.c

amdgpu_emit.c

main.c

データ構造はすべて固定サイズ配列で、ホットパスでは malloc を使わず、再帰もなく、ループはすべて境界付きです。

現状対応できない項目

• 単独の

  unsigned

  （

  unsigned int

  または

  int

  を使用）
• 代入演算子 (

  +=

  ,

  -=

  ,

  >>=

  など)

• const

  修飾子

• __constant__

  メモリ
• 2D シェアードメモリ配列 (

  __shared__ float a[16][16]

  ) – 1D にフラット化
• 整数リテラルサフィックス (

  0xFFu

  ,

  1ULL

  )

• __device__

  関数内でのパラメータ再代入（ローカル変数を使用）
• テクスチャとサーフェス
• 動的並列性（デバイス側カーネル起動）
• 複数トランスレーションユニット
• ホストコード生成（デバイスコードのみコンパイル）

これらはアーキテクチャ上の障害ではなく、まだ実装に手が回っていないだけです。

テストスイート

14 ファイル、35+ カーネル、約1,700 BIR 命令、約27 000 バイトのマシンコード：

• vector_add.cu

  – GPU 計算の「Hello World」

• cuda_features.cu

  – アトミック, ワープオペ, バリア, goto, switch, short-circuit

• test_tier12.cu

  – ベクトル, シェアードメモリ, 演算子オーバーロード

• notgpt.cu

  – AI 生成 CUDA（テイル SGEMM, 減算, ヒストグラム等）

• stress.cu

  – N-body, ネスト制御フロー, ビット操作, 構造体値渡し

• canonical.cu

  – NVIDIA サンプルから改変したパターン

• test_errors.cu

  – 故意の構文エラーで回復機能を確認

• test_launch_bounds.cu

  –

  __launch_bounds__

  と VGPR 上限

• test_coop_groups.cu

  – コーポレーティブグループ低減

また、前処理テスト, テンプレートテスト, unsigned 整数テストも含む。

ロードマップ

短期（ハーデニング）

既知のギャップ（代入演算子, 単独 unsigned, 整数リテラルサフィックス,

const

.cu

中期（最適化）

• 命令スケジューリング（メモリ遅延隠蔽）
• より良いレジスタ割り当て（現在は線形スキャン、グラフカラーを検討）
• 定数フォールディングとデッドコード除去
• ループ不変コード移動
• レジスタ圧力に基づく占有率調整

長期（多アーキテクチャ対応）

IR (BIR) はターゲット非依存です。新しいターゲットを追加するには、

isel

emit

候補：

• Tenstorrent – RISC‑V AI アクセラレータ（タイルベース、SIMT ではない）
• Intel Arc – Xe アーキテクチャ
• RISC‑V ベクトル拡張 – CUDA をソフトコアで動かす

GFX11 エンコーディングノート（勇敢な方へ）

独自 AMDGPU バックエンドを書こうとしているなら、以下の点が午後を台無しにします：

• SOP1 プレフィックスは

  0xBE800000

  （ドキュメントとは異なる）
• SOPC プレフィックスは

  0xBF000000

• VOP3 VDST はビット

  [7:0]

  に位置し、

  [15:8]

  ではない
• Null SADDR：グローバルメモリ用は

  0x7C

  、スクラッチは

  0xFC

• RDNA 3 のデフォルトは Wave32（Wave64 ではない）
• ISA マニュアルは 500 ページに及び、二度矛盾

amdgpu_emit.c

お問い合わせ

バグが見つかったら？ AMDGPU 命令エンコーディングについて議論したい？ GPU 計算で共感し合いたいなら：

zanehambly@gmail.com

議論したいことがあれば issue を立ててください。無理にではありません。私はあなたの母親ではありません。

ニュージーランド在住です。ここはすでに明日で、GPU は他所と同じくらい混乱しています。

ライセンス

Apache 2.0 – 何でも好きに使ってください。このコンパイラが本番環境で使用されたらぜひ教えてください（主に「面白いものを作った」という LinkedIn の更新のため）。