【2026年】Triton コンパイラ入門ガイド｜カスタムCUDAカーネル

Triton コンパイラ入門ガイド：はじめに

近年、大規模言語モデル（LLM）や深層学習の推論・学習速度は、ハードウェアの進化だけでなく、ソフトウェア側の最適化によっても劇的に向上しています。特に NVIDIA の GPU 上で動作する AI パフォーマンスを最大化するためには、従来の CUDA C++ による低レベルなコーディングから脱却し、より高水準で効率的な言語の使用が求められています。本稿では、OpenAI が開発・提供している「Triton コンパイラ」について、2026 年 4 月時点の最新バージョンである Triton 3.2 と PyTorch 2.5 を基盤に解説します。

特に、最新の RTX 5090 や H100（Hopper アーキテクチャ）のような次世代 GPU を活用した自作 AI ワークステーションやデータセンター環境において、Triton を用いてカスタム CUDA カーネルを Python で記述する手法は、開発効率と実行性能のバランスを劇的に改善します。Flash Attention 2 のような高性能なアテンションカーネルの実装例を通じて、共有メモリ（Shared Memory）の効率的な利用やブロック並列処理の設計思想を理解し、独自の最適化技術を習得していただくことが本記事の目的です。

Triton の基本コンセプトとアーキテクチャ

Triton は、「Python による GPU プログラミング」というパラダイムシフトを実現したコンパイラツールチェーンです。従来の CUDA C++ では、メモリ管理やスレッド同期など低レベルな細工を記述する必要がありましたが、Triton では Python の構文に近い DSL（ドメイン固有言語）でカーネルを定義し、コンパイラ側が GPU 命令列へ自動的に変換してくれます。これにより、開発者はアルゴリズムの実装に集中でき、コードの可読性と保守性が大幅に向上します。2026 年現在では PyTorch のインテグレーター（torch.compile）と深く統合されており、ネイティブな Python コードから直接高効率な GPU カーネルが生成される環境が整っています。

ハードウェア方面では、Hopper アーキテクチャを搭載した H100 や、2025 年末に発売された Blackwell アーキテクチャ搭載の RTX 5090 が主要なターゲットです。これらの GPU は従来のアーキテクチャとは異なるメモリ階層と計算ユニットを備えています。例えば、H100 では HBM3e（ハイバンド幅メモリー）により 3.35TB/s という圧倒的な帯域幅を提供し、Triton カーネルはこれを最大限に活用できるよう設計されています。RTX 5090 はコンシューマー向けでありながらプロ向け GPU に匹敵する性能を持ち、64GB の GDDR7 メモリと 1.5TB/s を超える帯域幅を誇ります。Triton コンパイラは、こうした最新のハードウェア特性を認識し、Tensor Core や Tensor Memory Accelerator（TMA）を活用した最適化コードを生成します。

CUDA C++ との最大の違いは「抽象度」にあります。CUDA ではスレッドブロックごとのインデックス計算や、メモリ転送の同期処理を手動で行う必要がありますが、Triton には tl.load や tl.store といった高機能な API が用意されており、複雑なメモリアクセスパターンも簡潔に記述できます。また、Triton 3.2 では自動チューニング機能が強化され、実行環境のハードウェア構成（メモリ容量やコア数）を自動的に検知し、最適なブロックサイズやタイルサイズを選択するようになりました。これにより、異なる GPU で実行する場合でも、手動でのパラメータ調整なしに高パフォーマンスを発揮できるようになっています。

Triton DSL 解説：Python ベースのカーネル記述

Triton の核心である DSL（ドメイン固有言語）は、Python の構文をベースとしつつ、GPU 特有の概念を組み込んだ拡張言語です。基本的なカーネル定義には @triton.jit デコレータを使用し、関数内で GPU スレッドが実行するロジックを記述します。このデコレータがあることで、Python コードはコンパイラによって解析され、GPU 用のアセンブリ語に変換されます。2026 年時点の Triton 3.2 では、型推論機能も強化されており、明示的な型注釈がなくても多くの場合で正確なデータ型を判別できるようになっていますが、安定性を期すため tl.dtype を用いた明示的な指定は推奨されます。

カーネル関数内で使用される変数は、レジスタや共有メモリ上の位置にマッピングされます。特に重要なのは、引数の型とメタデータの指定です。例えば、行列演算を行う場合、入力テンソルの形状（shape）やデータ型（dtype）をコンパイル時に固定するか、実行時ダイナミックにするかを指定する必要があります。tl.dot 関数を使用して行列乗算を実行する際、引数のデータ型が FP16 か BF16 か INT8 かによって計算ユニットの選択が変わります。また、Triton は SIMD（単一命令多重化）および SIMT（スレッドインストメンタル多重化）モデルに基づいて動作するため、同じコードが数千のスレッドで並列に実行されます。

データ型に関する詳細も DSL の重要な要素です。Triton 3.2 では FP8 形式のサポートが本格的に強化されており、tl.float8_e4m3fnz や tl.float8_e5m2 といった型が利用可能です。これは推論コストを削減し、RTX 5090 の Tensor Core でより高速な計算を実現するために不可欠です。さらに、メモリアクセスの粒度を制御する BLOCK_SIZE パラメータも DSL 内で定義されます。例えば @triton.autotune(configs=[...]) デコレータと共に、ブロックサイズを最適化してキャッシュヒット率を上げることが可能になります。DSL の記述は Python のインデント規則に則るため、Python プログラマにとって学習コストは低く抑えられています。

ブロック並列処理とメモリ階層の最適化

GPU 上でカーネルを実行する際、データがどのように配置され、どのように読み込まれるかが性能を決定づけます。Triton では「ブロックごとの並列処理」が基本単位となります。これは、巨大な行列やテンソルを小さなタイル（block）に分割し、複数のストリーミングマルチプロセッサ（SM）で同時に計算を行う手法です。例えば、1024x1024 の行列乗算を行いたい場合、これを 64x64 のブロックに分割し、各 SM が特定のブロックを担当して処理します。このアプローチにより、メモリ帯域幅のボトルネックを分散させ、全体のスループットを最大化できます。

メモリアクセスにおける「共有メモリ（Shared Memory）」の利用は、Triton 開発者が最も注力すべきポイントの一つです。これは GPU の SM 内に存在する高速な SRAM で、スレッドブロック内の全スレッドがアクセス可能です。グローバルメモリからデータを一度読み込み、共有メモリにキャッシュした後で計算を行うことで、メモリアクセスのレイテンシを隠蔽できます。Triton では tl.to_tensor_memory や tl.atomic_add などの機能を用いて、共有メモリへの書き込みと読み出しを効率的に行います。RTX 5090 の場合、SM あたりの共有メモリ容量が拡張されているため、より大きなデータセットをローカルに保持することが可能になっています。

キャッシュの整合性を保つための同期処理も重要です。スレッドブロック内のすべてのスレッドが共有メモリへの書き込みを終了し、そのデータを次の計算ステップで使用できるようになるまで、tl.static_barrier() や tl.sync() を用いて同期を行います。これを怠ると、一部のデータがまだ読み込まれていない状態で計算が行われ、不正な結果を招きます。また、メモリアクセスの効率化には「メモリ結合（Memory Coalescing）」という概念も不可欠です。隣接するスレッドが連続したメモリアドレスを同時にアクセスすることで、1 つのメモリ転送で複数のデータを取得でき、帯域幅の利用率を高めます。Triton のコンパイラは自動でこれらの最適化を試行しますが、コードを書き直すことでさらに効果的な結合を実現できる場合があります。

行列乗算（MatMul）カーネルの実装例

最も基本的かつ重要な演算である行列乗算（GEMM）のカーネルを Triton で実装する例を見てみましょう。これは AI モデルにおいて線形層の推論や学習に頻繁に使用される演算です。以下のコードは、2 つの入力行列 A と B を受け取り、結果の C を生成するシンプルな MatMul カーネルの skeleton です。この実装では、ブロックサイズを固定し、共有メモリを使用してデータをローディングしています。

import triton
import triton.language as tl
import torch

@triton.jit
def matmul_kernel(
    A_ptr, B_ptr, C_ptr,
    M, N, K,
    stride_am, stride_ak,  # A のストライド
    stride_bk, stride_bn,  # B のストライド
    stride_cm, stride_cn,  # C のストライド
    BLOCK_M: tl.constexpr,
    BLOCK_N: tl.constexpr,
    BLOCK_K: tl.constexpr,
):
    pid = tl.program_id(0)
    num_pid = (M * N + BLOCK_M * BLOCK_N - 1) // (BLOCK_M * BLOCK_N)
    
    # B のインデックス計算
    start_bk = pid % (N // BLOCK_N)
    start_am = pid // (N // BLOCK_N)
    
    a_block_ptr = tl.make_block_ptr(
        base=A_ptr, shape=(M, K),
        strides=(stride_am, stride_ak),
        offsets=(start_am * BLOCK_M, 0),
        block_shape=(BLOCK_M, BLOCK_K),
        order=(1, 0)
    )
    
    b_block_ptr = tl.make_block_ptr(
        base=B_ptr, shape=(K, N),
        strides=(stride_bk, stride_bn),
        offsets=(0, start_bk * BLOCK_N),
        block_shape=(BLOCK_K, BLOCK_N),
        order=(1, 0)
    )

    # ここにロード・計算・ストアのロジックが続く...

このコードの make_block_ptr 関数は、Triton の強力な機能の一つです。メモリポインタを直接扱う代わりに、ブロックごとのオフセットとストライドを指定してスケーラブルなアクセスを実現します。これにより、メモリ配置が変わってもコードを書き換える必要がなくなります。計算部分では、tl.dot 関数を使用して行列積を実行しますが、ここでは省略しています。実際の性能最適化では、A と B のデータを共有メモリにロードし、スレッド間で累算を行うループを記述します。

具体的な計算ロジックは以下のようになります。まず、A のブロックと B のブロックを各々 tl.load で読み込みます。次に、acc = tl.dot(a_tile, b_tile) を実行して部分和を計算します。これを K 次元のサイズ（BLOCK_K）でループし、最終結果を C に書き込みます。RTX 5090 のような Blackwell アーキテクチャでは、この dot 演算がハードウェアレベルで Tensor Core にマッピングされ、FP8 や FP16 で極めて高速に処理されます。また、C のストライド計算は、結果の次元に応じて正しく設定される必要があります。誤ったストライド指定は、メモリの重複書き込みや未定義動作を引き起こすため注意が必要です。

自動チューニング機能と性能プロファイリング

Triton コンパイラの最大の利点の一つに「自動チューニング」機能があります。手動でブロックサイズ（BLOCK_SIZE）やタイルサイズを調整するのは時間がかかる作業ですが、Triton の @triton.autotune デコレータは、指定された構成候補の中から実行時に最適な設定を選択します。これは、異なる GPU やメモリ容量を持つ環境でも、同じコードベースで高パフォーマンスを発揮することを可能にします。2026 年時点の Triton 3.2 では、このチューニングプロセスがさらに高速化され、コンパイル時のオーバーヘッドを大幅に削減しています。

自動チューニングでは、実行速度だけでなく、メモリ使用量やキャッシュ効率も考慮されます。具体的には、triton.testing.do_bench() を用いて複数の構成でベンチマークを行い、最小の実行時間を示す構成を選択します。例えば、BLOCK_M = 128, BLOCK_N = 128 の設定は H100 では最適ですが、メモリ容量の少ない RTX 4090 ではオーバーフローする可能性があります。Triton はこれらのハードウェア制約を認識し、自動的に安全かつ高速な構成を選定します。また、triton.testing.profiler を使用して、カーネルの実行時間とメモリアクセス量を可視化することもできます。

プロファイリング結果を元に、ボトルネックの特定も行えます。例えば、実行時間が長い場合、それが計算リソース不足か、メモリ帯域幅不足かで対策が変わります。Triton のビルトインツールは、カーネル内の各ステートメントの実行時間を計測し、どの部分が遅延の原因かを特定します。これにより、メモリアクセスの最適化や並列度の調整を効果的に行えます。さらに、PyTorch 2.5 と統合された場合、torch.compile を介してプロファイリングが行われ、コンパイラレベルでの最適化パスが可視化されます。このように、自動チューニングとプロファイリングは密接に連携し、開発者が性能を最大化するのを支援します。

Flash Attention の Triton 実装解説

Flash Attention は、従来のアテンション計算の O(n^2) 複雑度を低減し、メモリ使用量を劇的に削減したアルゴリズムです。Triton を用いることで、この Flash Attention 2 の高性能なカーネルを Python で実装することが可能になります。Flash Attention の核心は、出力テンソル（O）と softmax 統計量（L, M）を計算しながら、入力シーケンスの一部分のみを読み込む「I/O-aware」なアプローチです。これにより、巨大なシーケンス長でも HBM を頻繁にアクセスすることなく処理を進められます。

具体的な実装では、入力 Q, K, V のデータを読み込み、スレッドブロック内でローカル softmax を計算します。Triton 3.2 では、このプロセスを tl.reduce や tl.max といった関数で効率的に記述できます。重要なのは、シーケンス長が長い場合でもメモリバウンドを起こさないように、Q, K, V のブロックサイズを適切に調整することです。RTX 5090 のような GPU では、GDDR7 の帯域幅を活かしつつ、L1/L2 キャッシュの命中率を最大化するために、データのリファレンスパターンを工夫します。

Flash Attention の実装例では、以下のステップを踏みます。まず、Q, K, V をブロックごとに読み込みます。次に、各ブロック内での最大値（M）と和（L）を計算し、これらを用いて正規化を行います。最後に、結果の O に加算します。この際、tl.atomic_add を使用して複数のスレッドからの書き込みを同期させます。Triton 実装版 Flash Attention は、公式ライブラリよりも柔軟なカスタマイズが可能で、特定のアーキテクチャ向けに最適化することもできます。2026 年時点では、FlagAttention や Liger Kernel のような派生バージョンも Triton ベースで開発されており、これらとの互換性や比較検討が重要です。

PyTorch との統合および torch.compile

PyTorch は深層学習フレームワークのデファクトスタンダードであり、2026 年時点では PyTorch 2.5 のバージョンが主流です。この環境下で Triton カーネルを効果的に利用するには、torch.compile を介した統合が最も効率的です。torch.compile は NVIDIA のインテグレーター（torchinductor）を使用して、PyTorch の計算グラフを最適化し、Triton コードに変換します。これにより、開発者は明示的にカーネルを書くことなく、PyTorch の Python コードから自動的に高性能な GPU 実行コードが生成されます。

統合のプロセスはシンプルです。torch.compile(model, backend="triton") と指定するだけで、コンパイラが Triton ベックエンドを選択します。これにより、モデル内の全ての演算が最適化されたカーネルに置き換えられます。特に、ユーザー定義のオペレーション（Custom Ops）や特定の数学関数において、Triton の自動チューニング機能が発揮されます。また、torch.compile はグラフキャプチャ機能も提供しており、動的な形状を持つ入力テンソルに対しても効率的な実行計画を立てます。

この統合により、開発者は CUDA カーネルの複雑さを直接扱わずに済みますが、パフォーマンスを微調整したい場合は手動で Triton カーネルを実装し、torch.library.register_fake を用いて PyTorch に登録することも可能です。これにより、既存のモデル構造を変更せずに、特定の演算部分のみを高速化できます。2026 年時点では、PyTorch と Triton の連携は非常に緊密であり、エッジケースでの互換性も向上しています。ただし、最新の PyTorch バージョンを使用しないと、一部の新しいハードウェア機能（Blackwell 用など）が活用できないため、環境の構築には注意が必要です。

競合ライブラリとの比較（CUDA, xformers, Liger）

高性能な AI ランタイムを構築する際、Triton の他に CUDA C++ や他のライブラリも選択肢となります。特に Flash Attention 2 や Liger Kernel との比較は重要です。CUDA C++ は最も低レベルで制御力が高いですが、開発コストとデバッグ難易度が高くなります。一方、xformers や Liger Kernel は Triton を基盤としたライブラリであり、より高い抽象度を提供します。下表に、主要なライブラリと実装手法の性能比較を示します。

CUDA C++ はメモリ帯域幅利用率において依然として最高レベルを維持していますが、その分、コード記述に数日かかることもあります。Triton を手動で実装する場合、Python の利便性を保ちつつ 85-90% の性能を発揮できます。特に複雑なアテンション計算では、Triton の自動チューニングが CUDA の手動最適化に近い結果を出すことも珍しくありません。xformers はセットアップの容易さが特徴ですが、ハードウェア特性への細かい制御は制限されます。Liger Kernel は Triton ベースでありながら、PyTorch との統合がスムーズで、特に混合精度計算において優れたパフォーマンスを示します。

また、Flash Attention 2 の実装においても、公式ライブラリと Triton 版の実装では差異があります。Triton 3.2 では Flash Attention 2 のアルゴリズムを Python で再現可能であり、特定のレイアウトに最適化できます。例えば、H100 の TMA（Tensor Memory Accelerator）を活用する場合、Triton を用いて HBM と L1 キャッシュの間の転送を細かく制御することで、帯域幅効率をさらに向上させられます。2026 年時点では、各ライブラリのバージョンアップにより互いの性能差は縮まってきていますが、特定のワークロードにおいては Still CUDA C++ が頭一つ抜けているケースもあります。

デバッグとエラー処理の実践

Triton カーネルの開発において、最も難しい部分の一つがデバッグです。GPU での実行時エラーは、CPU のスタックトレースとは異なり、詳細な情報が得にくい場合があります。triton.testing モジュールには、カーネルの単体テストやプロファイリングを行うための関数が用意されています。特に do_bench() は、カーネルの実行時間を正確に計測し、初期化コストを除いた純粋な計算時間を評価するために使用されます。これにより、最適化前後の性能差を明確に定量化できます。

エラー処理では、コンパイル時の警告や実行時エラーを適切に処理することが重要です。Triton 3.2 では、型不一致によるエラーメッセージが改善されており、どの行で問題が発生したかがより明確に表示されます。ただし、メモリリークやデッドロックのような問題は依然として検出しにくい場合があります。このような場合、nvprof や nsys といった NVIDIA のプロファイリングツールを併用して、カーネルの実行履歴を確認することが有効です。また、Triton のコンパイラログは詳細な情報を出力するため、エラー発生時には必ずログを出力設定に切り替えて分析を行います。

具体的には、ブロックサイズやストライドの計算ミスが原因で生じるアライメントエラーやオーバーフローに注意します。これらは実行時にクラッシュせず、不正な値を返すことがあります。デバッグ用のブレイクポイント機能は GPU 上で直接動作しないため、CPU 側から呼び出すロジック内に print や torch.testing.assert_close を組み込んで、中間結果を確認する手法が一般的です。また、Triton のサードパーティ製デバッガーも存在し、複雑なグラフ構造の可視化や変数の監視が可能です。これらのツールを駆使することで、開発効率とコード品質を同時に向上させることができます。

よくある質問（FAQ）

Q1. Triton と CUDA C++ のどちらを使うべきですか？ A. 基本的には PyTorch や既存ライブラリで対応している場合はそちらを使用し、性能がボトルネックとなる場合や特殊な演算が必要な場合にのみ手動カーネル開発（CUDA または Triton）を検討してください。CUDA は制御性が高く、Triton は開発効率とパフォーマンスのバランスに優れています。

Q2. RTX 5090 で TMA を有効にするにはどうすればよいですか？ A. Triton 3.2 以降では tl.make_block_ptr の引数やコンパイラ設定を調整することで、Blackwell アーキテクチャの TMA（Tensor Memory Accelerator）を利用可能です。具体的には、block_order を指定し、メモリアクセスパターンが HBM と L1 キャッシュに最適化されるよう設定します。

Q3. Flash Attention 2 の実装で注意すべき点は？ A. 最大の注意点シーケンス長とメモリ容量のバランスです。ブロックサイズを小さくしすぎるとオーバーヘッドが増え、大きくしすぎると共有メモリ不足になります。Triton の自動チューニングを使用して、対象 hardware に最適な値を見つけ出すことを推奨します。

Q4. PyTorch 2.5 と Triton 3.2 の互換性は？ A. 2026 年 4 月時点では完全な互換性が保証されています。ただし、最新の機能を利用するには両方を最新版にアップグレードしてください。古いバージョンでは torch.compile のベックエンドとして指定できない場合があります。

Q5. Liger Kernel と xformers を同時に使用できますか？ A. はい、可能です。ただし、特定の演算パスで競合する可能性があります。通常はモデル全体を xformers で処理し、特定のレイヤーのみを Liger Kernel に置き換えるハイブリッド構成が推奨されます。

Q6. Triton カーネルのメモリ使用量を制限する方法は？ A. @triton.autotune の config 内でメモリの上限を指定できます。また、共有メモリへのデータロード量を抑えるため、ブロックサイズを動的に調整するロジックを実装することも可能です。

Q7. 自動チューニングのオーバーヘッドはどの程度ですか？ A. コンパイル時に行われるため、実行時のパフォーマンスには影響しません。ただし、コンパイル時間が数分になることがあり、頻繁なコード変更には注意が必要です。2026 年時点ではキャッシュ機能が強化され、この時間は短縮されています。

Q8. FP8 計算を有効にするには？ A. tl.float8_e4m3fnz などのデータ型を使用し、モデルの重みと活性化値を FP8 に変換する必要があります。PyTorch 2.5 では torch.compile が自動的に FP8 パスを選択するオプションを提供しています。

Q9. デバッグ時に print を使うのは問題ありませんか？ A. GPU 上での print はデバッグには有効ですが、パフォーマンスに悪影響を与えるため本番環境では使用しないでください。代わりにプロファイリングツールやログ出力関数を使用してください。

Q10. Triton コードのポート方法は？ A. 既存の CUDA カーネルを Triton に移植するには、まずメモリアクセスパターンと並列処理構造を把握し、tl.load/store API で置き換えることから始めます。Triton の公式ドキュメントには翻訳ガイドも用意されています。

まとめ

本記事では、2026 年 4 月時点の最新技術環境において、OpenAI Triton コンパイラを用いたカスタム CUDA カーネルの実装方法を詳しく解説しました。以下に本稿の要点をまとめます。

• Triton の基本: Python ベースで GPU コードを記述できる高機能コンパイラであり、PyTorch 2.5 との統合がスムーズ。
• アーキテクチャ対応: H100 や RTX 5090（Blackwell）などの最新 GPU を活用し、[メモリ帯域幅](/glossary/帯域幅)と計算ユニットを最大限に利用可能。
• DSL と最適化: tl.dot や make_block_ptr といった API を用いた行列演算やブロック並列処理の設計が重要。
• 自動チューニング: @triton.autotune により、環境に応じた最適なパラメータを自動的に選択し、開発コストを削減。
• Flash Attention: Triton を用いて [Flash Attention](/glossary/attention) 2 の実装が可能であり、メモリ効率と計算速度の両面で劇的な改善が見込める。
• ライブラリ比較: CUDA C++ との性能差は縮小しており、特に複雑なアテンション演算では Liger Kernel や xformers と併用する戦略が有効。
• デバッグ: triton.testing モジュールやプロファイリングツールを活用し、エラーとボトルネックを特定することが重要。

Triton を活用することで、AI ワークステーションの性能をソフトウェア側から最大化できる可能性が開かれます。特に自作 PC 愛好家にとっては、ハードウェア購入後のチューニングとして非常に有効な手段です。本ガイドが、高性能 AI システム構築への一助となることを願っています。