GPUチップの設計方法
Chapter 2 Gpu Rogramming Models

第2章: GPUプログラミングモデル

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

この章では、GPUのための並列プログラミングモデルの主要な概念と原理を探求します。特に、実行モデル、GPU命令セットアーキテクチャ (ISA)、NVIDIA GPU ISA、AMD のGraphics Core Next (GCN) ISAについて焦点を当てます。また、これらの概念がどのように実践されるかを示す例も提供します。

実行モデル

現代のGPUプログラミングモデルの実行モデルは、カーネルの概念に基づいています。カーネルとは、GPUの上で並列に実行される大量のスレッドによって実行される関数です。カーネルを起動する際、プログラマはスレッドの数と、それらがグリッド、ブロック (または協調スレッド配列 - CTA)、個々のスレッドといった階層に編成される方法を指定します。

  • グリッドは全体の問題領域を表し、1つ以上のブロックで構成されます。
  • ブロックは、共有メモリとバリアを通じて協調および同期できる一群のスレッドです。ブロック内のスレッドは、同じGPUコア (ストリーミングマルチプロセッサまたはコンピューティングユニット) で実行されます。
  • 各スレッドには、そのブロックとグリッド内での固有のIDがあり、これを使ってメモリアドレスの計算や制御フロー判断を行うことができます。

この階層的な組織化により、プログラマはデータ並列性 (同じ操作が複数のデータ要素に適用される) とタスク並列性 (異なるタスクが並行して実行される) の両方を表現できます。以下は、提供されたMarkdownファイルの日本語訳です。コードの部分は翻訳せず、コメントのみ翻訳しました。

Figure 2.1: GPUの実行モデルのスレッド階層

            Grid
       ________________
      /   /   /   /   /
     /   /   /   /   / 
    /   /   /   /   /
   /   /   /   /   /
  /__/__/__/__/__/
    |   |   |   |
    |   | Block |
    |   |   |   |
  Thread Thread ...

Figure 2.1: GPUの実行モデルのスレッド階層

SIMTの実行

CUDA 及びOpenCLのようなGPUプログラミングモデルは、Single-Instruction, Multiple-Thread (SIMT)実行モデルに従います。SIMTモデルでは、スレッドはワープ (NVIDIA用語) またはウェーブフロント (AMD用語)と呼ばれるグループで実行されます。ワープ内のすべてのスレッドは同じ命令を同時に実行しますが、それぞれのスレッドは異なるデータを扱います。

ただし、すべての処理要素が同期して実行される従来のSingle-Instruction, Multiple-Data (SIMD)モデとは異なり、SIMTではスレッドが独立した実行経路を持ち、分岐命令で分岐することができます。ワープが分岐命令に遭遇すると、GPUハードウェアはワープ内のそれぞれのスレッドの分岐条件を評価します。すべてのスレッドが同じ経路を辿る (収束する) 場合は、ワープは通常どおり実行を続けます。一部のスレッドが異なる経路を取る (発散する) 場合は、ワープは2つ以上のサブワープに分割され、それぞれが異なる経路を辿ります。GPUハードウェアは、発散した経路の実行をシリアル化し、各サブワープの非アクティブなスレッドをマスクします。すべての経路の実行が完了すると、サブワープは再び同期して、ロックステップでの実行を再開します。

Figure 2.2は、発散したコントロールフローでのSIMTの実行を示しています。

         Warp
   ________________
  /   /   /   /   /
 /   /   /   /   / 
/   /   /   /   /
   |   |   |
   | Branch |
   |   |   |
  / \ / \ / \
 /   X   \   \
/   / \   \   \
   /   \   \
  /     \   \
 /       \   \
/         \   \
           \
            \
             \
   Reconvergence

Figure 2.2: 発散したコントロールフローでのSIMTの実行

この発散ハンドリングメカニズムにより、SIMTはより柔軟なコントロールフローをサポートできるようになります。以下は上記のマークダウンファイルの日本語訳です。コードの部分については、コメントのみ日本語に翻訳しています。

ワープ内の発散を最小限に抑えることで、最適なパフォーマンスを達成することができます。

メモリ階層

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

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

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

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

      ____________
     |            |
     |   Global   |
     |   Memory   |
      ____________
           |
      ____________
     |            |
     |  Constant  |
     |   Memory   |
      ____________
           |
      ____________
     |            |
     |  Texture   |
     |   Memory   |
      ____________
           |
           |
      ____________
     |            |
     |   Shared   |
     |   Memory   |
      ____________
           |
      ____________ 
     |            |
     |   Local    |
     |   Memory   |
      ____________

図以下は、上記のマークダウンファイルの日本語翻訳です。コードについては、コメントのみ翻訳しています。

GPUインストラクション・セット・アーキテクチャ

GPUインストラクション・セット・アーキテクチャ(ISA)は、ソフトウェアとハードウェアの間の低レベルのインターフェースを定義します。命令、レジスタ、メモリアドレッシングモードなどをサポートしています。GPU ISAを理解することは、効率的なGPUコードを開発し、パフォーマンスを最適化するために不可欠です。

このセクションでは、主要なGPUベンダーであるNVIDIAとAMDのISAを探っていきます。NVIDIAのParallel Thread Execution(PTX)とSASSISA、およびAMDのGraphics Core Next(GCN)ISAに焦点を当てます。

NVIDIA GPU ISAs

NVIDIA GPUは、PTX(Parallel Thread Execution)とSASS(Streaming ASSembler)の2つのレベルのISAをサポートしています。PTXは仮想ISAで、CUDA コンパイラの安定したターゲットを提供し、SASSはNVIDIA GPUの native ISAです。

PTX (Parallel Thread Execution)

PTXは、NVIDIA GPUのための低レベルの仮想ISAです。LLVM IRやJavaバイトコードに似ており、コンパイラ用の安定したアーキテクチャ独立のターゲットを提供します。CUDA プログラムは通常、PTXコードにコンパイルされ、その後NVIDIAのGPUドライバによってネイティブのSASS命令に変換されます。

PTXは、算術、メモリ、制御フロー命令を幅広くサポートしています。無制限の仮想レジスタを持ち、述語を使った制御フローの効率的な実装をサポートしています。また、スレッド同期、アトミック操作、テクスチャサンプリングのための特別な命令も提供しています。

以下は、単純なベクトル加算カーネルのPTXコードの例です:

// PTXコードの例
.version 7.0
.target sm_70
.address_size 64

.visible .entry vecAdd(
    .param .u64 vecAdd_param_0,
    .param .u64 vecAdd_param_1,
    .param .u64 vecAdd_param_2,
    .param .u32 vecAdd_param_3
)
{
    .reg .b32 %r<4>;
    .reg .b64 %rd<6>;

    ld.param.u64 %rd1, [vecAdd_param_0]; // パラメータレジスタから値をロード
    ld.param.u64 %rd2, [vecAdd_param_1];
    ld.param.u64 %rd3, [vecAdd_param_2];
    ld.param.u32 %r1, [vecAdd_param_3];
    cvta.to.global.u64 %rd4, %rd1; // グローバルメモリへのアドレスに変換
    cvta
```Here is the Japanese translation of the provided markdown file, with the code sections left untranslated:

このPTXコードは、vecAddと呼ばれるカーネル関数を定義しています。このカーネルは4つのパラメータを受け取ります: 入力ベクトルと出力ベクトルのポインタ、およびベクトルのサイズ。このカーネルは、グローバルスレッドIDを計算し、対応する入力ベクトル要素をロードし、加算を実行し、その結果を出力ベクトルに格納します。

SASS (Streaming ASSembler)

SASSは、NVIDIAGPUのネイティブISAです。これは低レベルで機械に依存するISAで、GPUハードウェアに直接マッピングされます。SASS命令は、NVIDIAのGPUドライバーによってPTXコードから生成され、通常はプログラマーには見えません。

SASSの命令は、メモリ帯域幅とインストラクションキャッシュの容量を削減するために、コンパクトな形式でエンコードされています。レジスタ、即値、メモリアクセスのための様々なアドレッシングモードなど、さまざまな演算子タイプをサポートしています。

ベクトル加算カーネルのSASS命令の例を以下に示します:

code_version_number 90
                     @P0 LDG.E R2, [R8];
                     @P0 LDG.E R0, [R10];
                     @P0 FADD R0, R2, R0;
                     @P0 STG.E [R12], R0;
                         EXIT;

このSASS命令は、先ほどのPTXコードに対応しています。グローバルメモリからの入力ベクトル要素の読み込み (LDG.E)、加算の実行 (FADD)、結果のグローバルメモリへの書き込み (STG.E)、そしてカーネルの終了 (EXIT) を行っています。

AMDグラフィックスコアネクストISA

AMDのGPUはグラフィックスコアネクスト(GCN)アーキテクチャとISAを使用しています。GCNは、グラフィックスやコンピューティングのワークロードをサポートするRISCベースのISAです。高パフォーマンス、スケーラビリティ、低消費電力を目指して設計されています。

GCNは以下のような主要な機能を導入しています:


効率的なスカラー演算とフロー制御の実行のための ALAR ALU。

- 並列データ並列演算の実行のためのベクトルALU。
- アトミック演算とシェアードメモリへの低遅延アクセスをサポートする高帯域幅メモリシステム。
- ベース+オフセットとスカラー+ベクターアドレッシングをサポートする柔軟なメモリ演算アドレッシングモード。

GCN ISAコードのベクトル加算カーネルの例:

```asm
.text
.globl vecAdd
.p2align 2

.type vecAdd,@function
vecAdd:
    # DPTRレジスタを0に設定
    .set DPTR, 0

    # カーネル引数をメモリからロード
    s_load_dwordx4 s[0:3], s[4:5], 0x0
    s_load_dword s4, s[4:5], 0x10
    s_waitcnt lgkmcnt(0)

    # vメモリアドレスを計算
    v_lshlrev_b32 v0, 2, v0
    v_add_u32 v1, vcc, s1, v0
    v_mov_b32 v3, s3
    v_addc_u32 v2, vcc, s2, v3, vcc
    flat_load_dword v1, v[1:2]

    v_add_u32 v3, vcc, s0, v0
    v_mov_b32 v5, s3
    v_addc_u32 v4, vcc, s2, v5, vcc
    flat_load_dword v0, v[3:4]

    # ベクトル加算を実行
    v_add_f32 v0, v0, v1
    flat_store_dword v[3:4], v0
    s_endpgm

この GCN コードは、flat_load_dwordを使ってインプットベクトル要素をロードし、v_add_f32を使って加算を実行し、flat_store_dwordを使ってメモリに結果を書き戻しています。s_load_dwordx4とs_load_dword命令は、カーネル引数をメモリからロードするために使われています。

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

GPU アーキテクチャにアルゴリズムを効率的にマッピングすることは、高性能を達成するために重要です。考慮すべき主なポイントは以下の通りです。

十分な並列性の露出

アルゴリズムは、GPU の並列処理能力を最大限に活用するために、多数の細粒度スレッドに分解されるべきです。これには、異なるデータ要素で独立に実行できるデータ並列部分を特定することが含まれます。

ブランチ発散の最小化

ワープ/ウェーブフロント内の発散制御フローは、シリアル化とSIMD効率の低下につながります。アルゴリズムは、可能な限りブランチ発散を最小限に抑えるように構造化されるべきです。これは、データ依存の分岐の使用を減らすことで達成できます。### メモリ階層の活用

グローバルメモリへのアクセスは高コストです。アルゴリズムではシェアードメモリやレジスタの使用を最大化し、グローバルメモリアクセスを削減する必要があります。また、同一ワープ内のスレッドが連続したメモリ領域にアクセスできるよう、データレイアウトを最適化する必要があります。適切なメモリ階層の活用により、メモリ待ち時間やバンド幅のボトルネックを大幅に削減できます。

演算とメモリアクセスのバランス

アルゴリズムは、メモリ待ち時間を隠蔽し、高い演算スループットを達成するため、演算演算とメモリ操作の比率が高いことが重要です。これは、データの再利用の最大化、プリフェッチの活用、演算とメモリアクセスの重複などによって実現できます。

ホストとデバイス間のデータ転送の最小化

CPUとGPU間のメモリ転送は低速です。アルゴリズムでは、できるだけGPU上で演算を行い、転送オーバーヘッドを最小限に抑える必要があります。大きなバッチでデータをGPUに転送し、必要な間GPUに保持することが重要です。

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

  • マップ: 各スレッドが同じ演算を異なるデータ要素に適用し、大規模データセットの単純な並列処理を実現します。

  • リダクション: 並列リダクションを使用して、大規模入力データセットから単一の値(合計、最大値など)を効率的に計算します。スレッドが局所的なリダクションを実行し、最終結果を生成します。

  • スキャン: prefix sumとしても知られ、配列内の要素の累積和を計算するために使用されます。効率的な並列スキャンアルゴリズムは、多くのGPU加速アプリケーションの重要な構成要素です。

  • ステンシル: 各スレッドが隣接するデータ要素に基づいて値を計算します。ステンシル計算は、科学シミュレーションや画像処理アプリケーションで一般的に使用されます。以下は、提供されたマークダウンファイルの日本語翻訳です。コードの部分については、コメントのみ翻訳しています。

  • Gather/Scatter: スレッドはグローバルメモリの任意の場所から読み取り(Gather)、または書き込み(Scatter)を行います。効率性のためには、慎重なデータレイアウトとアクセスパターンが必要です。

図3.20は、mapパターンの例を示しています。各スレッドが入力配列の異なる要素に対して関数(例えば平方根)を適用しています。

入力配列:  
               |  |   |   |   |   |   |   |
               v  v   v   v   v   v   v   v
              ______________________________
スレッド:     |    |    |    |    |    |    |    |
             |____|____|____|____|____|____|____|
                |    |    |    |    |    |    |
                v    v    v    v    v    v    v
出力配列:

図3.20: GPUプログラミングにおけるmapパターンの例

結論

CUDA やOpenCLなどのGPUプログラミングモデルは、現代のGPUの並列処理機能をデベロッパーに露出し、幅広いアプリケーションの高速化を可能にしています。これらのプログラミングモデルは、微細な並列ワークロードをGPUハードウェアに効率的にマッピングするための抽象化を提供しています。

高性能なGPUコードを記述するためには、実行モデル、メモリ階層、同期プリミティブなどのプログラミングモデルの理解が不可欠です。スレッド構成、分岐分散、メモリアクセスパターン、アルゴリズム設計などの要因を慎重に検討する必要があります。

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

Chapter 1 Introduction to Gpu Chip DesignChapter 3 Parallel Programming Models