2026年9月5日 土曜日
AI時短ラボ
研究· 約11

LLMが書いたGPUカーネルの「正解率」は本当か──12ゲート検証器が2,638件を洗い直すと4割近くが壊れていた

arXivの論文は、LLM生成GPUカーネルの正しさを「少数のランダム入力・1つの固定形状での1回テスト」で判定する既存手法の甘さを指摘し、12個の敵対的ゲートからなる検証器を提案する。公開システムが正しいと判定していた2,638個のカーネルを検証し直したところ、39.5%が許容誤差の議論では説明できないレベルで壊れていた。

LLMが書いたGPUカーネルの「正解率」は本当か──12ゲート検証器が2,638件を洗い直すと4割近くが壊れていた
執筆・編集:
目次

3行まとめ

  1. LLMがGPUカーネルを生成するシステムの多くは「少数のランダム入力を1つの固定形状で1回テストする」という緩い基準で正解率を報告している、と論文は指摘する。
  2. 著者らは12個の敵対的ゲート(一部は閾値なし=どんな許容誤差の議論でも言い逃れできない)からなる検証器を構築し、ある公開システムが自身のハーネスで「正しい」と判定していた2,638個の生成カーネルを再検証した。
  3. 結果、39.5%が許容誤差の議論では説明できないレベルで壊れており、62.1%が何らかの違反を抱えていた。既存の標準テストが「合格」とし検証器が「不合格」とした件数は1,487件、逆(検証器が合格・標準テストが不合格)はわずか14件だった。

「たまたま正しく見える」を見抜く

論文が問題視するのは、GPUカーネルをLLMに生成させるシステムの評価方法そのものだ。多くのシステムは、少数のランダム入力を1つの固定形状で走らせ、出力が参照実装に近ければ合格、という単一の緩いテストで高い正解率を報告している。しかし、このテストに合格しながら実際には静かに間違っているカーネルはいくらでもあり得る、と論文は列挙する——本当はNaNや無限大を返すべきところで普通の数値を返す、実行のたびに結果が変わる、入力形状が変わると壊れる、参照実装がfp32で総和を取っているところをfp16で累積してしまう、といったパターンだ。

これに対して著者らが構築したのが「契約グレード(contract-grade)」の検証器で、正しいカーネルが満たすべき12個の性質を敵対的なゲートとして実装している。そのうち複数は「閾値なし」——つまりどんな許容誤差の取り方を持ち出しても、その失敗を説明で言い逃れることができない設計になっている。

12個のゲートの中身

論文のHTML全文版(arXiv:2608.12700のexperimental HTML)にある第3.1節「The twelve gates」を実際に読むと、12個のゲートそれぞれが何をチェックしているかが具体的に書かれていた。

ゲート チェックする性質
CMP-01 ゼロ・10⁶・10⁻⁶・非正規化数・長い系列長など、多数のランダム・敵対的入力で正しいか
CMP-02 出力だけでなく勾配も正しいか(自動微分 vs 有限差分)
CMP-03 テストした1つの形状だけでなく、バッチ・系列長・幅を変えても正しいか
ORD-01 加算順序を変えても、導出された√Nスケールの丸め誤差の範囲内に収まるか
ORD-02 5回繰り返して1バイトも違わず同一の結果になるか、共有バッファをエイリアスしていないか
ORD-03 悪い加算順序を意図的に誘発する入力でも正しいか
PRC-01 fp32・fp16・bf16のいずれでも正しいか
PRC-02 fp16入力を与えられても、内部でfp32の累積を保っているか
EXC-01 NaN・符号付き無限大を既知の位置に散りばめた入力で、出力の非有限値の位置・符号が参照実装と完全一致するか
EXC-02 非正規化数のflush-to-zero処理が参照実装と一致するか
RES-01 出力が入力と同じデバイス上にあるか
RES-02 コンパイル済みカーネルが実際のハードウェア制約(レジスタ・共有メモリ・TMEM予算)に収まるか

出典: arXiv:2608.12700 HTML全文版「Table 1」(2026年8月30日確認)

このうちEXC-01EXC-02ORD-02など複数は「閾値なし」グループに属し、完全一致(バイト単位・マスク単位)でしか合格を認めない。論文はこの検証器自身の健全性も検証しており、「入力をそのまま返す」「答えをキャッシュする」「fp16で累積する」「非決定的に振る舞う」「無限大を有限値にすり替える」「1つの形状でしか動かない」「こっそり別デバイスにデータを移す」という7種のトリックを含む、意図的に壊れた19個のカーネルを検証器に通し、すべて不合格・正しい参照実装だけ合格になることを確認した、としている。

監査対象の「ある公開システム」の正体:Dr. Kernel / KernelGYM

論文のHTML全文版・第3.2節「The audit」を読むと、abstractだけでは特定できなかった監査対象のシステム名が明記されていた。

"We ran the verifier over Dr. Kernel / KernelGYM [7] (hkust-nlp/drkernel-coldstart-8k, MIT license), a public corpus of 8,920 supervised-fine-tuning trajectories each ending in a Triton kernel that replaces a PyTorch module in the KernelBench Model/ModelNew convention."

(訳:この検証器を、Dr. Kernel / KernelGYM(hkust-nlp/drkernel-coldstart-8k、MITライセンス)に対して実行した。これは8,920件の教師ありファインチューニング用トラジェクトリからなる公開コーパスで、それぞれがKernelBenchのModel/ModelNew規約に従って、PyTorchモジュールを置き換えるTritonカーネルで終わる)

hkust-nlp/drkernel-coldstart-8kは、実際にHugging Face Datasets上で公開されているデータセットであることを確認した。論文は、この8,920件のうちSSM近縁の演算子クラス(matmul・attention・softmax・scan・norm・conv・reduction)に該当する3,134件を抽出し、そこから「速度向上(final_speedup > 0)が記録されている」「監査ハーネスが結果を読める」「対象範囲に含まれる」という3段階のフィルターを経て、最終的な監査対象2,638件に絞り込んだ、という経緯が記されていた。

2,638件を再検証:39.5%が壊れていた

この検証器を外向きに使い、ある公開システム自身のハーネスがすでに「正しい」と判定していた2,638個のLLM生成カーネルを監査した結果が論文の中心的な数字だ。

  • **39.5%**が、どんな許容誤差の議論でも説明できないレベルで壊れていた
  • **62.1%**が、何らかの違反を1つ以上抱えていた
  • 業界標準のテストが合格とし、この検証器が不合格とした件数は1,487件
  • 逆に、検証器が合格とし標準テストが不合格とした件数はわずか14件

つまり「標準テストは通ったが実は壊れている」というケースが「検証器が間違って厳しすぎる」というケースの100倍以上あった、という主張だ。この結果自体の信頼性についても、論文は4通りの独立した方法で裏付けを取っている——7件中7件が正しく通る陽性コントロール(positive control)、閾値を振った感度分析(threshold-calibration sweep)、参照ベンチマーク自身の正誤判定コードとの98.5%の一致、そして層別の手作業監査だ。

自分たちのカーネルにも同じ検証器を向ける

論文の後半では、この検証器を「内向き」にも使っている。対象は著者ら自身が実装した、Gated-Linear-Recurrence(GDN)ファミリー向けの、NVIDIA Blackwell世代のtcgen05命令を使うネイティブな学習用backward——業界がまだフォールバック実装に頼っている「reverse-state」ステージを含む。この自作カーネルの正しさを倍精度(double-precision)のオラクルに対して独立に検証し、GDNファミリーの5つのメンバーモデルをこのカーネルを通して実際に学習させた、としている。

本文は精読したが、独立した再現検証はしていない

本記事はarXiv:2608.12700のAbstractページに加え、HTML全文版(experimental)の第3.1節・第3.2節、そして著者自身が公開しているGitHubリポジトリ(github.com/RishiShah99/lethe)とHugging Face上のデータセット(hkust-nlp/drkernel-coldstart-8k)を情報源としている。論文は2026年8月13日提出、8月17日にv2へ改訂、20ページ・3図とarXivのメタデータにある。著者はRishi Shah氏・Rishav Shrestha氏の2名で、査読付き学会・論文誌への掲載情報はarXivページ上には見当たらなかった。

「39.5%」「62.1%」といった数字、監査対象コーパスの絞り込み経緯(8,920件→3,134件→2,638件)は、いずれも著者自身の論文本文の記述であり、この記事を書いている自分自身がgithub.com/RishiShah99/letheのコードを実際にクローンして12ゲートの検証器を動かし、hkust-nlp/drkernel-coldstart-8kのカーネルを自分の手元で再監査したわけではない。査読前の論文であり、この39.5%という数字自体がさらに外部から検証されたものかどうかは、本記事の調査範囲では確認できていない。

LLM生成GPUカーネルの検証手法という話題自体まだ薄い

Zenn・Qiitaともに実質的な言及はなかった(Qiitaの2件は無関係の可能性が高いが、タイトルベースの確認にとどまり本文までは確認していない)。LLMによるGPUカーネル生成というテーマ自体、日本語圏での解説記事はまだ少ない。

関連記事

シェア: ポスト はてブ

出典・参照資料

AIニュースの解説を動画でも

YouTubeでは注目ニュースの背景を解説し、Xでは新着記事をお知らせしています。

コメント

まだコメントはありません。最初のコメントを書いてみませんか?

AIについて聞きたいことはありますか?

質問箱で無料で受け付けています。回答は公開され、他の方の参考にもなります。

質問箱を見る →

新しい記事をメールで受け取る

AIの新しい発表を、出典付きで整理して届けます。

関連記事

「Intern-S2-Preview」とは何か──397Bの本体を凍結したまま、4Bの「記憶デコーダ」だけ足して専門特化させる科学AIの記事画像
研究09.05読了11

「Intern-S2-Preview」とは何か──397Bの本体を凍結したまま、4Bの「記憶デコーダ」だけ足して専門特化させる科学AI

出典 ─ Intern-S2-Preview: Sci
絶対に失敗しないテストは、実は壊れている──ソフトウェアテスト理論の『Oracle Anchoring』の記事画像
研究08.27読了7

絶対に失敗しないテストは、実は壊れている──ソフトウェアテスト理論の『Oracle Anchoring』

出典 ─ Oracles That Cannot Fa
論文をAIに再現実装させると、コードは「動くのに違うもの」になる──『Semantic Alignment Units』で測る意味的ドリフトの記事画像
研究08.27読了7

論文をAIに再現実装させると、コードは「動くのに違うもの」になる──『Semantic Alignment Units』で測る意味的ドリフト

出典 ─ SA-Bench: Evaluating S
「AI4AI at Test-Time」とは何か──強いモデルが弱いモデルのためにharnessを設計する、学習不要の能力転移の記事画像
研究09.05読了11

「AI4AI at Test-Time」とは何か──強いモデルが弱いモデルのためにharnessを設計する、学習不要の能力転移

出典 ─ AI4AI at Test-Time: St
全DRAMを殴れるパターンを見つけた研究「Blacksmith」──GPU向けRowhammer攻撃の源流をたどるの記事画像
研究09.05読了11

全DRAMを殴れるパターンを見つけた研究「Blacksmith」──GPU向けRowhammer攻撃の源流をたどる

出典 ─ Blacksmith(COMSEC公式研究ペ
「Agentic ESOpt」とは何か──強化学習ではなく「進化戦略」でエージェントを鍛える、GPUに優しいやり方の記事画像
研究09.04読了12

「Agentic ESOpt」とは何か──強化学習ではなく「進化戦略」でエージェントを鍛える、GPUに優しいやり方

出典 ─ Agentic ESOpt: Fine-Tu
強化学習環境を「作り直さず、着せ替える」──EnvHarnessは静的な訓練環境をラップして弱点を突くの記事画像
研究09.04読了12

強化学習環境を「作り直さず、着せ替える」──EnvHarnessは静的な訓練環境をラップして弱点を突く

出典 ─ EnvHarness: Awakening
『記憶』だけを6.9Bパラメータまで太らせたら、モデル本体を大きくするより効率が良かったの記事画像
研究09.04読了12

『記憶』だけを6.9Bパラメータまで太らせたら、モデル本体を大きくするより効率が良かった

出典 ─ Memory Decoder at Scal