
NVIDIA SASS命令辞書のリバースエンジニアリング、カーネル監査、およびGPUアーキテクチャ間のパターン認識。
制御されたカーネルからプロダクション監査まで、NVIDIA SASS をリバースエンジニアリングする。
記事1 · 記事2 · 知識ベース · パターンライブラリ · SM120 命令用語集 · エンコーディングノート · ここから始める · プロジェクト構成 · テンソルコア章 · コントリビューション
SASS King は、NVIDIA SASS(コンパイルされた CUDA バイナリ内部に発行されるネイティブ GPU 命令セット)の体系的なリバースエンジニアリングプロジェクトです。このプロジェクトは、SM120 / SM120a コンシューマ向け Blackwell ハードウェアから始まり、時間の経過とともに完全なクロスアーキテクチャ ISA およびパターンライブラリへと拡大していきます。
目的は実用的です。カーネルエンジニアが SASS ダンプを開き、コンパイラパターンを認識し、パフォーマンスに関連する構造を特定し、バイナリをソースレベルの最適化判断に結びつけることができるようにすることです。
プロジェクトは初期のフェーズ3パターンライブラリを完了しました。29 の再利用可能な SASS シグネチャが patterns/ の下で形式化され、knowledge/FINDINGS.md に完全な証拠の連鎖が保持されています。次の主要なステップはフェーズ4です。これらのパターンを実際のプロダクションカーネルに適用することです。
リポジトリは証拠パイプラインとして編成されています:
corpus/ 制御されたカーネルと生の SASS 証拠
knowledge/ プロジェクト全体の発見、命令ノート、エンコーディングノート
patterns/ 再利用可能なフェーズ3監査シグネチャ
production/ フェーズ4本番カーネル監査
同様の精神で最後に広く公開された SASS リバースエンジニアリングの取り組みは、2018年の Jia らによる Volta および Turing に関するものでした。Ampere、Hopper、Blackwell では命令ミックスが大幅に変更されています。非同期コピーパス、テンソルコアファミリ、行列ロード/ストア命令、スパースおよびスケーリングされた MMA 形式、新しいユニフォームレジスタフローです。
SASS King は、制御されたマイクロカーネル、生の SASS 読み取り、ランタイムプローブ、および本番カーネル監査を組み合わせることで、そのギャップを埋めます。
正式なパターンライブラリがフェーズ3の主要な成果です。章ごとの証拠を再利用可能な監査シグネチャに変換し、監査が毎回完全な研究の道筋を書き直す代わりに、名前付きパターンを引用できるようにします。
フェーズ3は完了と見なされます。なぜなら:
knowledge/FINDINGS.md のソース証拠に限定されたままであること;patterns/README.md から始まること;各パターンページには以下が含まれます:
patterns/README.md を監査向けのインデックスとして使用します。パターンの背後にあるより長い研究コンテキストが必要な場合は knowledge/FINDINGS.md を使用します。
フェーズ3は、すべての NVIDIA SASS 動作がデコードされたとは主張しません。手動のプロダクション監査を開始するのに十分な、再利用可能な SM120 / SM120a パターンレイヤーを確立します。ランタイムレイアウトデコード、完全な制御コードビット配置、自動化された cubin レポート、クロスアーキテクチャリプレイは将来の作業です。
公開記事:
制御された変動。 2つのカーネルは、dtype、オペランド順序、アンロール係数、メモリレイアウト、またはコンパイルターゲットの1つの変数のみが異なります。SASS の差分により、コンパイラの判断が分離されます。
厳格なクレームタグ。 すべての技術的クレームはタグを使用します:
トップダウンとボトムアップの併用。 マイクロカーネルは個々の命令とコンパイラの判断を分離します。プロダクションライクなカーネルは、どのパターンが実際のコードで重要かを示します。
パターン優先の監査。 プロダクション監査では、可視の SASS シグネチャと一致し、その信頼限界、アンチパターン、未解決のギャップを引き継いだ後にのみ、正式な PATTERN-NN ページを引用する必要があります。
最初のパスは、SM120 テンソルコアとメモリパイプラインに焦点を当てます:
HMMA、QMMA、OMMALDSM、STSMLDGSTS、LDGDEPBAR、DEPBARLDG、STG、LDS、STS、REDGBRA、EXIT、BSSY、、このプロジェクトは、ISA がまだ完全であると主張するものではありません。公開用語集は観測および説明されたものを追跡し、knowledge/encoding/ の下のより詳細なページは、マッチャースタイルのドキュメントに十分な証拠を持つファミリを追跡します。
SASS King はビットレベルの SASS 逆アセンブラと競合しません。このプロジェクトはローカルダンプを主要な証拠として使用し、命令フィールド、スケジューリングテーブル、述語、レジスタ追跡の相互チェックとして redplait/denvdis を使用することがあります。denvdis は低レベルのエンコーディング解釈を検証できます。SASS King は、制御された変動の証拠、意味パターンレイヤー、およびプロダクション監査の解釈を担当します。
flowchart LR
P1["Phase 1<br/>Teaching kernels<br/>01-12"] --> P2["Phase 2<br/>SM120 tensor-core corpus<br/>13-25"]
P2 --> P25["Phase 2.5<br/>denvdis cross-validation<br/>bit-level backend"]
P25 --> P3["Phase 3<br/>Pattern library<br/>compiler signatures"]
P3 --> P4["Phase 4<br/>Production audits<br/>real kernels"]
P4 --> P5["Phase 5<br/>Audit tool<br/>cubin reports"]
P5 --> P6["Phase 6<br/>Cross-architecture replay<br/>SM80/86/89/90a/100a/120"]
classDef done fill:#0b6d55,color:#fff,stroke:#0b6d55;
classDef active fill:#f4c95d,color:#111,stroke:#b89422;
classDef planned fill:#1f2937,color:#fff,stroke:#6b7280;
class P1,P2,P25,P3 done;
class P4 active;
class P5,P6 planned;
カーネル01-12は、ベースライン SASS の概念を確立します:FMA 融合、スコアボード動作、ループの低レベル化、共有メモリ、グローバルメモリ、ワーププリミティブ、スローパス算術、ローカルメモリスピル。
カーネル13-25は、現在の SM120 テンソルコアパスをカバーします:
プロダクション監査がパターンライブラリに依存する前に、redplait/denvdis を SM120 / SM120a のビットレベル相互チェックバックエンドとして検証します。このパスは、代表的なローカル cubin または HMMA、QMMA、QMMA.SF、QMMA.SP、OMMA、LDSM、STSM b16/b8、LDGSTS、DEPBAR、分岐マーカーをカバーするダンプに対して nvd -O、nvd -S、nvd -p、および必要に応じて nvd -T を実行します。
出力は knowledge/DENVDIS_INTEGRATION.md です:ファミリから denvdis 認識ステータス、修飾子カバレッジ、公開制御コードフィールド、SASS King アクションまでの事実に基づく互換性テーブル。denvdis 出力は補助的な証拠であり、ローカルダンプ観測の代替ではありません。
形式化された繰り返し構造を再利用可能なシグネチャに:
LDGSTS -> DEPBAR -> LDSM -> MMAHMMA / QMMA / OMMASTSM -> BAR -> LDS -> STG初期のフェーズ3ライブラリには、patterns/ の下に29のパターンページが含まれています。knowledge/FINDINGS.md は研究ログおよび真実の源泉として残り、patterns/ は監査向けのエントリポイントです。
フェーズ3は初期ライブラリレベルで完了しています。ランタイムレイアウトデコード、完全な制御コードビット配置、クロスアーキテクチャリプレイなどの残りの項目は、ギャップまたは将来のフェーズとして追跡され、フェーズ4プロダクション監査の開始を妨げるものではありません。
FlashAttention、CUTLASS、xFormers、Transformer Engine、FlashInfer、llama.cpp / ggml、tinygrad などのライブラリの実際のカーネルにパターンライブラリを適用します。目標は、カーネルごとに1つのマークダウンファイルではなく、アルゴリズムパターンによる代表的なカバレッジです。
フェーズ4の最初の成果物は、以下の手動監査レポートであるべきです:
PATTERN-NN ページを引用する;cubin を入力とし、既知のパターンを検出し、最適化指向のレポートを出力するパイプラインを構築します。
追加のターゲットで方法論を再現します:
.
├── corpus/ # 制御されたカーネル、ダンプ、章の解説
│ ├── basics/ # カーネル01-08:スカラー/ベクターとメモリの基礎
│ ├── warp_collectives/ # カーネル09-10:シャッフル、投票、リダクション
│ ├── math_and_spills/ # カーネル11-12:スローパスとスピル
│ └── tensor_cores/ # カーネル13-25:テンソルコア研究
├── knowledge/ # 発見、用語集、エンコーディングノート
│ ├── FINDINGS.md
│ ├── SASS_INSTRUCTIONS_SM120.md
│ └── encoding/
├── patterns/ # 正式なフェーズ3パターンライブラリ
├── production/ # フェーズ4プロダクションカーネル監査
├── docs/ # オンボーディング、構成、リリース関連ノート
└── guide/ # 外部 SASS 読み取りガイドサブモジュール
各章フォルダには、ソースカーネル、該当する場合のコンパイル済みアーティファクト、検証済み証拠セットの一部である場合の SASS ダンプ、および conclusion<N>.md 解説が含まれています。
各ディレクトリに何が属するかの詳細な説明については、Project Structure をお読みください。
cuobjdump --dump-sass 生の逆アセンブル用。gpuasm.com スコアボード、ストール、プレッシャー、依存矢印用。%clock マイクロベンチマーク命令レイテンシプローブ用。nvcc -Xptxas -v レジスタとスピルメタデータ用。SASS King はアルゴリズムパターンレイヤーで動作します。コンパイルされたカーネルがどのように構造化されているかを認識し、それらの構造をソースレベルの最適化判断に結び付けます。
コントリビューションを歓迎します。特に:
期待されるメタデータと執筆基準については CONTRIBUTING.md を参照してください。
Florian Mattana. florianmattana.com
| 目的 | ここから始める | 次に読む |
|---|
| プロジェクトを10分で理解する | docs/README.md | docs/START_HERE.md、次に docs/PROJECT_STRUCTURE.md |
| 証拠を再現する | corpus/README.md | 1つの章 conclusion*.md、次にその .sass ダンプ |
| 真実の源泉を見つける | knowledge/FINDINGS.md | knowledge/SASS_INSTRUCTIONS_SM120.md、knowledge/encoding/README.md |
| 新しいダンプでパターンを認識する | patterns/README.md | 対応する patterns/NN-*.md ページ |
| プロダクション監査を開始する | production/README.md | 対応する PATTERN-NN ページとソース証拠 |
| 修正やダンプをコントリビュートする | CONTRIBUTING.md | docs/START_HERE.md |
| 領域 | ステータス | 場所 |
|---|
| SM120 教育用カーネル | カーネル01-12 完了 | corpus/basics/01_vector_add/ から corpus/math_and_spills/12_register_spill/ |
| テンソルコア研究 | カーネル25 完了 | corpus/tensor_cores/ |
| 全体的な発見 | 現在の真実の源泉 | knowledge/FINDINGS.md |
| SM120 命令用語集 | 現在進行中、証拠に基づく | knowledge/SASS_INSTRUCTIONS_SM120.md |
| エンコーディングパイロット | LDSM、STSM、QMMA から開始 | knowledge/encoding/ |
| denvdis 相互検証 | 初期パス完了。制御コードの深いギャップが残る | knowledge/DENVDIS_INTEGRATION.md |
| パターンライブラリ | 初期フェーズ3ライブラリ完了 | patterns/ |
| プロダクション監査 | 次のフェーズ | production/ |
| パターンファミリ | 例 | 場所 |
|---|
| テンソルコア計算 | HMMA、QMMA、OMMA アキュムレータチェーン; スパースメタデータ; 狭断片 | patterns/02-* から patterns/04-*、patterns/10-*、patterns/21-* |
| 行列メモリとエピローグ | LDSM、STSM、非同期コピーパイプライン、REDG リダクションエピローグ | patterns/05-*、patterns/06-*、patterns/07-*、patterns/28-* |
| 制御フロー | 分岐/再収束、ループバックエッジ、述語付き出口、コールドトラップ、ローカル CALL | patterns/08-*、patterns/14-*、patterns/16-*、patterns/26-*、patterns/29-* |
| メモリとレジスタ | ベクトル化グローバルメモリ、スピル、共有メモリステージング、ディスクリプタ、ユニフォームレジスタフロー | patterns/09-*、patterns/11-*、patterns/17-*、patterns/19-*、patterns/20-* |
| 算術とスケジューリング | FFMA 融合、定数、MUFU スローパス、スコアボード、ライフタイムリサイクル | patterns/12-*、patterns/18-*、patterns/22-*、patterns/23-*、patterns/24-* |
| ワープコレクティブ | ワープリダクション、シャッフル/投票/マッチ/同期プリミティブ | patterns/01-*、patterns/25-* |
| タグ | 意味 |
|---|
[OBS] | ダンプ、ログ、ランタイム出力、またはプロファイルで直接観測。 |
[INF] | 観測された証拠から推論。 |
[HYP] | もっともらしいが未確認。 |
[RES] | 以前の仮説が後の証拠で解決された。 |
[GAP] | 未解決の質問が明示的に文書化されている。 |
BSYNCWARPSYNCSHFL、VOTE、REDUXS2UR、R2UR、UMOV、ULEA、LDCU| フェーズ | ステータス | 成果物 | なぜ重要か |
|---|
| 1. 教育用カーネル | 完了 | corpus/basics/、corpus/warp_collectives/、corpus/math_and_spills/ | 制御された CUDA から SASS への実験から読み取り語彙を確立します。 |
| 2. SM120 テンソルコアコーパス | 完了 | corpus/tensor_cores/13_hmma_fp16/ から 25_stsm_epilogue/ | 最初の SM120 / SM120a テンソルコア、行列メモリ、制御フロー、エピローグ証拠セットをキャプチャします。 |
| 2.5. denvdis 相互検証 | 初期パス完了 | knowledge/DENVDIS_INTEGRATION.md、knowledge/encoding/CONTROL_CODE.md | ローカルダンプ証拠を置き換えることなく、denvdis をビットレベル相互チェックとして使用します。完全なストール/イールドビット配置は未解決のままです。 |
| 3. パターンライブラリ | 初期ライブラリ完了 | patterns/ | 繰り返されるコンパイラ/SASS 構造を再利用可能なシグネチャに変換します。 |
| 4. プロダクション監査 | 次 | production/ | コーパスパターンがプロダクションライブラリの実際のカーネルを説明できるかテストします。 |
| 5. 監査ツール | 計画済み | cubin からレポートへのパイプライン | パターンレイヤーをスクリプト化可能で再現可能にします。 |
| 6. クロスアーキテクチャリプレイ | 計画済み | SM80、SM86、SM89、SM90a、SM100a、SM120 の比較 | アーキテクチャ固有の事実と一般的な NVIDIA SASS 動作を分離します。 |
| カーネル | トピック |
|---|
| 13 | HMMA ベースライン、レジスタ割り当て、アキュムレータチェーン |
| 14 | QMMA FP8 / FP6 / FP4 ベースライン |
| 15 | 狭 MMA バリアント |
| 16 | FP4 ピークとブロックスケーリング OMMA/QMMA |
| 17 | LDSM と行列ロード動作 |
| 18 | パイプラインマルチプレクサ MMA タイルと非同期コピーステージング |
| 19 | スパース MMA メタデータ |
| 20 | 制御フローとバックエッジ検出 |
| 21 | 分岐と再収束 |
| 22 | STSM 行列ストア動作 |
| 23 | FP4 / FP6 断片レイアウトプローブ |
| 24 | プロダクションミニ GEMM 監査 |
| 25 | STSM エピローグレイアウトとストアバックセマンティクス |
| アーキテクチャ | 代表的な GPU | 理由 |
|---|
| SM80 | A100 | データセンター Ampere ベースライン |
| SM86 | RTX 3090 | コンシューマ Ampere コーパス |
| SM89 | RTX 4090 | 一般的なコンシューマ推論カード |
| SM90a | H100 | TMA、WGMMA、ワープスペシャライゼーション、クラスター |
| SM100a | B200 | tcgen05.mma、TMEM |
| SM120 | RTX 5070 Ti / 5090 | コンシューマ Blackwell 開始点 |