cuda

CUDA プログラミングスキル

基本方針

推測する前に計測します。 GPU のパフォーマンスは、直感に反することがよくあります。まずプロファイルし、次に仮説を立て、その後変更し、最後に検証します。

変更は小さく、切り分けて行います。 CUDA のバグは重なり合うことがあります。変更を一つ加えたらテストして、コミットします。「一度にすべて直す」誘惑には抗いましょう。

printf は最も頼れるツールです。 デバッガーが使えないときや、ツールの出力が解読しにくいときは、デバイスコード内の printf が真相を明らかにします。大いに活用することをためらわないでください。

差分をじっくり見ることもあります。 原因の分かりにくいセグメンテーション違反はよく起きます。ツールが役に立たないこともあります。人間ならではの方法として、差分を最小限にして丁寧に読み、バグを見つけます。これは正当な方法であり、ツールを使うより速いこともよくあります。

デバッグのワークフロー

バグへの最初の対応

1. 最小限の条件で再現する — 失敗するカーネルを、可能な限り小さい入力で切り分けます
2. printf を追加する — ツールを使う前に、デバイスコードに printf を追加して実行を追跡します
3. compute-sanitizer を実行する — メモリエラーを対話操作なしで検出します：

   compute-sanitizer --tool memcheck ./your_program
   compute-sanitizer --tool racecheck ./your_program  # for race conditions
   compute-sanitizer --tool initcheck ./your_program  # uninitialized memory
   

4. まだ解決しない場合、バックトレースを取得するために cuda-gdb を対話操作なしで試します：

   cuda-gdb -batch -ex "run" -ex "bt" ./your_program
   

5. ツールが役に立たない場合 — 動作するコードと壊れたコードの差分を最小限にします。それを読みます。バグは差分の中にあります。

デバイスコード内の printf

__global__ void myKernel(float* data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx == 0) {  // Limit output
        printf("Kernel launched, n=%d, data[0]=%f\n", n, data[0]);
    }
    // ... kernel logic ...
    if (idx < 10) {  // Sample a few threads
        printf("Thread %d: result=%f\n", idx, someValue);
    }
}

主なパターン：

• 出力の大量発生を防ぐため、if (idx == 0) または if (idx < N) で条件を付けます
• 起動を確認するため、カーネルのエントリーポイントで出力します
• 問題が疑われる箇所で、中間値を出力します
• フラッシュはカーネル完了時に自動で行われます

compute-sanitizer クイックリファレンス

よくある落とし穴： 「Invalid shared write... out of bounds」は通常、配列のインデックスが誤っているのではなく、カーネル起動時の動的共有メモリ割り当てが不足していることを意味します。<<<grid, block, smem_size>>> を確認してください。

# Memory errors (most common)
compute-sanitizer --tool memcheck ./program

# Other tools: racecheck, initcheck, synccheck
# For detailed options, see references/debugging-tools.md

cuda-gdb の非対話モード

# Get backtrace on crash
cuda-gdb -batch -ex "run" -ex "bt" ./program

# For breakpoints, thread inspection, see references/debugging-tools.md

デバッグ情報を付けてコンパイルします：

nvcc -g -G -lineinfo program.cu -o program

バイナリ検査に使う cuobjdump

# Dump PTX and SASS
cuobjdump -ptx ./program
cuobjdump -sass ./program

# For resource usage, symbol listing, see references/debugging-tools.md

デバッグツールの完全なリファレンス： compute-sanitizer の詳細なオプション、cuda-gdb のワークフロー、cuobjdump の解析パターンについては、references/debugging-tools.md を参照してください。

パフォーマンス最適化のワークフロー

黄金律

まずプロファイルし、結果を確認してから最適化します。 GPU のボトルネックに関する直感は、ほとんどの場合、間違っています。プロファイル → 修正 → 検証のサイクルこそが最適化の作業そのものであり、単なる準備段階ではありません。

パフォーマンス調査の手順

1. ベースラインを確立する — 処理時間を計測して記録します
2. nsys でプロファイルする — タイムラインを取得し、重要なカーネルを特定します
3. ncu で詳細に調べる — 特定のボトルネックとなるカーネルを分析します
4. 仮説を立てる — メトリクスに基づいて、具体的な仮説を立てます
5. 一つだけ変更する — 狙いを定めた変更を一つ加えます
6. 検証する — 再度プロファイルし、改善を確認します
7. 繰り返す

nsys（Nsight Systems）— タイムラインプロファイリング

nsys は、「時間はどこに使われているか？」を調べるために使います。CPU/GPU の相互作用、カーネル起動パターン、メモリ転送、全体のタイムラインを確認できます。

# Basic profile
nsys profile -o report ./program
nsys stats report.nsys-rep --report cuda_gpu_kern_sum

# With NVTX markers
nsys profile --trace=cuda,nvtx -o report ./program

# Key reports: cuda_gpu_kern_sum, cuda_api_sum, cuda_gpu_mem_time_sum, nvtx_sum
# For detailed usage, see references/nsys-guide.md

nsys の詳細な解析パターン： タイムラインの解釈、よくあるボトルネックの特定、解析ワークフローについては、references/nsys-guide.md を参照してください。

ncu（Nsight Compute）— カーネル解析

ncu は、「このカーネルはなぜ遅いのか？」を調べるために使います。詳細なメトリクス、ルーフライン、メモリ解析、占有率を確認できます。

# Profile specific kernel
ncu --kernel-name "myKernel" -o report ./program

# Quick summary to stdout
ncu --set basic ./program

# Sets: basic, full, memory, launch, roofline
# Sections: ComputeWorkloadAnalysis, MemoryWorkloadAnalysis, Occupancy
# For detailed metrics and interpretation, see references/ncu-guide.md

警告： ncu のエキスパートシステムによる推奨は、誤解を招くことがあります。必ず実際のメトリクスと実験で検証してください。

規模が重要です： 大規模では効果がある最適化が、小規模では逆効果になることがあります。理論上の最大規模ではなく、実際の問題サイズで必ずプロファイルしてください。

ncu メトリクスの詳細な解釈： ルーフライン解析、メモリのボトルネック、占有率の上限、ワープスケジューリングについては、references/ncu-guide.md を参照してください。

カスタム計測に使う NVTX

カーネル単位より細かい粒度が必要な場合は、NVTX を使います：

#include <nvtx3/nvToolsExt.h>

nvtxRangePush("Operation Name");
// ... code to profile ...
nvtxRangePop();

コンパイル： -lnvToolsExt | プロファイル： nsys profile --trace=cuda,nvtx

完全なパターンについて： ネストした範囲、色、解析ワークフローについては、references/nvtx-patterns.md を参照してください。

一般的なパフォーマンスパターン

症状	考えられる原因	調査方法
GPU 使用率が低い	カーネル起動オーバーヘッド、CPU がボトルネック	nsys のタイムラインでギャップを確認
メモリがボトルネック	アクセスパターンの不備、キャッシュヒット率の低さ	ncu のメモリセクションで合体アクセスを確認
計算がボトルネックなのに遅い	占有率の低さ、レジスタの圧迫	ncu の占有率を確認し、レジスタ数を減らす
小さなカーネルが多数ある	起動オーバーヘッドが支配的	nsys のタイムラインを確認し、融合を検討
memcpy にかかる時間が長い	H2D/D2H 転送が過剰	nsys cuda_gpu_mem で転送を確認し、まとめて転送
サイクルの大部分が停止している	バンク競合、メモリ待ち	ncu SchedulerStatistics で共有メモリを確認
リクエストあたりのセクター数が多い	合体アクセスの不備（1 リクエストあたり 4 セクター超）	ncu のメモリメトリクスを確認し、ベクトル化ロードを使用

重大な落とし穴： バンク競合やメモリ合体アクセスの問題は、プロファイリングを行わないと気づきにくい一方で、パフォーマンスを大きく左右することがよくあります。詳しい診断と修正方法については、references/performance-traps.md を参照してください。

現実的な見通し： 予測できなかった問題に最適化時間の 80％を充てる必要があります。プロファイリングに基づく反復によって、実際のボトルネックを見つけられます。

コンパイルリファレンス

# Debug build
nvcc -g -G -lineinfo -O0 program.cu -o program_debug

# Release build
nvcc -O3 -lineinfo program.cu -o program

# Specific architecture
nvcc -arch=sm_80 program.cu -o program  # Ampere
nvcc -arch=sm_89 program.cu -o program  # Ada Lovelace
nvcc -arch=sm_90 program.cu -o program  # Hopper

# Generate PTX (inspect it)
nvcc -ptx program.cu

# Verbose compilation (see register usage)
nvcc --ptxas-options=-v program.cu

# With NVTX
nvcc program.cu -lnvToolsExt -o program

本番環境でプロファイリングする場合は、必ず -lineinfo を指定してコンパイルしてください — オーバーヘッドは最小限で、ソースとの対応付けが可能になります。

ローカル API ドキュメント

grep による検索が可能な、完全なリファレンスドキュメントを利用できます。

PTX ISA 9.1 — references/ptx-docs/（405 ファイル、2.3MB）

• 検索ガイド：references/ptx-isa.md
• 用途：命令レベルの最適化、インライン PTX、TensorCore 操作（WMMA、WGMMA、TMA）、メモリのスウィズリング

CUDA Runtime API 13.1 — references/cuda-runtime-docs/（107 ファイル、0.9MB）

• 検索ガイド：references/cuda-runtime.md
• 用途：エラーコード、API パラメーター、デバイスプロパティ（cudaDeviceProp）、メモリ管理、ストリームの動作

CUDA Driver API 13.1 — references/cuda-driver-docs/（128 ファイル、0.8MB）

• 検索ガイド：references/cuda-driver.md
• 用途：コンテキスト管理（cuCtxCreate）、モジュールの読み込み（cuModuleLoad）、仮想メモリ、ドライバーエラー（CUDA_ERROR_*）、高度な機能

各検索ガイドには、grep の例、ドキュメントの構成、一般的な使用パターンが記載されています。

検索方法： grep/ripgrep を使って *-docs/ ディレクトリを直接検索してください。検索ガイド（.md ファイル）には、ナビゲーションのパターンと一般的な検索クエリが記載されています。

その他のリファレンス

• references/performance-traps.md — バンク競合、メモリ合体アクセス、スケールに依存する最適化
• references/debugging-tools.md — compute-sanitizer、cuda-gdb、cuobjdump の詳しい使用方法
• references/nsys-guide.md — nsys のタイムライン分析とボトルネックの特定
• references/ncu-guide.md — ncu のメトリクス、ルーフライン、占有率の解釈
• references/nvtx-patterns.md — NVTX による計測とプロファイリングのパターン

最適化前のチェックリスト

• 再現可能なベースラインの計測時間を確立した
• nsys でプロファイリングしてホットスポットを特定した
• 実行時間の大部分を占めるカーネルを把握している
• 対象カーネルを ncu でプロファイリングした
• 具体的なボトルネック（メモリ、計算、レイテンシ）を特定した
• 具体的で検証可能な仮説を立てた
• 変更するのは 1 つだけにする計画を立てた