cuda-profiling

CUDA プロファイリング

目的

Nsight Systems（タイムラインレベル）と Nsight Compute（カーネルレベルのメトリクス）を使用して CUDA アプリケーションをプロファイリングする方法をエージェントに案内します。NCU CLI によるメトリクスの自動収集、ルーフラインモデルの解釈、カーネルがメモリバウンドかコンピュートバウンドかの診断を扱います。

使用する場面

• CUDA カーネルの速度が想定より遅く、ボトルネックを特定する必要がある場合
• カーネルのバリエーション（タイル戦略、ブロックサイズ）を比較する場合
• ncu のメトリクスを使用して CI の性能回帰チェックを構築する場合
• マルチストリームパイプラインで CPU と GPU のアクティビティを関連付ける場合
• タイムラインで処理を確認できるよう、NVTX を使ってアプリケーションのフェーズに注釈を付ける場合
• 占有率、メモリスループット、SM 利用率のメトリクスを解釈する場合

ワークフロー

1. プロファイリングツールを選択する

What do you need?
├── System-wide timeline (CPU+GPU+CUDA API) → Nsight Systems (nsys)
├── Per-kernel deep metrics (occupancy, memory) → Nsight Compute (ncu)
└── Quick metric from CLI in CI → ncu --metrics ...

2. Nsight Systems — タイムラインプロファイリング

# Profile entire application
nsys profile --trace=cuda,nvtx,osrt --output=report ./my_cuda_app

# Open report
nsys-ui report.nsys-rep

# CLI summary
nsys stats report.nsys-rep

タイムラインで確認する項目：

• カーネル起動間の空白（CPU ボトルネックまたは同期ポイント）
• cudaDeviceSynchronize による停滞
• ストリーム間での H2D コピーとカーネル実行のオーバーラップ
• CUDA API 呼び出しのオーバーヘッド

# Capture with CUDA graph info
nsys profile --capture-range=cudaProfilerApi ./my_cuda_app

3. NVTX 範囲の注釈

#include <nvtx3/nvToolsExt.h>

void pipeline(void) {
    nvtxRangePushA("H2D copy");
    cudaMemcpyAsync(d_in, h_in, size, cudaMemcpyHostToDevice, stream);
    nvtxRangePop();

    nvtxRangePushA("kernel");
    my_kernel<<<grid, block, 0, stream>>>(d_in, d_out, n);
    nvtxRangePop();

    nvtxRangePushA("D2H copy");
    cudaMemcpyAsync(h_out, d_out, size, cudaMemcpyDeviceToHost, stream);
    nvtxRangePop();
}

-lnvToolsExt を指定してコンパイルするか、ヘッダーのみの nvtx3 をリンクします。範囲は Nsight Systems で色付きの帯として表示されます。

4. Nsight Compute — カーネル分析

# Profile all kernels, save report
ncu -o kernel_report ./my_cuda_app

# Profile specific kernel by name
ncu --kernel-name regex:matmul_tiled ./my_cuda_app

# Launch UI
ncu-ui kernel_report.ncu-rep

NCU レポートの主なセクション：

• Speed of Light：SM スループット、メモリスループット、ピーク値の比較
• Occupancy：アクティブなワープ数とハードウェア上限
• Memory Workload Analysis：L1/L2 ヒット率、コアレッシング効率
• Warp State Statistics：停滞の理由（メモリ、バリア、ディスパッチ）

5. NCU CLI のメトリクス

# Essential metrics set
ncu --metrics \
  sm__throughput.avg.pct_of_peak_sustained_elapsed,\
  dram__throughput.avg.pct_of_peak_sustained_elapsed,\
  sm__warps_active.avg.pct_of_peak_sustained_active,\
  l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum,\
  smsp__sass_thread_inst_executed_op_ffma_pred_on.sum \
  ./my_cuda_app

# CSV export for CI
ncu --csv --metrics dram__bytes_read.sum,dram__bytes_write.sum ./my_cuda_app

# Set kernel replay mode for accurate counters
ncu --kernel-replay-mode application ./my_cuda_app

6. メモリバウンドかコンピュートバウンドかの診断

Roofline interpretation
├── dram__throughput near peak AND sm__throughput low → memory-bound
│   └── Fix: coalescing, shared mem tiling, reduce traffic
├── sm__throughput near peak AND dram low → compute-bound
│   └── Fix: tensor cores, loop unrolling, ILP
└── Both low → launch config, occupancy, or sync overhead

ルーフラインモデル（概念図）：

Performance (GFLOP/s)
    |     /\  compute roof
    |    /  \
    |   /    \____ memory roof (bandwidth-limited region)
    |  /
    +------------------ Arithmetic Intensity (FLOP/byte)

演算強度を測定します：smsp__sass_thread_inst_executed_op_ffma_pred_on.sum * 2 / dram__bytes.sum

7. 占有率の分析

ncu --metrics sm__warps_active.avg.pct_of_peak_sustained_active,\
launch__occupancy_limit_registers,\
launch__occupancy_limit_shared_mem,\
launch__occupancy_limit_block_size \
./my_cuda_app

制限要因	一般的な対処法
レジスター	-maxrregcount を指定する、カーネルを簡素化する
共有メモリ	タイルサイズを縮小する、フェーズを分割する
ブロックサイズ	512 以上ではなく、128 または 256 を試す

8. プロファイリングのワークフローチェックリスト

# 1. Build with line info (not -G unless debugging)
nvcc -lineinfo -O3 -arch=sm_80 -o app main.cu

# 2. Timeline first
nsys profile --trace=cuda,nvtx -o timeline ./app

# 3. Deep dive on hot kernel
ncu --kernel-name regex:hot_kernel --set full ./app

# 4. Compare before/after
ncu --csv --metrics sm__throughput.avg.pct_of_peak_sustained_elapsed ./app_v1 > v1.csv
ncu --csv --metrics sm__throughput.avg.pct_of_peak_sustained_elapsed ./app_v2 > v2.csv

よくある問題

症状	原因	対処法
ERR_NVGPUCTRPERM	プロファイリング権限が不足している	sudo で実行するか、NVreg_RestrictProfilingToAdminUsers=0 を設定する
すべてのメトリクスがゼロになる	プロファイリングが無効になっている、または GPU が正しくない	CUDA_VISIBLE_DEVICES を確認し、--target-processes all を使用する
NCU レポートが空になる	カーネルの実行時間が短すぎる、または起動されていない	ワークロードを増やし、cudaGetLastError() を確認する
プロファイリングのオーバーヘッドが非常に大きい	多数のカーネルに対してすべてのメトリクスセットを取得している	--kernel-name フィルターと --launch-skip を使用する
タイムラインにオーバーラップが表示されない	デフォルトストリームを 1 つだけ使用している	複数のストリームを作成し、非同期コピーを使用する
占有率に問題がないのにカーネルが遅い	メモリレイテンシを隠蔽できていない	メモリのコアレッシングを確認し、アクティブなワープ数を増やす

関連スキル

• skills/gpu/cuda — カーネルの記述、占有率の調整、nvcc フラグ
• skills/gpu/gpu-memory-model — コアレッシング、バンク競合、SIMT モデル
• skills/gpu/cuda-debugging — 性能調整の前に正しさを確認する
• skills/profilers/intel-vtune-amd-uprof — CPU 側のルーフライン分析とホットスポット分析
• skills/profilers/flamegraphs — nsys タイムラインを補完する CPU フレームグラフ
• skills/profilers/hardware-counters — perf stat の一般的な概念