cuda

CUDA

目的

NVIDIA CUDA C/C++ 開発において、カーネル起動構成、レジスターからグローバルメモリーまでのメモリー階層、ストリームによる非同期実行、nvcc のコンパイルフラグ、Thrust ライブラリーの使い方、ワープダイバージェンスやメモリーアクセスの非合体といった一般的なパフォーマンス上の問題の診断について、エージェントをガイドします。

使用する場面

• 行列積、リダクション、ステンシル演算用の CUDA カーネルを作成または最適化するとき
• ブロック／グリッドの次元を選び、オキュパンシーを見積もるとき
• メモリーアクセスパターンや低いオキュパンシーが原因でカーネルの動作が遅い場合にデバッグするとき
• cudaMemcpyAsync を使ったマルチストリームパイプラインを設定するとき
• nvcc でコンパイルし、アーキテクチャーフラグ（-gencode）を選択するとき
• 手書きのカーネルの代わりに、並列プリミティブとして Thrust を使うとき

ワークフロー

1. 最小限のカーネルと起動

CUDA では、スレッドをブロックに、ブロックをグリッドにまとめて処理を編成します。

Thread hierarchy
├── grid (1D/2D/3D)
│   └── block (1D/2D/3D, max 1024 threads)
│       └── thread (threadIdx, blockIdx, blockDim, gridDim)

// vector_add.cu
#include <cuda_runtime.h>
#include <stdio.h>

__global__ void vector_add(const float *a, const float *b, float *c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n)
        c[i] = a[i] + b[i];
}

int main(void) {
    const int n = 1 << 20;
    size_t bytes = n * sizeof(float);
    float *h_a, *h_b, *h_c, *d_a, *d_b, *d_c;

    cudaMalloc(&d_a, bytes);
    cudaMalloc(&d_b, bytes);
    cudaMalloc(&d_c, bytes);
    // ... host init and cudaMemcpy H2D ...

    int threads = 256;
    int blocks = (n + threads - 1) / threads;
    vector_add<<<blocks, threads>>>(d_a, d_b, d_c, n);
    cudaDeviceSynchronize();

    cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost);
    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
    return 0;
}

2. メモリー階層

メモリー	スコープ	レイテンシー	一般的な用途
レジスター	スレッドごと	約 1 サイクル	ローカルスカラー、ループインデックス
共有（__shared__）	ブロックごと	約 5 サイクル	タイルデータ、ハロー交換
グローバル	全スレッド	約 400 サイクル以上	大規模配列、合体アクセス
定数（__constant__）	読み取り専用、キャッシュあり	高速なブロードキャスト	カーネルパラメーター、ルックアップテーブル
テクスチャー	キャッシュされた 2D アクセス	キャッシュあり	画像サンプリング、不規則な読み取り

共有メモリーの例（行列タイル）：

#define TILE 16

__global__ void matmul_tiled(const float *A, const float *B, float *C, int N) {
    __shared__ float As[TILE][TILE];
    __shared__ float Bs[TILE][TILE];

    int row = blockIdx.y * TILE + threadIdx.y;
    int col = blockIdx.x * TILE + threadIdx.x;
    float sum = 0.0f;

    for (int t = 0; t < (N + TILE - 1) / TILE; t++) {
        As[threadIdx.y][threadIdx.x] = (row < N && t * TILE + threadIdx.x < N)
            ? A[row * N + t * TILE + threadIdx.x] : 0.0f;
        Bs[threadIdx.y][threadIdx.x] = (col < N && t * TILE + threadIdx.y < N)
            ? B[(t * TILE + threadIdx.y) * N + col] : 0.0f;
        __syncthreads();

        for (int k = 0; k < TILE; k++)
            sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];
        __syncthreads();
    }
    if (row < N && col < N)
        C[row * N + col] = sum;
}

3. ストリームと非同期コピー

cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);

cudaMemcpyAsync(d_a, h_a, bytes, cudaMemcpyHostToDevice, stream1);
cudaMemcpyAsync(d_b, h_b, bytes, cudaMemcpyHostToDevice, stream2);
vector_add<<<blocks, threads, 0, stream1>>>(d_a, d_b, d_c, n);
cudaMemcpyAsync(h_c, d_c, bytes, cudaMemcpyDeviceToHost, stream1);
cudaStreamSynchronize(stream1);

ページロックされたホストメモリー（cudaMallocHost）を使うと、カーネル実行と真の非同期 DMA をオーバーラップできます。

4. nvcc によるコンパイル

# Single architecture (local GPU)
nvcc -O3 -arch=sm_80 -o prog vector_add.cu

# Fat binary for multiple GPUs
nvcc -O3 \
  -gencode arch=compute_80,code=sm_80 \
  -gencode arch=compute_90,code=sm_90 \
  -o prog vector_add.cu

# Debug symbols for cuda-gdb
nvcc -G -g -O0 -arch=sm_80 -o prog_debug vector_add.cu

# Show PTX/SASS
nvcc -arch=sm_80 -ptx vector_add.cu
nvcc -arch=sm_80 -cubin vector_add.cu
cuobjdump -sass prog

よく使われるフラグ：

フラグ	効果
-O3	積極的な最適化
-G	デバッグ時に最適化を無効化
-lineinfo	プロファイル内でソース行と対応付け
-Xcompiler -fopenmp	CUDA と併用するホスト側 OpenMP
--use_fast_math	精度を下げて高速化する数学関数
-maxrregcount=N	レジスター数を制限してオキュパンシーを向上

5. オキュパンシーの見積もり

# CUDA Occupancy Calculator (spreadsheet) or programmatic:
./occupancy_tool --kernel vector_add --block-size 256 --regs 16 --smem 0

判断ツリー：

Kernel slow?
├── Low occupancy (< 25%) → reduce registers, shared mem, or block size
├── Memory-bound → check coalescing, use shared memory tiling
└── Compute-bound → increase arithmetic intensity, use tensor cores

cudaOccupancyMaxActiveBlocksPerMultiprocessor API、または Nsight Compute の sm__warps_active.avg.pct_of_peak_sustained_active メトリックを使います。

6. Thrust の基本

#include <thrust/device_vector.h>
#include <thrust/sort.h>
#include <thrust/reduce.h>

thrust::device_vector<int> d_vec(1000000);
thrust::sort(d_vec.begin(), d_vec.end());
int sum = thrust::reduce(d_vec.begin(), d_vec.end());

Thrust は一時ストレージとカーネル起動を内部で処理します。ソート、スキャン、リダクションには Thrust を優先し、用途に特化した融合演算にはカスタムカーネルを作成してください。

7. よくある落とし穴

ワープダイバージェンス：ワープ内のスレッド（32 個）は、SIMT のロックステップで実行されます。threadIdx に応じた分岐はシリアル化を引き起こします。

// Bad: divergent branch
if (threadIdx.x % 2 == 0) { heavy_a(); } else { heavy_b(); }

// Better: separate kernels or predication

非合体アクセス：連続するスレッドは、連続するアドレスにアクセスする必要があります。

// Bad: strided access
float val = data[threadIdx.x * stride];

// Good: coalesced
float val = data[blockIdx.x * blockDim.x + threadIdx.x];

よくある問題

症状	原因	対処法
cudaErrorIllegalAddress (700)	デバイスポインターの範囲外アクセス	compute-sanitizer --tool memcheck を実行し、境界チェックを追加する
cudaErrorLaunchTimeout (702)	カーネルの実行がウォッチドッグの制限を超過	処理量を減らす、カーネルを分割する、開発用 GPU で TDR を無効にする
オキュパンシーが低いという警告	レジスターまたは共有メモリーの使いすぎ	__shared__ のサイズを減らす、-maxrregcount を使う、ブロックを小さくする
CPU と GPU で結果が異なる	共有メモリーまたはグローバルメモリーの競合	__syncthreads() を追加する、リダクションにはアトミック操作を使う
no kernel image available	-arch=sm_XX の指定が誤っている	GPU のコンピュートケイパビリティーに合わせる：nvidia-smi --query-gpu=compute_cap
H2D/D2H コピーが遅い	ページ可能なホストメモリー	ページロックされたバッファーには cudaMallocHost を使う

関連スキル

• skills/gpu/cuda-profiling — Nsight Systems/Compute によるカーネルパフォーマンスの診断
• skills/gpu/cuda-debugging — 正しさの確認に使う cuda-gdb と compute-sanitizer
• skills/gpu/gpu-memory-model — SIMT 実行、合体ルール、バンク競合
• skills/gpu/hip-rocm — CUDA カーネルを AMD HIP に移植
• skills/gpu/triton-lang — カスタムカーネル向けの Python DSL 代替
• skills/hpc/openmp — CUDA と併用するホスト側の並列処理