GPUチップの設計方法
Chapter 3 Parallel Programming Models

第3章: GPUデザインにおける並列プログラミングモデル

グラフィックスプロセッシングユニット(GPU)は、固定機能のグラフィックスアクセラレータから、幅広いアプリケーションを高速化できる高度に並列化可能なプログラマブルコンピューティングエンジンへと進化してきました。プログラマがGPUの大規模な並列性を効果的に活用できるよう、NVIDIA CUDA、OpenCL、DirectComputeなどの並列プログラミングモデルとAPIが開発されています。これらのプログラミングモデルは、GPUハードウェアの低レベルの詳細を隠蔽しつつ、アプリケーションの並列性を表現できる抽象化を提供します。

この章では、GPUのための並列プログラミングモデルの主要な概念と原則を探求します。特に、SIMT(Single Instruction, Multiple Thread)実行モデル、CUDAプログラミングモデルとAPI、OpenCLフレームワークに焦点を当てます。また、高性能と効率を達成するためのアルゴリズムをGPUアーキテクチャにマッピングする手法についても議論します。

SIMT(Single Instruction, Multiple Thread)実行モデル

SIMTモデルは、現代のGPUが大規模な並列性を実現するための基本的なパラダイムです。SIMTモデルでは、多数のスレッドが同じプログラム(カーネルと呼ばれる)を並列に実行しますが、各スレッドは独自のプログラムカウンターを持ち、スレッドIDやオペランドデータに基づいて異なる実行パスを取ることができます。

カーネルとスレッド階層

GPUカーネルは、大量のスレッドによって並列に実行される関数です。カーネルを起動する際、プログラマはスレッドの数と、それらをグリッド、ブロック(または協調スレッド配列 - CTA)、個々のスレッドという階層に組織する方法を指定します。

  • グリッドは全体の問題空間を表し、1つ以上のブロックで構成されます。
  • ブロックは、共有メモリやバリアを使って協調および同期できるスレッドのグループです。ブロック内のスレッドは、同じGPUコア(ストリーミングマルチプロセッサと呼ばれる)で実行されます。以下は、提供されたマークダウンファイルの日本語翻訳です。コードの部分は翻訳せず、コメントのみ翻訳しています。

(スレッドまたはコンピューティングユニット)。

  • 各スレッドには、そのブロックとグリッド内で一意のIDがあり、これを使ってメモリアドレスを計算し、制御フローの決定を行うことができます。

この階層的な組織により、プログラマーは、データ並列性(同じ操作が複数のデータ要素に適用される)とタスク並列性(異なるタスクが並行して実行される)の両方を表現することができます。

図3.1は、SIMTの実行モデルにおけるスレッドの階層を示しています。

            グリッド
       ________________
      /   /   /   /   /
     /   /   /   /   / 
    /   /   /   /   /
   /   /   /   /   /
  /__/__/__/__/__/
    |   |   |   |
    |   | ブロック |
    |   |   |   |
  スレッド スレッド ...

図3.1: SIMTの実行モデルにおけるスレッドの階層

SIMT実行

SIMTの実行モデルでは、各スレッドが同じ命令を実行しますが、異なるデータを操作します。しかし、すべての処理要素が同期して実行されるSIMD(Single Instruction, Multiple Data)とは異なり、SIMTではスレッドが独立した実行パスを持ち、分岐命令で分岐することができます。

ワープ(NVIDIAのGPUでは32スレッド、AMDのGPUでは64スレッドのグループ)が分岐命令に遭遇すると、GPUハードウェアはワープ内の各スレッドの分岐条件を評価します。すべてのスレッドが同じパスを取る(収束する)場合、ワープは通常どおり実行を続けます。しかし、一部のスレッドが異なるパスを取る(分岐する)場合、ワープは2つ以上のサブワープに分割され、それぞれが異なるパスを追跡します。GPUハードウェアは、分岐したパスの実行をシリアル化し、各サブワープの非アクティブなスレッドをマスクします。すべてのパスが完了すると、サブワープは再び同期して実行を続けます。

図3.2は、分岐制御フローを持つSIMTの実行を示しています。

         ワープ
   ________________
  /   /   /   /   /
 /   /   /   /   / 
/   /   /   /   /
   |   |   |
   | 分岐 |
   |   |   |
  / \ / \ / \
 /   X   \   \
/   / \   \   \
   /   \   \
  /     \   \
 /       \   \
/         \   \
           \
```再収束

図3.2: 分岐制御フローを持つSIMT実行

この分岐処理メカニズムにより、SIMTはSIMDよりも柔軟な制御フローをサポートできますが、分岐が発生したときのSIMD効率の低下というコストがかかります。プログラマーは最適なパフォーマンスを得るために、ワープ内の分岐を最小限に抑えるよう努める必要があります。

メモリ階層

GPUには、並列ワークロードの高帯域幅と低レイテンシの要件をサポートするための複雑なメモリ階層があります。メモリ階層は通常以下で構成されます:

  • グローバルメモリ: 最大容量だが最も遅いメモリ領域で、カーネル内のすべてのスレッドからアクセス可能。グローバルメモリは通常、高帯域幅のGDDRやHBMメモリを使って実装されています。
  • シェアードメモリ: ブロック内のすべてのスレッドで共有される高速なオンチップメモリ領域。シェアードメモリはスレッド間の通信とブロック内のデータ共有に使用されます。
  • 定数メモリ: 読み取り専用のメモリ領域で、すべてのスレッドにデータをブロードキャストするために使用されます。
  • テクスチャメモリ: 空間的な局所性に最適化された読み取り専用のメモリ領域で、テクスチャキャッシュを介してアクセスされます。テクスチャメモリはグラフィックスワークロードでより一般的に使用されます。
  • ローカルメモリ: レジスタスピルやサイズの大きなデータ構造に使用される、各スレッドのプライベートメモリ領域。ローカルメモリはグローバルメモリにマッピングされます。

メモリ階層の効果的な活用は、GPUでの高パフォーマンスを達成するために不可欠です。プログラマーはシェアードメモリの使用を最大化し、グローバルメモリへのアクセスを最小限に抑えることで、メモリレイテンシとバンド幅のボトルネックを軽減する必要があります。

図3.3はGPUのメモリ階層を示しています。

|   共有     |
|   メモリ   |
 ____________
      |
 ____________ 
|            |
|   ローカル |
|   メモリ   |
 ____________

図3.3: GPUメモリ階層

## CUDA プログラミングモデルとAPI

CUDA (Compute Unified Device Architecture) は、NVIDIAが開発したGPUを使った一般目的コンピューティングのためのパラレルコンピューティングプラットフォームとプログラミングモデルです。CUDAは、C、C++、Fortranなどの標準プログラミング言語に拡張機能を提供し、プログラマがパラレリズムを表現し、NVIDIA GPUの計算能力を活用できるようにしています。

### CUDAプログラミングモデル

CUDAプログラミングモデルは、カーネルと呼ばれる関数をGPU上の多数のスレッドで並列に実行するという概念に基づいています。プログラマは、起動するスレッドの数とそれらのスレッドブロックへの編成を指定します。

CUDAは、パラレルプログラミングを容易にするための以下のような主要な抽象化概念を導入しています:

- スレッド: CUDAの基本的な実行単位。それぞれのスレッドは独自のプログラムカウンタ、レジスタ、ローカルメモリを持っています。
- ブロック: 協調して同期できるスレッドのグループ。ブロック内のスレッドは同じストリーミングマルチプロセッサ上で実行され、共有メモリを介して通信できます。
- グリッド: 同じカーネルを実行するスレッドブロックの集合体。グリッドは全体の問題領域を表し、1次元、2次元、3次元のいずれかで構成できます。

CUDAはまた、スレッドが自身の位置を識別し、スレッド階層に基づいてメモリアドレスを計算できるようにする組み込み変数(threadIdx、blockIdx、blockDim、gridDim など)も提供しています。

図3.4にCUDAプログラミングモデルを示します。

グリッド

/ / / / / / / / / / / / / / / / / / / / /////__/ | | | | | | ブロック | | | | | スレッド スレッド ...

図3.4: CUDAプログラミングモデル

### CUDAメモリ階層CUDA は、プログラマーにGPUメモリ階層を公開し、データの配置と移動を明示的に制御できるようにしています。CUDAの主なメモリ領域は以下のとおりです:

- グローバルメモリ: カーネル内のすべてのスレッドからアクセス可能で、カーネル起動間で持続します。グローバルメモリは最も遅延が高く、大規模なデータ構造に使用されます。
- シェアードメモリ: ブロック内のすべてのスレッドで共有される高速なオンチップメモリ。スレッド間の通信とブロック内でのデータ共有に使用されます。
- 定数メモリ: 読み取り専用のメモリ領域で、すべてのスレッドにブロードキャストされる読み取り専用データに使用されます。定数メモリはキャッシュされ、低遅延アクセスを提供します。
- テクスチャメモリ: 空間的な局所性に最適化された読み取り専用のメモリ領域で、テクスチャキャッシュを介してアクセスされます。テクスチャメモリは主にグラフィックスワークロードで使用されます。
- ローカルメモリ: レジスタのスピルアウトや大規模なデータ構造に使用される、各スレッドのプライベートメモリ領域。ローカルメモリは通常グローバルメモリにマッピングされます。

プログラマーは、cudaMalloc、cudaMemcpy、cudaFreeなどのCUDAランタイムAPIを使用して、ホスト(CPU)とデバイス(GPU)メモリ間でデータを割り当てて転送できます。

図3.5はCUDAメモリ階層を示しています。

| | | Global | | Memory |

|

| | | Constant | | Memory |

|

| | | Texture | | Memory |

| |

| | | Shared | | Memory |

|

| | | Local | | Memory |

図3.5: CUDAメモリ階層

### CUDA同期とコーディネーション

CUDAは、スレッド間の協力と通信を可能にする同期とコーディネーションのプリミティブを提供しています:

- バリア同期: __syncthreads()s()関数は、ブロック内のすべてのスレッドが同じ地点に到達するまで待機し、その後に処理を続行することを保証するバリアとして機能します。

- 原子演算: CUDAは、他のスレッドの干渉なしに共有メモリまたはグローバルメモリ上で読み取り-修正-書き込み操作を実行できる原子演算(atomicAdd、atomicExchなど)をサポートしています。
- ワープレベルのプリミティブ: CUDAは、ワーフ内での効率的な通信と同期を可能にするワープレベルの組み込み関数(例: __shfl、__ballot)を提供しています。

同期と調整のプリミティブを適切に使用することは、CUDAでの正しく効率的な並列プログラムを記述するために不可欠です。

例3.1は、ベクトル加算を実行する単純なCUDAカーネルを示しています。

```c
__global__ void vectorAdd(int *a, int *b, int *c, int n) {
    // インデックスの計算
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    // 範囲内の場合のみ処理を実行
    if (i < n) {
        c[i] = a[i] + b[i];
    }
}

int main() {
    int *a, *b, *c;
    int n = 1024;
    
    // ホスト上のメモリ確保
    a = (int*)malloc(n * sizeof(int));
    b = (int*)malloc(n * sizeof(int));
    c = (int*)malloc(n * sizeof(int));
    
    // 入力ベクトルの初期化
    for (int i = 0; i < n; i++) {
        a[i] = i;
        b[i] = i * 2;
    }
    
    // デバイス上のメモリ確保
    int *d_a, *d_b, *d_c;
    cudaMalloc(&d_a, n * sizeof(int));
    cudaMalloc(&d_b, n * sizeof(int));
    cudaMalloc(&d_c, n * sizeof(int));
    
    // ホストからデバイスへの入力ベクトルのコピー
    cudaMemcpy(d_a, a, n * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, b, n * sizeof(int), cudaMemcpyHostToDevice);
    
    // カーネルの起動
    int blockSize = 256;
    int numBlocks = (n + blockSize - 1) / blockSize;
    vectorAdd<<<numBlocks,blockSize>>>(d_a, d_b, d_c, n);
    
    // デバイスからホストへの結果ベクトルのコピー
    cudaMemcpy(c, d_c, n * sizeof(int), cudaMemcpyDeviceToHost);
    
    // デバイスメモリの解放
    cudaFree(d_a);
    cudaFree(d_b); 
    cudaFree(d_c);
    
    // ホストメモリの解放
    free(a); 
    free(b);
    free(c);
    
    return 0;
}
// ベクトル加算カーネルを、numBlocks個のブロックと、blockSize個のスレッドで起動します。
// カーネルは、入力ベクトルaとbの要素ごとの加算を行い、その結果をベクトルcに格納します。
// <<<...>>>構文は、カーネルを起動する際のグリッドとブロックの次元を指定するために使用されます。

CUDA ストリームとイベント

CUDAストリームとイベントは、カーネルと memory操作の並行実行と同期のためのメカニズムを提供します:

  • ストリーム: 順序に従って実行される一連の操作(カーネル起動、メモリコピー)。異なるストリームは並行して実行できるため、計算とメモリ転送の重複が可能です。
  • イベント: 特定の操作の完了を記録するマーカー。同期と計測の目的で使用できます。

ストリームとイベントにより、プログラマーはCUDAアプリケーションのパフォーマンスを最適化できます。計算とメモリ転送の重複や、GPUハードウェアの機能を最大限に活用することができます。

例3.2は、CUDAストリームを使ってカーネル実行とメモリ転送を重複させる方法を示しています。

// 2つのストリームを作成
cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);
 
// 非同期にデバイスにデータをコピー
cudaMemcpyAsync(d_a, a, size, cudaMemcpyHostToDevice, stream1);
cudaMemcpyAsync(d_b, b, size, cudaMemcpyHostToDevice, stream2);
 
// 異なるストリームでカーネルを起動
kernelA<<<blocks, threads, 0, stream1>>>(d_a);
kernelB<<<blocks, threads, 0, stream2>>>(d_b);
 
// 非同期にホストにデータをコピーバック
cudaMemcpyAsync(a, d_a, size, cudaMemcpyDeviceToHost, stream1);
cudaMemcpyAsync(b, d_b, size, cudaMemcpyDeviceToHost, stream2);
 
// ストリームを同期
cudaStreamSynchronize(stream1);
cudaStreamSynchronize(stream2);

この例では、2つのCUDAストリームを作成しています。入力データは各ストリームを使って非同期にデバイスにコピーされます。その後、異なるストリームでカーネルが起動され、結果がホストにコピーバックされます。最後に、ストリームの同期が行われます。以下は、提供されたマークダウンファイルの日本語翻訳です。コードについては、コメントのみ翻訳しています。

OpenCLフレームワーク

OpenCL (Open Computing Language) は、CPUやGPU、FPGAなどの異種プラットフォームにわたる並列プログラミングのためのオープンで使用料無料の標準規格です。OpenCLは、開発者が移植性の高い効率的な並列コードを書くことができるようなユニファイドプログラミングモデルとAPIを提供しています。

OpenCLプログラミングモデル

OpenCLプログラミングモデルはCUDAに似ていますが、用語と抽象化にいくつかの違いがあります:

  • カーネル: 大量のワークアイテム(スレッド)によって並列に実行されるOpenCLデバイス上の関数。
  • ワークアイテム: OpenCLにおける基本的な実行単位で、CUDAのスレッドに相当します。
  • ワークグループ: 同期とローカルメモリを介したデータ共有が可能なワークアイテムの集合体。CUDAのスレッドブロックに相当します。
  • NDRange: カーネル実行のためのインデックス空間とワークアイテムの構成を定義します。1次元、2次元、または3次元にできます。

OpenCLはCUDAに似た階層的なメモリモデルも定義しています:

  • グローバルメモリ: すべてのワークグループのすべてのワークアイテムがアクセス可能で、CUDAのグローバルメモリに相当します。
  • ローカルメモリ: 同一ワークグループ内のすべてのワークアイテムが共有できるメモリで、CUDAのシェアードメモリに相当します。
  • プライベートメモリ: 単一のワークアイテムにのみ私有のメモリで、CUDAのレジスタに相当します。
  • 定数メモリ: 読み取り専用で、すべてのワークアイテムがアクセス可能なメモリ。

OpenCLカーネルはOpenCLランタイムによってランタイムにコンパイルされます。ホストプログラムは利用可能なOpenCLデバイスを問い合わせ、適切なデバイスを選択し、コンテキストを作成し、そのデバイス用にカーネルをビルドできます。これにより、OpenCLアプリケーションは異なるハードウェアプラットフォーム間で高い移植性を持つことができます。

例3.3は、例3.1のCUDAの例と同様のベクトル加算を行うOpenCLカーネルを示しています。

// グローバルメモリ上の定数ポインタaとbを受け取り、結果をグローバルメモリ上のcに書き込む
__kernel void vectorAdd(__global const int *a, __global const int *b, __global int *c)
{
    // ワークアイテムのグローバルインデックスを取得
    int gid = get_global_id(0);
 
    // a[gid] + b[gid]の結果をc[gid]に格納
    c[gid] = a[gid] + b[gid];
}
__kernel void vector_add(const __global int *a, const __global int *b, __global int *c, int n) {
    int i = get_global_id(0);
    if (i < n) {
        c[i] = a[i] + b[i];
    }
}

The __kernel キーワードは OpenCL カーネル関数を定義します。__global キーワードは、ポインタがグローバルメモリを指していることを示します。get_global_id 関数は、現在のワークアイテムのグローバルインデックスを返し、入力および出力ベクトルのメモリアドレスを計算するために使用されます。

GPUアーキテクチャへのアルゴリズムのマッピング

アルゴリズムをGPUアーキテクチャに効率的にマッピングすることは、高パフォーマンスを達成するために重要です。主な考慮事項は以下のとおりです:

  • 十分な並列性の露出: アルゴリズムは、GPUの並列処理機能を最大限に活用するために、多くの細粒度のスレッドに分解される必要があります。

  • ブランチ発散の最小化: ワープ/ウェーブフロント内の発散制御フローは、直列化と SIMD 効率の低下につながる可能性があります。アルゴリズムはできる限りブランチ発散を最小限に抑えるように構造化される必要があります。

  • メモリ階層の活用: グローバルメモリへのアクセスは高コストです。アルゴリズムは、グローバルメモリアクセスを削減するために、共有メモリやレジスタの使用を最大化する必要があります。データはまた、コアレスメモリアクセスを可能にするようにメモリ上に配置される必要があります。

  • 計算とメモリアクセスのバランス: アルゴリズムは、メモリ待ち時間を効果的に隠蔽し、高い計算スループットを達成するために、メモリ操作に対する算術演算の比率が高くなるようにする必要があります。

  • ホストデバイス間のデータ転送の最小化: ホストとデバイスメモリ間のデータ転送は遅いです。アルゴリズムは、できる限りGPU上で計算を行うことで、このような転送を最小限に抑える必要があります。

一般的に、GPUカーネルの開発時に使用される並列アルゴリズムのデザインパターンには以下のようなものがあります:

  • マップ: 各スレッドが同じ操作を異なるデータ要素に適用することで、大規模なデータセットの単純な並列処理を可能にします。

  • 縮小: 並列縮小は、大規模な入力データセットから単一の値(例: 合計、最大値)を効率的に計算するために使用されます。スレッドはローカルな縮小を実行し、それらを組み合わせて最終的な結果を生成します。

  • スキャン: プレフィックス和としても知られ、配列の要素の実行中の和を計算するのに使用されます。効率的な並列スキャンアルゴリズムは、多くのGPU加速アプリケーションの重要な構成要素です。

  • ステンシル: 各スレッドは隣接するデータ要素に基づいて値を計算します。ステンシル計算は、科学的なシミュレーションや画像処理アプリケーションでよく見られます。

  • ギャザー/スキャッター: スレッドは任意の場所からデータを読み取り(ギャザー)、または任意の場所にデータを書き込み(スキャッター)ます。効率のためには、慎重なデータレイアウトとアクセスパターンが必要です。

結論

CUDA やOpenCLのようなGPUプログラミングモデルは、開発者にモダンGPUの並列処理機能を公開し、幅広いアプリケーションの高速化を可能にします。これらのプログラミングモデルは、細かい粒度の並列ワークロードをGPUハードウェアに効率的にマッピングできる抽象化を提供します。

GPUコードの高パフォーマンスを書くには、これらのプログラミングモデルが提供する実行モデル、メモリ階層、同期プリミティブの理解が不可欠です。開発者は、スレッド構成、分岐発散、メモリアクセスパターン、アルゴリズム設計などの要因を慎重に検討する必要があります。

GPUアーキテクチャが進化し続けるにつれ、プログラミングモデルとツールも新しいハードウェア機能と機能を効果的に活用できるよう進化していく必要があります。プログラミング言語設計、コンパイラ最適化、オートチューニングなどの分野での継続的な研究が、ヘテロジニアス コンピューティングの時代におけるプログラマ生産性とパフォーマンスポータビリティの向上に不可欠です。

Chapter 2 Gpu Rogramming ModelsChapter 4 Gpu Memory System Design