- はじめに — プロファイラが示したカーネルは誰も手で書いていなかった
- 共通する一つの問題 — グラフを受け取りカーネルを作る
- 最下層 — NVCC、PTX、SASS、そしてLLVMのGPUバックエンド
- MLIRとその上に築かれたもの — IRを作るための枠組み、そしてTriton
- torch.compileとInductorが実際にやっていること
- XLA、IREE、TVM — グラフを丸ごと受け取る側
- どの層を、いつ知るべきか
- おわりに — 層を分けるのは性能ではなく決定権です
- 参考資料
はじめに — プロファイラが示したカーネルは誰も手で書いていなかった
Nsight Computeで学習ステップをプロファイルすると、最も時間を使っているカーネルの名前が出てきます。ところがその名前がtriton_poi_fused_add_mul_native_layer_norm_7のようなものだと困ったことになります。リポジトリを検索しても見つかりません。誰も書いていないからです。機械が作ったのです。
それが何によって作られたのかを突き止めるには、次のような問いに答えられなければなりません。そのカーネルはどの段階で生成されたのか。PyTorchが作ったのか、Tritonが作ったのか、それともさらに下のLLVMが作ったのか。名前を変えたい、いや、コードを直したいとなったとき、どの層に手を入れればよいのか。
この問いが難しいのは層が多いからではなく、層ごとに名前が違い、所属組織が違い、ドキュメントも別々にあるせいで、互いに無関係に見えるからです。NVCC、LLVM、MLIR、Triton、Inductor、XLA、IREE、TVMは、それぞれまったく別の世界のもののように見えます。
実際には、すべて同じ問題を解いています。この記事はまずその問題が何かを述べ、各層がその問題のどの部分を切り取ったのかを整理します。このシリーズの前の記事がカーネルを直接書き換える作業、カーネル記述の三つの層、推論エンジンのチューニング、二つのベンダースタックの違いを扱ったのだとすれば、この記事はそれらが置かれている座標系にあたります。
バージョンと事実関係は2026年8月2日に各プロジェクトの公式ドキュメントとリリースノートで直接確認しました。この領域は6か月もあれば地形が変わるため、確認できたこととできなかったことを本文中で区別して記しています。
共通する一つの問題 — グラフを受け取りカーネルを作る
名前が何であれ、この記事に出てくるツールは入力と出力が同じです。
入力は演算グラフです。行列積、加算、活性化関数、正規化といったノードがテンソルで結ばれたものです。出力はGPU上で実際に実行される機械語です。そしてその間でやるべきことも、驚くほど似通っています。
- フュージョンの決定です。どのノードを一つのカーネルにまとめるか。まとめれば中間テンソルをグローバルメモリに書かずに済むため、メモリ帯域幅を節約できます。XLAのドキュメントはフュージョンを「XLAにおける単一最重要の最適化」と呼んでいます。
- レイアウトの決定です。テンソルをメモリ上にどの順序で並べるか。同じ演算であってもレイアウト次第でコアレッシングされることもあれば、されないこともあります。
- タイリングの決定です。大きな行列をどのサイズの断片に切り、各スレッドブロックに割り当てるか。断片が大きすぎれば共有メモリとレジスタが足りなくなり、小さすぎれば再利用が減ります。
- ライブラリ呼び出しとコード生成のどちらを選ぶかです。この演算はcuBLASを呼ぶほうが速いのか、それとも自前でカーネルを生成したほうが速いのか。
- ロワリングです。決定した内容を実際の命令へと落とし込む段階です。
層が複数存在するのは、これらの決定の性質がそれぞれ異なるからです。フュージョンはグラフ全体を見なければ判断できず、レジスタ割り当ては命令レベルでしか判断できません。一つの抽象化で両方をうまくこなすのは困難です。そのため上位の層はグラフを見て、下位の層は命令を見て、その間に中間層が挟まっています。
まとめると次のようになります。
| 層 | 見る単位 | 行う決定 | 代表的なツール |
|---|---|---|---|
| グラフ | モデル全体 | フュージョン、レイアウト、分散分割 | XLA、Inductor、IREE、TVM |
| カーネル | カーネル一つ | タイルサイズ、パイプライニング、共有メモリ配置 | Triton、CUTLASS、Helion |
| 命令 | 基本ブロック | 命令選択、スケジューリング、レジスタ割り当て | LLVM、ptxas |
| 機械語 | チップ | 実際のエンコーディング | SASS、AMDGCN |
この表を頭に置いて以下を読むと、各ツールの位置がはっきり定まります。
最下層 — NVCC、PTX、SASS、そしてLLVMのGPUバックエンド
まず一番下から見ていきます。ここは20年間、構造がほとんど変わっていない唯一の場所です。
.cuファイルをnvccに渡すと、ホストコードとデバイスコードに分かれます。デバイスコードはPTXとCUBINという二つの形式で出力されます。CUDA 13.3のコンパイラドキュメントが説明する手順はこうです。デバイス関数がPTXまたはCUBINにコンパイルされ、それらがfatbinaryコンテナに格納され、元のソースはそのfatbinaryを内包する標準C++に変換されてホストコンパイラに渡されます。実行時には「埋め込まれたfatbinaryをCUDAランタイムが検査」し、今のGPUに合ったイメージを選びます。
ここで理解しておくべき概念が、仮想アーキテクチャと実アーキテクチャの区別です。
# compute_90 = 仮想アーキテクチャ、PTXを生成する命令セット世代
# sm_90 = 実アーキテクチャ、実チップ向けバイナリ
nvcc -gencode arch=compute_90,code=sm_90 \
-gencode arch=compute_90,code=compute_90 \
kernel.cu -o kernel
# 前の行はHopper向けのSASSを、後ろの行はPTX自体をfatbinaryに入れます。
# 後ろの行があることで、将来世代のGPUでもドライバがJITで蘇らせられます。
ドキュメントの表現どおり、PTXを一緒に入れておくと「現在のGPUに合ったバイナリロードイメージが存在しない場合、埋め込まれたPTXコードをCUDAランタイムが動的にコンパイル」します。これが前方互換性のすべてです。そしてこれこそが、PTXが消えない理由でもあります。PTXは人間が読めるアセンブリのように見えますが、実際には配布フォーマットであり、安定した契約です。SASSは世代ごとにエンコーディングが変わり文書化もされていませんが、PTXは文書化されており下位互換性が保たれています。
実務でこの区別が効いてくる場面があります。PyTorch 2.13のリリースノートを見ると、CUDA 13ビルドではptxasがもはやバイナリに同梱されていません。PTXをSASSに変換するツールがホイールの中にないということであり、PTX JITパスに依存する環境や、ptxasを直接呼び出すカスタムビルドスクリプトが突然壊れる可能性があります。ある層を知らなくてもよいということと、その層が存在しなくてもよいということは別の話です。
PTXからSASSへ至る最後の区間はptxasが担当します。この段階でレジスタ割り当てと命令スケジューリングが行われ、占有率を左右するレジスタ数もここで決まります。カーネルのソースは変えていないのにCUDAツールキットのマイナーバージョンだけを上げたら性能が変わった、という現象はほぼ常にptxasのスケジューリングが変わったことが原因です。以下のツールで確認できます。
# 実際に生成されたSASSを見ます
cuobjdump -sass ./kernel | head -40
# レジスタと共有メモリの使用量(占有率計算の入力値)
nvcc -Xptxas -v -arch=sm_90 -c kernel.cu
# ptxas info : Used 64 registers, 8192 bytes smem, 384 bytes cmem[0]
# PTXだけを取り出してみる
nvcc -ptx -arch=compute_90 kernel.cu -o kernel.ptx
ですが、GPU用の機械語を作るのはCUDA C++だけではありません。LLVMにはNVPTXと呼ばれるバックエンドがあり、LLVM IRをPTXへとロワリングします。Tritonも、XLAも、その他ほぼすべてのコード生成器が最終的にはこのバックエンドを通過します。つまりPTXは、CUDAの産物であると同時にLLVMの産物でもあるのです。
AMD側も構造は同じで、名前が違うだけです。LLVMのAMDGPUバックエンドがLLVM IRをAMDGCNアセンブリへとロワリングし、それがコードオブジェクトになります。ただしAMDはアップストリームのLLVMをそのまま使わず、フォークを維持しています。ROCmのドキュメントは自社コンパイラを「llvm/llvm-projectのフォーク」であると明記し、amdclang++をデフォルトコンパイラとして、hipccを「clangまたはnvccを呼び出し、適切なincludeとlibraryのオプションを渡す」ドライバとして説明しています。確認時点のドキュメントに記載されたllvm-projectのバージョンは22.0.0でした。
フォークを維持すること自体は悪い兆候ではありません。新しいチップのサポートや機能がアップストリームのリリース周期よりも速く必要になるためであり、AMDの変更分もかなりの部分がアップストリームに取り込まれています。ただし実務上は意味を持ちます。AMD GPUを対象とするサードパーティのコンパイラプロジェクトは、自分のLLVMバージョンとROCmのLLVMバージョンの間で調整をとらなければならず、これが二つのベンダースタックを比較した記事で扱った移植摩擦の一つの系統です。
MLIRとその上に築かれたもの — IRを作るための枠組み、そしてTriton
ここから先は地形が変わります。
LLVM IRはCPUのために設計されています。スカラーとベクトル、基本ブロック、SSA。テンソルという概念はなく、ループのネストという概念もありません。行列積をLLVM IRで表現すると、すでにループが展開された低レベルのコードになっているため、「このループをタイリングしよう」といった決定を下すための情報がもう残っていません。
そのため2010年代後半のあらゆるMLコンパイラが、同じ作業を繰り返すことになりました。テンソルレベルのIRを新たに定義し、パス基盤を新たに組み、パーサーと検証器と出力器を新たに作り、最後にLLVM IRへロワリングするコードを新たに書く。XLAも、TVMも、各ハードウェアベンダーも、それぞれが個別に行っていました。
MLIRの提案は、その繰り返しをなくそうというものです。IRをもう一つ定義する代わりに、IRを定義するためのインフラを提供します。演算と型のまとまりをダイアレクトと呼び、プロジェクトごとに自分のダイアレクトを定義しつつ、パス管理・検証・シリアライズ・ロワリングのフレームワークは共有します。
重要なのは、ダイアレクトが共存するという点です。一つのモジュールの中に抽象度の異なる演算が混在していてもよく、ロワリングはその混在を段階的に解消していく過程になります。
// 同じ関数の中で抽象度の異なる二つのダイアレクトが共存しています。
// linalgは「これは行列積である」という意味を保持し、
// scf/arithはすでにループとスカラー演算にまで落とし込まれています。
func.func @mm(%A: tensor<128x256xf32>, %B: tensor<256x64xf32>)
-> tensor<128x64xf32> {
%init = tensor.empty() : tensor<128x64xf32>
%C = linalg.matmul
ins(%A, %B : tensor<128x256xf32>, tensor<256x64xf32>)
outs(%init : tensor<128x64xf32>) -> tensor<128x64xf32>
return %C : tensor<128x64xf32>
}
linalg.matmulのたった一行が「これは行列積である」という情報をそのまま保持しています。タイリングパスはこの情報を見てループを切ることができます。同じものを先にLLVM IRまでロワリングしてしまうとその情報は失われ、以降はループの再構成がはるかに難しくなります。情報を手放すのを遅らせるほど、より良い決定を下せる — これがMLIRの存在理由を一文に凝縮したものです。
MLIRはそれ自体では何もコンパイルしません。その上に何を築くかがすべてです。そして実際に、多くのものが築かれてきました。TritonのIRも、IREEのIRも、XLAの一部のコンポーネントも、複数のハードウェアベンダーの社内コンパイラも、MLIRの上に乗っています。名前の異なるプロジェクトたちが同じ基盤を共有しているという事実は、この地形を理解するうえで重要です。
MLIRの上に築かれたもの — Tritonがカーネルを作る経路
Tritonはカーネル記述言語であると同時に、MLIRベースのコンパイラでもあります。カーネル記述の三つの層を比較した記事で言語としてのTritonは扱ったので、ここではコンパイラとしての位置づけを見ていきます。
パイプラインは明確に定義されています。PyTorchチームが2026年7月に公開したTritonプラグイン拡張のドキュメントは、各段階を次のように記しています。
Pythonカーネル関数
↓ (ASTトラバーサル)
TTIR — Triton IR。タイル単位の演算。ハードウェア非依存。
↓ (レイアウト決定、スレッドマッピング)
TTGIR — TritonGPU IR。ワープ配置と共有メモリが決定済み。
↓
LLVM IR
↓
PTX (NVIDIA) または AMDGCN (AMD)
この四段階は、この記事の前半で整理した「層」の構造を凝縮して示しています。TTIRはカーネル層、TTGIRはカーネル層と命令層の境界、LLVM IRより下は命令層です。ユーザーはおおよそTTIRのレベルでコードを書き、残りはコンパイラに任せます。これがTritonの取引条件です。
リポジトリのREADMEによれば、サポートされているバックエンドはNVIDIA(Compute Capability 8.0以上)とAMD(ROCm 6.2以上)で、CPUは「開発中」と表示されています。CPUバックエンドは複数の方向で試みられていますが、確認時点では実験段階であり、実務で期待できる対象にはまだなっていません。
2026年にこの層で実際に起きた変化のうち、注目に値するのがプラグイン拡張の仕組みです。PyTorchが配布するpytorch-triton 3.7に組み込まれ、PyTorch 2.13はTriton 3.7.1をピン留めしています。以前はカスタムパスやダイアレクトを追加しようとすると、Tritonをフォークする必要がありました。フォークはすぐにアップストリームに後れを取り、マージの衝突が積み重なり、最終的には新しいハードウェアサポートを受けられない古いリリースに縛られてしまいます。プラグインの仕組みは、ランタイムに共有ライブラリを載せることで、パイプラインの任意の地点にパスを挿入したり、特定のパスを無効にしたり、段階そのものを置き換えたりできるようにします。
# プラグインは環境変数で読み込まれます。importの順序が重要です。
import os, sysconfig
dist_packages = sysconfig.get_paths()["purelib"]
os.environ["TRITON_PLUGIN_PATHS"] = os.path.join(
dist_packages, "utlx_plugin", "libutlx.so"
)
import triton
import triton.language as tl
import utlx_plugin as tlx # TLX拡張演算がここで読み込まれます
最初の利用者がMetaのTLX(Triton Language Extensions)です。tlx.local_alloc、tlx.async_load、tlx.async_dotといった演算により、共有メモリバッファと非同期パイプライニングをカーネルの中で明示的に扱います。つまり、Tritonが自動で行っていた決定の一部を、人間の手に取り戻す方向性です。発表された数値では、H100のパーシステントGEMMでcuBLAS比2〜4パーセント上回り、MI350のパイプラインGEMMでrocBLAS比12〜15パーセント上回るとされています。ベンダーライブラリより速いカーネルがコンパイラの言語から生まれているというのが要点です。
ここに一つの流れが見えます。抽象化層は上へ向かうだけではありません。Tritonがスレッドを隠すことで上へ上がったのだとすれば、TLXは非同期命令を再び露出させることで下へ降りてきます。抽象化の高さを決めるのは性能ではなく、どの決定を人間が握るかです。
torch.compileとInductorが実際にやっていること
PyTorchユーザーが最もよく出会い、そして最も理解していない層です。torch.compile(model)というたった一行の裏で、三つの段階が順番に動いています。
TorchDynamoはPythonのバイトコードを横取りして、FXグラフを抜き出します。純粋なテンソル演算はグラフに入り、Pythonのリスト操作やprintのようなものはグラフの外に押し出されます。その地点がグラフブレイクです。グラフが複数の断片に切られると、それぞれの断片の前後でPythonに戻ってまた出るというコストが発生し、断片が小さいとフュージョンの機会も失われます。torch.compileが期待したほど速くならない事例の大半は、カーネルが悪いからではなく、グラフが切られてしまっているからです。
AOTAutogradはそのグラフを順伝播と逆伝播、二つのATenレベルのグラフに分割し、何を保存し何を再計算するかを決めます。
TorchInductorが実際のコードを生成します。GPUが対象ならTritonカーネルを、CPUが対象ならC++とOpenMPを生成します。そのため、先ほど見たtriton_poi_fused_...のような名前が出てくるのです。名前そのものが情報を含んでいます。poiはpointwiseの略で、fused_の後ろに続くものが、まとめられた演算のリストです。
生成されたコードは目で見ることができます。これがこの層を理解する一番の近道です。
import os
os.environ["TORCH_COMPILE_DEBUG"] = "1" # 生成コードをディスクに残します
os.environ["TORCH_LOGS"] = "output_code,graph_breaks"
import torch
def block(x, w, b):
return torch.nn.functional.gelu(x @ w + b)
x = torch.randn(1024, 512, device="cuda", dtype=torch.bfloat16)
w = torch.randn(512, 512, device="cuda", dtype=torch.bfloat16)
b = torch.randn(512, device="cuda", dtype=torch.bfloat16)
compiled = torch.compile(block)
out = compiled(x, w, b)
# グラフブレイクだけを個別に確認する方法
explanation = torch._dynamo.explain(block)(x, w, b)
print("グラフ数:", explanation.graph_count)
print("ブレイク数:", explanation.graph_break_count)
TORCH_LOGS="output_code"を有効にすると、Inductorが生成したTritonのソースが標準出力に出てきます。行列積はたいていcuBLAS呼び出しのまま残り、GELUとバイアス加算が一つのpointwiseカーネルにまとめられているのが見て取れます。そのカーネルこそが、先ほどプロファイラが示していた名もなきカーネルの正体です。
2026年、この層で起きた重要な変化がもう一つあります。PyTorch 2.13から、InductorがTritonに加えてCuTeDSL経路を追加で持つようになりました。NVIDIAのCuTeの上に構築されたDSLで、リリースノートによればトランスフォーマーのGEMMとRMSNormを狙い、「Tritonなしでもより高品質な行列積コードを生成する」第二の高性能経路です。カーネルのコンパイルがスレッドプールからサブプロセスプールへと移されたのも同じリリースで、これはPythonのGILがコンパイル時間のボトルネックになっていたためです。
同じリリースにはtorch.compiler.set_default_backendも入りました。torch.set_default_dtypeと同じ方式でプロセス全体のデフォルトバックエンドを設定するAPIで、すべてのtorch.compile呼び出しにbackend=を付ける必要がなくなります。ツリー外のバックエンドを構築するハードウェアベンダーにとって意味の大きい変化です。
まとめると、torch.compileはグラフコンパイラであり、カーネル生成は下位層に委譲します。その下位層はかつてTriton一つでしたが、今は二つになりました。
XLA、IREE、TVM — グラフを丸ごと受け取る側
InductorがPyTorchに付随するグラフコンパイラだとすれば、フレームワークから独立してグラフを受け取ろうとする系統が別に存在します。
XLAが最も古く、最も広く使われています。JAXの唯一の実行経路であり、TPUの唯一の実行経路でもあり、PyTorchもPyTorch/XLAを通じて接続します。入力はHLOで、フレームワークとコンパイラの間の契約はStableHLOが担います。StableHLOは後方5年、前方2年の互換性を約束する演算集合であるため、保存しておいたグラフを何年か後のコンパイラで開くということが成立します。
XLA:GPUのコード生成方式は、この記事のテーマとぴったり噛み合っています。公式ドキュメントの表現をそのまま訳すと、XLA:GPUは「ネイティブ(LLVM経由のPTX)エミッターとTritonIRエミッターの組み合わせ」を使います。これは三つの経路に分かれます。
- よくある演算は、cuBLAS、cuDNN、NCCLをそのまま呼び出します。
- リダクションや転置のようにパターンが認識できる演算は、LLVM IRを直接生成してPTXへとロワリングします。
- 行列積やソフトマックスが絡む高度なフュージョンは、HLOフュージョンをTritonIRへ変換し、タイルパラメータを選んだうえでTritonを呼び出してPTXを受け取ります。
三つ目が重要です。XLAとTritonは競合関係ではなく、呼び出す側と呼び出される側の関係にあります。XLAはグラフ層の決定を行い、カーネル層の決定はTritonに委ねます。先ほどの表で分けた層の構造が、そのまま実装されているわけです。最後に、XLAランタイムはカーネル呼び出しとライブラリ呼び出しの並びをRuntimeIRという独自のMLIRダイアレクトへと移し、そこからCUDAグラフを抽出します。
IREEは同じ問題を、はるかに広いターゲット範囲で解こうとします。READMEの表現では、「MLモデルを統合IRへとロワリングし、データセンターの要求まで拡張でき、モバイルとエッジの制約まで縮小できる」MLIRベースのコンパイラ兼ランタイムです。リポジトリのトピックタグには、フロントエンドとしてJAX、PyTorch、ONNX、TensorFlowが、ターゲットとしてCUDA、ROCm、Vulkan、SPIR-Vが挙げられています。2024年5月にLF AI & Data財団のサンドボックス段階のプロジェクトとして参加しました。
IREEの特徴は、コンパイラとランタイムを一緒に設計している点です。成果物がPythonプロセスを必要とするオブジェクトではなく自己完結した実行モジュールであるため、Pythonの存在しない組み込みターゲットまで同じパイプラインで到達できます。データセンターでIREEに出会うケースはまだ稀で、AMDが自社スタックの一部として推し進めていることが、確認時点における最大の採用事例です。
TVMは、この地形の中で立ち位置が最も大きく変わったプロジェクトです。2018年頃には唯一の本格的なオープンソース深層学習コンパイラであり、オートチューニングでベンダーライブラリに勝つという主張で注目を集めました。現在はRelaxをグラフレベルの表現、TensorIRをテンソルレベルの表現として使うクロスレイヤー設計であり、ドキュメントが掲げる設計目標は「ほとんどの変換をPythonからカスタマイズ可能にすることで、MLコンパイラをアクセスしやすくする」ことです。
率直に言えば、フロンティアLLMの学習とサービングの現場でTVMに出会うことは、今ではほとんどありません。その座はInductorとXLAが奪いました。TVMが生き残っている場所は別にあります。ベンダーライブラリの存在しないハードウェア、Pythonランタイムを載せられないエッジターゲット、そしてMLC LLMのようにブラウザやモバイルでLLMを動かすデプロイ経路です。オートチューニングがチューニング済みのライブラリに勝つという当初の命題は、cuBLASとCUTLASSが存在する場所では成立しにくく、それらが存在しない場所では今も成立しています。
一つの表にまとめると、次のようになります。
| XLA | IREE | TVM | Inductor | |
|---|---|---|---|---|
| 主なフロントエンド | JAX、TF、PyTorch/XLA | PyTorch、ONNX、JAX、TF | 複数フレームワークからのインポート | PyTorch専用 |
| グラフIR | HLO / StableHLO | MLIRダイアレクト | Relax | FX → ATen |
| カーネル生成 | LLVM直接生成 + Triton呼び出し | MLIRロワリング | TensorIR + オートチューニング | Triton、CuTeDSL、C++/OpenMP |
| ランタイム | XLAランタイム、PJRT | 独自の軽量ランタイム | 独自ランタイム | Pythonプロセス内部 |
| 強い領域 | TPU、JAX、大規模学習 | エッジからデータセンターまで | 非主流ハードウェア、エッジ展開 | PyTorchの標準経路 |
どの層を、いつ知るべきか
層が多いからといって、すべてを知る必要があるわけではありません。必要なときにだけ降りればよいのです。その「必要なとき」を示すサインが何かこそ、この記事で最も実用的な部分です。
| 今起きている症状 | 降りるべき層 | 具体的にやること |
|---|---|---|
| 学習がただ遅い | まだどの層でもない | まずプロファイラから。データローダーと同期ポイントが原因の大半です |
torch.compileを有効にしたのに速くならない | グラフ層 | torch._dynamo.explainでグラフブレイクを確認 |
| コンパイル自体に時間がかかりすぎる | グラフ層 | 動的shapeによる再コンパイルの有無、mode="reduce-overhead"の検討 |
| 生成されたカーネル名がプロファイルに出るが正体がわからない | カーネル層 | TORCH_LOGS="output_code"でTritonソースを読む |
| Inductorが作ったカーネルがライブラリより遅い | カーネル層 | Tritonで直接書くか、カスタムopに置き換える |
| Tritonカーネルが理論帯域幅の近くまで届かない | カーネル・命令の境界 | TTGIRダンプでレイアウトと共有メモリ配置を確認 |
| レジスタスピルが疑われる | 命令層 | -Xptxas -vまたはncuでレジスタ数を確認 |
| ツールキットのバージョンだけ上げたら性能が変わった | 命令層 | SASSの逆アセンブルを比較 |
| 新しいハードウェアにフレームワークを対応させる必要がある | 全層 | MLIRダイアレクトとPJRTプラグインの設計から |
経験上、実務者の90パーセントは上の二行で完結します。グラフブレイクを一つなくすことは、たいていカーネルを手で最適化するよりも大きな利益をもたらします。下へ降りるほど得られる利益は大きくなりますが、かかる時間ははるかに急な傾きで増えていきます。
Triton IRを直接見たい場合、ダンプの方法は次のとおりです。
# 各段階のIRをファイルとして残します
export TRITON_KERNEL_DUMP=1
export TRITON_DUMP_DIR=/tmp/triton_ir
python train.py
ls /tmp/triton_ir/*/
# *.ttir TTIR — タイル単位、ハードウェア非依存
# *.ttgir TTGIR — レイアウトとワープ配置が決定済みの状態
# *.llir LLVM IR
# *.ptx PTX
# *.cubin SASSを含む最終バイナリ
同じカーネルの五つのファイルを並べて開いてみるほうが、この記事全体を読むよりも理解が早いはずです。上から下へ降りていくにつれて、何が決定され何が消えていくのかが目に見えてきます。
最後に、今回の調査で確認できなかったことを明らかにしておきます。
- Tritonの安定版リリースのバージョンと日付です。リポジトリのREADMEはバージョンを明記しておらず、確認できたのはPyTorchがピン留めしている3.7.1のみでした。Triton自体のリリース番号と、PyTorchが配布するpytorch-tritonの番号は一致しない可能性があります。
- IREEの最新リリースタグと日付です。READMEには表示されておらず、別途の確認も行っていません。
- Apache TVMの最近のリリース時期と活動レベルです。プロジェクトのドキュメントから設計の方向性は確認しましたが、リリース履歴そのものは直接調べていません。上記の記述は採用状況についての判断であり、統計に裏付けられたものではありません。
- Helionです。PyTorchが2026年に推し進めているカーネルDSLであり、TPUバックエンドがPallasにコンパイルされることまでは確認できましたが、NVIDIAおよびAMDの経路でTritonとどのような関係にあるのかは、参照したドキュメントからは確認できませんでした。
おわりに — 層を分けるのは性能ではなく決定権です
この記事の地図を一文に畳み込むと、こうなります。層が分かれる基準は速度ではなく、誰がどの決定を握るかです。
グラフ層はフュージョンとレイアウトを引き受けます。その代わりに、個々のカーネルの細部は手放します。カーネル層はタイリングとパイプライニングを引き受けます。その代わりに、グラフ全体は見えません。命令層はレジスタとスケジューリングを引き受けます。その代わりに、何を計算しているのかを知りません。どの層も他の層を代替できない理由はここにあり、30年前にCPUコンパイラが経験した層の分化と同じ構造です。
だから、新しいプロジェクト名が出てきたときに投げるべき問いも決まっています。何が新しいかではなく、どの決定を誰から奪ったのかです。Tritonはスレッドマッピングを人間から奪いました。TLXは非同期命令をコンパイラから取り戻しました。Inductorはフュージョンをユーザーから奪いました。MLIRはIRインフラを各プロジェクトから奪いました。
この問い一つで、新しく出てくるフレームワークをこの地図のどこに置くべきか、たいてい決めることができます。そしてそれこそが、6か月ごとに変わる名前を全部追いかけるよりも長く使える知識です。
参考資料
- NVIDIA CUDA Compiler Driver NVCC 13.3 — コンパイル軌跡と仮想・実アーキテクチャ
- ROCm LLVM Projectドキュメント — amdclang、hipcc、comgr
- Tritonリポジトリ README — サポートされるバックエンドとMLIRベースの書き直し
- Triton Plugin Extensions: TLX and Custom Compiler Passes — PyTorchブログ、2026-07-15
- PyTorch 2.13リリースノート — CuTeDSLバックエンド、Triton 3.7.1ピン留め、CUDA 13の変更
- XLA:GPU Architecture Overview — ネイティブエミッターとTritonIRエミッター
- StableHLO — フレームワークとコンパイラの間の互換レイヤー
- IREEリポジトリ README — MLIRベースのコンパイラとランタイム
- Apache TVM — RelaxとTensorIRのクロスレイヤー設計
- MLIRプロジェクト — ダイアレクトと段階的ロワリング
현재 단락 (1/152)
Nsight Computeで学習ステップをプロファイルすると、最も時間を使っているカーネルの名前が出てきます。ところがその名前が`triton_poi_fused_add_mul_native_la...