tilegym-cutile-autotuning

CuTile の自動チューニング

exhaustive_search API を使用し、1 回だけのチューニング／キャッシュ／直接起動のパターンで CuTile カーネルに自動チューニングを追加します。

手順

判断ツリーに従ってカーネルを分類し、探索空間を設計して、1 回だけのチューニング／キャッシュ／起動のパターンを実装し、パフォーマンスを検証します。

1. 分類 — 判断ツリーを使用して、探索する次元（占有率のみ、またはタイル全体の探索）を決定します。
2. 探索空間の設計 — references/kernel-type-templates.md から該当するテンプレートを選択します。アーキテクチャフィルターを使用して、最終コードの構成数を 30 以下に絞り込みます（方向付けられた探索のプローブは一時的にこの数を超えてもかまいません。詳しくは設計方針を参照してください）。
3. 実装 — ステップバイステップのワークフローに従って、exhaustive_search、キャッシュ、ct.launch を追加します。必要に応じて、インプレース書き込みには分割バッファーを使用します。
4. テスト — 自動チューニングを有効にした場合と DISABLE_AUTOTUNE=1 の場合の両方で、正しさをテストします。
5. 検証 — 固定された既知の最良構成と A/B ベンチマークを実施します。references/search-strategies.md を参照してください。
6. 削減 — 一度も最良にならない不要な構成を削除し、コンパイルコストを最小限にするため、アーキテクチャごとに 8 構成以下を目標にします（ステップ 10）。

タスクルーター — 必要な項目に移動

やりたいこと	移動先
新しいカーネルに自動チューニングを追加する（最も一般的）	以下のクイックリファレンス → ワークフロー：自動チューニングの追加 → references/kernel-type-templates.md（カーネルの種類に応じて選択：T1=要素単位、T2=インプレース、T3=行列積、T4=永続的、T5=FMHA、T6=FP8、T7=グループ化 GEMM、T8=可変長アテンション、T9=デュアル GEMM 融合）
デバッグ：初回実行後にデータが破損する／結果が正しくない	落とし穴 #1（インプレースカーネル）
デバッグ：自動チューニングに 5 分以上かかる	落とし穴 #2（コンパイルのタイムアウト）
デバッグ：探索空間ジェネレーターが構成を 1 つも返さない	まず落とし穴 #5 を確認します。あわせて、アーキテクチャフィルター、サイズガード、num_ctas の制約も確認してください。
既存の自動チューニング構成を最適化する	ワークフロー：既存構成の最適化

クイックリファレンス — Occupancy のみの自動チューニング（チューニングは一度だけ／キャッシュ／起動）

多くの CuTile カーネル（要素ごとの演算、リダクション、LayerNorm）で必要なのは Occupancy のチューニングだけです。次のパターンをコピーしてください。

from types import SimpleNamespace
from cuda.tile.tune import exhaustive_search
import cuda.tile as ct
import torch

def _my_autotune_configs():
    for occ in [1, 2, 4, 8]:
        yield SimpleNamespace(occupancy=occ)

# Module-level cache: tune once, launch fast forever after
_autotune_cache = {}

def my_op(x, output):
    stream = torch.cuda.current_stream()
    NUM_SM = torch.cuda.get_device_properties(x.device).multi_processor_count

    # Cache key: anything that affects optimal config (use str() for device)
    cache_key = (x.shape, x.dtype, str(x.device))

    if cache_key not in _autotune_cache:
        configs = list(_my_autotune_configs())
        result = exhaustive_search(
            configs,
            stream,
            grid_fn=lambda cfg: (min(NUM_SM * cfg.occupancy, M), 1, 1),
            kernel=my_kernel,
            args_fn=lambda cfg: (x, output, ...),
            hints_fn=lambda cfg: {"occupancy": cfg.occupancy},
        )
        best_cfg = result.best.config
        tuned_kernel = my_kernel.replace_hints(occupancy=best_cfg.occupancy)
        _autotune_cache[cache_key] = (best_cfg, tuned_kernel)  # cache BOTH

    cfg, tuned_kernel = _autotune_cache[cache_key]
    grid = (min(NUM_SM * cfg.occupancy, M), 1, 1)
    ct.launch(stream, grid, tuned_kernel, (x, output, ...))

重要なルール：

• 一度だけチューニングし、キャッシュして、直接起動する — exhaustive_search は形状ごとに初回の呼び出しでのみ実行されます。それ以降の呼び出しでは、設定がキャッシュされ、オーバーヘッドなしで ct.launch が実行されます
• インプレースカーネルでは、探索中に分割バッファを使用します（入力テンソルと出力テンソルを分けます）
• 最終コードの設定は 30 個以下にしてください（一時的な指向性プローブについては「設計思想」を参照）
• exhaustive_search には Sequence（リスト／タプル）が必要です。ジェネレーターは list() で変換してください
• 探索空間には元の固定設定を含めてください — これにより、自動チューニングによってパフォーマンスが低下することを防げます

このパターンを使う場合：カーネルのブロックサイズが固定で、タイルサイズを調整できない場合です。要素ごとの演算（SwiGLU、GeGLU）、リダクション（RMSNorm、LayerNorm）、RoPE、ヒューリスティックなブロックサイズを使う永続カーネル（grouped GEMM）などが該当します。

より複雑なカーネル（タイルサイズを使う行列積、FMHA、num_ctas を使う FP8）については、以下の完全なガイドと kernel-type-templates.md を参照してください。

⚠️ ほとんどの人が陥る 3 つの落とし穴 — 提出前に確認してください：

• ホットパスで replace_hints を使っていませんか？ → exhaustive_search から設定とカーネルオブジェクトの両方をキャッシュしてください。呼び出しのたびに replace_hints() を実行すると再コンパイルが発生し、100～500 倍遅くなります → 落とし穴 #7
• インプレースカーネル（入力テンソルに書き戻すカーネル）ですか？ → 探索中に分割バッファパターンを必ず使用してください → 落とし穴 #1
• 探索空間が空ですか？ → アーキテクチャフィルターと num_ctas の制約を確認してください → 落とし穴 #5

最低限のカバレッジ：sm100 以降では、FMHA／行列積／可変長の探索空間に num_ctas=1 と num_ctas=2 の両方を含める必要があります。主要な次元（タイルサイズ、Occupancy）については、どちらが優れているか確信がなくても、異なる値を少なくとも 2 つ含めてください。exhaustive_search に判断させましょう。

チューニングを終了するタイミング：平均高速化率が [0.98, 1.02] の範囲にある場合、現在の探索空間が効果を発揮していないことを意味します。ただし、効果のある設定が存在しないという意味ではありません。終了する前に、このカーネルタイプの主要な次元を網羅できているか確認してください（references/kernel-type-templates.md を参照）。探索空間がすでにテンプレート推奨の次元を網羅していて、最良の結果も依然としてノイズの範囲内であれば、そこで終了してください。さらに細かな調整をしても効果はありません。主要な次元が抜けている場合（たとえば、2 つの GEMM を使うカーネルで num_ctas=2 を試していない場合）は、諦めるのではなく探索空間を広げてください。

正当性テストに合格し、自動チューニングしたカーネルが固定設定のベースラインより高速になったら、そこで終了してください。「確認」のために再実行しないでください。 GPU カーネルの計測時間は、クロックのスケーリングや OS のスケジューリングにより、呼び出し間で ±5～10％変動します。その後の計測値が低下しても、コードが誤っているわけではありません。

高速化率を改善するには、自動チューニングの探索空間（設定、タイルサイズ、Occupancy、num_ctas）のみを変更してください。高速化を追求するために、それ以外のコード（Python ラッパー、ストリーム管理など）を変更しないでください。カーネルのパフォーマンスはホスト側のコードではなく、設定の選択によって決まります。

読み方ガイド

• 占有率のみを調整するカーネル（要素単位、リダクション、ブロックサイズ固定の永続カーネル）：クイックリファレンスと落とし穴チェックリストで十分です。references/ のドキュメントは読み飛ばしてください。インプレースカーネルでは、落とし穴 #1 も読んでください。
• 複雑なカーネル（タイルサイズを調整可能な matmul、FMHA、num_ctas を使う FP8）：クイックリファレンス → 決定木 → API リファレンス → ステップごとのワークフロー → 関連する references/ ドキュメントの順に読みます。

5 ステップの概要：カーネルを分類 → 探索空間を設計（parameter-space-design.md）→ テンプレートを使って実装（kernel-type-templates.md）→ A/B テストで検証 → 落とし穴チェックリストを確認。

リファレンスの読み方：カーネルの種類に関連するリファレンスだけを読みます。たとえば FMHA では、references/kernel-type-templates.md のテンプレート 5 のセクションを読みます。ハードウェア制約については、対象アーキテクチャのセクションだけを読んでください。必要な箇所を調べれば十分な場合は、すべてのリファレンスを最初から最後まで読むのは避けてください。

設計方針

大きな探索空間を削っていくのではなく、小さく正確な探索空間をボトムアップで構築します。 CuTile のコンパイルは Triton よりはるかに重いため（構成ごとに約 0.5～1 秒）、最終コードに含める構成は 30 以下にしてください。まずカーネルの種類を分類し、その種類とアーキテクチャに関連する構成だけを作成します。

開発中の方向性を定めた探索：初期テンプレート構成で得られる高速化が 1.0 未満の場合は、どの次元が重要かを特定するため、bash + python3 -c で 一時的に より大規模なプローブ（30～100 構成）を実行してもかまいません。ただし、このプローブは方向性を定めたものにし、無差別な直積にしてはいけません。カーネルの種類の分類に基づき、どの次元を変化させるかを決めます（たとえば dual-GEMM ではタイルサイズを固定して num_ctas × occupancy を調べ、FMHA では TILE_N を固定して TILE_M × num_ctas を調べます）。プローブで有望な領域が特定できたら、最終コードの探索空間は有力候補の上位 8 個以下に絞り込みます。大規模なプローブをソースファイルに書き込まないでください。これは一度だけ使う診断ツールです。

決定木：このカーネルで必要な探索次元は？

すべてのカーネルに自動チューニングを追加してください。問題は自動チューニングを行うかどうかではなく、どの次元を探索するかです。

What type of kernel is this?
├── Compute-bound (matmul, GEMM, FMHA) → Does it have multiple tunable dimensions (tile sizes)?
│   ├── YES → Is it a fused multi-GEMM kernel (dual-GEMM, e.g. Linear+GLUAct)?
│   │   ├── YES → Template 9: low occupancy (1–2), conservative tiles (2× SHMEM/register pressure)
│   │   └── NO  → Full search: TILE_M × TILE_N × (TILE_K) × occupancy × num_ctas
│   │             (see matmul/FMHA templates in kernel-type-templates.md)
│   └── NO  → Occupancy-only search: [1, 2, 4, 8]
│             (see Quick Reference above)
├── Balanced (LayerNorm, reduction + compute) →
│   Occupancy-only search: [1, 2, 4, 8]
│   Expected benefit: 2-15%
└── Memory-bound (CE Loss, pure elementwise) →
    Occupancy-only search: [1, 2, 4, 8]
    Expected benefit: 0-15% (varies by kernel; zero-cost after tuning)

メモリ律速カーネルが num_ctas やタイルサイズではなく、占有率のみを探索する理由：

• num_ctas による利点はありません：num_ctas > 1 にすると TMA マルチキャストが有効になり、複数の CTA が共有メモリ内のタイルデータを共有します（たとえば、CTA 間で再利用される matmul の A/B タイル）。メモリ律速カーネルでは、タイルを再利用しない要素ごとの ct.gather／ct.scatter を使うため、複数 CTA の協調にはデータ共有の利点がなく、オーバーヘッドが増えるだけです。
• タイルサイズは事前に決まっています：メモリ律速カーネルの BLOCK_SIZE はオフラインスイープで決定されます（たとえば B200 では [256, 512, 1024, 2048, 4096, 8192] の中で 1024 が全体を通して最適です）。これは定数であり、実行時に調整するものではありません。
• 有効な調整項目は占有率のみです：占有率が高いと、メモリー要求で停止している CTA の代わりに GPU が別の CTA に切り替わり、メモリーのレイテンシーを隠せます。

根拠 — CE Loss の実験：Cross-Entropy Loss で 12 構成の探索（occupancy × num_ctas）を行ったところ、改善はわずか 2.5% でした（Triton 比で 0.79x → 0.81x）。num_ctas の次元は何の効果もありませんでした。コンパイルコストがわずかな利点を上回ったため、この結果は取り消されました。occupancy のみ（4 構成）でも、コンパイル時間を 3 分の 1 に抑えながら同じ結果が得られます。

メモリ律速カーネルに関する注記：occupancy のみの自動チューニングを追加する価値は常にあります。その理由は次のとおりです。

• 一度チューニングしてキャッシュし、起動するパターンでは、初回の呼び出し後に実行時オーバーヘッドが発生しません。
• 探索空間が非常に小さい（4 構成、コンパイルは約 2～4 秒）です。
• わずかな改善でも、大規模に適用すれば価値があります。

占有率の選択ガイド

占有率は、SM ごとに同時実行する CTA の数を制御します。占有率の探索範囲を設計する際の出発点として、以下を参考にしてください。

占有率の範囲	最適な用途	カーネルの例
1–4	計算律速（複雑な演算）	複雑な変換、行列積
4–8	バランス型（GEMM、TMA）	行列積、FMHA
8–16	メモリ律速（リダクション）	Softmax、LayerNorm
16–32	非常に軽量な処理（コピー、キャスト）	型変換、要素単位の処理

これらの範囲を初期探索空間の候補として使用してください。占有率のみを対象とするカーネルでは、[1, 2, 4, 8] でほとんどのケースをカバーできます。上記のクイックリファレンスを参照してください。

exhaustive_search API リファレンス

exhaustive_search API の全容については、references/api-reference.md を参照してください。現在のシグネチャ、TuningResult、一度だけチューニングしてキャッシュし起動するパターン、replace_hints、カーネルヒント、search_space の設計、grid_fn のパターンを説明しています。

ステップごとのワークフロー

エンドツーエンドのワークフローについては、references/workflow.md を参照してください。新しいカーネルへの自動チューニングの追加、既存のマルチアーキテクチャ設定の処理、torch.autograd.Function との統合、バックエンド間での設定の移行（Triton → CuTile）、既存設定の最適化を説明しています。

落とし穴のチェックリスト

よくある落とし穴の全一覧については、references/pitfalls.md を参照してください。インプレース操作によるデータ破損、コンパイルのタイムアウト、コールドキャッシュによる性能の偏り、NCU プロファイリングによる干渉、search_space ジェネレーターの枯渇、FP8 の精度低下、ホットパスでの replace_hints による再コンパイルを取り上げています。

対象範囲と境界

このスキルの対象は、探索空間の設計、exhaustive_search の呼び出し、キャッシュ、チューニング済みヒントを使用した ct.launch など、自動チューニングの設定のみです。カーネルコードは変更しません。

対象範囲内（自動チューニング設定）：

• 探索空間のジェネレーター関数
• exhaustive_search() の呼び出しと結果の処理
• チューニング済みヒントを適用する kernel.replace_hints()
• キャッシュのロジック（キーの設計、辞書の管理）
• チューニング済みカーネルを使用する ct.launch()
• DISABLE_AUTOTUNE のフォールバックパス

対象範囲外（カーネルコードの変更 — これらの変更は行わないでください）：

• 数学演算フラグ（flush_to_zero、rounding_mode）
• パフォーマンスヒント（slice_hint、buffer_depth、copy_config）
• メモリアクセスパターン（2D→1D の gather/scatter 変換）
• コード生成の最適化（safe_offs → padding_value）
• アルゴリズムの変更（K ループの分割、負荷分散）

さらなる最適化の提案

自動チューニングを追加した後、以下のカーネルレベルの最適化によって、さらに性能が向上する可能性があります。これらはこのスキルの対象範囲外です。次のステップの候補としてユーザーに伝えてください。ただし、自動チューニングの一環として実装しないでください。

• 数学演算フラグ：flush_to_zero=True と rounding_mode=APPROX を設定すると、FMHA クラスのカーネルで 34～72% の性能向上が得られる場合があります（環境変数 TILEIR_ENABLE_FTZ=1 TILEIR_ENABLE_APPROX=1 を設定するか、カーネルコード内で設定します）。因果関係：大きなタイルは、非正規化数の処理オーバーヘッドによって当初は性能を 18～43% 低下させますが、FTZ+APPROX を有効にするとこの低下を解消でき、結果は +34～72% の向上に転じます。そのため、FMHA クラスのカーネルで大きなタイル設定を効果的にするには、数学演算フラグが前提条件となります。
• パフォーマンスヒント：slice_hint、buffer_depth、copy_config — カーネルの IR コードの変更が必要です
• メモリアクセスパターン：ct.gather の代わりに TMA ロード（ct.load）を使用する、不要な境界チェックを削除する（安全な場合は check_bounds=False を指定）
• コード生成の品質：手動の ct.where マスキングの代わりに padding_value パラメーターを使用する、safe_offs を削除する
• アルゴリズムの再構成：K ループの分割、負荷分散、代数的単純化

Triton Autotune との違い

主な違いは次のとおりです。Triton では @triton.autotune デコレーターと Config(...) オブジェクトを使用しますが、CuTile では exhaustive_search() と SimpleNamespace の設定、別個のキャッシュ、ct.launch を使用します。CuTile には num_warps／num_stages がなく（コンパイラーが決定します）、タイルサイズ、occupancy、num_ctas のみを指定します。CuTile のコンパイルは負荷が大きいため、最終コードでは設定を30個以下にしてください。CuTile のキャッシュはユーザーがメモリー上で管理するもので、自動的には永続化されません。CuTile では args_fn（カーネル引数）と hints_fn（コンパイラーヒント）が分かれています。

参考ドキュメント

カテゴリー	ドキュメント	内容
API リファレンス	api-reference.md	exhaustive_search のシグネチャ、TuningResult、一度だけのチューニング／キャッシュ／起動のパターン、replace_hints、カーネルヒント、search_space の設計、grid_fn のパターン
ワークフロー	workflow.md	エンドツーエンドのワークフロー：新しいカーネルへの自動チューニングの追加、複数アーキテクチャ向けの設定、torch.autograd.Function との統合、Triton から CuTile への移行、既存設定の最適化
落とし穴	pitfalls.md	よくある落とし穴：インプレース破損、コンパイルタイムアウト、コールドキャッシュによる偏り、NCU の干渉、search_space の枯渇、FP8 の精度、replace_hints による再コンパイル
パラメーター設計	parameter-space-design.md	カーネルタイプごとのパラメーター空間、アーキテクチャ間のパターン、grid_fn のパターン、プルーニングルール
探索戦略	search-strategies.md	網羅的探索、A/B テストの方法、DISABLE_AUTOTUNE のパターン
テンプレート	kernel-type-templates.md	8種類のカーネルタイプ向けのコピー＆ペースト可能な自動チューニングテンプレート
ハードウェア	hardware-constraints.md	アーキテクチャごとの制約、タイルサイズの範囲、num_ctas のルール、TMA の要件

ソースコードの参照先

主なファイル：ops/cutile/matmul.py（matmul の自動チューニング）、ops/cutile/attention.py（FMHA の自動チューニング）、suites/unsloth/cutile/ct_ops.py（共有の autotune_configs()、occupancy=[1,2,4,8]）、suites/unsloth/cutile/swiglu.py（要素単位の例）、suites/unsloth/cutile/rope_embedding.py（分割バッファーのパターン）、suites/unsloth/cutile/grouped_gemm.py（永続 GEMM、occupancy のみ）。

実例

各例では変更前 → 変更後のパターンを示します。fixed_launch.py（ハードコードされた ct.launch）と autotuned_launch.py（一度だけのチューニング／キャッシュ／起動にリファクタリング済み）を比較します。

ディレクトリ	カーネル	自動チューニングのパターン	複雑さ	主なポイント
assets/examples/01_rmsnorm_occupancy_only/	RMSNorm（リダクション）	occupancy のみ [1,2,4,8]	低	最も一般的なパターンです。タイルのチューニングは行わず、最適な occupancy を見つけます。Grid = NUM_SM * cfg.occupancy。インプレースではありません。
assets/examples/02_matmul_full_search/	GEMM C=A@B	全項目：TILE_M/N/K + occupancy + num_ctas（sm90+）	高	調整可能な次元が複数ある、演算負荷の高いカーネルです。args_fn はタイルサイズを ct.Constant[int] として渡します。grid_fn は cfg に依存します。設定は30個以下にしてください。
assets/examples/03_rope_inplace_splitbuffer/	RoPE エンベディング（インプレース）	分割バッファーを使用した occupancy のみ	中	破損を防ぐため、インプレースカーネルでは探索中に必ず分割バッファーを使用してください。探索ではスクラッチ領域に書き込み、最終的な ct.launch では実際のインプレース引数を使用します。