AMDとIntelのGPUとIntelのXeon Phiの実行方式

AMDも以前はSIMD方式を使っていたが、2011年12月に発表された 「GCN(Graphic Core Next)アーキテクチャ」のGPU(最初の製品はHD7000)からはSIMT方式に変わった。なお、AMDは64スレッドがひとまとまりであり、これをWave Front(波頭)と呼んでいる。これも感じの分かる命名である。

GPUに詳しい後藤弘茂氏が執筆された記事では、IntelのCore CPUに内蔵されているIntel HDグラフィックスユニットは、ハードウェア的にはSIMDとSIMTの両方の実行方式をサポートし、頂点処理はSIMD、その他の処理はSIMTを使っていたという。しかし、Skylakeでは頂点処理もSIMTで行うようにして、すべての処理でSIMT実行だけを使うように変更されたとのことである。このように、GPUの実行方式はSIMTが事実上の標準という状況に成りつつある。

なお、IntelのXeon PhiはXeon CPUとの互換性を重視しており、512ビット長のベクトル型のデータを扱うSIMD方式となっている。しかし、データごとに命令の実行の可否を指定するマスクを設けたり、Scatter/Gatherという飛び飛びのアドレスをアクセスする機能を設けたりして、SIMTに近い実行ができるようになっている。

このように動作させるためにはマスクデータを作ってマスクレジスタに書き込んで置く、あるいはScatter/Gatherでアクセスするアドレスのリストを作って置くという準備が必要であり、SIMTと比べると、ひと手間、余計にかかるという感じである。しかし、このようにXeon PhiもSIMT実行に必要なほとんどのハードウェアを持っており、SIMT実行もできるようになっているのかも知れない。

NVIDIAのCUDAによる超マルチスレッドの実行

NVIDIAのKepler GPUでは、SM(Streaming Multiprocessor)と呼ばれる単位が、独立したプログラムを実行できるものであり、CPUで言うコアに対応する。なお、NVIDIAはSMの中の個々の演算ユニットをCUDA Coreと呼び、GK110チップは2880コアというような言い方をするが、これはマーケティングで、実際は15コア(SM)で、各SMは192個の演算ユニットを含んでいるというのが本当である。

Kepler GPUのそれぞれのSMは、64エントリのワープバッファを持ち、最大64Warpを同時に(実際は、サイクルごとに切り替えて)処理することができるようになっている。そして、各ワープは32スレッドを含む。したがって、1個のSMで最大2048スレッドを同時に実行することができる。そして、SMを15個搭載するGK110チップを満員にしようとすると、2,048×15=30,720スレッドを実行させる必要がある。このような多数のスレッドをpthreadsで生成したり、OpenMPで生成したりして動作させるプログラムを作るのは容易ではない。

そこで、NVIDIAが考えたのが、CUDAというプログラミング言語である。CUDA(Compute Unified Device Architecture)は、本来はGPUで科学技術計算を行わせる体系全体を指すものであるが、最近では、CUDA言語の意味で使われる方が多いようである。

CUDA言語は、C言語をベースにSIMT方式の超多数のマルチスレッド実行を記述するために最低限の拡張を行ったと説明されており、主要な拡張は、マルチスレッドの実行を指示する文の追加と変数の定義の拡張である。

C言語の関数呼び出しは、関数名(引数)と書くが、ホストCPUで動作するプログラムからGPUで動作するプログラム(Kernel:カーネルと呼ぶ)を呼び出す場合には関数名<<>>(引数)のように、関数名と(引数)の間に3重のカギかっこで囲んだブロック数(nBlocks)とスレッド数(nThreads)を指定する。CUDAでは変数は最大4要素を含むベクトル変数が使えるように拡張されており、このブロック数とスレッド数として、最大3要素のベクトル変数を指定することができる。nBlocksを3要素のベクトルで指定すると、ブロックは3次元の配列で、ベクトルの3要素は、配列のX方向、Y方向、Z方向の大きさを示す。nThreadsはそれぞれのブロックの中にいくつのスレッドが含まれるかを示すもので、これもnBlocksと同様に3要素のベクトル変数で指定できる。

つまり、スレッドの3次元配列であるブロックを3次元配列にまとめた超多数のスレッド群を、この3重カギかっこの関数呼び出しで一括して生成することができる。この3重カギかっこで呼び出されるブロックの集合をNVIDIAはグリッドと呼んでいる。

なお、この3重のカギかっこの表記は警察や軍隊の山形袖章に似ているのでシェブロン(Chevron)と呼ばれているが、これは通称で、NVIDIAの正式の文書には出てこない。

図3-34 2次元のブロックの、2次元のグリッドを起動する例 (出典:NVIDIA CUDA C Programming Guide)

図3-34の例では、nBlocksとnThreadsは2要素のベクトルで、<<>>全体がグリッドである。この例では、グリッドは3×2の合計6個のブロックからなっており、各ブロックは4×3の12スレッドからなっている。したがって、このグリッド指定では、指定された関数が12×3×2＝72スレッド起動されることになる。ただし、これは紙面に収まるブロック数、スレッド数にした結果で、実際には、図3-34よりずっと大きなnBlocksとnThreadsが使われる。

そして、SIMT実行であるので関数を記述するプログラムは1つであり、72スレッドすべてが同じプログラムを実行する。同じプログラムでどのように処理を分担して並列処理を行うのかのイメージを掴んで戴くためには、次の例を見て戴くのが手っ取り早い。

次の例は、行列の加算を行うMatAddという関数(出典:CUDA C Programming Guide)である。なお、GPUで実行する関数をKernel(カーネル)と呼ぶので、最初にKernel definitionと書かれている。

// Kernel definition 
__global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N]) 
{ 
int i = threadIdx.x; 
int j = threadIdx.y; 
C[i][j] = A[i][j] + B[i][j]; 
} 
int main() 
{ 
    ... 
    // Kernel invocation with one block of N * N * 1 threads 
int numBlocks = 1; 
dim3 threadsPerBlock(N, N); 
MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C); 
    ... 
}

最初のKernel definitionに続く部分がGPUで実行されるプログラム(kernel)の定義で、行列AとBの1要素の加算を行って行列Cに格納するMatAdd関数を定義している。threadIdxは組み込み変数で、threadIdx.xは、そのスレッドの2次元のブロックの中のX方向の位置、threadIdx.yはY方向の位置を与えるようになっている。

最初の__global__はMatAdd関数がGPUで動作する関数であり、ホスト(CPU)側のプログラムから呼び出しが可能であることを示している。

そしてKernel invocationからがCPU側で実行するホストプログラムの定義で、ここではグリッドのブロック数は1で、ブロックのスレッド数はN×NとしてMadAdd関数を呼び出している。

その結果、N2のカーネルスレッドがGPUで起動されるのであるが、それぞれどの位置のスレッドであるかは、各スレッドがthreadIdx.xとthreadIdx.yを読めばわかる。そして、これをiとjとして各スレッドは、C[i][j] = A[i][j] + B[i][j];を計算している。

普通、2次元の行列の加算というと2重ループで記述することを考えるが、GPUで並列処理を行う場合は、各スレッドは行列A、Bの1つの要素の加算しか行わない。それでもN2のすべてのスレッドの動作を考えれば、行列AとBの全要素の加算が実行される。この計算パラダイムの変更は、最初は戸惑うのであるが、慣れれば問題はない。なお、ここでは必要がないのでループの構文は使っていないが、CUDA Cではループが使えないわけではない。