Zaneham/BarraCUDA: AMD GPU をターゲットとしたオープンソース CUDA コンパイラー (将来的にはさらに多くのものも!)。 .cu を GFX11 マシンコードにコンパイルします。

AMD GPU をターゲットとするオープンソース CUDA コンパイラーであり、より多くのアーキテクチャが計画されています。 C99 の 15,000 行で書かれています。 LLVM 依存性はありません。コンパイルします .cu ファイルを GFX11 マシンコードに直接変換し、ELF を吐き出します .hsaco AMD GPU が実際に実行できるバイナリ。

NVIDIA の壁に囲まれた庭園を見て、「それはどれほど難しいことだろう?」と考えると、このようなことが起こります。答えは、実際にはかなり難しいですが、とにかくやり遂げました。

注: 現在の tenstorrent 実装をテストするためにここにいる場合は、それぞれのブランチをクローンする必要があります 🙂

CUDA C ソースコードを使用します。同じです .cu フィードするファイル nvcc、それらを AMD RDNA 3 (gfx1100) バイナリにコンパイルします。 LLVMはありません。 HIP 変換層はありません。 「最初に CUDA を他のものに変換する」ことはできません。レクサー、パーサー、IR、そしてコンパイラーの教科書が泣くような約 1,700 行の手書きの命令選択だけです。

┌──────────────────────────────────────────────────────────────┐
│                     BarraCUDA Pipeline                        │
├──────────────────────────────────────────────────────────────┤
│  Source (.cu)                                                │
│       ↓                                                      │
│  Preprocessor → #include, #define, macros, conditionals      │
│       ↓                                                      │
│  Lexer → Tokens                                              │
│       ↓                                                      │
│  Parser (Recursive Descent) → AST                            │
│       ↓                                                      │
│  Semantic Analysis → Type checking, scope resolution         │
│       ↓                                                      │
│  BIR (BarraCUDA IR) → SSA form, typed instructions           │
│       ↓                                                      │
│  mem2reg → Promotes allocas to SSA registers                  │
│       ↓                                                      │
│  Instruction Selection → AMDGPU machine instructions         │
│       ↓                                                      │
│  Register Allocation → VGPR/SGPR assignment                  │
│       ↓                                                      │
│  Binary Encoding → GFX11 instruction words                   │
│       ↓                                                      │
│  ELF Emission → .hsaco ready for the GPU                     │
│       ↓                                                      │
│  Your kernel runs on ya silicon                              │
└──────────────────────────────────────────────────────────────┘

すべてのエンコーディングが検証されています llvm-objdump デコード失敗はゼロです。コンパイルには LLVM を使用しませんでしたが、宿題をチェックするために使用しました。

# It's C99. It builds with gcc. There are no dependencies.
make

# That's it. No cmake. No autoconf. No 47-step build process.
# If this doesn't work, your gcc is broken, not the Makefile.

• C99 コンパイラ (gcc、clang など、何でも)
• 生きる意志（任意ですが推奨）
• LLVM は必要ありません。 BarraCUDA は、大人と同じように独自の命令エンコーディングを行います。

# Compile to AMD GPU binary
./barracuda --amdgpu-bin kernel.cu -o kernel.hsaco

# Dump the IR (for debugging or curiosity)
./barracuda --ir kernel.cu

# Just parse and dump the AST
./barracuda --ast kernel.cu

# Run semantic analysis
./barracuda --sema kernel.cu

次の CUDA 機能は、動作する GFX11 マシン コードにコンパイルされます。

• __global__、 __device__、 __host__ 関数修飾子
• threadIdx、 blockIdx、 blockDim、 gridDim ビルトイン
• 構造体、列挙型、typedef、名前空間
• ポインタ、配列、ポインタ演算
• すべて C の制御フロー: if/else、 for、 while、 do-while、 switch/case、 goto/label
• 短絡 && そして ||
• 三項演算子
• テンプレート (基本的なインスタンス化)
• 複数のリターンパス、 continue、 break

• __shared__ メモリ (LDS から割り当てられ、適切に追跡される)
• __syncthreads() → s_barrier
• アトミックな操作: atomicAdd、 atomicSub、 atomicMin、 atomicMax、 atomicExch、 atomicCAS、 atomicAnd、 atomicOr、 atomicXor
• ワープ組み込み関数: __shfl_sync、 __shfl_up_sync、 __shfl_down_sync、 __shfl_xor_sync
• ワープ投票: __ballot_sync、 __any_sync、 __all_sync
• ベクトルの種類: float2、 float3、 float4、 int2、 int3、 int4 と .x/.y/.z/.w アクセス
• 半精度: __half、 __float2half()、 __half2float()
• __launch_bounds__ (解析、伝播、VGPR キャップの適用)
• 協力団体： cooperative_groups::this_thread_block() と .sync()、 .thread_rank()、 .size()
• 演算子のオーバーロード

• フル C プリプロセッサ: #include、 #define/#undef、関数のようなマクロ、 #ifdef/#ifndef/#if/#elif/#else/#endif、 #pragma、 #error、 -I/-D フラグ
• エラー回復 (ハングせずに複数のエラーを報告)
• IRダンプでのソース位置追跡
• 構造体の値渡し

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

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

LLVM は必要ありません:-)

すべてのデータ構造は、事前に割り当てられた固定サイズの配列を使用します。ホットパスにはmallocはありません。再帰はありません。どこにでもある境界ループ。 JPL のコーディング標準委員会が火星への着陸に戻る前に、納得してうなずくような種類のコードです。

制限について正直であることが重要です。足りないものは次のとおりです。

• 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、スイッチ、ショートサーキット
• test_tier12.cu – ベクター、共有メモリ、演算子のオーバーロード
• notgpt.cu – 非常に皮肉なコメントを含む AI 生成の CUDA (タイル化された SGEMM、リダクション、ヒストグラム、プレフィックス スキャン、ステンシル、半精度、協力グループ、および「キッチン シンク」カーネル)
• stress.cu – N ボディ シミュレーション、ネストされた制御フロー、ビット操作、値渡しの構造体、連鎖関数呼び出し
• canonical.cu – NVIDIA サンプルの正規パターンをパーサーに適合させたもの
• test_errors.cu – エラー回復を検証するための意図的な構文エラー
• test_launch_bounds.cu – __launch_bounds__ 解析と VGPR キャップの強制
• test_coop_groups.cu – 協力グループの引き下げ
• さらに、プリプロセッサ テスト、テンプレート テスト、符号なし整数テスト

既知のギャップを修正: 複合代入演算子、裸 unsigned、整数リテラル接尾辞、 const、パラメータの再割り当て。これらはすべてパーサー/下位の小さな変更です。目標は現実世界をコンパイルすることです .cu 変更を加えていないファイル。

生成されたコードは機能しますが、ベンチマークには勝っていません。優先順位:

• 命令スケジューリング (メモリ遅延を非表示)
• レジスタ割り当ての改善 (現在はリニアスキャン、グラフの色付けを考慮)
• 定数フォールディングとデッドコードの除去
• ループ不変コードの動作
• レジスタ圧力に基づいた占有率の調整

IR (BIR) はターゲットに依存しません。バックエンドはきれいに分離されています。新しいターゲットを追加するということは、新しいターゲットを作成することを意味します isel + emit ペア。候補者:

• テントレント – RISC-V ベースの AI アクセラレータ。 ISAを開きます。非常に異なる実行モデル (SIMT ではなくタイルベース) ですが、IR は適切にマップされます。
• インテル アーク – Xe アーキテクチャ。 BarraCUDA は 3 つの主要 GPU ベンダーすべてをカバーします。
• RISC-V ベクトル拡張 – GPU が主流すぎて、ソフトコアで CUDA を実行したい場合。

独自の AMDGPU バックエンドを作成することを検討している場合、午後の時間を台無しにするものは次のとおりです。

• SOP1 プレフィックスは 0xBE800000、ドキュメントから期待されるものではありません
• SOPC プレフィックスは 0xBF000000
• VOP3 VDST はビット単位です [7:0]、 ない [15:8] 賢明な人ならそう思うだろう
• ヌル SADDR は 0x7C グローバルメモリの場合、 0xFC スクラッチ用
• RDNA 3 はデフォルトで Wave32 であり、GCN のような Wave64 ではありません
• ISA マニュアルは 500 ページあり、少なくとも 2 回矛盾しています

全1,735行 amdgpu_emit.c これらのページを読んだ証拠なので、読む必要はありません。

バグが見つかりましたか? AMDGPU 命令エンコーディングのより詳細な点について議論したいですか? GPU コンピューティングの現状について共感できる人が必要ですか?

zanehambly@gmail.com

議論したいことがあれば、問題を開いてください。あるいは、しないでください。私はあなたのお母さんではありません。

ニュージーランドに拠点を置いていますが、すでに明日を迎えており、GPU は他の場所と同様に混乱しています。

アパッチ2.0。やりたいことは何でもしてください。このコンパイラがどういうわけか実稼働環境に入ることができたら、ぜひそのことについて聞きたいです。それは主に、遊びで CUDA コンパイラを書くよりももっと興味深い内容で LinkedIn を更新できるようにするためです。