cuda-auto-tune

NCU 駆動の反復的カーネル最適化（CUDA／CUTLASS／Triton／CuTe DSL）

ゲートチェック（最適化の前に必ず実施）

STOP — Do you have NCU profile data for this kernel?
  NO  → Go to Step 1. Do NOT touch any kernel code.
  YES → Go to Step 2.

厳守事項 — いずれかに違反すると、最適化全体が無効になります：

• NCU のデータなしに、カーネルコード、起動設定、テンプレートパラメーターを絶対に変更しないでください。
• 推奨事項にはすべて、根拠となる具体的な NCU メトリクス値を必ず示してください。
• 各反復では、ルーフライン、メモリ階層、ワープストール、占有率を最低限確認してください。
• 最適化のプレイブックは、カーネルの実装タイプに一致させてください。
• コードを変更するたびに、再プロファイリングを行い、--diff で比較してください。
• 改善が頭打ちになった場合、またはメトリクスがハードウェアの上限に近づいた場合は、反復を終了してください。

---

必須の最適化ループ

┌─────────────────────────────────────────────────────────────────────┐
│  Step 1: Profile (NCU --set full)                                   │
│      ↓                                                              │
│  Step 2: Multi-dimensional analysis + identify kernel type          │
│      ↓                                                              │
│  Step 3: Apply type-specific playbook (one change per iteration)    │
│      ↓                                                              │
│  Step 4: Re-profile + diff → improved? → loop or stop               │
│      ↑                                           │                  │
│      └───────────────────────────────────────────┘                  │
└─────────────────────────────────────────────────────────────────────┘

---

ステップ 1：NCU でプロファイリングする（必須 — データなしで最適化しない）

オプション A：プロファイリングスクリプト（推奨）

# Native CUDA / CUTLASS binaries
bash cuda-auto-tune/scripts/ncu_profile.sh ./kernel report_v1

# Triton / Python
bash cuda-auto-tune/scripts/ncu_profile.sh "python your_kernel.py" report_v1

# CuTe DSL / Python
bash cuda-auto-tune/scripts/ncu_profile.sh "python your_cutedsl_kernel.py" report_v1

このスクリプトは --set full を収集し、CSV にエクスポートした後、詳細分析を実行してレポートを生成します。

オプション B：手動プロファイリング

ncu --set full -o report_v1 --target-processes all ./your_kernel
ncu --import report_v1.ncu-rep --page raw --csv > report_v1.csv
python3 cuda-auto-tune/scripts/ncu_analyse.py report_v1.csv

カーネル名フィルター（ノイズを低減）

# CUTLASS only
ncu --set full -o report_v1 --target-processes all \
    --kernel-name "cutlass_\|sm90_\|ampere_" ./cutlass_program

# Triton only
ncu --set full -o report_v1 --target-processes all \
    --kernel-name "triton_" "python triton_kernel.py"

# CuTe DSL (kernel name often generic — use --type override in analysis)
python3 cuda-auto-tune/scripts/ncu_analyse.py report_v1.csv --type cutedsl

想定される出力

ncu_reports/
├── report_v1.ncu-rep           # Full binary report
├── report_v1.csv               # Raw metrics CSV
├── report_v1_analysis.md       # Deep analysis report
└── report_v1_summary.txt       # Per-kernel summary

---

ステップ 2：多面的な分析

2.1 実装タイプを特定する

NCU の「Function Name」とソースのコンテキストから、カーネルのタイプを判定します：

タイプ	判定の手がかり
Native CUDA	ライブラリの接頭辞がない。手書きの __global__ 関数
CUTLASS	cutlass_ 接頭辞、smXX_xmma_、tensorop または cutlass を含む
Triton	triton_ 接頭辞、triton を含む、エンコードされた接尾辞（例：_0d1d...e）
CuTe DSL	@cute.kernel による汎用的な名前。ソースのインポート（cutlass.cute、cute.compile）または --type cutedsl で確認
ライブラリ	cublas*、cudnn* — ベースライン／参照用のみ。最適化対象外

2.2 共通診断（すべてのカーネルタイプで必ず実行）

分析項目	主な NCU メトリクス	出力
ルーフライン	SM スループット、メモリスループット	計算律速／メモリ律速／レイテンシ律速／バランス
メモリ階層	L1／L2 ヒット率、コアレス化率、DRAM スループット	キャッシュ効率＋帯域幅のサブボトルネック（DRAM／L2／L1）
ワープストール	PC サンプリングによるストール理由（long_scoreboard、wait、barrier、など）	ストール理由の上位項目と割合
命令構成	パイプラインの FMA／ALU／LSU／Tensor 使用率	パイプラインの不均衡、Tensor Core の使用状況
占有率	アクティブワープ率、制限要因の内訳（レジスタ／smem／ワープ／ブロック）	制限要因＋レジスタ数＋smem サイズ
メモリハザード	バンク競合、レジスタスピル（ローカルストアセクター）	深刻度と根本原因
分岐	平均実行スレッド数と平均アクティブスレッド数（実数）	分岐率

2.3 タイプ別の重点項目

タイプ	主な重点項目
Native CUDA	起動設定（ブロックサイズ、グリッド）、メモリアクセスパターン、非同期コピー（cp.async／TMA）、Tensor Core の活用可能性
CUTLASS	ThreadblockShape、WarpShape、ステージ数、アライメント、スケジュールポリシー、エピローグ融合、CTA スウィズル
Triton	num_warps、num_stages、BLOCK_* サイズ、コンパイラーヒント（tl.multiple_of、tl.max_contiguous）、tl.dot の設定
CuTe DSL	threads_per_cta、elems_per_thread、CopyAtom（num_bits_per_copy）、tiled_copy のレイアウト、smem ステージング、cta_reduce パターン

2.4 ボトルネック分類の判断ツリー

SM% > MEM% + 20  →  COMPUTE_BOUND
MEM% > SM% + 20  →  MEMORY_BOUND
  ├─ DRAM throughput > 70%        → DRAM-Bound (near HBM ceiling)
  ├─ L2 hit < 50%, DRAM > 40%    → DRAM-Bound (L2 miss driven)
  ├─ L1 hit < 20%, L2 hit >= 50% → L2-Bound
  └─ L1 hit < 20%                → L1-Bound
SM% < 40 AND MEM% < 40           →  LATENCY_BOUND
SM% > 60 AND MEM% > 60           →  BALANCED (near peak)

2.5 結論のテンプレート（分析のたびに必須）

=== Conclusion ===
Kernel:    {kernel_name}
Type:      {Native CUDA | CUTLASS | Triton | CuTe DSL}
Arch:      SM_{arch}
Overall:   {COMPUTE_BOUND | MEMORY_BOUND | LATENCY_BOUND | BALANCED}
Duration:  {duration_us} us
Roofline:  SM {sm}%, MEM {mem}%, DRAM {dram}%
Occupancy: {occ}% (theoretical: {theo}%), limited by {limiter}
Regs/Thread: {regs}, Smem/Block: {smem} KB

Findings (sorted by severity):
  [CRITICAL] {finding}: {NCU evidence with numbers} -> {specific action}
  [WARNING]  {finding}: {NCU evidence with numbers} -> {specific action}
  [INFO]     {finding}: {NCU evidence with numbers}

Optimization priorities:
  1. {highest_priority} (expected gain: Nx, evidence: {metric}={value})
  2. {second_priority}  (expected gain: Nx, evidence: {metric}={value})
  3. {third_priority}   (expected gain: Nx, evidence: {metric}={value})

---

ステップ 3：タイプ別プレイブックを適用する

直感だけに頼った変更は禁止です。変更はすべて、NCU で判明した問題に直接対処する必要があります。 反復ごとに変更は 1 つだけ行い、その後再プロファイリングします（ステップ 4）。

---

3.1 プレイブック：Native CUDA

3.1.1 起動構成

NCU の検出結果	対応	コードパターン
占有率 < 50%、ブロックサイズ < 128	ブロックサイズを 128–256 に増やします	kernel<<<grid, 256>>>
レジスターが占有率の制限要因です	__launch_bounds__ でレジスター数を制限します	__global__ void __launch_bounds__(256, 2) kernel()
グリッドが小さすぎます（< SM 数）	すべての SM をカバーできるだけのブロック数を確保します	十分な N を指定して grid = (N + block - 1) / block
占有率が低く、ブロック数が制限要因です	ブロックサイズを小さくして、SM あたりのブロック数を増やします	256 の代わりに 128 を試します

3.1.2 メモリアクセスの最適化

NCU の検出結果	対応	コードパターン
ロードの合体率 > 8	ワープ内で連続したアドレス指定を確保し、AoS→SoA にします	data[threadIdx.x + blockIdx.x * blockDim.x]
ストアの合体率 > 8	スキャッター書き込みには共有メモリのステージングを使用します	まず smem に書き込み、その後、合体した書き戻しを行います
L1 ヒット率 < 20%	頻繁に再利用されるデータには __shared__ を使用します	__syncthreads() を使って共有メモリにタイル化します
L2 ヒット率 < 50%	L2 永続化ヒント（Ampere 以降）を使用します	ホットデータ範囲に cudaAccessPolicyWindow を使用します
DRAM スループット > 80%	混合精度や圧縮を使ってデータ移動を減らします	帯域幅が重要な処理には half / __nv_bfloat16 を使用します
バンク競合 > 100K	共有メモリにパディングを追加するか、レイアウトをスウィズルします	__shared__ float smem[32][33];（+1 のパディング）
レジスターのスピル > 0	スレッドごとの状態を減らし、__launch_bounds__ を使用します	アキュムレーターを単純化するか、サブカーネルに分割します

3.1.3 レイテンシ隠蔽とパイプライン処理

NCU の検出結果	対応	コードパターン
stall_long_scoreboard > 30%（SM>=80）	cp.async とダブルバッファリングを使用します	__pipeline_memcpy_async(&smem, &gmem, size)
stall_long_scoreboard > 30%（SM>=90）	バルクの非同期転送には TMA を使用します	cute::copy(tma_load, ...) または CuTe TMA アトム
stall_barrier > 25%	同期の頻度を減らし、ワーププリミティブを使用します	__shfl_sync()、cooperative_groups
stall_wait > 30%、long_scoreboard < 15%	パイプラインのバッファーが過剰です。深さを減らします	バッファーステージを 1 つ削除します
stall_math_pipe_throttle > 20%	計算が飽和しています（好ましい兆候です）	Tensor Core の使用、または FLOP 数の削減を検討します

3.1.4 Tensor Core の使用率

NCU の検出結果	対応
pipe_tensor < 5%、FP16/BF16 ワークロードで GEMM に似たパターン	WMMA（wmma::mma_sync）またはインライン PTX（mma.sync）を使用します
pipe_tensor < 5% だが、データは FP32	nvcuda::wmma::precision::tf32 を指定した wmma::mma_sync で TF32 パスを使用します
pipe_fma_fp16 > 10%、pipe_tensor < 5%	スカラー FP16 FMA から Tensor Core パスに切り替えます

3.1.5 ベクトル化メモリアクセス

// NCU evidence: coalescing ratio > 4 for 32-bit loads
// Before: scalar loads
float val = input[idx];

// After: vectorized 128-bit load (4x float)
float4 val = reinterpret_cast<const float4*>(input)[idx / 4];

---

3.2 プレイブック：CUTLASS

3.2.1 カーネル構成の解析

CUTLASS のカーネル名には構成がエンコードされています。次の情報を抽出します。

• アーキテクチャ：sm80_、sm90_、ampere_、hopper_
• 計算タイプ：tensorop と simt のどちらか
• タイル形状：128x128x32、256x128x64
• パイプラインのステージ数：末尾の x3、x5
• アラインメント：align8
• スケジュール（3.x）：WarpSpecialized、WarpSpecializedCooperative、WarpSpecializedPingpong

3.2.2 タイル形状と占有率

NCU の検出結果	対応
占有率 < 40%、smem が制限要因です	ThreadblockShape を小さくする（例：256x128→128x128）か、ステージ数を減らします
占有率 < 40%、レジスターが制限要因です	WarpShape を小さくして（例：64x64→32x32）、スレッドあたりのレジスター数を減らします
SM スループット < 30%、グリッドが小さい	ThreadblockShape を大きくして、CTA あたりの処理要素数を増やします
SM スループット > 80%、MEM < 40%	すでに計算律速です。重なりを増やすため、パイプラインのステージ数を増やします

3.2.3 パイプラインのステージ数

NCU の検出結果	対応
stall_long_scoreboard > 30%	ステージ数を増やします（Ampere：3→5、Hopper：2→3）
stall_wait > 30%、long_scoreboard < 15%	パイプラインのバッファーが過剰です。smem を節約するため、ステージ数を減らします
smem が制限要因で、ステージ数 > 3	ステージ数を減らして smem を空け、占有率を高めます

3.2.4 アラインメントとベクトル化

NCU の検出結果	対応
ロードの結合が 4 を超え、アラインメントが 8 未満	CUTLASS のアラインメントを 8（128 バイト）に引き上げます。行列の先頭次元をアラインメントの倍数にパディングします
データが TensorOp に対応しているのに SIMT パスが使用されている	CUTLASS の設定を tensorop に切り替えます（2～8 倍の高速化）
TensorOp が設定されているのに pipe_tensor が 5% 未満	アラインメント要件を確認します。LD は InstructionShape::kK の倍数である必要があります

3.2.5 スケジュールとアーキテクチャ

NCU の検出結果	対応
SM>=90 で CUTLASS 2.x を使用	WarpSpecialized + TMA を備えた CUTLASS 3.x にアップグレードします（1.2～1.5 倍の高速化）
大規模な GEMM で L2 ヒット率が 50% 未満	ThreadblockSwizzle を追加します（2.x：GemmIdentityThreadblockSwizzle<N>、3.x：StreamK またはタイルスウィズル）
Hopper で stall_long_scoreboard が 30% を超える	TMA ロードを使用する WarpSpecializedCooperative スケジュールに切り替えます

3.2.6 エピローグの融合

NCU の検出結果	対応
CUTLASS カーネルが連続して複数実行されている（例：GEMM + バイアス + 活性化）	CUTLASS エピローグビジターツリーを使用して、単一のカーネルに融合します
DRAM トラフィックが多い（読み書きの GB が想定値を超える）	GEMM 後の処理をエピローグに移し、中間テンソルをなくします

---

3.3 プレイブック：Triton

3.3.1 カーネルの分類

Triton カーネルのサブタイプ（カーネル名から判定）：

• triton_poi_：Inductor のポイントワイズ（自動生成）
• triton_red_：Inductor のリダクション（自動生成）
• triton_per_：Inductor の永続リダクション（自動生成）
• カスタムの @triton.jit：手書きカーネル（全面的にチューニング可能）

Inductor 生成カーネル：PyTorch レベル（torch._inductor.config）で最適化するか、ホットパスの場合はカスタムの @triton.jit として書き直します。

3.3.2 num_warps のチューニング

NCU の検出結果	対応
レジスター数が 128 以上、num_warps が 8 以上	重大：num_warps を減らします（4 または 2 を試します）
レジスター数が 64 以上、num_warps が 8 以上	num_warps を 4 に減らします
占有率が 40% 未満で、レジスターが制限要因	num_warps を減らす、および／または BLOCK_* タイルサイズを小さくします
SM スループットが 30% 未満で、ワープ数が少ない	レイテンシー隠蔽を改善するため、num_warps を増やします

3.3.3 num_stages のチューニング

NCU の検出結果	対応
stall_long_scoreboard が 30% を超える	num_stages を増やします（Ampere では 2→3→4、Hopper では 2→3）
stall_wait が 30% を超え、long_scoreboard が 15% 未満	num_stages を減らす（バッファー過多）か、タイルあたりの処理量を増やします
Smem が占有率の制限要因	num_stages を減らします（各ステージで smem バッファーが 2 倍になります）
Hopper で long_scoreboard が高い	TMA ベースのロードには tl.make_block_ptr() の使用も検討します

3.3.4 BLOCK_* タイルサイズのチューニング

NCU の検出結果	対応
レジスターの負荷が高い	BLOCK_M、BLOCK_N、または BLOCK_K を小さくします
SM スループットが低く、計算律速を改善できる余地がある	タイルあたりの計算量を増やすため、BLOCK_M/BLOCK_N を大きくします
DRAM 帯域幅が上限に近い	書き戻し前のデータ再利用を増やすため、BLOCK_K を大きくします

3.3.5 メモリアクセスの最適化

NCU の検出結果	対応	コードパターン
ロードの結合が 8 を超える	ストライドのヒントを追加します	tl.multiple_of(stride, 16) および tl.max_contiguous(offsets, BLOCK)
転置された入力でアクセスが結合されていない	構造化ポインターを使用します	tl.make_block_ptr(base, shape, strides, offsets, block_shape, order)
L1 ヒット率が低い	アクセスパターンの連続性を確認します	最内次元のストライドが 1 であることを確認します

3.3.6 Tensor Core の使用率

NCU の検出結果	対応
カーネルが tl.dot を使用しているのに pipe_tensor が 5% 未満	1）fp32 では allow_tf32=True を設定します。2）BLOCK_K を 16 の倍数にします。3）データ型が fp16/bf16/tf32/fp8 であることを確認します
コードに tl.dot がなく、pipe_tensor が 5% 未満	GEMM に似たパターンがありません。tl.dot を使うように構造を変更します

3.3.7 Triton autotune の統合

@triton.autotune(
    configs=[
        triton.Config({'BLOCK_M': 128, 'BLOCK_N': 128, 'BLOCK_K': 32}, num_warps=4, num_stages=3),
        triton.Config({'BLOCK_M': 64,  'BLOCK_N': 64,  'BLOCK_K': 64}, num_warps=4, num_stages=4),
        triton.Config({'BLOCK_M': 128, 'BLOCK_N': 64,  'BLOCK_K': 32}, num_warps=8, num_stages=3),
    ],
    key=['M', 'N', 'K'],
)
@triton.jit
def kernel(...):
    ...

NCU でボトルネックが判明したら、手当たり次第に検索するのではなく、有望な領域に絞って autotune の設定を調整します。

---

3.4 プレイブック：CuTe DSL

3.4.1 主要な調整パラメーター

パラメーター	効果	一般的な範囲
threads_per_cta	CTA あたりのワープ数です。占有率、バリアコスト、リダクションコストに影響します	128–512
elems_per_thread	スレッドあたりの要素数です。レジスター使用量、データ再利用に影響します	4–32
num_bits_per_copy	CopyAtom の幅です。ベクトル化されたロード／ストアの帯域幅に影響します	32, 64, 128
Smem ステージングバッファー	パイプラインの深さ × タイルサイズです。smem の使用量に影響します	占有率を高めるため最小化

3.4.2 占有率の最適化

NCU の検出結果	対処
占有率 < 40%、レジスターが制限要因	elems_per_thread または threads_per_cta を減らします。cute.compile() に --maxrregcount=128 を追加します
占有率 < 40%、smem が制限要因	threads_per_cta を減らします（ワープが少ないほどリダクションバッファーが小さくなります）。またはステージングバッファーの数を減らします
レジスター >= 128、ワープ >= 8	重大：threads_per_cta を 128 または 256 に減らします

3.4.3 メモリーアクセス（TiledCopy）

NCU の検出結果	対処
ロードのコアレッシング > 8	1) num_bits_per_copy を 128 に増やします。2) t_layout が連続するアドレスに沿ってスレッドを配置していることを確認します。3) from_dlpack() で assumed_align=16 が使われていることを確認します
stall_long_scoreboard > 30%	1) num_bits_per_copy を 128 に増やします。2) 再利用を増やすため elems_per_thread を増やします。3) SM>=80 では CpAsyncOp コピーアトムを使用します。4) ダブルバッファリングを追加します
stall_wait > 30%、long_scoreboard < 15%	パイプラインのバッファーが過剰です。1 ステージあたりの計算量を増やすため elems_per_thread を増やすか、パイプラインの深さを減らします

3.4.4 同期とリダクション

NCU の検出結果	対処
stall_barrier > 25%	1) threads_per_cta を減らします（バリアで待機するワープが少なくなります）。2) ワープ数が 32 以下の場合、2 回目の sync_threads をシャッフルブロードキャストに置き換えます。3) 複数の cta_reduce 呼び出しをまとめます
バリアコストが高く、リダクションが小さい	要素数が少ない場合は、smem を使わないワープのみのリダクションを使用します
反復ごとに sync_threads が複数回ある	同期ポイントを最小限にします。非同期パイプラインのコミット／待機パターンを使用します

3.4.5 スレッドの分岐

NCU の検出結果	対処
分岐 > 20%	threads_per_cta * elems_per_thread が問題の次元 N にできるだけ近くなるよう調整し、述語によって無効化されるスレッドを減らします
述語付きコピーで分岐が多い	N が threads_per_cta * elems_per_thread で割り切れることを確認するか、末尾処理の方法を使用します

3.4.6 計算の最適化

NCU の検出結果	対処
pipe_tensor < 5%、FP16 の GEMM 型演算	Tensor Core パスに MmaOp を指定して cute.make_mma_atom() を使用します
pipe_fma が高く pipe_tensor が低い（RMSNorm／LayerNorm などの非 GEMM 演算）	リダクションには Tensor Core を適用できません。代わりにメモリーとバリアの最適化に注力します

3.4.7 再プロファイリングのためのキャッシュ無効化

CuTe DSL は JIT を介して Python を CUDA にコンパイルします。コードを変更した後：

# Clear compilation cache to ensure re-compilation
rm -rf __pycache__/ .cache/ /tmp/cutlass_cute_cache/
# Then re-profile
bash cuda-auto-tune/scripts/ncu_profile.sh "python your_cutedsl_kernel.py" report_v2

---

ステップ 4：再プロファイリングと検証（変更のたびに必須）

4.1 再プロファイリング

# Clear JIT caches first
rm -rf ~/.triton/cache            # Triton
rm -rf __pycache__/ .cache/       # CuTe DSL

# Profile updated version
bash cuda-auto-tune/scripts/ncu_profile.sh ./kernel_v2 report_v2
# or
bash cuda-auto-tune/scripts/ncu_profile.sh "python kernel_v2.py" report_v2

4.2 ベースラインとの比較

python3 cuda-auto-tune/scripts/ncu_analyse.py ncu_reports/report_v2.csv --diff ncu_reports/report_v1.csv

4.3 検証チェックリスト

チェック項目	基準
実行時間は改善しましたか？	gpu__time_duration.sum が減少
対象のボトルネックは改善しましたか？	変更のきっかけとなった特定の指標が改善
新たなボトルネックはありませんか？	差分レポートに新たな重大な検出結果がない
ハードウェアの上限に達していますか？	SM スループット > 80% または DRAM スループット > 85% は、ピーク近くを意味します

4.4 反復記録のテンプレート

各反復を記録し、説明責任を果たします。

=== Iteration {N} ===
Change:  {what was changed and why}
NCU evidence: {metric}={before_value} -> {finding}
Report: report_v{N}.csv

Result:
  Duration: {before} us -> {after} us ({delta}%)
  Target metric: {metric}={before} -> {after}
  New findings: {any new issues introduced}

Decision: {CONTINUE to next bottleneck | STOP — at ceiling | ROLLBACK — regression}

---

クイックリファレンス：NCU の重要指標

指標	NCU キー
実行時間	gpu__time_duration.sum [us]
SM スループット	sm__throughput.avg.pct_of_peak_sustained_elapsed [%]
メモリスループット	gpu__compute_memory_throughput.avg.pct_of_peak_sustained_elapsed [%]
DRAM スループット	gpu__dram_throughput.avg.pct_of_peak_sustained_elapsed [%]
L1 ヒット率	l1tex__t_sector_hit_rate.pct [%]
L2 ヒット率	lts__t_sector_hit_rate.pct [%]
ロードのコアレス化	l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum / l1tex__t_requests_pipe_lsu_mem_global_op_ld.sum
バンクコンフリクト	l1tex__data_bank_conflicts_pipe_lsu_mem_shared.sum
レジスタスピル	l1tex__t_sectors_pipe_lsu_mem_local_op_st.sum [sector]
占有率	sm__warps_active.avg.pct_of_peak_sustained_active [%]
ワープの実行可能性	smsp__warps_eligible.avg.per_cycle_active [warp]
スレッドあたりのレジスタ数	launch__registers_per_thread [register/thread]
ブロックあたりの共有メモリ	launch__shared_mem_per_block [Kbyte/block]

---

概要

このスキルでは、厳密な プロファイル → 分析 → 変更 → 検証 のサイクルを徹底します。 NCU データがなければ最適化は行いません。指標による根拠がなければコードは変更しません。 カーネルの種類（Native CUDA／CUTLASS／Triton／CuTe DSL）ごとに、NCU 指標と対応するアクションをまとめた専用のプレイブックを用意しています。 すべての変更を追跡し、再プロファイリングによって検証します。