• 1. GPUとCPUのアーキテクチャの違い
  • 1.1 CPU: 逐次処理に最適化されたプロセッサ
  • 1.2 GPU: 大規模並列処理に最適化されたプロセッサ
  • 1.3 SIMT実行モデル
• 2. CUDAプログラミングモデル: Grid、Block、Thread階層
  • 2.1 Thread
  • 2.2 Thread Block（Block）
  • 2.3 Grid
  • 2.4 Thread Block Cluster（Compute Capability 9.0以降）
• 3. CUDAカーネルの開発と実行設定
  • 3.1 カーネル関数の定義
  • 3.2 実行設定
• 4. スレッドインデックス: threadIdx、blockIdx、blockDim、gridDim
  • 4.1 組み込み変数
  • 4.2 グローバルスレッドIDの計算
• 5. GPUメモリの種類
  • 5.1 レジスタ
  • 5.2 ローカルメモリ
  • 5.3 Shared Memory
  • 5.4 Global Memory
  • 5.5 Constant Memory
  • 5.6 Texture Memory
  • 5.7 メモリ階層のまとめ
• 6. メモリ管理API
  • 6.1 明示的メモリ管理
    • cudaMalloc: デバイスメモリの割り当て
    • cudaMemcpy: ホスト-デバイス間のデータ転送
    • cudaFree: デバイスメモリの解放
    • 典型的なCUDAプログラムの流れ
  • 6.2 Unified Memory
• 7. ワープ実行とワープダイバージェンス
  • 7.1 ワープの概念
  • 7.2 ワープダイバージェンス
    • ワープダイバージェンスを最小化する戦略
• 8. 占有率の概念と最適化
  • 8.1 占有率とは
  • 8.2 占有率に影響する要因
  • 8.3 占有率の最適化方法
• 9. 実践例
  • 9.1 ベクトル加算
  • 9.2 行列積
    • ナイーブバージョン
    • タイルバージョン（Shared Memory使用）
• 10. NVIDIA GPU世代の特徴（Compute Capability）
  • アーキテクチャ別の主なプログラミング関連の変更
• 11. デバッグとプロファイリングツール
  • 11.1 cuda-gdb
  • 11.2 Compute Sanitizer
  • 11.3 NVIDIA Nsight
    • デバッグ/プロファイリングのワークフロー
• 12. エラーハンドリングのベストプラクティス
• 結論
• 参考文献
  • クイズ

1. GPUとCPUのアーキテクチャの違い

CUDAプログラミングを理解するには、まずGPUとCPUの根本的なアーキテクチャの違いを把握する必要があります。

1.1 CPU: 逐次処理に最適化されたプロセッサ

CPU（Central Processing Unit）は、複雑な制御フロー、分岐予測、大容量キャッシュを備え、逐次的なタスクに最適化されています。一般的な高性能CPUは8〜64個のコアを持ち、各コアが複雑な命令を高速で独立して実行できます。CPUのトランジスタの大部分は制御ロジックとキャッシュに割り当てられ、シングルスレッドの実行速度の最大化に重点を置いています。

1.2 GPU: 大規模並列処理に最適化されたプロセッサ

一方、GPU（Graphics Processing Unit）は数千の小さなコアを備え、超並列計算に特化しています。NVIDIA GPUはStreaming Multiprocessor（SM）と呼ばれるユニットで構成され、各SMには数十から数百のCUDA Coreが含まれています。GPUのトランジスタの大部分は演算論理ユニット（ALU）に割り当てられ、数千のスレッドの同時実行を可能にしています。

1.3 SIMT実行モデル

NVIDIA GPUはSIMT（Single Instruction, Multiple Threads）実行モデルを使用します。SIMTはSIMD（Single Instruction, Multiple Data）と類似していますが、重要な違いがあります。SIMDではベクトル幅がソフトウェアに公開されるのに対し、SIMTは個々のスレッドの実行と分岐の挙動を規定します。各スレッドは独自のプログラムカウンタとレジスタ状態を持ち、論理的には独立した実行パスを辿ることができます。

SIMTの核心はワープベースの実行です。GPUは32スレッドを1つのワープとしてグループ化し、同じ命令を同時に実行します。ワープ内のすべてのスレッドが同じコードパスを辿る場合に最大のパフォーマンスが得られます。スレッドが異なる分岐を取ると、パフォーマンスが低下します（ワープダイバージェンスと呼ばれ、後述で詳しく説明）。

2. CUDAプログラミングモデル: Grid、Block、Thread階層

CUDAプログラミングモデルの最も基本的な概念はスレッド階層です。CUDAでカーネル関数が呼び出されると、多数のスレッドが作成されて並列に実行されます。これらはGrid、Block、Threadの3段階の階層で構成されます。

2.1 Thread

Threadは、CUDA実行の最も基本的な単位です。各スレッドはカーネルコードの1つのインスタンスを実行し、独自のレジスタとローカルメモリを持ちます。各スレッドは一意のIDを通じて、自身が処理すべきデータを決定します。

2.2 Thread Block（Block）

Thread Blockは、スレッドのグループです。同じブロック内のスレッドは以下の特徴を共有します:

• Shared Memoryを通じてデータを共有できる
• __syncthreads()による同期が可能
• 単一のSM上で実行され、実行中に他のSMに移動しない
• 最大1024スレッドを含むことができる（Compute Capabilityにより異なる場合がある）

Thread Blockは1D、2D、3Dで構成でき、ベクトル、行列、ボリュームデータに対する自然なインデックス付けが可能です。

2.3 Grid

Gridは、Thread Blockの集合です。1回のカーネル呼び出しで1つのGridが作成されます。Gridも1D、2D、3Dで構成できます。異なるブロック間ではShared Memoryによる直接的なデータ共有はできず、同期も制限されています（Cooperative Groupsなどの特殊なAPIを使用する必要があります）。

2.4 Thread Block Cluster（Compute Capability 9.0以降）

NVIDIA Hopperアーキテクチャ（Compute Capability 9.0）から、Thread Block Clusterというオプション層が追加されました。クラスタは複数のThread Blockで構成され、同じクラスタ内のブロックは同じGPC（GPU Processing Cluster）上で実行されるため、Distributed Shared Memoryを通じて互いのShared Memoryにアクセスできます。

Grid
 +-- Block Cluster（オプション、CC 9.0以降）
      +-- Thread Block（最大1024スレッド）
           +-- Thread（個々の実行単位）

3. CUDAカーネルの開発と実行設定

3.1 カーネル関数の定義

CUDAでは、GPU上で実行される関数をカーネルと呼びます。カーネル関数は__global__修飾子を使用して定義され、戻り値の型はvoidでなければなりません。

__global__ void myKernel(int *data, int n) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < n) {
        data[idx] = data[idx] * 2;
    }
}

CUDAには3つの関数修飾子があります:

__host__と__device__を併用すると、ホストとデバイスの両方で関数がコンパイルされます。

3.2 実行設定

カーネルを呼び出す際、**<<<gridDim, blockDim>>>**構文でGridとBlockの次元を指定します。

// 1D設定
int N = 1024;
int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
myKernel<<<blocksPerGrid, threadsPerBlock>>>(d_data, N);

// 2D設定
dim3 blockDim(16, 16);      // 16x16 = ブロックあたり256スレッド
dim3 gridDim(64, 64);       // 64x64ブロック
matMulKernel<<<gridDim, blockDim>>>(d_A, d_B, d_C, N);

// Shared Memoryサイズとストリームの指定
myKernel<<<gridDim, blockDim, sharedMemBytes, stream>>>(args);

<<<>>>構文の完全な形式は<<<gridDim, blockDim, sharedMemBytes, stream>>>です。第3引数は動的に割り当てるShared Memoryのバイト数、第4引数はCUDA Streamです。省略した場合、それぞれ0とデフォルトストリームが使用されます。

重要: ブロックあたりのスレッド数にはハードウェア制限があります。現在、すべてのNVIDIA GPUはブロックあたり最大1024スレッドをサポートしています。この値を超えるとカーネルの起動が失敗します。

4. スレッドインデックス: threadIdx、blockIdx、blockDim、gridDim

カーネル内部では、各スレッドが組み込み変数を使用して自身の位置を決定します。これらの変数はuint3またはdim3型で、.x、.y、.zメンバを持ちます。

4.1 組み込み変数

4.2 グローバルスレッドIDの計算

1Dグリッドと1Dブロックの場合:

int globalIdx = blockIdx.x * blockDim.x + threadIdx.x;

2Dグリッドと2Dブロックの場合:

int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
int globalIdx = row * width + col;

3Dにも同じパターンが拡張されます:

int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int z = blockIdx.z * blockDim.z + threadIdx.z;

境界チェックは不可欠です。 データサイズがブロックサイズの正確な倍数でない場合、最後のブロックの一部のスレッドが有効な範囲外に出る可能性があるため、必ず境界チェックを行う必要があります:

if (globalIdx < N) {
    // 有効なインデックスに対してのみ計算を実行
    output[globalIdx] = input[globalIdx] * 2;
}

5. GPUメモリの種類

CUDA GPUは複数のメモリ空間を提供し、それぞれアクセス速度、サイズ、可視性（スコープ）が異なります。適切なメモリを選択することが、CUDAプログラムの最適化の鍵です。

5.1 レジスタ

• 場所: オンチップ（SM内部）
• 可視性: 各スレッドにプライベート
• 速度: 最速（1サイクルのレイテンシ）
• サイズ: SMあたり64K個の32ビットレジスタ、スレッドあたり最大255個

カーネル内のローカル変数は、デフォルトでレジスタに割り当てられます。レジスタ使用量が多いと、SM上に同時に常駐できるスレッド数が減少し、占有率が低下します。

5.2 ローカルメモリ

• 場所: オフチップ（Device Memory、Global Memoryと同じ物理的な場所）
• 可視性: 各スレッドにプライベート
• 速度: Global Memoryと同じ（遅い、数百サイクル）
• 用途: レジスタに収まらないローカル変数（レジスタスピル）、大きな配列

「ローカル」という名前にもかかわらず、実際にはオフチップに配置されるため、アクセス速度は遅いです。レジスタが不足するとコンパイラが自動的にローカルメモリにスピルします。

5.3 Shared Memory

• 場所: オンチップ（SM内部、L1キャッシュと物理的な空間を共有）
• 可視性: 同じThread Block内のすべてのスレッド
• 速度: レジスタに近い（バンクコンフリクトなしで約5サイクル）
• サイズ: 通常SMあたり48KB〜164KB（アーキテクチャにより異なる）

__global__ void sharedMemExample(float *data) {
    __shared__ float sharedData[256];  // 静的割り当て

    int tid = threadIdx.x;
    sharedData[tid] = data[blockIdx.x * blockDim.x + tid];

    __syncthreads();  // ブロック内のすべてのスレッドを同期

    // sharedDataを使用した演算
    float result = sharedData[tid] + sharedData[255 - tid];
    data[blockIdx.x * blockDim.x + tid] = result;
}

Shared Memoryはバンク（通常32個）に分割されており、複数のスレッドが同時に同じバンクにアクセスするとバンクコンフリクトが発生し、シリアル化されます。バンクコンフリクトを回避するアクセスパターンの設計が重要です。

5.4 Global Memory

• 場所: オフチップ（Device DRAM、すなわちHBMまたはGDDR）
• 可視性: すべてのスレッド + ホスト
• 速度: 最も遅い（数百サイクルのレイテンシ）
• サイズ: 最大（数GB〜数十GB）

cudaMalloc()で割り当てられるメモリがGlobal Memoryです。L1/L2キャッシュによりアクセス速度を改善でき、コアレスドアクセス（隣接するスレッドが連続するメモリアドレスにアクセスすること）により帯域幅を最大化する必要があります。

5.5 Constant Memory

• 場所: オフチップ（Global Memory領域）、専用キャッシュを介してキャッシュ
• 可視性: すべてのスレッド（読み取り専用）
• 速度: キャッシュヒット時は非常に高速
• サイズ: 64KB

__constant__ float constData[256];

// ホストから値を設定
cudaMemcpyToSymbol(constData, hostData, sizeof(float) * 256);

ワープ内のすべてのスレッドが同じアドレスを読み取る場合に最高のパフォーマンスを発揮します。ブロードキャストメカニズムにより、1回のメモリ読み取りでワープ内のすべてのスレッドに値が配信されます。

5.6 Texture Memory

• 場所: オフチップ、専用のTexture Cacheを介してキャッシュ
• 可視性: すべてのスレッド（読み取り専用）
• 特徴: 2D空間局所性に最適化、ハードウェア補間をサポート

Texture Memoryは、画像処理などの2D/3Dデータセットで空間的に隣接するデータへのアクセスパターンに有利です。モダンなCUDAではSurface Objectsと併用されます。

5.7 メモリ階層のまとめ

高速 <--------------------------------------------> 低速
レジスタ > Shared Memory > L1/L2キャッシュ > Constant/Textureキャッシュ > Global Memory
(オンチップ)  (オンチップ)    (オンチップ)    (キャッシュ済み)              (オフチップ)

6. メモリ管理API

6.1 明示的メモリ管理

従来のCUDAメモリ管理アプローチは、ホストメモリとデバイスメモリを明示的に分離して管理します。

cudaMalloc: デバイスメモリの割り当て

float *d_array;
cudaMalloc((void **)&d_array, N * sizeof(float));

cudaMallocはGPUのGlobal Memoryにメモリを割り当てます。割り当てられたポインタ（d_array）はデバイス上でのみ有効であり、ホスト上で直接デリファレンスすることはできません。

cudaMemcpy: ホスト-デバイス間のデータ転送

// ホスト → デバイス
cudaMemcpy(d_array, h_array, N * sizeof(float), cudaMemcpyHostToDevice);

// デバイス → ホスト
cudaMemcpy(h_result, d_result, N * sizeof(float), cudaMemcpyDeviceToHost);

// デバイス → デバイス
cudaMemcpy(d_dest, d_src, N * sizeof(float), cudaMemcpyDeviceToDevice);

cudaMemcpyは同期関数であり、転送が完了するまでホストスレッドをブロックします。非同期転送には、CUDA Streamと共にcudaMemcpyAsyncを使用します。

cudaFree: デバイスメモリの解放

cudaFree(d_array);

典型的なCUDAプログラムの流れ

// 1. ホストメモリの割り当てとデータの初期化
float *h_input = (float *)malloc(N * sizeof(float));
float *h_output = (float *)malloc(N * sizeof(float));
initializeData(h_input, N);

// 2. デバイスメモリの割り当て
float *d_input, *d_output;
cudaMalloc(&d_input, N * sizeof(float));
cudaMalloc(&d_output, N * sizeof(float));

// 3. データ転送 ホスト → デバイス
cudaMemcpy(d_input, h_input, N * sizeof(float), cudaMemcpyHostToDevice);

// 4. カーネルの起動
int blockSize = 256;
int gridSize = (N + blockSize - 1) / blockSize;
myKernel<<<gridSize, blockSize>>>(d_input, d_output, N);

// 5. 結果の転送 デバイス → ホスト
cudaMemcpy(h_output, d_output, N * sizeof(float), cudaMemcpyDeviceToHost);

// 6. メモリの解放
cudaFree(d_input);
cudaFree(d_output);
free(h_input);
free(h_output);

6.2 Unified Memory

CUDA 6.0で導入されたUnified Memoryは、ホストとデバイスが単一のアドレス空間を共有できるようにします。データ移動はCUDAランタイムにより自動的に管理されます。

float *data;
cudaMallocManaged(&data, N * sizeof(float));

// ホストからアクセス可能
for (int i = 0; i < N; i++) {
    data[i] = (float)i;
}

// 同じポインタをデバイスで使用
myKernel<<<gridSize, blockSize>>>(data, N);
cudaDeviceSynchronize();

// ホストから結果にアクセス（同期後）
printf("Result: %f\n", data[0]);

cudaFree(data);  // cudaMallocManagedもcudaFreeで解放

Unified Memoryの利点:

• cudaMemcpyの呼び出しが不要で、より簡潔なコード
• ホストとデバイスで同じポインタを使用
• ランタイムがオンデマンドでデータ移行を処理
• デバイスメモリ容量を超えるデータの処理が可能（オーバーサブスクリプション）

ただし、パフォーマンスの面では、StreamsとcudaMemcpyAsyncを使用してカーネル実行とデータ転送をオーバーラップさせるように最適化されたプログラムの方が、Unified Memoryのみを使用するプログラムよりも高いパフォーマンスを達成できます。cudaMemPrefetchAsyncを使用してUnified Memoryのパフォーマンスを向上させることもできます。

7. ワープ実行とワープダイバージェンス

7.1 ワープの概念

ワープは、CUDAにおける実行のスケジューリングの基本単位です。SMのワープスケジューラは32スレッドをグループ化してワープを形成し、ワープレベルで命令を発行します。同じワープ内の32スレッドは同じプログラムカウンタ（PC）を共有し、同じ命令を同時に実行します。

Thread Block内では、スレッドIDの順にワープが形成されます:

• ワープ0: スレッド0〜31
• ワープ1: スレッド32〜63
• ワープ2: スレッド64〜95
• ... 以下同様

7.2 ワープダイバージェンス

ワープダイバージェンスは、ワープ内のスレッドが条件分岐（if、switch、forなど）で異なるパスを辿る場合に発生します。

__global__ void divergentKernel(int *data, int *result) {
    int tid = threadIdx.x;

    // ワープダイバージェンスが発生!
    if (tid % 2 == 0) {
        result[tid] = data[tid] * 2;      // 偶数スレッド
    } else {
        result[tid] = data[tid] + 10;     // 奇数スレッド
    }
}

上記のコードでは、ワープ内の偶数番号と奇数番号のスレッドが異なる分岐を取ります。GPUはこれを各分岐パスを順次実行し、現在のパスに属さないスレッドを無効化することで処理します。結果として、両方の分岐の実行時間が合算され、パフォーマンスが低下します。

ワープダイバージェンスを最小化する戦略

// 悪い例: 同じワープ内で分岐
if (threadIdx.x % 2 == 0) { ... }

// 良い例: ワープ境界で分岐するように設計
if (threadIdx.x / 32 % 2 == 0) { ... }
// または
if (blockIdx.x % 2 == 0) { ... }

核心原則は、同じワープ内のスレッドが同じコードパスを辿るようにコードを設計することです。分岐条件がワープ境界に揃っていれば、各ワープ全体が1つのパスのみを実行し、ダイバージェンスが防止されます。

8. 占有率の概念と最適化

8.1 占有率とは

**占有率（Occupancy）**は、SM上で同時にアクティブにできる最大ワープ数に対する、実際のアクティブワープ数の比率です。

占有率 = アクティブワープ数 / SMあたりの最大ワープ数

例えば、SMが最大64ワープをサポートし、実際に32ワープがアクティブである場合、占有率は50%です。

8.2 占有率に影響する要因

占有率は3つのリソースによって決まります:

1. レジスタ使用量: スレッドあたりのレジスタ使用量が多いと、SM上に常駐できるスレッドの総数が減少します。SMあたり64Kの32ビットレジスタの場合、スレッドあたり128レジスタを使用すると、SM上には512スレッド（16ワープ）しか配置できません。

2. Shared Memoryの使用量: ブロックあたりのShared Memory使用量が多いと、SM上に同時に配置できるブロック数が減少します。

3. ブロックサイズ（スレッド数）: SMあたりの最大ブロック数の制限があり（例: Ampere CC 8.0で32、CC 8.6で16）、ブロックが小さすぎるとこの制限に達し、SMのスレッド容量を埋めることができない場合があります。

8.3 占有率の最適化方法

• ブロックサイズは128、256、512を選択: 経験的に256が良い出発点です。小さすぎる場合（32、64）はSMあたりのブロック数制限に達し、大きすぎる場合（1024）はリソース要件が増加して占有率が低下する可能性があります。

• レジスタ使用量の制御: __launch_bounds__修飾子やコンパイラオプション-maxrregcountを使用して、カーネルごとのレジスタ数を制限します。

__global__ void __launch_bounds__(256, 4)  // maxThreadsPerBlock, minBlocksPerSM
myKernel(float *data) {
    // ...
}

• CUDA Occupancy Calculatorの使用: NVIDIAが提供するOccupancy CalculatorスプレッドシートやcudaOccupancyMaxPotentialBlockSize() APIを使用して、最適なブロックサイズを自動的に決定します。

int blockSize;
int minGridSize;
cudaOccupancyMaxPotentialBlockSize(&minGridSize, &blockSize, myKernel, 0, 0);

注意: 高い占有率が常に高いパフォーマンスを保証するわけではありません。メモリアクセスパターン、命令レベル並列性（ILP）、Shared Memoryの活用など、複数の要因が総合的にパフォーマンスに影響します。

9. 実践例

9.1 ベクトル加算

最も基本的なCUDAの例であるベクトル加算です。2つの配列の対応する要素を加算し、結果を格納します。

#include <stdio.h>
#include <cuda_runtime.h>

__global__ void vectorAdd(const float *A, const float *B, float *C, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        C[idx] = A[idx] + B[idx];
    }
}

int main() {
    int N = 1 << 20;  // 約100万要素
    size_t size = N * sizeof(float);

    // ホストメモリの割り当て
    float *h_A = (float *)malloc(size);
    float *h_B = (float *)malloc(size);
    float *h_C = (float *)malloc(size);

    // データの初期化
    for (int i = 0; i < N; i++) {
        h_A[i] = 1.0f;
        h_B[i] = 2.0f;
    }

    // デバイスメモリの割り当て
    float *d_A, *d_B, *d_C;
    cudaMalloc(&d_A, size);
    cudaMalloc(&d_B, size);
    cudaMalloc(&d_C, size);

    // ホスト → デバイス転送
    cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);

    // カーネルの起動
    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
    vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

    // デバイス → ホスト転送
    cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);

    // 検証
    for (int i = 0; i < N; i++) {
        if (fabs(h_C[i] - 3.0f) > 1e-5) {
            fprintf(stderr, "Verification failed at index %d!\n", i);
            return -1;
        }
    }
    printf("Vector addition successful!\n");

    // クリーンアップ
    cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);
    free(h_A); free(h_B); free(h_C);

    return 0;
}

重要なポイント:

• 各スレッドが1つの要素を処理（1:1マッピング）
• (N + threadsPerBlock - 1) / threadsPerBlockは切り上げ除算で、すべての要素をカバー
• 境界チェック（if (idx < N)）で範囲外アクセスを防止

9.2 行列積

行列積は、CUDA最適化の代表的な例です。まずナイーブバージョンを見て、次にShared Memoryを使用したタイルバージョンを取り上げます。

ナイーブバージョン

__global__ void matMulNaive(float *A, float *B, float *C, int N) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;

    if (row < N && col < N) {
        float sum = 0.0f;
        for (int k = 0; k < N; k++) {
            sum += A[row * N + k] * B[k * N + col];
        }
        C[row * N + col] = sum;
    }
}

// 起動
dim3 blockDim(16, 16);
dim3 gridDim((N + 15) / 16, (N + 15) / 16);
matMulNaive<<<gridDim, blockDim>>>(d_A, d_B, d_C, N);

ナイーブバージョンでは、各スレッドが結果行列の1要素を計算しながら、行列Aの行全体と行列Bの列全体をGlobal Memoryから読み込みます。これにより O(N^3) のGlobal Memoryアクセスが発生し、非常に非効率です。

タイルバージョン（Shared Memory使用）

#define TILE_SIZE 16

__global__ void matMulTiled(float *A, float *B, float *C, int N) {
    __shared__ float tileA[TILE_SIZE][TILE_SIZE];
    __shared__ float tileB[TILE_SIZE][TILE_SIZE];

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

    // タイルを反復処理
    for (int t = 0; t < (N + TILE_SIZE - 1) / TILE_SIZE; t++) {
        // タイルをShared Memoryにロード
        if (row < N && (t * TILE_SIZE + threadIdx.x) < N)
            tileA[threadIdx.y][threadIdx.x] = A[row * N + t * TILE_SIZE + threadIdx.x];
        else
            tileA[threadIdx.y][threadIdx.x] = 0.0f;

        if ((t * TILE_SIZE + threadIdx.y) < N && col < N)
            tileB[threadIdx.y][threadIdx.x] = B[(t * TILE_SIZE + threadIdx.y) * N + col];
        else
            tileB[threadIdx.y][threadIdx.x] = 0.0f;

        __syncthreads();

        // タイル内の積和演算
        for (int k = 0; k < TILE_SIZE; k++) {
            sum += tileA[threadIdx.y][k] * tileB[k][threadIdx.x];
        }

        __syncthreads();
    }

    if (row < N && col < N) {
        C[row * N + col] = sum;
    }
}

タイル行列積の重要なアイデア:

1. Global MemoryからのデータをTILE_SIZE x TILE_SIZEのタイルに分割し、Shared Memoryにロード
2. Shared Memoryからの読み取りはGlobal Memoryから約100倍高速
3. ブロック内の複数スレッドが同じデータを再利用し、Global Memoryアクセスを大幅に削減
4. __syncthreads()により、すべてのスレッドがタイルのロードを完了してから計算を開始することを保証

10. NVIDIA GPU世代の特徴（Compute Capability）

Compute Capability（CC）は、GPUのハードウェア機能と仕様を示すバージョン番号です。メジャーバージョンはアーキテクチャ世代を表し、マイナーバージョンは世代内の改良を表します。

アーキテクチャ別の主なプログラミング関連の変更

Volta（CC 7.0）: Independent Thread Schedulingが導入され、ワープ内のスレッドがより柔軟に分岐できるようになりました。以前のアーキテクチャではワープ内の暗黙的なロックステップ同期に依存していましたが、Volta以降は明示的な同期（__syncwarp()）が必要です。

Hopper（CC 9.0）: Thread Block Clusterの概念が追加され、複数のThread BlockがDistributed Shared Memoryを通じて連携できるようになりました。非同期データ移動用のTMA（Tensor Memory Accelerator）ユニットも導入されました。

Blackwell（CC 10.0）: 第5世代Tensor CoreがFP4とFP6精度をネイティブにサポートし、マイクロテンソルフォーマットと動的レンジスケーリングを使用します。AI計算性能がHopperに比べて大幅に向上し、最大20 PFLOPSのAI計算性能を実現しています。

11. デバッグとプロファイリングツール

11.1 cuda-gdb

cuda-gdbは、GNU GDBのCUDA拡張であり、実際のGPUハードウェア上で動作するCUDAアプリケーションをデバッグできます。

主要機能:

• GPUカーネル内でのブレークポイントの設定
• スレッド、ブロック、ワープレベルでの状態のクエリ
• デバイスメモリの内容の検査
• ホストコードとデバイスコードの同時デバッグ

# デバッグ情報を含めてコンパイル
nvcc -g -G -o myapp myapp.cu

# cuda-gdbでデバッグを開始
cuda-gdb ./myapp

# cuda-gdb内のコマンド
(cuda-gdb) break myKernel
(cuda-gdb) run
(cuda-gdb) cuda thread       # 現在のスレッド情報
(cuda-gdb) cuda block        # 現在のブロック情報
(cuda-gdb) info cuda threads # すべてのCUDAスレッドをリスト

11.2 Compute Sanitizer

Compute Sanitizerは、CUDAプログラムの機能的正確性をチェックするためのツール群で、CUDA Toolkitに含まれています。4つのサブツールを提供します:

# memcheckの実行
compute-sanitizer --tool memcheck ./myapp

# racecheckの実行
compute-sanitizer --tool racecheck ./myapp

# initcheckの実行
compute-sanitizer --tool initcheck ./myapp

11.3 NVIDIA Nsight

NVIDIA Nsightは、統合的な開発、デバッグ、プロファイリング環境を提供するツール群です。

• Nsight Systems: システム全体のパフォーマンスを分析します。CPU-GPUタイムライン、カーネル実行時間、メモリ転送、API呼び出しを視覚的に表示します。全体的なボトルネックを特定する最初のステップとして使用されます。

• Nsight Compute: 個々のCUDAカーネルの詳細なパフォーマンスメトリクスを分析します。占有率、メモリ帯域幅の利用率、命令スループット、ワープの状態などを細部にわたって調べます。特定のカーネルの最適化に使用されます。

• Nsight Visual Studio Edition / VS Code Extension: IDEに統合されたCUDAデバッグとプロファイリングサポートを提供します。

# Nsight Systemsでプロファイリング
nsys profile --stats=true ./myapp

# Nsight Computeでカーネルを分析
ncu --set full ./myapp

デバッグ/プロファイリングのワークフロー

1. 機能的正確性の確認: Compute Sanitizer（memcheck、racecheck）を使用して、まずメモリエラーと競合状態を検出
2. システムレベルの分析: Nsight Systemsを使用して全体的なボトルネックを特定（CPU-GPU同期、メモリ転送など）
3. カーネルレベルの最適化: Nsight Computeを使用してボトルネックカーネルの詳細なパフォーマンスメトリクスを分析し、最適化

12. エラーハンドリングのベストプラクティス

ほとんどのCUDA API呼び出しはcudaError_t型のエラーコードを返します。本番コードでは、常にエラーをチェックする必要があります。

#define CUDA_CHECK(call)                                                    \
    do {                                                                    \
        cudaError_t err = call;                                             \
        if (err != cudaSuccess) {                                           \
            fprintf(stderr, "CUDA error at %s:%d: %s\n",                    \
                    __FILE__, __LINE__, cudaGetErrorString(err));            \
            exit(EXIT_FAILURE);                                             \
        }                                                                   \
    } while (0)

// 使用例
CUDA_CHECK(cudaMalloc(&d_array, size));
CUDA_CHECK(cudaMemcpy(d_array, h_array, size, cudaMemcpyHostToDevice));

// カーネル起動後のエラーチェック
myKernel<<<gridSize, blockSize>>>(d_array, N);
CUDA_CHECK(cudaGetLastError());       // カーネル起動エラー
CUDA_CHECK(cudaDeviceSynchronize());  // カーネル実行中のエラー

カーネル呼び出しは非同期であるため、cudaGetLastError()は起動時のエラーをチェックし、cudaDeviceSynchronize()はカーネル実行中に発生したエラーをチェックします。

結論

CUDAプログラミングは、GPUの大規模並列処理能力を活用するためのコア技術です。以下に解説した内容をまとめます:

• GPUアーキテクチャ: 数千の軽量コアがSIMTモデルで動作し、データ並列処理に最適化
• スレッド階層: 並列処理をGrid、Block、Threadの3段階の階層で構成
• メモリ階層: レジスタ、Shared、Globalなど様々なメモリ空間の特性を理解し、適切に活用する必要がある
• ワープと占有率: ワープダイバージェンスを最小化し、占有率を考慮して実行設定を最適化
• デバッグツール: Compute SanitizerおよびNsight Systems/Computeを使用した体系的なデバッグと最適化

CUDAプログラミングは、単にカーネルを書くだけにとどまらず、GPUハードウェアの特性を深く理解し、それに応じてコードを最適化することが求められます。NVIDIAの公式CUDA Programming Guideを継続的に参照しながら、実際のプロジェクトにこれらの概念を適用することをお勧めします。

参考文献

• CUDA Programming Guide - NVIDIA Official Documentation
• CUDA C++ Programming Guide (Legacy)
• CUDA Programming Guide - Programming Model
• CUDA Programming Guide - Writing CUDA SIMT Kernels
• CUDA Programming Guide - Advanced Kernel Programming
• CUDA Programming Guide - Unified and System Memory
• CUDA Programming Guide - Unified Memory
• CUDA Programming Guide - Compute Capabilities
• CUDA Runtime API - Memory Management
• CUDA GPU Compute Capability - NVIDIA Developer
• CUDA Refresher: The CUDA Programming Model - NVIDIA Technical Blog
• Using Shared Memory in CUDA C/C++ - NVIDIA Technical Blog
• Using CUDA Warp-Level Primitives - NVIDIA Technical Blog
• CUDA-GDB - NVIDIA Developer
• Compute Sanitizer - NVIDIA Documentation
• Nsight Developer Tools - NVIDIA Developer
• NVIDIA Blackwell Tuning Guide
• NVIDIA Ampere GPU Architecture Tuning Guide
• NVIDIA Ada GPU Architecture Tuning Guide
• CUDA Samples - Matrix Multiplication (GitHub)

クイズ