第11章: GPU上のスカラー化とアフィン実行に関する研究の方向性
第2章で説明したように、CUDA やOpenCLなどのGPUコンピューティングAPIは、プログラマがGPU上に大量のスカラースレッドを起動できるMIMD型のプログラミングモデルを提供しています。これらのスカラースレッドは、それぞれ独自の実行パスを辿り、任意のメモリ位置にアクセスできますが、一般的には少数の実行パスを辿り、類似した操作を行います。
ほとんどの現代のGPUでは、SIMT実行モデルを利用してGPUスレッドの収束制御フローを活用しています。ここでスカラースレッドはワープにグループ化され、SIMD ハードウェア上で実行されます(3.1.1節参照)。本章では、スカラー化とアフィン実行を通じてこれらのスカラースレッドの類似性をさらに活用する一連の研究を要約します。
この研究の鍵となる洞察は、同じコンピュートカーネルを実行するスレッド間の値の構造[Kim et al., 2013]に関するものです。一様(uniform)とアフィン(affine)の2種類の値の構造は、例11.1のコンピュートカーネルで示されています。
一様変数
コンピュートカーネル内のすべてのスレッドで同じ定数値を持つ変数。アルゴリズム11.1では、変数a、定数THRESHOLD、Y_MAX_VALUEがすべてのスレッドで一様な値を持っています。一様変数は単一のスカラーレジスタに格納され、カーネル内のすべてのスレッドで再利用できます。
アフィン変数
コンピュートカーネル内のスレッドIDに対して線形関数で表される値を持つ変数。アルゴリズム11.1では、y[idx]の メモリアドレスがスレッドID threadIdx.xのアフィン変換で表現できます:
&(y[idx]) = &(y[0]) + sizeof(int) * threadIdx.x;このアフィン表現は、ベースとストライドの2つのスカラー値として格納できるため、完全に展開したベクトルよりもはるかに効率的です。
__global__ void vsadd( int y[], int a ) {
// コードは変更せず、コメントのみ翻訳
// スレッドIDに基づいてyの要素にaを加算する
}
```以下は、提供されたマークダウンファイルの日本語翻訳です。コードの部分は翻訳せず、コメントのみ翻訳しています。
int idx = threadIdx.x; y[idx] = y[idx] + a; if ( y[idx] > THRESHOLD ) y[idx] = Y_MAX_VALUE; }
アルゴリズム11.1: GPUコンピュートカーネルにおけるスカラーおよびアフィン演算の例 ([Kim et al., 2013]より).
GPUでの均一または線形変数の検出と活用に関する多くの研究提案があります。本章の残りでは、これらの2つの側面における提案を要約します。
## 均一または線形変数の検出
GPUコンピュートカーネルにおける均一または線形変数の存在を検出するには、主に2つのアプローチがあります: コンパイラ主導の検出と、ハードウェアによる検出です。
### コンパイラ主導の検出
GPUコンピュートカーネルにおける均一または線形変数の存在を検出する1つの方法は、特別なコンパイラ解析を行うことです。これは可能です。なぜなら、既存のGPUプログラミングモデルであるCUDAやOpenCLでは、プログラマがコンピュートカーネル全体を通して変数を定数として宣言したり、スレッドIDの特別な変数を提供したりすることができるためです。コンパイラは制御依存性の解析を行い、定数やスレッドIDのみに依存する変数を検出し、それらを均一/線形として印付けすることができます。均一/線形変数のみを操作する命令は、スカラー化の候補となります。
AMD GCN [AMD, 2012]は、コンパイラによる均一変数の検出と、専用のスカラープロセッサによるスカラー演算の活用に依存しています。
Asanovic et al. [2013]は、収束と変異の解析を組み合わせたものを紹介しています。これにより、コンパイラは任意のコンピュートカーネルにおいて、スカラー化やアフィン変換の対象となる演算を特定できます。コンパイルカーネルの収束領域内の命令は、スカラー/アフィン命令に変換できます。収束領域から発散領域への遷移時には、コンパイラは`syncwarp`命令を挿入して、2つの領域間の制御フロー依存レジスタを処理します。Asanovic et al. [2013]は、この手法をGPUアーキテクチャに組み込んでいます。以下は、提供された Markdown ファイルの日本語翻訳です。コードについては、コメントのみ翻訳しています。
Temporal-SIMT アーキテクチャ [Keckler et al., 2011, Krashinsky, 2011] のためのスカラー演算を生成するコンパイラ分析を提供します。
Decoupled Affine Computation (DAC) [Wang and Lin, 2017] は、同様のコンパイラ分析を使用して、スカラーおよびアフィンの候補を抽出し、別のワープにデカップルします。Wang and Lin [2017] は、発散アフィン分析プロセスを拡張し、命令の連続を抽出することを目的としています。