- HipKittens は、AMD GPU向けの 高性能カーネルとC++ベースのプログラミングプリミティブ を提供し、AI演算の効率を高めるプロジェクト
- 既存のAMDエコシステムである AITER、PyTorch、Triton、TileLang、Composable Kernel などは、性能の不安定さと未成熟なサポート により限界を露呈
- HipKittensは タイルベース抽象化 (tile abstraction) を中心に、NVIDIAとAMDの間で共通インターフェース を保ちつつ、ハードウェアごとの実装を分離
- 約 500行以下のコードで書かれたカーネル が、既存のAMDの手書きアセンブリカーネルより高速な性能を達成
- AI演算を マルチシリコン環境へ拡張 するための実践的な基盤を示し、オープンなハードウェアエコシステム への移行可能性を提示
AMD GPUエコシステムの現状と課題
- AI演算はこれまで 単一ハードウェアベンダー中心 に発展してきたが、AMD GPUは 最新世代で最高水準の演算性能とメモリ帯域幅 を提供
- 例: AMD MI355X OAMは BF16 2.5 PFLOPs、MXFP8 5.0 PFLOPs、メモリ 288GB、帯域幅 8.0TB/s
- しかし、成熟したソフトウェアスタックの欠如 により、実際のAIワークフローでこの性能が活用されていない
- AMDの主要なソフトウェア構成要素は AITER、PyTorch、Triton、Mojo、TileLang、Composable Kernel (CK) などで構成
- AITERとPyTorch のSDPA Llama GQA逆伝播カーネルは、それぞれ SoTA比で30%、24%程度の性能 にとどまる
- Mojo はバンク競合 (bank conflict) により 50%程度の性能 にとどまる
- TileLang はCDNA3までしか対応しておらず、中核機能の不足とCK依存 により複雑性が増大
- Triton はレジスタ寿命追跡とメモリアクセス最適化に課題がある
- 結果として、最高性能のAMDカーネルは専門家がアセンブリで直接記述しなければならない状況 であり、これは拡張性と保守性に大きな制約となっている
CUDA中心エコシステムとの比較
- AIコミュニティでは CUDAエコシステムの参入障壁 (CUDA moat) が存在すると評価されている
- 過去のNVIDIAカーネル開発でも、低レベルのCUDA/CUTLASSベースで何年にもわたる努力 が必要だった
- 最近では DSL (ドメイン固有言語) と AI支援ツール の発展により、NVIDIAカーネル開発は簡素化されている
- 例: ThunderKittens (TK)、CuTe DSL、TinyGrad “tinykittens”、TileLang、Gluon など
- こうした流れを踏まえ、AMDでも 新しいプログラミングプリミティブの必要性 が模索されている
HipKittensの設計原理
- HipKittensは ThunderKittens (NVIDIA)、ThunderMittens (Apple Silicon) に続く AMD向けに開発されたバージョン
- 中核概念: タイルベース抽象化 (tile abstraction)
- タイル抽象化の汎用性 – NVIDIAで効果的だったタイルベース演算モデルはAMDでも自然に適用できる
- バックエンド実装のアーキテクチャ特化 – メモリアクセスパターンとレジスタスケジューリングはハードウェアごとに異なる設計
- スケジューリング戦略の適応性 – AMDのCDNA3・CDNA4ではwaveベースのスケジューリングは非効率だが、タイル単位のスケジューリングは依然として単純さと性能を維持
- インターフェース (タイルおよび演算) と 実装 (ハードウェアマッピング) を分離することで、多様なGPUアーキテクチャに共通適用できるモデル を提示
性能結果
- Attention Forwardカーネル: 約 500行のコード で書かれており、AITERのアセンブリカーネルより優れた性能 を達成
- さまざまな ヘッド次元 (D) と シーケンス長 (N) で、Causal/Non-Causal 設定の両方において優位
- GEMMカーネル: 100行未満の中核ループ で構成され、BF16・FP8演算の両方で最高性能 を達成
- Attention Backward、Rotary、Fused Dropout-Residual-LayerNormカーネル でも 既存最高性能を上回る改善 を確認
- すべてのカーネルが 可読性と修正のしやすさ を維持しながら、手書きアセンブリ並みの性能 を確保
マルチシリコンAIへの拡張
- AIの潜在力を実現するには、多様でオープンなハードウェアの活用 が必要
- HipKittensはAMD GPUを AI開発者が実用的に利用できるプラットフォーム にすることを目指す
- AMDの オープンソースソフトウェアスタック と組み合わせることで、マルチシリコンベースのAIエコシステム への移行可能性を示す
- 続編ではHipKittensの 技術的な詳細構造 を扱う予定 (part two予告)
1件のコメント
Hacker Newsのコメント
私たちはAMDと契約し、Llama 405BをMI350XでMLPerf向けに学習させています。
AMDの状況は明らかに良くなっています。AMD GPUを持っているなら、pytorch.orgでLinux+ROCmを選んでPyTorchをインストールしてみることを勧めます。3年前は絶望的でしたが、今では主要機能のほとんどがきちんと動きます。MI300Xでnanochatを動かしましたが、そのまますぐに動きました。MI350Xも同様に安定しています。
もちろんNVIDIAにはまだ後れを取っており、ソフトウェアエコシステムとコンパイラ、ドライバへの多くの投資が必要です。しかし2年前の絶望的な状況と比べれば、今は希望が見えます。
AMDのLLVMバックエンドの足りない部分を見たければ、HipKittensのコードをCUDA Kittensと比較してみるとよいです。
学習用途ではNVIDIAとGoogleが1位、AMDが2位で、他はほぼ存在感がありません。IntelとTenstorrentはまだ遠く、Huaweiはサンプルがsegfaultで落ち、Groqはチップ販売を断念し、Cerebrasはそもそも入手できません。Trainiumはインスタンスを1つ受け取るのに5日かかり、興味を失いました。
初期設定はまだ少し荒いですが、以前と比べればずっと良くなりました。たとえば1か月前まではnanochatがあまりうまく動きませんでしたが、今は動きます。重要なのは、人々が今やAMDエコシステムに関心を持っていることです。
AIにはハードウェアの多様性が必要です。すべてのAIハードウェアとソフトウェアを1社が独占するのは、株主には良いかもしれませんが、技術の進歩には有害です。
NVIDIAの企業価値が理解できません。今はTransformerのような少数のアルゴリズムが勝ち残り、データの方が重要になっています。昔のように多様なHPCコードが必要ではないので、今や競合は狭いアルゴリズム群だけを最適化すれば済みます。
vLLMとTransformerを効率よく動かせるなら巨大な市場が開けます。だとすればAMDやHuaweiも簡単に追いつけそうですが、NVIDIAの**本当の堀（moat）**は何なのか気になります。InfiniBandだけで十分なのでしょうか?
データセンター向けGPUでは依然としてNVIDIAが強いですが、GoogleはTPUを大規模に使っており、OpenAIもAMDハードウェアを発注しています。
CUDAエコシステムは依然として重要ですが、誰もがそこから抜け出そうとする動きが活発です。AMD、Intel、Qualcommなどもこの競争に参入しています。HipKittensは中立的なソフトウェアエコシステムへ向かう重要な一歩に見えます。
AMDがこうしたプロジェクトに資金提供していてもおかしくないのに、私の考えではまた機会を逃したのだと思います。GPUやAIではいつもそうでした。
HipKittensは改善ですが、AMD内部にはカーネル性能を追跡する能力が不足しています。DevOpsはインドのTCSに外注されており、先方は状況をよく把握していません。
良いリーダーがいるチームは独自にシャドーITチームを運営しています。ROCmにはそうしたチームがなく、結局は大手クラウド顧客が抗議して改善が始まりました。
人材を採用しても、提示する給与は市場以下です。QualcommやWalmart水準と比較しているからです。
この4年間、ボーナスが満額支給されたことはありません。
過去にはRadeon VIIのようにサポートを打ち切ったり、エコシステムを頻繁に作り直したりして、成熟できませんでした。CUDA互換性の不足と投資不足のため、ほとんどの組織は結局NVIDIAを選びます。
それでも最近はROCmとInstinctエコシステムに継続して投資しており、少しずつ改善しています。ただしネットワーキング分野では、依然としてNVIDIAがはるかに先を行っています。
ThunderKittensの上に構築されたプロジェクトがあるのか気になります。
これがHIPに移植された版なら、単なるCUDA移植よりずっと実用的価値があるように思えます。
「raw assembly vs cooked assembly」という表現が興味深いです。
昔はCPU向けにアセンブリコードを直接書くのが一般的でした。GPUでも必要以上に恐れることはありません。結局はコンパイラがそのコードを生成するのです。
数学的に見れば、推論（inference）は基本的な線形代数/BLAS演算です。
dtypeとsparsityを考慮し、ユースケースの80%をカバーする簡潔な推論APIが可能なのではないかとも思います。CUDAのように複雑である必要はなさそうです。
composable-kernelは、私のGentooシステムで**OOM（メモリ不足）**を最も頻繁に引き起こしたソフトウェアです。ClangがCKをコンパイルするとき、スレッドあたり2.5GB以上を使います。
公式ビルドサーバーではアーキテクチャ数を減らす必要がありそうです。依存パッケージの大半ではヘッダーだけが必要だと聞いています。ビルド時間短縮のための改善作業が進行中です。
device_grouped_conv2d_fwd_xdl_ngchw_gkcyx_ngkhw_f16_instance、いったいclangに何をさせているのか分かりません。こうした進展によって、コンシューマー向けAMDチップでもLLMをうまく動かせるようになるのか気になります。
たとえばAMD CPU/iGPUとRTX 5080を一緒に積んだHP OMEN MAX 16ノートPCを検討していますが、AMD側はRTXと競争できるでしょうか?
このブログ記事や高性能Macの結果を見ると興味深いです。
Sonnet 4 APIより少し遅いですが、ローカルでこの程度の性能なら十分満足です。
Opencode + Ollama + Qwen3 Coderの組み合わせは、ClaudeCode + Sonnet4の優れた代替です。
ただし、AIにコーディングをすべて代行させたいなら話は別かもしれません。それでも個人用アシスタントとしては完璧です。
AMDがB300を完全に無視した理由が分かりません。