SankalpがCodex loopを使用して、GPU ModeのB200 QRコンテストで12位に入賞

SankalpはCodex、プロファイリングツール、そして1,500件以上の提出を用いてB200 QR kernelを最適化し、183人の参加者中12位に入った。

By · Published

Primary source: Sankalp's blog

Why it matters

Sankalp's result shows how a strict benchmark can turn a coding agent into an experimental system. Engineers still have to choose the architecture, interpret profiles and change the search policy when the agent stalls.

The focused labor of a programmer optimizing code (Oil painting in the manner of Edward Hopper)

独立して検証可能な成果は、SankalpがNVIDIA B200で1,804.779マイクロ秒、参加者183名中12位を記録したことだ。彼はOpenAI Codexを使用して14日間で1,500件以上の提出を行った後その結果を達成した。プロジェクトはGPU ModeのLinear Algebra Kernels in the Age of Researchシリーズとそのqr_v2 competitionの一部だった。

7月8日のプロジェクト報告でSankalpはこう書いている。「14日間の間に、1500回以上の提出を行った。」Sankalpのブログは約419,000マイクロ秒のPyTorchベースラインに対して232倍の改善を報告している;その比較はこのベンチマークワークロードに特有のものであり、提供された資料では独立して再現されていない。

コンペは6月29日に終了し、本記事の約6週半前だった。Sankalpの最終結果は優勝の1,220.774マイクロ秒の提出より概ね48%遅かった。彼の成績が注目に値するのは、[Tritonを中心にGPUカーネル最適化を学んで約1年の経験しかなく、プロとしてこの分野で働いたことはないと本人が述べている](https://sankalp.bearblog.dev/autoresearch/?ref=runtimewire#learning-enough-to-ask-better-questions)からだ。

Sankalpはコード生成、ベンチマーク、プロファイリング、リーダーボード提出を組み合わせた最適化ループを記述した。彼がアーキテクチャを選び、結果を解釈した。以下のワークフローの詳細は彼の報告に基づくものであり、独立して測定可能な成果は彼のリーダーボードでの時間と順位である。

エージェント向けに設計されたベンチマーク

GPU Mode、GPUプログラミングの競技会とカーネルベンチマークのプラットフォーム、及びCore Automationは、競技者にバッチ化された正方行列のcompact-Householder QR分解を実装するよう求めた。各提出はFP32のCUDA行列を受け取り、PyTorchのtorch.geqrfが使うのと同じコンパクト表現を返す必要があった:上三角のR結果と格納されたHouseholderベクトルを含むH行列と、反射子係数を入れたtauベクトルだ。

チェッカーはQを再構成し、直交性と残差誤差をテストし、正しい提出を行列形状と入力条件全体での幾何平均実行時間でランク付けした。ワークロードは4,096×4,096までの行列をカバーしており、条件の悪いケースも含まれていた。競技者は内部で低精度を使うことが許されたが、出力はFP32スタイルのチェックに合格しなければならなかった。

その組み合わせはCodexにとって進捗を高速かつ定量的に定義する手段を与えた。GPU Modeのpopcorn command-line interfaceにより、エージェントは候補をテスト、ベンチマーク、提出できた。形状ごとのタイミングは各変更がどこで改善や悪化をもたらしたかを示し、プロファイリングは別の証拠の層を提供した。

Sankalpの報告によれば、彼はAGENTS.mdファイルに運用手順、問題の定義、実験記録、タイムスタンプ付き提出ログを維持していた。後のCodexセッションは既に失敗したアプローチを読み取って再発見せずに済んだ。彼はまたエージェントに数値目標を与え、いくつかの最適化ランを一晩中継続させ、2〜3時間ごとに何を変えたか、どのボトルネックを追っているかを確認した。

最終的に生成されたディレクトリには560の名前付き提出バリアント、119のModal B200プローブと比較スクリプト、68の実験別ドキュメントが含まれていた。Modal、クラウドインフラプロバイダーはコンペ用のGPUクレジットを提供したとSankalpは述べている。記録は失敗したアプローチ、プロファイリング結果、提出履歴を保存し、後のCodexセッションが以前の作業に基づいて構築できるようにした。

逐次処理を行列乗算に変える

SankalpはまずClaudeと教育資料を使ってHouseholder QRを理解したと報告している。彼はトレイリングWY更新を伴うブロック化Householder設計に落ち着き、その後プロファイリング、正確性チェック、繰り返しのリーダーボード提出でさらなる最適化を導いた。

コアの性能問題は逐次依存だった。従来のHouseholder実装は列を順に処理し、各反射子が前のステップで生成された行列に依存する。これにより、B200のテンソルコアが待機する一方で、より遅い行列ベクトル演算にかなりの作業が残ることになる。

ブロック化アプローチは逐次作業を狭いパネルに限定し、より大きなトレイリング更新を行列乗算に変換する。GPUはそこではるかに多くの並列作業を実行できる。Sankalpはそのアーキテクチャを使用した最初の日のうちに、重み付けの大きい512×512ケースで約5,000マイクロ秒に到達したと報告している。

さらなる改善はスタック全体での変更を必要とした。彼の投稿の提出履歴は、カスタムTritonパネル、グループ化されたWY更新、CUDAグラフのリプレイ、融合されたレイアウト組み立て、固定形状カーネルの特殊化、大規模行列向けのカスタムCholesky経路へと進んだ。Sankalpは追跡された全表結果が108,803マイクロ秒から約1,805マイクロ秒に低下したと報告している。この進展は彼の232倍計算に使われた概算の419,000マイクロ秒のPyTorchベースラインとは別のものだ。

Sankalpの後期プロファイルに関する報告では、起動オーバーヘッドとパネル処理が支配的であり、カーネルが生のメモリ帯域幅や計算能力で制限されることは稀だったとされる。彼の最適化テーブルには、起動回数の削減、削減操作の融合、固定形状の特殊化、V/T組み立ての結合、コピー・連結・一時表現の除去に関する作業が記録されている。

SankalpはCodexにModalのプロファイリングアクセスを与え、Torchプロファイリング、NVIDIA Nsight Systems、後にNCUを使用したと書いている。これらは自己申告の手法の詳細である。彼の説明では、Codexが変更を実装し測定する一方で、彼は結果を検査し、目標を設定し、探索を再指示した。

人間が探索プロセスを変えた

Sankalpの報告によれば、結果が3,000マイクロ秒を下回った後、最適化はより困難になった。彼はCodexがパラメータ調整や既に試したアイデアの小さな変種に多くの時間を費やすようになったと述べている。Sankalpは実験が注目を集める方法を変更して応じた。

彼はCodexに、単一の現行最良解を保持して遅い試行をすべて排除する代わりに、3〜5の候補ファミリのビームを維持するよう指示した。候補には現在の最良解に近い変更、有望なニアミス、高リスクの構造的アイデアが含まれていた。

同じ報告でSankalpは、headlessなclaude -p呼び出しをアドバイザーとして使い、サブエージェントに最適化アイデアを探索させ、新しい実行の前に蓄積されたコンテキストをクリアしたと述べている。彼のAGENTS.mdの指示は、タイムアウトを結論未定と扱い、タイムスタンプ付きログを保存し、完成した出力を証拠として要求し、なぜ候補が昇格または却下されたかを記録するよう求めていた。ファイルはまた、実質的な変更なしに却下されたアイデアを繰り返すことへの警告も含んでいた。

Codexは持続性、コード生成、大きな探索予算を供給した。Sankalpはアーキテクチャを選び、フィードバックループを改訂し、反復的な行動を認識し、実験方針を変更した。

In 彼のトップ10エントリーのレビュー、Sankalp は、より高速な競合者が入力分布検出器を使用し、より多くのライブラリ呼び出しを削除し、カスタムの三角行列逆行列を実装し、より低精度のデータをより積極的に扱っていたと述べた。これらは独立して監査された所見ではなく、Sankalp の観察に基づくものであった。同じ投稿の結論では、B200 上でそのテンソルコアをさらに活用するために tcgen05 命令を「使用できなかった」と書いている。

GPU Mode turns contests into an agent test bed

Core Automation の AI が書いたシステムコードに関するエッセイで、GPU Mode 共同創業者の Mark Saroufim は、彼と共同創業者の Andreas Kopf が 2023 年後半に GPU プログラミングの勉強会を始めたと書いている。グループはその後 YouTube チャンネル、カーネルコンペティション、ハッカソンへと拡大した。

Saroufim は、コーディングモデルを訓練するためにオンラインで利用できる GPU カーネル資料の不足を指して「データ不足(data starvation)」という表現を用いた。GPU Mode のコンテストは、そのまれな種類のシステム問題に対してエージェントと人間が取り組む計測可能な例を生み出している。

Sankalp の12位の成績はその活動の現実的な見方を提供する。新人がエージェントを使って、狭く測定可能なタスクで競争的なリーダーボードの上位に迫った。作業はそれでも問題を学び、実験を構成し、反復的な振る舞いに挑戦し、プロファイリング結果を解釈することを彼に要求した。彼の成果は、各実行後に信頼できる証拠を生成できる評価者の価値を示している。

Reader comments

Conversation for this story loads after sign-in.