BarraCUDA - AMD GPU向けオープンソースCUDAコンパイラ

 (github.com/Zaneham)
5 ポイント 投稿者 GN⁺ 2026-02-18 | 1件のコメント | WhatsAppで共有
  • CUDA Cソース（.cu） を AMD RDNA3（GFX11） 向けの機械語へ直接変換する スタンドアロンのオープンソースコンパイラ
  • LLVMやHIPレイヤーなし に、独自の 字句解析器、パーサー、中間表現（BIR） を通じて ELF .hsaco バイナリを生成
  • 約15,000行のC99コード で書かれており、単一の make コマンド でビルド可能
  • CUDAの スレッド組み込み変数、共有メモリ、アトミック演算、ワープ演算、協調グループ など主要機能をサポート
  • Apache 2.0ライセンス で公開されており、今後は Tenstorrent、Intel Arc、RISC-V など追加アーキテクチャへの拡張を目指す

BarraCUDA 概要

  • BarraCUDAは AMD GPU向けCUDAコンパイラ で、.cu ファイルを GFX11機械語コード に変換
    • 出力物はAMD GPU上で実行可能な ELF .hsaco バイナリ 形式
    • LLVM依存なし で完全に独立して動作
  • コード全体は C99で書かれた約15,000行 で、単一のMakefile でビルド可能
  • このプロジェクトは ニュージーランド拠点の開発者 が個人で開発

動作方式

  • 入力された .cu ファイルを 前処理 → 字句解析 → パース → 意味解析 → BIR生成 → 命令選択 → レジスタ割り当て → バイナリエンコード → ELF出力 の順で処理
  • BIR（BarraCUDA IR） はSSA形式の内部表現で、アーキテクチャ非依存の設計
  • すべてのエンコードは llvm-objdump を通じて検証され、デコードエラーは0件

対応機能

  • CUDAコア構文: __global__, __device__, __host__, threadIdx, blockIdx など
  • CUDA機能: __shared__ メモリ、__syncthreads(), アトミック演算、ワープシャッフル／投票、ベクター型、half精度、協調グループ など
  • コンパイラ機能: 完全なCプリプロセッサ、エラー回復、ソース位置追跡、構造体の値渡しをサポート

未対応項目

  • unsigned 単独使用、複合代入演算子（+=, -= など）、const, __constant__ メモリ、2D共有配列、テクスチャ・サーフェス、動的並列実行などは未実装
  • 複数の翻訳単位およびホストコード生成には未対応

テストとロードマップ

  • 14個のテストファイル、35個以上のカーネル、約 27KBの機械語コード で検証
  • 短期目標: 構文補完と実用的な .cu ファイルとの互換性強化
  • 中期目標: 命令スケジューリング、レジスタ割り当て改善、定数畳み込み、ループ不変コード移動などの最適化
  • 長期目標: Tenstorrent、Intel Arc、RISC-V Vector Extension など新しいバックエンドの追加

ライセンス

  • Apache 2.0ライセンス

関連記事

1件のコメント

 
GN⁺ 2026-02-18
Hacker Newsのコメント

  • プロジェクトのREADMEに書かれたニュージーランド風ユーモアが印象的だった
    LLVMに依存せず独自に命令エンコーディングを実装している点が新鮮
    こうしたプロジェクトを始めるには並外れた低レベルの知識が必要だと感じさせる
    AMD陣営ではCUDA非対応がNVIDIA独占の言い訳になりがちだが、こうした試みは市場の均衡に役立ちそう

    • 元文で言っていたのはLLMではなくLLVMだと指摘
    • このプロジェクトにも明らかにAI関連の貢献は含まれていると思う。ただ、うまく使えばAIは優れた道具になり得る
    • 「オセアニア風ユーモア」という話から、Beached Whaleアニメーションを思い出した
    • ソースコードのコメント「/* 80 keywords walk into a sorted bar */」が気が利いていた (lexer.cリンク)
    • 「AI slope」という表現を見て、科学的にはgradient descentと呼ぶという冗談を付け加えた

  • 最初の外部Issueをgeohotが投稿していたことに驚いた (Issueリンク)
    こういう人物たちが力を合わせてNVIDIAのGPU市場独占を崩すところを見てみたい

  • AI推論ワークロードをNVIDIA GPUで回すのはコスト負担が大きい
    こうしたプロジェクトは、スタートアップが手の届く代替手段を作るうえで重要
    conv2dやattentionのような演算での性能が気になる

  • 「# It’s C99. It builds with gcc. There are no dependencies.」
    make一発で終わる単純さが本当に美しい

    • こういうシンプルで直感的なアプローチがとても好き

  • 投稿タイトルの大文字表記を見て、会社のGPUファーム名が「barracuda」な理由を今さら理解した。かなり笑えた

  • もし情熱的な開発者たちがAMDにできなかったことをやってのけるなら、笑えると同時に悲しくもあるだろう

    • AMDがCUDAをサポートしなかったのは「できなかった」からではなく戦略的選択だったと思う
      CUDAをサポートすると、かえってNVIDIAのエコシステムを強化することになるからだ
      CUDA代替を望むならZLUDAのほうが実用的かもしれない
    • オープンソースは株主の顔色をうかがう必要がないため、しばしば商用より優れた結果を出す
      ただ、大きくなると結局大企業に買収されて消えていくのが残念だ
      Linusとgitの例のように、意志と知識さえあれば壁を破れる
    • AMDは「できなかった」のではなく「やらなかった」ことが多い
      たとえばFSR4を旧世代カードで公式サポートしないのもその一例だ
    • 「うちにはHIPがある」という冗談で締めくくった

  • 旧世代AMDアーキテクチャ（GFX1010など）までサポートできるのか気になる

    • 難しいが十分可能だと思う。以前はROCmがひどくて自分でパッチを当てていたが、今はずっと良くなった
    • 自分もRX5700mを使っているので、旧世代AMDサポートの作業を直接進めている
      ISAドキュメントを読みながらバイナリエンコーディングを調整中
      ただしこの領域はLLMがあまり得意ではない分野なので、自分で理解して修正しなければならない

  • CUDAはC++をサポートしているのに、Clang/LLVMなしで純粋なCで進めるのは制約が大きいのではと気になった

    • 実際にはCUDAが使うC++機能の一部だけをパースしている。コンパイラ自体はC99で書かれている
    • テストを見るとテンプレートのような一部のC++機能も処理できそうだ
    • LLVMのAMD GX11対応は非常に限定的なので、依存しないほうがよい
    • 本物の開発者はAIに依存せずコードを書く
      最近はオープンソースにAI生成のPRがあふれているが、このプロジェクトはそうした依存を避けている点が印象的だった

  • AMDはこうしたプロジェクトを公式に支援すべきだと思う
    世界はNVIDIAの独占から抜け出す必要がある

  • 「OpenCLはもう終わったのか？」という疑問が湧いた
    よく分からないが、それでもこのプロジェクトは印象的だ

    • CUDAのほうがはるかに大きな市場シェアを取っており、OpenCLは基盤が弱い
      昔はAppleもサポートしていたが、今は違うようだ
      まるでUnixとWindowsの関係のように、技術的には優れていても市場は一方に偏った