ハクソク

世界を動かす技術を、日本語で。

AMD GPUを対象としたオープンソースのCUDAコンパイラ「BarraCUDA」

97日前原文(github.com)

概要

BarraCUDAは AMD GPU向け のオープンソースCUDAコンパイラ。 LLVMやHIPに依存せず、.cuファイルから直接GFX11バイナリを生成。 約 15,000行のC99 で実装され、シンプルなビルド手順を実現。 主要なCUDA機能とC言語機能に対応し、今後さらに 他アーキテクチャも対応予定。 制限事項や今後のロードマップも明示されており、開発の透明性を確保。

BarraCUDAとは

  • AMD GPU向け のオープンソースCUDAコンパイラ
  • LLVM非依存、HIP変換レイヤーも不要
  • .cuファイル を直接GFX11マシンコード（ELF .hsacoバイナリ）に変換
  • 15,000行以上の C99 による実装
  • 依存ライブラリなし、gcc等でmake一発ビルド可能
  • 現状、 AMD RDNA 3（gfx1100） に対応

パイプライン概要

  • プリプロセッサ ：#includeや#define等の前処理
  • 字句解析 ：トークン分割
  • 構文解析 ：再帰下降パーサでAST生成
  • 意味解析 ：型チェック・スコープ解決
  • BIR（独自IR） ：SSA形式による中間表現
  • mem2reg ：allocaをSSAレジスタに昇格
  • 命令選択 ：BIR→AMDGPU命令
  • レジスタ割り当て ：VGPR/SGPR割り当て
  • バイナリエンコード ：GFX11命令語化
  • ELF出力 ：.hsaco生成
  • カーネル実行 ：GPU上で実行可能

特徴・強み

  • LLVMやHIP不要、独自で命令エンコードを実装
  • エラー回復機能 （複数エラーを一度に報告）
  • 全データ構造は固定長配列、ヒープ割り当てなし
  • 再帰なし・有界ループのみ、堅牢な実装
  • IRやASTのダンプ機能 （デバッグ・解析用）
  • 全エンコーディングをllvm-objdumpで検証済み

対応機能（CUDA/C言語）

  • global, device, __host__関数修飾子
  • threadIdx, blockIdx, blockDim, gridDim等の組込変数
  • 構造体、enum、typedef、namespace
  • ポインタ、配列、ポインタ演算
  • C言語制御構文（if/else, for, while, do-while, switch/case, goto/label）
  • 論理演算子（&&, ||）、三項演算子
  • テンプレート（基本的なインスタンス化）
  • CUDA共有メモリ（shared）、バリア（__syncthreads→s_barrier）
  • アトミック演算（atomicAdd, atomicSub, atomicMin, 他）
  • ワープ命令（__shfl_sync, __ballot_sync等）
  • ベクタ型（float2, int3等）と.x/.y/.z/.wアクセス
  • 半精度型（__half, __float2half等）
  • __launch_bounds__対応（VGPR制限も反映）
  • Cooperative Groups（cooperative_groups::this_thread_block()等）
  • 演算子オーバーロード
  • Cプリプロセッサ全対応（#include, #define, #ifdef, #pragma, #error, -I/-D等）
  • ソース位置情報付きIRダンプ

使い方

  • AMD GPUバイナリ生成 ./barracuda --amdgpu-bin kernel.cu -o kernel.hsaco
  • IRダンプ ./barracuda --ir kernel.cu
  • ASTダンプ ./barracuda --ast kernel.cu
  • 意味解析のみ実行 ./barracuda --sema kernel.cu

ビルド方法

  • gcc等のC99コンパイラ でmake実行のみ
  • make失敗時はgcc側の問題

ソース構成（主なファイルと役割）

  • lexer.c（字句解析、747行）
  • preproc.c（プリプロセッサ、1,370行）
  • parser.c（構文解析、1,500行）
  • sema.c（意味解析、1,725行）
  • bir.c + bir_lower.c（SSA IR、3,032行）
  • bir_mem2reg.c（mem2reg、965行）
  • amdgpu_isel.c（命令選択、1,788行）
  • amdgpu_emit.c（レジスタ割当・バイナリエンコード・ELF、1,735行）
  • main.c（CLIドライバ、317行）

未対応・制限事項

  • unsigned単体指定（unsigned intで代用）
  • 複合代入（+=, -=, >>=等、明示的に記述）
  • const修飾子、__constant__メモリ
  • 共有メモリの2次元配列宣言（1次元にフラット化推奨）
  • 整数リテラルサフィックス（0xFFu, 1ULL等）
  • __device__関数でのパラメータ再代入（ローカル変数で代用）
  • テクスチャ・サーフェス、動的並列処理、複数翻訳単位
  • ホストコード生成（デバイスコードのみ対応）

テストスイート

  • 14ファイル、35+カーネル、約1,700 BIR命令、27KBマシンコード
  • vector_add.cu（GPU計算のHello World）
  • cuda_features.cu（アトミック、ワープ命令、バリア等）
  • test_tier12.cu（ベクタ型、共有メモリ、演算子オーバーロード）
  • notgpt.cu（AI生成CUDA、半精度、協調グループ等）
  • stress.cu（N-body、ビット演算、構造体値渡し等）
  • canonical.cu（NVIDIAサンプル準拠）
  • test_errors.cu（エラー回復テスト）
  • その他、launch_bounds・cooperative groups等のテスト

今後のロードマップ

  • 短期 ：未対応構文（複合代入、unsigned単体、const等）の対応
  • 中期 ：最適化（命令スケジューリング、レジスタ割当、定数畳み込み、ループ最適化等）
  • 長期 ：新アーキテクチャ対応（Tenstorrent、Intel Arc、RISC-V Vector Extension等）

GFX11エンコーディング注意点

  • SOP1プリフィックスは0xBE800000
  • SOPCプリフィックスは0xBF000000
  • VOP3 VDSTは[7:0]ビット
  • Null SADDRはグローバル0x7C、スクラッチ0xFC
  • RDNA 3はWave32がデフォルト
  • ISAマニュアルは500ページ、矛盾あり
  • amdgpu_emit.c（1,735行）はこれらの検証の証

ライセンス・連絡先

  • Apache 2.0ライセンス、商用利用含め自由
  • バグ報告・議論歓迎（zanehambly@gmail.com、New Zealand在住）
  • GitHub Issueも利用可能
  • 実運用事例の報告も歓迎

BarraCUDAは、 NVIDIAの独自エコシステムに対する挑戦 として生まれた、 シンプルかつパワフルなAMD GPU向けCUDAコンパイラ。 今後の拡張や最適化にも積極的で、 GPUコンパイラに興味のある開発者に最適なプロジェクト。

Hackerたちの意見

CUDAの開発には詳しくないけど、CUDAってC++をサポートしてるんじゃなかったっけ？Clang/LLVMを飛ばして「純粋な」Cに行くのは、ちょっと制限が多い気がするな。

└

本当の開発者は、良質なコードを書くのにAIに依存しないんだ。実際、左右に飛び交うスロープコードの量はLLMのせいだよ。オープンソースプロジェクトはAIからのPRで溢れてるけど、それに依存しないことがプロジェクトを制限するわけじゃない。このプロジェクトのオーナーは、何が起こっているのかをしっかり理解していて、依存関係をなくすのは簡単なスキルじゃないよ。多くの開発者は、依存関係だらけのコードを書いたり、LLMからコピペしたりすることが多いけどね。後者のことを「コーディング」って呼ぶ人もいるよ :)

└

正直、LLVMがAMDのGX11マシンコードをどれだけサポートしてるのかはわからない。かなりニッチなバックエンドだし、存在しても理想的な出力が得られないかもしれない。しかも、大きな依存関係があるしね。

└

それについても気になってる。コンパイラ自体はC99で書かれてるけど、テストを見る限り、テンプレートみたいなC++の機能も解析できるみたい。

これはC99だよ。gccでビルドできるし、依存関係もない。 > make Beautiful.

└

これ、好きだわ。シンプルでストレート。

AMDができなかったことを、熱心な人たちがやり遂げたら面白いし、悲しいよね :)

└

多くのプロジェクトは、オープンソースだからこそ株主を喜ばせる必要がなく、独自のものよりもはるかに良くなることが多いんだ。残念なのは、そういうプロジェクトが大きくなりすぎて、騒がしくなってしまい、大手企業が買収してしまうこと。結局、みんなが何も得られないことになる。独自の囲い込みを打破するには、知識と実行する意志を持った誰かが必要なんだ。LinusがgitとLinuxを作ったのがその完璧な例だよ。面白い事実として、BitKeeperは2005年にLinuxコミュニティに「くたばれ」と言ったけど、Linusは10日以内にgitを作ったんだ。BitKeeperは2016年にコードをオープンソースにしたけど、その頃には誰も彼らが誰か知らなかったよ、笑。だから、時間を与えてあげてね :)

└

AMDがCUDAをサポートしていないのは、単に「できなかった」わけじゃないよね（確かにソフトウェアは全体的にイマイチだけど）。明らかに戦略的な決定だと思う。AMDでCUDAをサポートしたら、NVidiaの優位性がさらに強まるだけだし、競合にGPUプログラミング環境を全部譲る理由なんてないよ。実際、これはいい賭けだったと思う。時が経つにつれて、CUDAはどんどん重要性が薄れてきてるしね。それに、CUDAの代替を実用的に考えるならZLUDAがいいよ。このプロジェクトは面白いし、ちょっとクールだけど、Cのサブセットに制限されてて、置き換えライブラリ（BLASやDNNなど）がないから、比較的役に立たないかな。

└

家にはHIPがあるよ。

└

できなかった というより「しなかった」って感じかな、ほとんどの場合。これって他のいくつかのことにも当てはまるよね？今のところ、古いカードでのFSR4がその一例だし、AMDはまだ公式にはサポートしてない。でも、サポートすると思うよ。周りのネガティブな意見が多すぎるし。r/AMDの投稿の半分はそのことで文句言ってる人たちだし。

HIPの翻訳レイヤーはなし。ストレージ容量がどこでも喜んでる。

要件 > 生きる意志（オプションだけど推奨） > LLVMは必要ない。BarraCUDAは大人のように自分で命令エンコーディングをする。 > 話したいことがあればイシューを開いてね。そうじゃなければ、別にいいよ。私はあなたのお母さんじゃないから。 > ニュージーランドに拠点を置いてるから、オセアニアのユーモアは他にはないよ、ハハ。このプロジェクトのオーナーは、LLMに依存しないことを強調してる。AIが流行ってる中で、これは本当に新鮮だね。こんなプロジェクトを始めるのに必要な知識の量は、ほんとに別次元だし、機械語レベルでマニュアルを覆すのは全く別の話だよ。AMDに関して言えば、「CUDAサポートなし」がNVIDIAの囲い込みに参加する最大の「言い訳」だね。このプロジェクトに幸運を祈るよ。競争が増えれば、NVIDIAがPCパーツの価格を壊し続けることも難しくなるから。

└

プロジェクトのオーナーは、LLMに依存しないことを強調してる。AIが流行ってる中で、これは本当に新鮮だよ。プロジェクトのオーナーが言ってるのはLLVM、コンパイラツールキットのことだから、LLMじゃないよ。

└

このプロジェクトには確実にAIへの貢献があるね。でも、あんまり気にしてない。AIは熟練した手で驚くべきことができるし、このプロジェクトを使うのが楽しみだよ。

└

マシン言語レベルでマニュアルが間違っていることを証明するなら、ここで水を差す役になりそうだな。マニュアルはまだ正しいし、リバースエンジニアリングでAMDが選んだシリコンのアーキテクチャを修正することはできないよ。ラスタGPUでCUDAの機能のサブセットを実装することは絶対可能だけど、OpenCLからずっとそうしてきたし、CUDAはまだ王様だね。業界がクソみたいなGPGPU計算標準に収束するのが一番いいことだと思う。でも、Intel、AMD、Appleはみんな対立してるから、CUDAの業界に対するヘッジベットは今後も利益を生むだろうね。

└

オセアニアのユーモアは他に類を見ないね、笑った！ビーチに打ち上げられたクジラのアニメーション短編を思い出したよ。[1]

└

LLVMは必要ないよ。BarraCUDAは大人みたいに自分で命令エンコーディングをやってる。 > プロジェクトのオーナーはLLMに依存しないことを強調してるし、AIの流れの中ではすごく新鮮だね。「シリコンバレーのテックインサイダーサイトから技術リテラシーは消えたのか？そんなことは信じられない。あなたたちの中に本物のエンジニアはいないのか？！」

└

関係ないけど、NZから帰ってきたばかり。素晴らしい人たちだったよ。

└

AMDがこれを最優先にしてないのがまだ信じられない。何年も言ってるけど、もし俺がAMDだったら、CUDA互換レイヤーを作るために何十億も使うよ。絶対に元は取れるし、そんなにお金はかからないはず。

└

LLVMは、LLMとは関係ないよ。

└

「これがうまくいかないなら、あなたのgccが壊れてるだけで、Makefileのせいじゃない。」... いや、マジで自信満々だな。

└

プロジェクトのオーナーは、LLMに依存しないことを強調してる。AIのスロープの世界では、これはすごく新鮮だね。え？ これはREADMEから明らかにAIのスロープだよ。行の終わりにずれた「ASCIIアート」の図を見てみて。あれは明らかにAIのスロープの証拠だよ。手で編集する人なら、余分なスペースを消して整列させるのが普通だから。

いいね！誰かがNvidiaのソフトウェアの壁を破るのは時間の問題だったよ。Nvidiaの弁護士が君の住所を知らないことを願うよ。

└

でも、これはプロダクション向けのものじゃないね。

これが古いAMDアーキテクチャ、例えばGFX1010、つまり初代RDNAシリーズにまで対応するのはどれくらい現実的なのかな？一番貧弱なGPUだし。

└

誰にも諦めさせないで。面倒くさいけど、できるから。ディフュージョンが新しかった頃、rocmがまだぐちゃぐちゃだった時に、手動でたくさんパッチ当てて、viiや1030、1200をなんとか動かしてたんだ。今は昔よりずっとマシだし、AMDは本当に評価されるべきだよ。環境が整えばCodexはそれを圧倒できるはず。

コードを見てみると、翻訳はかなり複雑そうだね。ここでも人気のあるhttps://github.com/vosen/ZLUDAに感謝！Zludaはcomfyuiと一緒にうまく動くようになったよ。

作者以外の誰かが作った最初の問題は、geohot本人からのものだよ.. やばい人だね: https://github.com/Zaneham/BarraCUDA/issues/17 こういう人たちが協力して、NVIDIAのGPU市場に対する独占を打破してくれるといいな。ネットによると、彼らは70%もの利益率を得てるらしいから、全ユーザーにとってコストが上がっちゃうんだよね。