mcpField
D
design-cute-dsl-kernel

design-cute-dsl-kernel

github.com/pepperu96開発084 回インストール32 週間前に更新
リポジトリ

design-cute-dsl-kernel について

CuTe Python DSL を使う GPU カーネル開発者向けの設計スキルです。cuTile から移行する条件を判断し、カーネル構造、自動チューニング、アーキテクチャの適合性確認を案内します。

主な機能

  1. 01CuTe DSL へ移行する条件の評価
  2. 02ホストとデバイスのカーネル構造の案内
  3. 03CuteKernel の自動チューニングパターン
  4. 04GPU アーキテクチャとリソース制限の確認
  5. 05同期とメモリ使用上の注意点

活用シーン

  1. 01warpgroup 制御を要する最適化方式の選定
  2. 02明示的な TMA パイプラインを持つカーネルの設計
  3. 03Blackwell GPU 間の命令互換性の確認

インストール

コマンドをコピーしてターミナルで実行します。

$npx skills add https://github.com/pepperu96/hyper-mla --skill design-cute-dsl-kernel

インストール後、Claude Code で/design-cute-dsl-kernelを入力して読み込まれたか確認します。

詳細

CuTe Python DSL カーネル設計

共通の命名、バージョン管理、ワークフローについては、必ず /design-kernel も読み込んでください。 API リファレンス、実行モデル、アーキテクチャ操作については、/cute-dsl-ref も読み込んでください。

CuTe Python DSL を使用する場面

cuTile の公開制御機能では不十分になったものの、Python でカーネルを記述するワークフローが適切な場合は、CuTe Python DSL（cute-dsl）を使用してください。

適合性の判断

次の最適化に、cuTile では公開されていない以下の制御機能が必要な場合は、CuTe DSL が適しています。

  • スレッド／ワープ／ワープグループの識別 — どのスレッドが何を行うかを明示的に制御します
  • CTA 内の同期 — バリア、名前付きバリア、到着／待機パターン
  • ワープグループのスケジューリング — プロデューサー／コンシューマーのワープグループの役割、永続的なワープグループループ
  • TMA パイプライン制御 — バリア同期を伴う明示的な多段階非同期コピーパイプライン
  • クラスターのプログラミング — CTA 間の共有メモリアクセス、分散共有メモリ
  • レジスタレベルのデータ移動 — レジスタ間シャッフル、ワープレベルプリミティブの明示的な使用
  • カスタムエピローグ — 細かな制御を伴う融合後処理
引き続き cuTile を使用する場面

タイルサイズ、CTA の再マッピング、occupancy／num_ctas ヒント、レイテンシヒント、allow_tma フラグで最適化を表現できる場合は、cuTile を使い続けてください。CuTe DSL は複雑さを増すため、その複雑さが不可欠な場合に限って使用してください。

設計上の厳守事項

プロファイリングで 1 CTA/SM、対象となるワープ数の少なさが確認され、解決に明示的なワープグループまたはバリアのスケジューリングが必要な場合、cuTile では性能差を埋められません。これが CuTe DSL に切り替える典型的なきっかけです。

命名とレイアウト

  • 公開言語キー：cute-dsl
  • Python パッケージパス：cute_python
  • カーネルレイアウト：src/mla_var3/kernel/cute_python/<layer>/<design>/<design>[_vN]/
  • モジュール：<design>[_vN].py
  • ラッパー：CuteKernel

CuteKernel のランタイムパターン

ランタイムラッパーは src/mla_var3/runtime/cute_kernel.py にあります。

CuTe DSL カーネルは、ホスト／デバイスの 2 層パターンを使用します。

  1. @cute.jit ホスト関数が TMA ディスクリプターを設定し、グリッドを計算して、カーネルを起動します
  2. @cute.kernel デバイス関数に GPU コードを記述します

CuteKernel データクラスは、このパターンをラップします。

from mla_var3.runtime.cute_kernel import CuteKernel

# In KernelPlan.plan():
def plan(self, *inputs):
    # Build a closure that captures inputs and calls the host function
    def launch_fn(tiling):
        # Convert tensors, set up TMA descriptors, compute grid
        host_fn(tiling, *converted_inputs)

    return CuteKernel(
        kernel_fn=device_kernel,       # The @cute.kernel function (for naming)
        launch_fn=launch_fn,           # Closure: (tiling) -> launches kernel
        input_tensors=list(inputs),
        output_tensors=[output],
        tiling=self.tiling,
        autotune_configs=self._autotune_configs(),
        algorithmic_flops_bytes_fn=self._algorithmic_flops_bytes,
    )
CtKernel との主な違い
項目CtKernel（cuTile）CuteKernel（CuTe DSL）
起動方法grid_fn + args_fn + cuTile コンパイラー@cute.jit ホストをラップする launch_fn クロージャー
コンパイルcuTile バイトコード + MLIRCuTe DSL JIT（自動キャッシュ）
自動チューニングcuda.tile_experimental の autotune_launch()構成ごとの triton.testing.do_bench_cudagraph
グリッド設定grid_fn(cfg) -> (x, y, z)launch_fn／ホスト関数内
TMA ディスクリプター暗黙的（cuTile が処理）ホスト関数内で明示的に設定
自動チューニング

CuteKernel は autotune_configs を順に処理し、それぞれを do_bench_cudagraph でベンチマークして、最速の構成を選択します。失敗した構成は捕捉され、スキップされます。

コンパイル

CuTe DSL JIT はアーティファクトを自動的にキャッシュします。compile() メソッドは出力ディレクトリを作成しますが、アーティファクトを明示的に抽出することはありません（カーネルごとに拡張できます）。

タイリング用データクラスの指針

CuTe DSL のタイリングには通常、cuTile のタイリングにはないフィールドが含まれます。

@dataclass
class Tiling:
    # Standard tile dimensions
    tile_m: int
    tile_n: int
    tile_k: int

    # CuTe DSL-specific fields
    num_warpgroups: int          # Warpgroups per CTA (typically 2-4)
    num_pipeline_stages: int     # Async copy pipeline depth
    num_tma_buffers: int         # TMA double/triple buffering
    cluster_m: int = 1           # Cluster shape (M dimension)
    cluster_n: int = 1           # Cluster shape (N dimension)

    def validate(self, pd) -> bool:
        # Validate against problem dimensions and device limits
        ...

具体的なフィールドはカーネル設計によって異なります。validate メソッドでは、タイルの各次元が問題の各次元を割り切れること、およびリソース使用量（レジスタ、SMEM）がデバイスの制限内に収まることを確認してください。

アーキテクチャ互換性の確認

カーネルは、特定の SM バージョンを必要とするアーキテクチャ固有の命令を対象にする場合があります。「Blackwell」というマーケティング上の名称は、命令セットが異なる複数の SM バージョンにまたがります。

  • SM100（B200、B300）— データセンター向け GPU。tcgen05 MMA、TMEM、すべてのクラスタ機能をサポートします
  • SM120（GeForce RTX 5090、RTX 5080）— コンシューマー向け Blackwell。sm_120a を使用し、tcgen05 ops はサポートしません

tcgen05 ops を使用するカーネルには SM100 が必要です。「Blackwell」であっても、RTX 5090（SM120）では実行できません。

CuTe DSL カーネルの実行、コンパイル、プロファイリングの前に、必ず利用可能なデバイスを確認してください。

nvidia-smi --query-gpu=name --format=csv,noheader

次に、GPU 名と SM バージョンを照合してください。カーネルの対象 SM バージョンが一致しない場合は、実行／プロファイリングの手順をスキップしてください。実行を試みないでください。コンパイル時に失敗するか、誤解を招く結果が出る可能性があります。カーネルの devlog にアーキテクチャ要件を記録し、スキップしたことも明記してください。

よくある落とし穴

  • CuTe DSL API を決してでっち上げないでください。/cute-dsl-ref スキル、docs/cute-dsl/ のドキュメント、または CUTLASS のサンプルカーネルで確認してください
  • @cute.kernel 関数名はモジュールのファイル名と一致させる必要があります（cuTile と同じルールです）
  • TMA ディスクリプターの設定はホスト側のみで行います。適切な @cute.jit によるホスト側の設定なしに、@cute.kernel 内で TMA 操作を行わないでください
  • レジスター予算は最大 255 個／スレッドです。docs/devices/ と src/mla_var3/conf/devices.json で、使用中のデバイスの SM 制限を確認してください
  • 共有メモリの容量は SM バージョンによって異なります。SM100（B200／B300）は最大 228 KB／SM、オプトイン時は 227 KB／ブロックです。SM120（RTX 5090）は 96 KB です。「Blackwell」だからといって、データセンター向けの SMEM 制限を想定しないでください
  • パイプラインのステージ数は SMEM 使用量に影響します（各ステージに専用バッファが必要です）。ステージ数を増やす前に、SMEM の合計を確認してください
  • バリア同期のエラーは表面化せず、クラッシュではなく誤った結果を引き起こします。必ず --check を使ってテストしてください
  • クラスタプログラミングでは、起動時にクラスタ互換のグリッド次元を使用する必要があります

参考リソース

  • API リファレンス：コア API の表、実行モデル、アーキテクチャ操作については /cute-dsl-ref を読み込んでください
  • 公式ドキュメント：docs/cute-dsl/（CuTe Python DSL）と docs/cutlass-cpp/cute/（CuTe C++ の概念）
  • サンプルカーネルと学習パス：分類された CUTLASS サンプルカーネルの一覧については /learn-cute-dsl を読み込んでください

ナレッジリンク

  • 共有最適化ナレッジ：docs/knowledge/optimizations/
  • 共有アンチパターン：docs/knowledge/anti-patterns/
  • CuTe Python DSL オーバーレイ：docs/knowledge/languages/cute-dsl/

開発ログの記録

docs/kernels/<kernel>.md を使用し、実装場所を Python パッケージパス形式で記録してください。

src/mla_var3/kernel/cute_python/mla/<kernel>/<kernel>[_vN]/

コメント

Markdown 対応