tilegym-converting-cutile-to-triton

cuTile → Triton の変換

@ct.kernel カーネルを @triton.jit に変換します。API マッピング：references/api-mapping.md（cuTile → Triton）。

このスキルの Markdown では、Triton の起動構文 kernel［grid］(…) に Unicode の角括弧を使用しています。リンクチェッカーが [grid](…) をハイパーリンクとして解析しないようにするためです。実際の Triton コードでは通常の ASCII 角括弧を使用してください。

手順

translations/workflow.md のフェーズゲート方式のワークフローに従ってください。各変換では、分析 → 変換 → 検証 → テスト → ベンチマークの順に進め、次のフェーズに移る前に明示的なゲートを通過する必要があります。タスクが特殊なケース（エラー、レイアウトフラグ、パフォーマンス）に該当する場合は、ワークフローの選択にあるドキュメントを使用してください。

0. 最適化戦略（パフォーマンス重視／attention） — 対象の演算がattention、FMHA、sliding window、soft cap、または GQA（例：Gemma の gemma_attention）の場合は、内部ループを変換する前に references/optimization-strategy.md を読み、その後 §4 Gemma FMHA チェックリスト を適用してください。その他の GEMM／BMM／attention 関連カーネルについても、TMA の実装後にこのファイルの §2～§3 に目を通してください。

1. 変換パスを選択 — 既存の TileGym 演算の場合は、translations/workflow.md の標準モードを使用してください。cuTile ソースで transpose／transpose_v、デュアルレイアウト、または MLA 形式のパスを使用している場合は、Triton を記述する前に translations/advanced-patterns.md を読んでください（1 つのカーネルと tl.trans ではなく、2 つのカーネルと META グリッドを使用します）。

2. 事前確認 — cuTile ソースに対して事前確認分析の grep コマンドを実行してください。@ct.kernel の定義数を数え、TMA に関連する ct.load／ct.store、ct.launch、Constant、レイアウトフラグを記録してください。

3. マッピングを確認 — cuTile → Triton の API 対応を確認できるよう、references/api-mapping.md を開いたままにしてください。実行時エラー（不正なアドレス、dtype、ストライド）については、references/debugging.md を参照してください。

4. 変換 — 変換チェックリストを todo リストにコピーし、順番に実行してください。構成とファイルの配置については、translations/file-structure.md を参照してください。必須： 2D 以上のブロック形状を持つタイルのロード／ストアには、TMA の tl.make_tensor_descriptor を使用してください。フルタイルに対して tl.load(ptr+offs, mask=…) を直接使ってはいけません。これを省略することが、大幅な性能低下の最も一般的な原因です。ホスト側では、タプルまたは自動チューニング用の lambda META: (…) を使って Triton の角括弧による起動 kernel［grid］(args) を記述してください。ct.launch は使いません。

5. 検証 — 新しい Triton モジュールの構文を確認し、対象の演算に対応する TileGym の pytest ターゲットを実行してください：pytest tests/ops/test_<op>.py -k "triton" -vs。ベンチマークの前に失敗を修正してください。

6. ベンチマーク — パフォーマンステストで Triton と cuTile を比較してください。Triton が明らかに遅い場合は、GEMM／BMM／attention について translations/workflow.md の パフォーマンス分析（フェーズ c2t-5） と references/optimizing-reference.md に従ってください。順序立てたチェックリストとして references/optimization-strategy.md を使用してください。10～50 倍の性能低下が見られる場合は、まず同じワークフローファイルの重大なパフォーマンスパターンを読んでください。

実行ルール（必須）：

• カーネルコードを編集する前に変換チェックリストを作成して進捗を記録してください（例：TodoWrite）。事前確認や TMA の判断を省略せず、手順を順番に完了してください。
• attention／FMHA／Gemma／GQA／soft cap／sliding windowの場合は、変換を最適化済みと見なす前に references/optimization-strategy.md を読み、§4を適用してください。
• TMA が適用される 2D 以上のタイルロードで、生のポインターとマスクを使う方法を採用しないでください。意図的に例外を設ける場合は、その理由を記載してください。
• テストまたはベンチマークのゲートに失敗した場合は、変換完了を宣言せず、いったん止まって修正してください。未検証の変更を重ねてはいけません。

ワークフローの選択

• 既存の TileGym op → 標準モード：translations/workflow.md
• エラー（cudaErrorIllegalAddress、形状の不一致、数値の不一致）→ references/debugging.md
• 高度なパターン（TMA、デュアルレイアウトフラグ transpose、オートチューニング + META グリッド、Array.slice、ct.gather().item()）→ translations/advanced-patterns.md（MLA スタイルの 2 カーネル、transpose=False での 3～15 倍の性能低下を回避）。
• パフォーマンス（Triton カーネルが cuTile より遅い、オートチューニング、プロファイリング）→ translations/workflow.md（**パフォーマンス分析（フェーズ c2t-5）**のセクション）
• 最適化戦略のハブ（順序付きチェックリスト：advanced-patterns + optimizing-reference）→ references/optimization-strategy.md — attention/FMHA/Gemma では最初に読み、その後必要に応じて 2 つのソースドキュメントを参照してください
• GEMM/BMM/attention の最適化（TMA の導入後、または Triton が 10～20％遅い場合）→ references/optimizing-reference.md — EVEN_K 高速パス、ポインター演算による transpose、グリッドレイアウト、オートチューニングの探索範囲、エピローグのサブタイル化。変換時およびパフォーマンス承認前にこれらのパターンを使用してください（optimization-strategy §2～§3に要約されています）
• Gemma attention / GQA FMHA の変換 → references/optimization-strategy.md §4
• Blackwell の最適化（反復アルゴリズム、レジスター負荷、ループ展開を伴う複雑なカーネル）→ references/optimizing-reference.md §9 — TMA ディスクリプター、loop_unroll_factor、占有率のオートチューニング、TMEM に適したブロックサイズ、スラブアロケーター、デュアルパスのカーネル設計
• ⚠️ 10～50 倍の性能低下（変換後に壊滅的な速度低下が発生）→ translations/workflow.md — **重大なパフォーマンスパターン（10～50 倍の性能低下を回避）**のセクション
• ⚠️ transpose=True では良好なパフォーマンスだが、transpose=False では大幅に低下する（またはその逆）→ translations/advanced-patterns.md — §1 デュアルレイアウトフラグ：2 つの @triton.jit カーネル + grid = lambda META: (... META["BLOCK_H"] ...)

事前分析（変換前に実行）

# Count kernels (only main kernel gets @triton.jit, helpers stay plain def)
grep "@ct\.kernel" source.py | wc -l

# Check for patterns needing special handling
grep "ct\.transpose\|ct\.permute" source.py   # → use tl.trans/tl.permute
grep "ct\.astype" source.py                    # → use .to(dtype)
grep "ct\.load\|ct\.store" source.py          # → TMA for 2D+ (tl.make_tensor_descriptor), NOT raw tl.load(ptr+offs)
grep "ct\.launch" source.py                    # → bracket launch: kernel then [grid] then (args)
grep "ct\.Constant\|ct\.ConstInt" source.py    # → tl.constexpr
grep "ct\.cdiv" source.py                      # → triton.cdiv (host) or Python (a+b-1)//b
grep "ct\.bid\|ct\.num_blocks" source.py       # → tl.program_id/tl.num_programs
grep "1 << .*\.bit_length" source.py           # → triton.next_power_of_2 if needed
grep "transpose\|transpose_v" source.py       # → if hit, read translations/advanced-patterns.md (dual kernels + META grid)

変換チェックリスト

このチェックリストをコピーして、進捗を記録してください。

Conversion Progress:
 [ ] Step 0 (attention / Gemma FMHA / GQA / soft cap / sliding window): Read [references/optimization-strategy.md](./references/optimization-strategy.md) and apply §4 checklist before inner-loop Triton
 [ ] Step 1: Pre-flight — run grep commands above, note special patterns and 2D+ loads (→ TMA)
 [ ] Step 2: Analyze source cuTile kernel (identify patterns, shapes, dtypes)
 [ ] Step 3: Create Triton file with correct structure (see translations/file-structure.md)
 [ ] Step 4: Convert kernel signature (tensor args → pointer args, Constant → constexpr)
 [ ] Step 4b: TMA (MANDATORY for 2D+ loads) — use tl.make_tensor_descriptor for every 2D+ tile load/store; do NOT ship raw tl.load(ptr+offs,mask) for block-shaped access (see workflow.md § TMA OPTIMIZATION)
 [ ] Step 5: Convert kernel body (apply gotchas table below + API mapping)
 [ ] Step 6: Convert host wrapper (grid tuple/lambda, bracket-style launch: kernel, grid, then arguments; no ct.launch); call triton.set_allocator(alloc_fn) if using TMA
 [ ] Step 7: Validate — run pytest or syntax check on Triton file
 [ ] Step 8: Test — run pytest, verify X passed 0 failed
 [ ] Step 9: If test fails → fix → re-validate → re-test (loop until green)
 [ ] Step 10: Benchmark — run perf test, compare vs cuTile (see workflow.md § PERFORMANCE ANALYSIS)
 [ ] Step 10b: If GEMM/BMM/attention and Triton >20% slower → walk [references/optimization-strategy.md](./references/optimization-strategy.md) §2–§3 then [references/optimizing-reference.md](./references/optimizing-reference.md) (EVEN_K, transpose, grid, autotune, epilogue subtile), then re-benchmark
 [ ] Step 10c: If op has `transpose` / layout flag → read [translations/advanced-patterns.md](./translations/advanced-patterns.md); verify **separate kernels** per layout (not transpose-kernel + `tl.trans`); **autotuned** launches use `lambda META: (triton.cdiv(..., META["BLOCK_H"]), ...)` — no fixed `BLOCK_H`/`BLOCK_N` through `apply()` unless autotune is disabled

Post-conversion Verification (TMA is mandatory for 2D+ loads):
 [ ] TMA: All 2D+ tile loads use tl.make_tensor_descriptor(...).load([...]); no raw ptr+mask for block-shaped 2D+ access (else 5x-20x regression)
 [ ] Grid uses tuple or lambda (not 3-tuple required like cuTile)
 [ ] Triton autotune added if cuTile op used kernel_configs/autotune (see workflow § PERFORMANCE ANALYSIS)
 [ ] Host grid uses triton.cdiv where appropriate (not (a+b-1)//b only)
 [ ] Pointer/offset indexing: Triton uses element offsets (ptr + offs), not block index in tl.load (or use TMA descriptor)
 [ ] ct.astype(x, dtype) → x.to(dtype) in Triton
 [ ] ct.mma(a, b, acc=acc) → tl.dot(a, b, acc) (no keyword in Triton)
 [ ] Optional/None args: Triton allows None in kernel args if desired (cuTile required dummy+flag)
 [ ] Masking applied when BLOCK_SIZE > actual dimension (same as cuTile); with TMA, masks can often be removed for full tiles
 [ ] Reduction divisor uses actual_size, NOT BLOCK_SIZE
 [ ] fp32/tf32: Triton defaults allow_tf32=True; match cuTile behavior if you had explicit tf32 cast
 [ ] If any 2D+ load uses raw ptr+mask (exception only): document WHY TMA was not used
 [ ] tl.assume() alignment hints added for strides and pointers

注意点（よくある変換エラー） {#gotchas-most-common-translation-errors}

@ct.kernel を @triton.jit に移植する際に、頻繁に問題や性能低下を引き起こすパターンを網羅した表です — mma アキュムレーター、型キャスト、グリッド、TMA の使用、dtype の処理、レイアウトフラグ、バッチ行列積など。

参照： references/gotchas.md — Triton カーネルを書く前に読んでください。

パフォーマンス上の注意点（10～50 倍の性能低下のリスク） {#performance-gotchas-10-50x-regression-risk}

⚠️ これらは壊滅的な速度低下を引き起こします。ベンチマークの前に確認してください。

パターンとその影響：TMA と raw ptr+mask の比較（5～20 倍）、オートチューニングと固定タイルサイズの比較（2～3 倍）、broadcast_to + tl.dot（10～50 倍）、extract_slice の連鎖（2～5 倍）など。

参照： references/performance-gotchas.md — 性能低下のリスクをまとめた表。

詳細： translations/workflow.md — **重大なパフォーマンスパターン（10～50 倍の性能低下を回避）**のセクション。

API マッピング全体：references/api-mapping.md。

Triton の数学関数の dtype（erf/erfc/exp/log/sqrt）と「erf を tanh で代用しない」パターン：references/debugging.md — **Triton 数学関数の dtype 要件（重要）**のセクション。

最適化戦略（ハブ）

ファイル： references/optimization-strategy.md

translations/advanced-patterns.md（レイアウトフラグ、デュアルカーネル、autotune+META、バッチ起動、Blackwell ポインター）と references/optimizing-reference.md（TMA 後のマイクロ最適化、§9）を §1～§3 と 必須の §4 Gemma FMHA チェックリストにまとめています。

ルール： attention / FMHA / Gemma 形式の変換では、workflow と同じセッションで optimization-strategy を開いてください。パフォーマンスの承認に TMA だけを頼らないでください。

参照ドキュメント {#reference-documents}

cuTile → Triton の観点から読みます。主要ファイルはこのスキル内の `` にあります。

カテゴリ	ドキュメント	内容
戦略	optimization-strategy.md	順序付きハブ： advanced-patterns + optimizing-reference、§4 Gemma FMHA 必須チェックリスト
ワークフロー	translations/workflow.md	標準的な c2t 変換（フェーズ + チェックリスト）
	translations/file-structure.md	cuTile から変換する際に Triton ファイルを配置する場所
	translations/advanced-patterns.md	デュアルレイアウトフラグ（転置）、autotune + META グリッド、MLA 形式の 2 つのカーネル
API	api-mapping.md	cuTile → Triton のマッピング
	optimizing-reference.md	GEMM/BMM/attention の最適化（EVEN_K、転置、グリッド、autotune、エピローグのサブタイル）
注意事項	gotchas.md	cuTile→Triton 翻訳時によくあるエラー（mma、dtype、グリッド、TMA、レイアウトフラグ）
	performance-gotchas.md	10～50 倍の性能低下リスク表（TMA と ptr+mask の比較、broadcast_to、extract_slice の連鎖、autotune）
テストとエラー	references/debugging.md	Triton のランタイムエラー（cudaErrorIllegalAddress、ポインター型、stride のオーバーフロー）

実例

cutile_kernel.py をソース、triton_kernel.py をターゲットとして使用します。

例	ディレクトリ	複雑さ
ベクトル加算	examples/01_vector_add/	基本
Softmax	examples/02_softmax/	中級
LayerNorm	examples/03_layernorm/	中級
MatMul	examples/04_matmul/	上級
Attention	examples/05_attention/	上級

逆方向のマッピングを確認するには、まず cutile_kernel.py を読み、次に triton_kernel.py を読みます。

⚠️ 必須の完了チェックリスト（省略しないでください）

すべての項目にチェックが付くまで、変換は完了していません。コピーして記入してください：

MANDATORY COMPLETION GATES:
 [ ] 1. CORRECTNESS: pytest passes with 0 failures
     Command: python -m pytest {test_path} -k "test_op and triton" -vs --tb=short
     Gate: "X passed, 0 failed"

 [ ] 2. TMA OPTIMIZATION: All 2D+ tile loads use tl.make_tensor_descriptor
     Verify: grep -n "tl.load.*mask" triton_file.py | wc -l  # Should be 0 for 2D+ ops
     Skip = 5-20x performance regression

 [ ] 3. PERFORMANCE TEST: Triton within 20% of cuTile baseline
     Command: python -m pytest {test_path} -k "test_perf" --print-record -v
     OR: Run benchmark script: cd tests/benchmark && python bench_{op}.py
     Gate: Triton TFLOPS >= 0.8 * CuTile TFLOPS

 [ ] 4. PERFORMANCE COMPARISON RECORDED:
     Document results:
     | Config | Triton (TFLOPS) | CuTile (TFLOPS) | Ratio |
     |--------|-----------------|-----------------|-------|
     | [fill] | [fill]          | [fill]          | [fill]|

CONVERSION COMPLETE: All 4 gates passed? → YES / NO

これが重要な理由：

• ゲート 1 は機能上のバグを検出します
• ゲート 2 は壊滅的な 5～20 倍の性能低下を防ぎます（最もよくあるミス）
• ゲート 3 は最適化が効果を発揮したことを検証します
• ゲート 4 は説明責任の記録を作成します

いずれかのゲートに失敗した場合： 完了と宣言する前に修正し、再検証してください。