この記事の対象読者
- WindowsでStable Diffusion、ComfyUI、FramePackなどのAIツールを使っている方
- 「A matching Triton is not available」というエラーに遭遇して困っている方
- PyTorchやxformersを使った開発でTriton関連の警告が気になっている方
- WSLを使わずにWindowsネイティブでAI開発をしたい方
この記事で得られること
- Tritonとは何か、なぜWindowsで問題になるのかの理解
- triton-windowsを使った完全な解決手順
- 環境別(ComfyUI/FramePack/一般Python環境)のセットアップ方法
- よくあるエラーの対処法とトラブルシューティング
この記事で扱わないこと
- WSL(Windows Subsystem for Linux)を使った回避策
- AMD GPU向けのTritonセットアップ(NVIDIA GPU専用)
- Tritonカーネルの自作・チューニング
1. Tritonとの出会い
「A matching Triton is not available, some optimizations will not be enabled.」
ComfyUIを起動するたびに表示されるこの警告。Stable Diffusion WebUIでxformersを有効にしようとしたときのあのエラー。FramePackで動画生成しようとしたら突然のクラッシュ。
私もWindowsでAI開発を始めた頃、このTritonという謎の存在に何度も心を折られました。ググっても「WSLを使え」「Linuxに移行しろ」という回答ばかり。でも、メモリ16GBのノートPCでWSLのスワップ管理なんて地獄でしかないし、そもそもWindowsで動かしたいからWindowsを使ってるんだよ、と。
そんな中で出会ったのがtriton-windowsプロジェクト。これがあれば、WSLなしでTritonをWindowsネイティブで動かせます。
Tritonとは、簡単に言えば「GPUプログラミングを劇的に簡単にするコンパイラ」です。CUDAを書かなくても、Pythonライクな記法で高速なGPUカーネルが書ける。PyTorchのtorch.compileやxformersの最適化機能は、このTritonに依存しています。
ここまでで、Tritonがどんな存在か、なんとなくイメージできたでしょうか。次は、この技術が生まれた背景と、なぜWindowsでサポートされてこなかったのかを見ていきましょう。
2. 前提知識の確認
本題に入る前に、この記事で使う用語を整理しておきます。
2.1 GPUカーネルとは
GPU上で実行される計算処理の単位です。AIモデルの推論や学習では、行列演算などの重い計算をGPUカーネルとして実行することで高速化しています。従来はCUDA(NVIDIA独自の言語)で書く必要がありました。
2.2 PTX(Parallel Thread Execution)とは
NVIDIAのGPU向け中間言語です。高級言語(Pythonなど)からPTXにコンパイルされ、さらにGPUが直接実行できるバイナリ(cubin)に変換されます。人間が読めるアセンブリに近い形式です。
2.3 LLVM(Low Level Virtual Machine)とは
様々なプログラミング言語のコンパイラを作るための基盤技術です。Tritonは内部でLLVMを使用しており、Pythonコードを最適化されたGPUコードに変換します。
2.4 xformersとは
Meta(Facebook)が開発したTransformerモデル向けの最適化ライブラリです。メモリ効率の良いAttention計算を提供し、Stable Diffusionなどで広く使われています。内部でTritonを使用しています。
2.5 JIT(Just-In-Time)コンパイルとは
プログラムの実行時にコンパイルを行う方式です。Tritonは@triton.jitデコレータを付けた関数を、実行時にGPUコードにコンパイルします。事前コンパイルと違い、実行環境に最適化されたコードを生成できます。
これらの用語が押さえられたら、次に進みましょう。
3. Tritonが生まれた背景
3.1 OpenAIによる開発
Tritonは2019年にHarvard大学のPhilippe Tilletによって論文が発表され、2020年からOpenAIが開発を引き継ぎました。
公式サイトの説明を引用します。
Triton makes it possible to reach peak hardware performance with relatively little effort; for example, it can be used to write FP16 matrix multiplication kernels that match the performance of cuBLAS—something that many GPU programmers can't do—in under 25 lines of code.
つまり、「多くのGPUプログラマーでも難しいcuBLAS並みの性能を、たった25行のPythonで実現できる」というのがTritonの価値です。
3.2 なぜWindowsがサポートされなかったのか
Tritonの内部アーキテクチャを簡単に説明すると、以下のような流れになっています。
Python AST → Triton-IR → LLVM-IR → PTX → cubin(GPUバイナリ)
この変換パイプラインは、LinuxのPOSIX API(ファイルシステム、プロセス管理など)に依存した実装になっていました。具体的には以下の問題がありました。
-
aligned_alloc関数がWindowsに存在しない(_aligned_mallocに置き換えが必要) - ファイルパスの扱いがLinux前提(
/と\の違い) - 一部のシステムコールがWindowsで異なる動作をする
OpenAIとしては、主要なユースケースがLinuxのサーバー/クラウド環境だったため、Windowsサポートの優先度が低かったのです。
3.3 triton-windowsプロジェクトの誕生
この状況を変えたのが、コミュニティの有志によるtriton-windowsプロジェクトです。andreigh、wkpark、woct0rdhoらの開発者が、Windows固有の問題を一つずつ解決していきました。
現在では、PyPIからpip install triton-windowsでインストールでき、公式Tritonとほぼ同等の機能がWindowsネイティブで動作します。
背景がわかったところで、抽象的な概念から順に、具体的な仕組みを見ていきましょう。
4. Tritonの基本概念
4.1 ブロックベースのプログラミングモデル
Tritonの最大の特徴は「ブロック」という概念です。CUDAでは個々のスレッドを意識する必要がありますが、Tritonではデータのブロック(塊)単位で処理を記述します。
import triton
import triton.language as tl
@triton.jit
def add_kernel(x_ptr, y_ptr, output_ptr, n_elements, BLOCK_SIZE: tl.constexpr):
# プログラムID = 何番目のブロックか
pid = tl.program_id(axis=0)
# このブロックが担当するデータの開始位置
block_start = pid * BLOCK_SIZE
# ブロック内のオフセット(0, 1, 2, ..., BLOCK_SIZE-1)
offsets = block_start + tl.arange(0, BLOCK_SIZE)
# 境界チェック用マスク
mask = offsets < n_elements
# データのロード・計算・ストア
x = tl.load(x_ptr + offsets, mask=mask)
y = tl.load(y_ptr + offsets, mask=mask)
output = x + y
tl.store(output_ptr + offsets, output, mask=mask)
このコードでは、長大なベクトルをBLOCK_SIZE個ずつのブロックに分割し、各ブロックを並列に処理しています。
4.2 自動最適化
Tritonコンパイラは、ユーザーが書いたコードに対して自動的に以下の最適化を適用します。
| 最適化 | 説明 |
|---|---|
| メモリ・コアレッシング | 連続したメモリアクセスをまとめて効率化 |
| 共有メモリの活用 | 頻繁にアクセスするデータをGPUのSRAMにキャッシュ |
| ループ展開 | ループを展開して分岐オーバーヘッドを削減 |
| レジスタ割り当て | 限られたレジスタを効率的に使用 |
CUDAでこれらを手動実装すると数百行のコードが必要ですが、Tritonなら数十行で同等の性能を実現できます。
4.3 PyTorchとの統合
PyTorch 2.0以降では、torch.compileによるモデル高速化機能が追加されました。この機能は内部でTritonを使用しています。
import torch
# 通常の関数定義
def my_function(x, y):
return torch.relu(x @ y)
# torch.compileでTritonカーネルに自動変換
compiled_fn = torch.compile(my_function)
torch.compileを使うと、PyTorchの計算グラフが分析され、Tritonカーネルとして最適化されたコードが生成されます。
基本概念が理解できたところで、これらの抽象的な概念を具体的な実装に落とし込んでいきましょう。
5. 実際に使ってみよう
5.1 環境構築
以下の3種類の環境設定を用意しています。用途に応じて選択してください。
環境設定(基本・Python venv環境用)
# setup_basic.ps1 - 基本環境用(このままコピーして使える)
# 対象: 通常のPython環境(venv/システムワイド)
# 1. Python環境の確認
python --version # Python 3.10〜3.13が必要
# 2. PyTorchのインストール(CUDA 12.x版)
pip install torch torchvision torchaudio --index-url https://download.pytorch.org/whl/cu124
# 3. triton-windowsのインストール
# PyTorch 2.9 + Triton 3.5の組み合わせ
pip install -U "triton-windows<3.6"
# 4. 動作確認
python -c "import triton; print('Triton version:', triton.__version__)"
環境設定(ComfyUI Portable用)
# setup_comfyui.ps1 - ComfyUI Portable用(このままコピーして使える)
# 対象: ComfyUI_windows_portableを使用している場合
# 1. ComfyUIのディレクトリに移動
cd C:\path\to\ComfyUI_windows_portable
# 2. embeded Pythonのパスを確認
# .\python_embeded\python.exe が存在することを確認
# 3. include/libsフォルダの追加(重要)
# https://github.com/woct0rdho/triton-windows/releases/v3.0.0-windows.post1/
# からPythonバージョンに合ったzipをダウンロード
# 例: python_3.13.2_include_libs.zip
# 解凍して python_embeded フォルダに include と libs を配置
# 4. triton-windowsのインストール
.\python_embeded\python.exe -m pip install -U "triton-windows<3.6"
# 5. 動作確認
.\python_embeded\python.exe -c "import triton; print('Triton version:', triton.__version__)"
環境設定(FramePack用)
# setup_framepack.ps1 - FramePack用(このままコピーして使える)
# 対象: FramePackを使用している場合
# 1. FramePackのディレクトリに移動
cd C:\path\to\FramePack
# 2. Pythonパスの確認(FramePackはsystem\pythonにある)
# .\system\python\python.exe が存在することを確認
# 3. include/libsフォルダの追加(重要)
# https://github.com/woct0rdho/triton-windows/releases/v3.0.0-windows.post1/
# からPython 3.10用のzipをダウンロード: python_3.10.11_include_libs.zip
# 解凍して system\python フォルダに include と libs を配置
# 4. triton-windowsのインストール
.\system\python\python.exe -m pip install -U "triton-windows<3.6"
# 5. 動作確認
.\system\python\python.exe -c "import triton; print('Triton version:', triton.__version__)"
5.2 PyTorchとTritonのバージョン対応表
triton-windowsは、PyTorchのバージョンに合わせて適切なバージョンを選ぶ必要があります。
| PyTorch | Triton | インストールコマンド |
|---|---|---|
| 2.9 | 3.5 | pip install -U "triton-windows<3.6" |
| 2.8 | 3.4 | pip install -U "triton-windows<3.5" |
| 2.7 | 3.3 | pip install -U "triton-windows<3.4" |
| 2.6 | 3.2 | pip install -U "triton-windows<3.3" |
5.3 基本的な動作確認
"""
Triton動作確認スクリプト
使い方: python test_triton.py
必要なパッケージ: pip install torch triton-windows
"""
import torch
import triton
import triton.language as tl
@triton.jit
def add_kernel(x_ptr, y_ptr, output_ptr, n_elements, BLOCK_SIZE: tl.constexpr):
"""ベクトル加算カーネル"""
pid = tl.program_id(axis=0)
block_start = pid * BLOCK_SIZE
offsets = block_start + tl.arange(0, BLOCK_SIZE)
mask = offsets < n_elements
x = tl.load(x_ptr + offsets, mask=mask)
y = tl.load(y_ptr + offsets, mask=mask)
output = x + y
tl.store(output_ptr + offsets, output, mask=mask)
def add(x: torch.Tensor, y: torch.Tensor) -> torch.Tensor:
"""Tritonカーネルを使ったベクトル加算"""
output = torch.empty_like(x)
n_elements = output.numel()
grid = lambda meta: (triton.cdiv(n_elements, meta["BLOCK_SIZE"]),)
add_kernel[grid](x, y, output, n_elements, BLOCK_SIZE=1024)
return output
def main():
print(f"PyTorch version: {torch.__version__}")
print(f"Triton version: {triton.__version__}")
print(f"CUDA available: {torch.cuda.is_available()}")
if not torch.cuda.is_available():
print("ERROR: CUDA is not available")
return
print(f"GPU: {torch.cuda.get_device_name(0)}")
# テスト実行
a = torch.rand(3, device="cuda")
b = a + a # PyTorch標準の加算
b_compiled = add(a, a) # Tritonカーネルでの加算
# 結果比較
diff = b_compiled - b
print(f"Difference: {diff}")
if torch.allclose(b, b_compiled):
print("SUCCESS: Triton is working correctly!")
else:
print("ERROR: Results do not match")
if __name__ == "__main__":
main()
5.4 実行結果
上記のコードを実行すると、以下のような出力が得られます。
PS C:\dev> python test_triton.py
PyTorch version: 2.9.0+cu124
Triton version: 3.5.1
CUDA available: True
GPU: NVIDIA GeForce RTX 4090
Difference: tensor([0., 0., 0.], device='cuda:0')
SUCCESS: Triton is working correctly!
5.5 よくあるエラーと対処法
| エラー | 原因 | 対処法 |
|---|---|---|
ModuleNotFoundError: No module named 'triton.language' |
テストスクリプトをtriton.pyという名前で保存した |
ファイル名をtest_triton.pyなどに変更 |
AttributeError: module 'pkgutil' has no attribute 'ImpImporter' |
setuptoolsが古い |
pip install -U pip setuptoolsを実行 |
ImportError: DLL load failed while importing libtriton |
vcredistが古い | 最新版vcredistをインストール |
PermissionError: [WinError 5] Access is denied: 'C:\\Users\\...\.triton' |
キャッシュフォルダの権限問題 | ユーザーフォルダの権限を確認、または管理者権限で実行 |
FileNotFoundError: ... too long |
Windowsのパス長制限(260文字) | レジストリで長いパスを有効化(後述) |
ImportError: DLL load failed while importing cuda_utils |
キャッシュの破損またはPythonバージョン不一致 |
.triton\cache\フォルダを削除して再実行 |
基本的な使い方をマスターしたので、次はより実践的なユースケースを見ていきましょう。
6. ユースケース別ガイド
6.1 ユースケース1: ComfyUIでの高速化
想定読者: ComfyUIでStable Diffusionの画像生成を行っている方
推奨構成: ComfyUI Portable + triton-windows + SageAttention
サンプルコード:
"""
ComfyUI用SageAttentionインストールスクリプト
使い方: python_embeded\python.exe install_sage.py
"""
import subprocess
import sys
import os
def get_python_path():
"""embeded Pythonのパスを取得"""
# スクリプトの場所からpython_embededを探す
script_dir = os.path.dirname(os.path.abspath(__file__))
candidates = [
os.path.join(script_dir, "python_embeded", "python.exe"),
os.path.join(script_dir, "..", "python_embeded", "python.exe"),
]
for path in candidates:
if os.path.exists(path):
return path
return sys.executable
def main():
python = get_python_path()
print(f"Using Python: {python}")
# triton-windowsのインストール
print("\n[1/3] Installing triton-windows...")
subprocess.run([python, "-m", "pip", "install", "-U", "triton-windows<3.6"])
# SageAttentionのインストール
print("\n[2/3] Installing SageAttention...")
subprocess.run([python, "-m", "pip", "install", "sageattention"])
# 動作確認
print("\n[3/3] Verifying installation...")
result = subprocess.run(
[python, "-c", "import triton; import sageattention; print('OK')"],
capture_output=True,
text=True
)
if result.returncode == 0:
print("SUCCESS: SageAttention is ready to use!")
print("Restart ComfyUI to enable the optimization.")
else:
print(f"ERROR: {result.stderr}")
if __name__ == "__main__":
main()
6.2 ユースケース2: torch.compileによるモデル高速化
想定読者: 自作のPyTorchモデルを高速化したい方
推奨構成: Python 3.12 + PyTorch 2.9 + triton-windows 3.5
サンプルコード:
"""
torch.compileによるモデル高速化の例
使い方: python torch_compile_example.py
必要なパッケージ: pip install torch triton-windows
"""
import torch
import torch.nn as nn
import time
class SimpleModel(nn.Module):
"""シンプルなMLPモデル"""
def __init__(self, input_dim=1024, hidden_dim=4096, output_dim=1024):
super().__init__()
self.fc1 = nn.Linear(input_dim, hidden_dim)
self.fc2 = nn.Linear(hidden_dim, hidden_dim)
self.fc3 = nn.Linear(hidden_dim, output_dim)
self.activation = nn.GELU()
def forward(self, x):
x = self.activation(self.fc1(x))
x = self.activation(self.fc2(x))
return self.fc3(x)
def benchmark(model, x, name, warmup=10, iterations=100):
"""モデルのベンチマーク"""
# ウォームアップ
for _ in range(warmup):
_ = model(x)
torch.cuda.synchronize()
# 計測
start = time.perf_counter()
for _ in range(iterations):
_ = model(x)
torch.cuda.synchronize()
elapsed = time.perf_counter() - start
print(f"{name}: {elapsed/iterations*1000:.2f} ms/iteration")
return elapsed / iterations
def main():
device = "cuda"
batch_size = 32
# モデルとデータの準備
model = SimpleModel().to(device)
x = torch.randn(batch_size, 1024, device=device)
# 通常モデルのベンチマーク
eager_time = benchmark(model, x, "Eager mode")
# torch.compileでコンパイル
# mode="reduce-overhead"はTritonを使用した最適化を行う
compiled_model = torch.compile(model, mode="reduce-overhead")
# コンパイル済みモデルのベンチマーク
compiled_time = benchmark(compiled_model, x, "Compiled mode")
# 高速化率の計算
speedup = eager_time / compiled_time
print(f"\nSpeedup: {speedup:.2f}x")
if __name__ == "__main__":
main()
6.3 ユースケース3: カスタムTritonカーネルの作成
想定読者: 独自の最適化カーネルを作りたい上級者
推奨構成: Python 3.12 + PyTorch 2.9 + triton-windows 3.5
サンプルコード:
"""
カスタムTritonカーネル: Fused Softmax
使い方: python fused_softmax.py
必要なパッケージ: pip install torch triton-windows
"""
import torch
import triton
import triton.language as tl
@triton.jit
def softmax_kernel(
output_ptr,
input_ptr,
input_row_stride,
output_row_stride,
n_cols,
BLOCK_SIZE: tl.constexpr,
):
"""Fused Softmaxカーネル"""
# 行のインデックス
row_idx = tl.program_id(0)
# 入力/出力の開始位置
row_start_ptr = input_ptr + row_idx * input_row_stride
out_start_ptr = output_ptr + row_idx * output_row_stride
# 列のオフセット
col_offsets = tl.arange(0, BLOCK_SIZE)
mask = col_offsets < n_cols
# データのロード
row = tl.load(row_start_ptr + col_offsets, mask=mask, other=-float('inf'))
# Softmax計算(数値安定性のためmax減算)
row_max = tl.max(row, axis=0)
row = row - row_max
numerator = tl.exp(row)
denominator = tl.sum(numerator, axis=0)
softmax_output = numerator / denominator
# 結果の保存
tl.store(out_start_ptr + col_offsets, softmax_output, mask=mask)
def softmax(x: torch.Tensor) -> torch.Tensor:
"""Triton版Softmax"""
n_rows, n_cols = x.shape
# BLOCK_SIZEは2のべき乗である必要がある
BLOCK_SIZE = triton.next_power_of_2(n_cols)
output = torch.empty_like(x)
# カーネル起動
softmax_kernel[(n_rows,)](
output,
x,
x.stride(0),
output.stride(0),
n_cols,
BLOCK_SIZE=BLOCK_SIZE,
)
return output
def main():
# テストデータ
x = torch.randn(128, 512, device="cuda")
# PyTorch版とTriton版の比較
torch_result = torch.softmax(x, dim=1)
triton_result = softmax(x)
# 結果の検証
if torch.allclose(torch_result, triton_result, atol=1e-5):
print("SUCCESS: Results match!")
else:
max_diff = (torch_result - triton_result).abs().max()
print(f"WARNING: Max difference = {max_diff}")
# ベンチマーク
import time
iterations = 1000
# PyTorch版
torch.cuda.synchronize()
start = time.perf_counter()
for _ in range(iterations):
_ = torch.softmax(x, dim=1)
torch.cuda.synchronize()
torch_time = time.perf_counter() - start
# Triton版
torch.cuda.synchronize()
start = time.perf_counter()
for _ in range(iterations):
_ = softmax(x)
torch.cuda.synchronize()
triton_time = time.perf_counter() - start
print(f"PyTorch: {torch_time/iterations*1000:.3f} ms")
print(f"Triton: {triton_time/iterations*1000:.3f} ms")
print(f"Speedup: {torch_time/triton_time:.2f}x")
if __name__ == "__main__":
main()
ユースケースが把握できたところで、この記事を読んだ後の学習パスを確認しましょう。
7. 学習ロードマップ
この記事を読んだ後、次のステップとして以下をおすすめします。
初級者向け(まずはここから)
- triton-windows公式README - インストールとトラブルシューティング
- Triton公式チュートリアル - ベクトル加算から始める基礎
- PyTorch torch.compileガイド - 既存コードの高速化
中級者向け(実践に進む)
- SageAttention - 実用的なAttention最適化
- Triton Puzzles - GPUプログラミングの練習問題
- xformers - Transformer最適化ライブラリ
上級者向け(さらに深く)
- Tritonソースコード - コンパイラ内部の理解
- ROCm版Triton - AMD GPU対応の仕組み
- NVIDIA CUDA Programming Guide - 低レベルGPUプログラミング
8. まとめ
この記事では、Tritonについて以下を解説しました。
-
Tritonとは: OpenAIが開発したGPUカーネルコンパイラで、PyTorchの
torch.compileやxformersの基盤技術 - Windowsでの問題: 公式TritonはLinux専用だったが、triton-windowsプロジェクトによりネイティブ動作が可能に
- セットアップ方法: 環境別(基本/ComfyUI/FramePack)の具体的な手順
- 実践的な活用: SageAttention導入、torch.compile高速化、カスタムカーネル作成
私の所感
正直なところ、「WSLを使え」という回答が蔓延する中で、woct0rdhoさんをはじめとするコミュニティの方々がWindowsサポートを地道に実装し続けてくれていることに、私は深く感謝しています。
triton-windowsプロジェクトのREADMEにある「Catgirl matters」という一文。これは「些細に見えるユーザーのニーズも大切にする」という開発哲学を表現したものだと思います。Windowsユーザーは決して少数派ではないし、全員がLinuxに移行できるわけでもない。そんな現実を見据えて、誰もが最先端のAI技術を使えるようにする。それがオープンソースの素晴らしさだと改めて感じました。
みなさんも、ぜひtriton-windowsを活用して、WindowsでのAI開発をより快適なものにしてください。
参考文献
- triton-windows GitHub - Windows版Tritonのメインリポジトリ
- Triton公式サイト - OpenAI Tritonの公式ドキュメント
- OpenAI Triton発表ブログ - Tritonの設計思想と技術的背景
- PyTorch torch.compile - PyTorchのコンパイル機能
- SageAttention - 実用的なAttention最適化ライブラリ