Skip to content
KitploitKITPLOIT
ツールブログ
提出
ツールブログ
提出

ハッキング、侵入テスト、サイバーセキュリティツールをあなたのセキュリティアーセナルに!

Kitploitはハッキング、サイバーセキュリティ、ペネトレーションテストのツールディレクトリです。最新のプロジェクトアップデートを見つけて、脆弱性の発見、システム分析、テストの自動化、セキュリティの強化を行いましょう。

··フィード·お問い合わせ·プライバシー·© 2026 Kitploit

ツールディレクトリ

カテゴリ

すべてのカテゴリを見る
Loading categories
sass-king — NVIDIA SASS命令辞書のリバースエンジニアリング、カーネル監査、およびGPUアーキテクチャ間のパターン認識。 | Kitploit
ツール/GitHubGitHub/florianmattana/sass-king
組み込みシステムセキュリティ静的分析コード分析リバースエンジニアリングハードウェアセキュリティハードウェアとIoTセキュリティバイナリ解析論文と研究学習と教育厳選リソースファームウェア解析
317153ヶ月前Kitploit レビュー済み

人気

すべて見る →

コミュニティで最も使われているツールを見つけましょう。

すべてのツールを探索

ツールコレクションを閲覧

すべてのツールを見る →
共有
GitHub
florianmattana/sass-king

sass-king

NVIDIA SASS命令辞書のリバースエンジニアリング、カーネル監査、およびGPUアーキテクチャ間のパターン認識。

リポジトリを見るウェブサイト

SASS King ロゴ

SASS King

制御されたカーネルからプロダクション監査まで、NVIDIA SASS をリバースエンジニアリングする。

記事1 · 記事2 · 知識ベース · パターンライブラリ · SM120 命令用語集 · エンコーディングノート · ここから始める · プロジェクト構成 · テンソルコア章 · コントリビューション

アーキテクチャ ステータス ライセンス

SASS King は、NVIDIA SASS(コンパイルされた CUDA バイナリ内部に発行されるネイティブ GPU 命令セット)の体系的なリバースエンジニアリングプロジェクトです。このプロジェクトは、SM120 / SM120a コンシューマ向け Blackwell ハードウェアから始まり、時間の経過とともに完全なクロスアーキテクチャ ISA およびパターンライブラリへと拡大していきます。

目的は実用的です。カーネルエンジニアが SASS ダンプを開き、コンパイラパターンを認識し、パフォーマンスに関連する構造を特定し、バイナリをソースレベルの最適化判断に結びつけることができるようにすることです。

プロジェクトは初期のフェーズ3パターンライブラリを完了しました。29 の再利用可能な SASS シグネチャが patterns/ の下で形式化され、knowledge/FINDINGS.md に完全な証拠の連鎖が保持されています。次の主要なステップはフェーズ4です。これらのパターンを実際のプロダクションカーネルに適用することです。

クイックナビゲーション

リポジトリは証拠パイプラインとして編成されています:

root@kitploit:~
corpus/      制御されたカーネルと生の SASS 証拠
knowledge/   プロジェクト全体の発見、命令ノート、エンコーディングノート
patterns/    再利用可能なフェーズ3監査シグネチャ
production/  フェーズ4本番カーネル監査

存在理由

同様の精神で最後に広く公開された SASS リバースエンジニアリングの取り組みは、2018年の Jia らによる Volta および Turing に関するものでした。Ampere、Hopper、Blackwell では命令ミックスが大幅に変更されています。非同期コピーパス、テンソルコアファミリ、行列ロード/ストア命令、スパースおよびスケーリングされた MMA 形式、新しいユニフォームレジスタフローです。

SASS King は、制御されたマイクロカーネル、生の SASS 読み取り、ランタイムプローブ、および本番カーネル監査を組み合わせることで、そのギャップを埋めます。

現在の状態

フェーズ3 成果物

正式なパターンライブラリがフェーズ3の主要な成果です。章ごとの証拠を再利用可能な監査シグネチャに変換し、監査が毎回完全な研究の道筋を書き直す代わりに、名前付きパターンを引用できるようにします。

フェーズ3は完了と見なされます。なぜなら:

  • 章01-25で見つかった繰り返しの構造が29の名前付きパターンページに昇格されたこと;
  • すべてのパターンに平易な英語の説明、SASS シグネチャ、バリアント、アンチパターン、未解決のギャップ、信頼レベルが含まれていること;
  • クレームタグは knowledge/FINDINGS.md のソース証拠に限定されたままであること;
  • 監査向けナビゲーションは現在 patterns/README.md から始まること;
  • 未解決の項目は、パターンテキスト内に隠されるのではなく、ギャップとして明示的に先送りされていること。

各パターンページには以下が含まれます:

  • 平易な英語での意味;
  • SASS シグネチャ;
  • 観測されたバリアント;
  • 解釈の境界;
  • アンチパターン;
  • 未解決のギャップ;
  • 信頼レベル。

patterns/README.md を監査向けのインデックスとして使用します。パターンの背後にあるより長い研究コンテキストが必要な場合は knowledge/FINDINGS.md を使用します。

フェーズ3は、すべての NVIDIA SASS 動作がデコードされたとは主張しません。手動のプロダクション監査を開始するのに十分な、再利用可能な SM120 / SM120a パターンレイヤーを確立します。ランタイムレイアウトデコード、完全な制御コードビット配置、自動化された cubin レポート、クロスアーキテクチャリプレイは将来の作業です。

ここから始める

  • プロジェクトが初めての方: Start Here をお読みください。
  • プロジェクト全体のマップをご希望の場合: knowledge base index をお読みください。
  • 現在の命令マップをご希望の場合: SASS instructions on SM120 / SM120a をお読みください。
  • エンコーディングノートをご希望の場合: encoding notes をお読みください。
  • 生の真実の源泉をご希望の場合: findings をお読みください。
  • 再利用可能な監査シグネチャをご希望の場合: pattern library をお読みください。
  • テンソルコアの証拠をご希望の場合: tensor-core chapters から始めてください。
  • ダンプや修正をコントリビュートしたい場合: contributing をお読みください。
  • v0.1 の境界をご希望の場合: release notes をお読みください。

公開記事:

  • Part 1 - Reading NVIDIA SASS from First Principles
  • Part 2 - Reading the Compiler's Mind

方法論

制御された変動。 2つのカーネルは、dtype、オペランド順序、アンロール係数、メモリレイアウト、またはコンパイルターゲットの1つの変数のみが異なります。SASS の差分により、コンパイラの判断が分離されます。

厳格なクレームタグ。 すべての技術的クレームはタグを使用します:

トップダウンとボトムアップの併用。 マイクロカーネルは個々の命令とコンパイラの判断を分離します。プロダクションライクなカーネルは、どのパターンが実際のコードで重要かを示します。

パターン優先の監査。 プロダクション監査では、可視の SASS シグネチャと一致し、その信頼限界、アンチパターン、未解決のギャップを引き継いだ後にのみ、正式な PATTERN-NN ページを引用する必要があります。

カバー範囲

最初のパスは、SM120 テンソルコアとメモリパイプラインに焦点を当てます:

  • HMMA、QMMA、OMMA
  • LDSM、STSM
  • LDGSTS、LDGDEPBAR、DEPBAR
  • LDG、STG、LDS、STS、REDG
  • BRA、EXIT、BSSY、、

このプロジェクトは、ISA がまだ完全であると主張するものではありません。公開用語集は観測および説明されたものを追跡し、knowledge/encoding/ の下のより詳細なページは、マッチャースタイルのドキュメントに十分な証拠を持つファミリを追跡します。

SASS King はビットレベルの SASS 逆アセンブラと競合しません。このプロジェクトはローカルダンプを主要な証拠として使用し、命令フィールド、スケジューリングテーブル、述語、レジスタ追跡の相互チェックとして redplait/denvdis を使用することがあります。denvdis は低レベルのエンコーディング解釈を検証できます。SASS King は、制御された変動の証拠、意味パターンレイヤー、およびプロダクション監査の解釈を担当します。

ロードマップ

root@kitploit:~
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;

フェーズ1 - 教育用カーネル

カーネル01-12は、ベースライン SASS の概念を確立します:FMA 融合、スコアボード動作、ループの低レベル化、共有メモリ、グローバルメモリ、ワーププリミティブ、スローパス算術、ローカルメモリスピル。

フェーズ2 - テンソルコアおよび SM120 カバレッジ

カーネル13-25は、現在の SM120 テンソルコアパスをカバーします:

フェーズ2.5 - denvdis 相互検証

プロダクション監査がパターンライブラリに依存する前に、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 出力は補助的な証拠であり、ローカルダンプ観測の代替ではありません。

フェーズ3 - パターンライブラリ

形式化された繰り返し構造を再利用可能なシグネチャに:

  • LDGSTS -> DEPBAR -> LDSM -> MMA
  • チェーンされた HMMA / QMMA / OMMA
  • STSM -> BAR -> LDS -> STG
  • ワープリダクションとクロスレーンコレクティブ
  • レジスタスピルシグネチャ
  • スカラーおよびユニフォーム制御フローパターン

初期のフェーズ3ライブラリには、patterns/ の下に29のパターンページが含まれています。knowledge/FINDINGS.md は研究ログおよび真実の源泉として残り、patterns/ は監査向けのエントリポイントです。

フェーズ3は初期ライブラリレベルで完了しています。ランタイムレイアウトデコード、完全な制御コードビット配置、クロスアーキテクチャリプレイなどの残りの項目は、ギャップまたは将来のフェーズとして追跡され、フェーズ4プロダクション監査の開始を妨げるものではありません。

フェーズ4 - プロダクション監査

FlashAttention、CUTLASS、xFormers、Transformer Engine、FlashInfer、llama.cpp / ggml、tinygrad などのライブラリの実際のカーネルにパターンライブラリを適用します。目標は、カーネルごとに1つのマークダウンファイルではなく、アルゴリズムパターンによる代表的なカバレッジです。

フェーズ4の最初の成果物は、以下の手動監査レポートであるべきです:

  • 実際のカーネルを SASS 領域に分割する;
  • 一致する PATTERN-NN ページを引用する;
  • 各結論に信頼レベルを割り当てる;
  • 説明されていない領域を新しいギャップとして記録する;
  • 少なくとも1つの手動レポートが安定するまで監査ツールを構築しない。

フェーズ5 - 監査ツール

cubin を入力とし、既知のパターンを検出し、最適化指向のレポートを出力するパイプラインを構築します。

フェーズ6 - クロスアーキテクチャ

追加のターゲットで方法論を再現します:

リポジトリマップ

root@kitploit:~
.
├── 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 スコアボード、ストール、プレッシャー、依存矢印用。
  • Nsight Compute プロファイリングとストール属性用。
  • %clock マイクロベンチマーク命令レイテンシプローブ用。
  • nvcc -Xptxas -v レジスタとスピルメタデータ用。

関連研究

  • Jia et al. 2018, "Dissecting the NVIDIA Volta GPU Architecture via Microbenchmarking" は、レイテンシ、スループット、依存関係検証の背後にある経験的マイクロベンチマーク手法に関するもの。
  • kuterdinel.com/nv_isa は、ファジングされた NVIDIA ISA エンコーディング研究、特に逆アセンブラの動作から機械可読なエンコーディングルールを導出するアイデアに関するもの。
  • redplait/denvdis は、オペコードテーブル、ビットレベル逆アセンブル、エンコーディングフィールド検査、スケジューリング分析、レジスタ追跡、cubin 操作に関するもの。抽出された SM120 data12 テーブル は、ローカルダンプを主要な証拠としたまま低レベル相互チェックとして使用される。
  • Redplait ツールとノート: ced cubin editor、SASS disassembly Perl bindings、SASS latency analysis、libcuda/nvasm_internal notes。
  • Huerta et al. 2025 は、コンパイラガイド付きスケジューリング、制御コード、依存カウンタ、再利用フラグ、イールド動作のリバースエンジニアリングに関するもの。
  • Yan et al. 2026 は、SASS 以下のドライバレイヤーの起動とプッシュバッファ分析に関するもの。
  • MaxAS および TuringAS は、古い NVIDIA アーキテクチャ向けの以前の公開 SASS アセンブラの取り組み。
  • NVIDIA CUDA Binary Utilities ドキュメントは、公式の cubin、fatbin、逆アセンブルツールに関するもの。

SASS King はアルゴリズムパターンレイヤーで動作します。コンパイルされたカーネルがどのように構造化されているかを認識し、それらの構造をソースレベルの最適化判断に結び付けます。

コントリビューション

コントリビューションを歓迎します。特に:

  • ここでは直接利用できないハードウェアからの生の SASS ダンプ;
  • コンパイラの1つの判断を分離する制御されたカーネル研究;
  • 既存の観測の修正;
  • 新しいプロダクションカーネルパターンの提案;
  • クロスアーキテクチャ比較。

期待されるメタデータと執筆基準については CONTRIBUTING.md を参照してください。

著者

Florian Mattana. florianmattana.com

ツールをダウンロード
目的ここから始める次に読む
プロジェクトを10分で理解するdocs/README.mddocs/START_HERE.md、次に docs/PROJECT_STRUCTURE.md
証拠を再現するcorpus/README.md1つの章 conclusion*.md、次にその .sass ダンプ
真実の源泉を見つけるknowledge/FINDINGS.mdknowledge/SASS_INSTRUCTIONS_SM120.md、knowledge/encoding/README.md
新しいダンプでパターンを認識するpatterns/README.md対応する patterns/NN-*.md ページ
プロダクション監査を開始するproduction/README.md対応する PATTERN-NN ページとソース証拠
修正やダンプをコントリビュートするCONTRIBUTING.mddocs/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-*
制御フロー分岐/再収束、ループバックエッジ、述語付き出口、コールドトラップ、ローカル CALLpatterns/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]未解決の質問が明示的に文書化されている。
BSYNC
WARPSYNC
  • SHFL、VOTE、REDUX
  • ユニフォームレジスタフロー:S2UR、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 動作を分離します。
    カーネルトピック
    13HMMA ベースライン、レジスタ割り当て、アキュムレータチェーン
    14QMMA FP8 / FP6 / FP4 ベースライン
    15狭 MMA バリアント
    16FP4 ピークとブロックスケーリング OMMA/QMMA
    17LDSM と行列ロード動作
    18パイプラインマルチプレクサ MMA タイルと非同期コピーステージング
    19スパース MMA メタデータ
    20制御フローとバックエッジ検出
    21分岐と再収束
    22STSM 行列ストア動作
    23FP4 / FP6 断片レイアウトプローブ
    24プロダクションミニ GEMM 監査
    25STSM エピローグレイアウトとストアバックセマンティクス
    アーキテクチャ代表的な GPU理由
    SM80A100データセンター Ampere ベースライン
    SM86RTX 3090コンシューマ Ampere コーパス
    SM89RTX 4090一般的なコンシューマ推論カード
    SM90aH100TMA、WGMMA、ワープスペシャライゼーション、クラスター
    SM100aB200tcgen05.mma、TMEM
    SM120RTX 5070 Ti / 5090コンシューマ Blackwell 開始点