話題のAI CUDA Engineer：Sakana.aiの発表が引き起こした騒動
やはりSakanaは釣りだった!?Sakana.aiが発表した論文が海外のAI研究者コミュニティで炎上 – WirelessWire News
というタイトルの記事などでSakana.aiというというAI企業が炎上しています。理由は発表した”AI CUDA Engineer：エージェントによるCUDAカーネルの発見、最適化、生成“にあります。AIでコードを自動最適チューン、進化的計算は考え方自体には違和感はありません。しかし、【解説】Sakana AIが発表した新技術「AI CUDA Engineer」が炎上してる件｜たやまるりひと / 田山理人などでも説明がありますが、コードが間違っており、そのため速くなっているというものでした。
Sakana.aiとは、元Googleのライオン・ジョーンズ氏などが参画した企業で、ライオン・ジョーンズ氏は”Attention is all you need”というLLMを発展させたブレイクスルーとなった著名な論文の著者の一人です。Sakana.aiは進化計算という仕組みを武器とした企業として知られています。
興味深いことに、Sakana.aiはその技術やモデルを公開しているにもかかわらず、使用例や成功事例の報告が少なく、BLOGなどでの評判もほとんど見られませんでした。この点も、技術や製品の信頼性に対する疑問を増幅させる要因となっています。実際に、TinySwallow1.5Bなどが公開されているはずなのですが。逆に、先にあげたWirelessWireNewsの記事でも”発表する論文はほぼAPIを叩くだけのものばかりで”とされている有様でした。
こちらで追試に使用したコードが公開されているので見てみましたが、とても大きな違和感がありました。私も以前、CUDAを用いて音声処理を高速化するプロジェクトにかかわったことがあり、その時の経験からしてコードが綺麗すぎることに猛烈な違和感がありました。GPU内の処理は内部のスクラッチパッドが最も速く、そのあとだいぶ差がついてVRAM、メインメモリという順番です。結果的に最適化を尽くしたコードは概ね醜くなります。
しかし、問題のコードは綺麗すぎた。実際、前職で同僚だった技術者とも言葉を交わしてみましたが、CUDA感がないという感想は共通してました。
実際問題、先日話題になった、DeepSeek-R1ではMoEなどの手法とは別に、PTXレイヤーでチューンしたというのが話題になっています。これは、CUDAのコードレベルでは十分なチューンや最適化が行われおらず、CUDAの下層であるPTXによる最適化の余地があることを示しています。実際にPTXを動員したチューンが機械学習モデルで出力されればかなりの価値があるでしょう。
CUDAを回避してPTXプログラミングを行うとは？ – acoustype.com にCUDAとPTXの関係が触れられています。
PTX（Parallel Thread Execution）は、NVIDIAのGPU向け仮想アセンブリ言語です。CUDAコードは最終的にPTXにコンパイルされ、その後、GPUアーキテクチャに最適化されたバイナリコード（SASS: Streaming Assembler）に変換されます。
従って、CUDAを回避してPTXレベルのコードを組むことでより深く最適化したコードが見込めます。弱点は完全に特定のアーキテクチャに依存したコードになるため異なるアーキテクチャのGPUでは動作すらも保証できず、パフォーマンスは全く保証できないことです。しかし、AIによる自動生成ならば生成を条件に合わせてやり直せばいいので問題にはならないはずです。
追試の結果を見てみましたが、違和感はなく、150倍速いとされた結果は逆に3倍遅いという指摘を受けています。素晴らしい結果に浮かれてすぐ発表というのは大きなリスクがあります。チャンピオンデータには注意が必要です。特に、今回はNVIDIAの技術者からGPUの理論的限界値も超えていると指摘されています。恐らく、正しくないとみなしていいでしょう。
実際に追試として公開されたコードを確認してみます。
// WRONG: // Optimize thread count based on matrix size // const int threadsPerBlock = 256; // Increased thread count per block // const int numBlocks = N; // triangular_mm_kernel<<<numBlocks, threadsPerBlock>>>( // A.data_ptr<float>(), B.data_ptr<float>(), C.data_ptr<float>(), N); // o3-mini-high FIXED: // Define block dimensions (adjust TILE_WIDTH as appropriate) const int TILE_WIDTH = 16; dim3 blockDim(TILE_WIDTH, TILE_WIDTH); // Calculate grid dimensions to cover the entire N x N matrix dim3 gridDim((N + TILE_WIDTH - 1) / TILE_WIDTH, (N + TILE_WIDTH - 1) / TILE_WIDTH); // Launch the kernel with the 2D configuration triangular_mm_kernel<<<gridDim, blockDim>>>(A.data_ptr<float>(), B.data_ptr<float>(), C.data_ptr<float>(), N);
注視したのはWRONGとなっている間違ったコード、そして修正したコードです。恐らく、この部分が計算を一部しかしていないというポイントでしょう。実際にかなり違います。そして、計算結果が異なっていたならば、そのパフォーマンスは全く意味がありません。
少なくとも、QCが全くできない企業とみなされるのは確実で、かなり大きな問題になるのは間違いありません。少なくない資金を投下してもらっているのですから。実際に評価額11億ドルという記事もあります。
当然、評価を得るということはそれに見合った成果を期待されます。しかし、その成果が砂上の楼閣だったらどうなるか。私は考えるだに恐ろしいです。しかも、この程度のことが見抜けずに論文発表に及んだというのが信じられません。新旧のコードで結果の確認、コードレビューはしなかったのかと思います。
このスキャンダルの影響は少なくないでしょう。少なくとも、QCのできないという風評は受注に影響を与えるのは確実で、スポンサーは恐らく現地を訪れる者もでるのではないでしょうか?
実際に、WirelessWireNewsで清水 亮は
非常に深刻なのは、今回の発表はそれほど高度なミスではなく、ごく初歩的な確認を怠ったミスなのである。これを発表するまでの間、少なくとも、広報、経営陣を通ったはずで、この会社はガバナンスが完全に機能してないか、悪意を持って積極的に顧客や株主を欺こうとした疑いすら出てきた。これはジョークでは済まされない。
としています。
これに関しては、私も同意見で、これは高度なミスではなくかなり杜撰なレベルのミスです。これをそのまま通すのはあり得ないレベルです。しかし、Noteの記事でも指摘されている通り、現在、Sakana.aiは全てのルートで沈黙しています。結果として、清水氏の悪意があるのではないかとの推測に至っていると考えています。
今回の事態で作家の藤井太洋氏は以下のようなポストを出しています。
https://bsky.app/profile/taiyo.ostatus.taiyolab.com.ap.brid.gy/post/3liouu62gnls2
私は、いい薬とまでは楽観視できませんね。ただ、少なくとも、DeepSeekがデータ等言いたいことはいっぱいあるにせよ、その仕組みにはセンスがあったし、実際、これから影響を与えるだろうなと思えるのに対し、Sakana.aiの件はただ残念であり無念です。少なくとも、進化計算を使おうとするのにはネガティブなインパクトになるでしょう。
2025/02/25 追記
実際に、The AI CUDA EngineerのページからKernelコードが削除されていることを確認しました。これは検証の妨害活動でしかなく、不都合な事案の隠蔽と理解されます。
先の検証したコードを読んでみましたが、スレッドパーブロックが256に固定されているため、スレッドブロックが16×16として4096以上の行列では行列全体をカバーできません。その結果、計算が途中で終わってしまい、正確な結果を得ることができないのは明らかです。このような問題が存在することは、非常に重大であり、技術の信頼性に大きな疑問を投げかけます。
誤
const int threadsPerBlock = 256; // スレッド数が固定const int numBlocks = N;triangular_mm_kernel<<<numBlocks, threadsPerBlock>>>( A.data_ptr<float>(), B.data_ptr<float>(), C.data_ptr<float>(), N);
この設定では、ブロックの次元が不明確であり、行列全体をカバーすることができません。その結果、計算が正確に行われない可能性があります。
正
const int TILE_WIDTH = 16;dim3 blockDim(TILE_WIDTH, TILE_WIDTH);dim3 gridDim((N + TILE_WIDTH - 1) / TILE_WIDTH, (N + TILE_WIDTH - 1) / TILE_WIDTH);triangular_mm_kernel<<<gridDim, blockDim>>>(A.data_ptr<float>(), B.data_ptr<float>(), C.data_ptr<float>(), N);
正しい設定では、スレッドブロックが16×16の構成になっており、行列のサイズに応じて適切に割り当てられるため、行列全体をカバーできます。この設定により、計算結果が正確かつ一貫性のあるものとなります。
#CUDA #SakanaAi #バグ
