CUDA
訪問者数 10880 最終更新 2008-10-25 (土) 10:36:24
＞ ＧＰＵ
CUDA Programming Guide Version 2.0 †
以下すべては、山内研の学生の勉強のため、原文（CUDA Programming Guide Version 2.0） を読んで、要約したものです。
原文の著作権はNVIDIA社にあります。
Chapter 2 Programming Model †
- CUDAはC言語の拡張
- kernel: Cの関数で、呼び出されると、N個の異なるスレッド(CUDAスレッド）上で、N回並列に実行される
- kernelを�＼觚世垢襪箸�、__global__で指定し、��呼出す時に何個の並列スレッドを走らせるかを、<<< ... >>>記法で指定する。
__global__ void vecAdd(float* A, float* B, float* C)
{
int i = threadIdx.x;
C[i] = A[i] + B[i];
}
int main()
{
// Kernel invocation
vecAdd<<<1, N>>>(A, B, C);
}
ここで、threadIdxは組込みの変数で、それぞれのスレッドに与えられるユニークなthread IDの値である。
関数vecAddの引数A, B, Cがいずれもポインタなのは、多分、配列とポインタの互換を利用しただけのことだろうと思う。
- スレッドの階層
スレッドを１次元・２次元・３次元で使えるように、threadIdxは３要素のベクトルになっている。
たとえば２次元の配列の例：
__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
dim3 dimBlock(N, N);
matAdd<<<1, dimBlock>>>(A, B, C);
}
まず、引数A, B, Cはいずれも２次元でNxNとしてある。iとjはthreadIdxのベクトルの
x要素とy要素。
mainの側は、dim3型のdimBlockを宣言しておき、スレッドの数は<<<1. dimBlock>>>と
いている。
dim3型はこのような次元を指定するのに用いられる、整数ベクトルの型で、指定されなかった要素は１に初期化される（4.3.1.12、34ページ）
- 多次元のスレッドIDは単純に、２次元：(x, y)は(x + y*Dx)、３次元：(x, y, z)は(x + y*Dx + z*Dx*Dy)で計算される。但しDx, Dy, Dzは各次元での大きさ。
- ブロック内のスレッドは、共有メモリによるデータ共有と、同期機構__syncthreads()を使って協調動作する。
__synctrehds()は全てのスレッドをその点で同期させるバリヤの働きをする。
- 効率的な実行のためには、共有メモリについては遅延が少ないこと、だからプロセッサコアの近くに置かれたＬ１キャッシュのような振る舞いが望まれるし、__syncthreads()についてはは軽量であること、だから同一ブロックのスレッドは同じプロセッサコア内にあることが望まれる。これによって、１ブロック内のスレッドの個数はプロセッサコアのメモリ資源によって制限される。NVIDIA Teslaのアーキテクチャでは512スレッドが上限になる。
- しかしながら、カーネルは、同じ形をした複数の（スレッド）ブロック上で実行することができ、スレッド総数は（ブロック当たりのスレッド数）×（ブロック数）にできる。個の複数ブロックは、１次元又は２次元の「グリッド」の形におくことができる。
グリッドの数は、<<< ... >>>の第１項目に書くことで指定できる。（今までの例では１だった）
グリッド内のブロックは、組込み変数blockIdxで示されるインデックス番号（１次元もしくは２次元）で表される。
__global__ void matAdd(float A[N][N], float B[N][N], float C[N][N])
{
int i = blockIdx.x * blockDim.x + threadIdx.x;
int j = blockIdx.y * blockDim.y + threadIdx.y;
if (i < N && j < N)
C[i][j] = A[i][j] + B[i][j];
}
int main()
{
// Kernel invocation
dim3 dimBlock(16, 16);
dim3 dimGrid( (N + dimBlock.x - 1) / dimBlock.x,
(N + dimBlock.y - 1) / dimBlock.y );
matAdd<<<dimGrid, dimBlock>>>(A, B, C);
}
ここで、組込み変数blockDimは、ブロックの大きさを示している。
また、個々のブロックのブロックサイズ16x16=256は適当に選んであって、グリッド内のブロック数は行列全体NxNの要素１つ１つがスレッドに割当たるように計算している。
- メモリの階層
- CUDAのメモリ構造は、３種類の空間（スレッドごとのローカルメモリ、ブロックごとの共有メモリ、全体のグローバルメモリ）からなる。
スレッドブロックの共有メモリは、ブロック内の全てのスレッドから参照可能であり、ブロックのライフタイムと同じライフタイムを持つ。グローバルメモリは存在する全てのスレッドからアクセス可能である。
- これに加えて、全てのスレッドからアクセス可能な、読出し専用のメモリが２種類存在する。定数メモリ空間とテクスチャメモリ空間である。それぞれ異なった使い方に最適化されている。また、テクスチャメモリは、データフィルタリングとともに、ある特定のデータ形式のためのいくつかの異なるアドレシングモードを備えている。
- グローバルメモリ空間・定数メモリ空間・テクスチャメモリ空間は、いずれも、同じアプリケーション内でのカーネル起動に対して持続的（カーネルが終了し次に別のカーネルが起動されたときでも、変わらない情報がアクセスできる）である。
- ホスト（要するにCPU）とデバイス（要するにGPU）
- CUDAのプログラム実行は、ホスト（CPU）とデバイス（GPU）という異なる２箇所で行われる。実行の様子の例は以下の図。
- このとき、メモリもホストメモリとデバイスメモリの２箇所に置かれることになる。
プログラムは、グローバル・定数・テクスチャメモリをカーネルに見えるようにするため、CUDAランタイムを呼び出す。ランタイムでは、デバイスメモリの確保と解放や、
ホスト・デバイスメモリ間の転送を行う。
- ソフトウェアスタック
- CUDAのソフトウェアスタックは、図にあるように、デバイスドライバ、ＡＰＩとそのためのランタイム、及び２つのよく使われる高レベル数学ライブラリ（CUFFTとCUBLAS、いずれも別のマニュアルに解説されている）から成っている。
- （デバイスの）Compute Capability 〜〜 省略
Chapter 3 GPU Implementation †
- (複数の）オンチップ共有メモリを持つSIMT (Single-Instruction Multiple-Thread) マルチプロセッサ
- Teslaアーキテクチャでは、ホストCPUでCUDAプログラムがカーネルを起動すると、グリッド内のブロックが数えられマルチプロセッサに分配される。ブロック内のスレッドは１つのマルチプロセッサ内で実行される。１つのブロックが終了すると、新しいブロックが空いたマルチプロセッサ上で起動される。
- １つのマルチプロセッサは、８つのスカラープロセッサ(SP)コア、２つの特殊な関数ユニット（超越関数のための）、多スレッド命令ユニット、オンチップ共有メモリから成る。マルチプロセッサは並列スレッドを生成し、管理し、スケジューリングオーバーヘッドなしにハードウェア上で実行する。また、__syncthreads()バリア同期を内部機能として１命令で実行できる。高速なバリア同期、軽量なスレッド生成、オーバーヘッドなしのスレッドスケジューリングはともに、非常に細粒度の並列を実現しており、たとえば、２次元画像のピクセル、３次元のボクセル、グリッドコンピューティングのセルなどの、個々のデータを１つのスレッドに割り当てるといった非常に細かい粒度の分割を可能にしている。
- 数個の異なるプログラムを実行している数百のスレッドを管理するため、マルチプロセッサは新しいSIMT (single-instruction, multiple-thread)アーキテクチャを採用している。マルチプロセッサは、それぞれのスレッドを１つのスカラプロセッサコアに割当て、個々のスカラスレッドはそれぞれ個別の命令アドレスとレジスタを持って独立に実行する。マルチプロセッサのSIMTユニットは、３２個の並列スレッドをグループとして（これをwarpと呼ぶ）、生成し、スケジュールし、実行する。１つのSIMT warpを構成するそれぞれのスレッドは、同じプログラムアドレスから開始するが、それ以外は自由に分岐し、独立に実行する。
- １つのマルチプロセッサが１つかそれ以上のブロックを実行するときは、SIMTユニットでスケジュールされるwarpに分割する。ブロックがwarpに分割されるやりかたは、いつも同じで、それぞれのwarpは連続したスレッドIDを持つスレッドを持つ（最初のwarpはスレッド0番を持つ）。
- 命令が発行されるときは、SIMTユニットは実行準備のできたwarpを選択し、warpのアクティブなスレッドに対してその命令を発行する。１つのwarpは１つの共通命令を１つずつ実行する。だからwarp上の３２個全てのスレッドが同じexecution pathを実行するとき、フルの効率が得られる。もしwarp上のスレッドがデータに依存する条件分岐で異なる実行パスを取る場合、warpはそれぞれの分岐パスを１つ１つ直列に（そのパスに乗らないスレッドの実行を停止しながら）実行する。その後、全てのパスが終わった段階で、スレッドは同じ実行パスに集合する。このような分岐による分裂は同じwarp内にのみ起こる。warpが異なれば、命令パスが共通か異なるかにかかわらず、実行は独立である。
- SIMTアーキテクチャはSIMD(Single-Instruction, Multiple-Data)のベクトル型の構成に似ている。大事な違いは、SIMDではSIMDの幅がソフトウェアに見えているのに対して、SIMTでは命令は１つのスレッドの実行と分岐の振舞いを指定しているに過ぎない。SIMTではプログラマが、独立したスカラースレッドからなるスレッド並列コードも、協調して動くスレッドからなるデータ並列のコードも、書くことができる。プログラムの論理的な正しさという点からはプログラマはSIMTの振舞いを無視することができる一方で、実行性能という点からは滅多に異なる分岐をしないようにコードを作成することによって性能を桁違いに向上できる。これはキャッシュの役割と同じように考えることが出来る。つまり、キャッシュはプログラムの論理的正しさという点からは存在を無視できるが、最高性能を得ようとデザインするときにはコードの構造を考えなければならない。ちなみに、ベクトルアーキテクチャの場合は、ソフトウェアがベクトルのロードをしたり分岐の異なる場合の管理を自分で（ソフトウェアで）しなければならない。
- 図のように、１つのマルチプロセッサは４種類のオンチップメモリを持っている
- ３２ビットのローカルレジスタを１プロセッサに付き１セット
- パラレルデータキャッシュ（共有メモリ）。全てのスカラプロセッサコアに共有されている。共有メモリ空間が置かれる場所。
- リードオンリー定数キャッシュ。全てのスカラプロセッサコアに共有されている。定数メモリ空間（デバイスメモリのリードオンリー領域）の読出しを高速化する。
- リードオンリーテクスチャキャッシュ。全てのスカラプロセッサコアに共有されている。テクスチャメモリ空間（デバイスメモリのリードオンリー領域）の読出しを高速化する。それぞれのマルチプロセッサは、テクスチャキャッシュを「テクスチャユニット」を経由してアクセスする。テクスチャユニットはいろいろなアドレシングモードとデータフィルタリングを実現している。
- ローカルとグローバルのメモリ空間は、読み書き可能の領域で、キャッシュされていない。
- １つのマルチプロセッサがいくつのブロックをプロセスできるかは、スレッドあたりに必要なレジスタの個数と、ブロックごとに必要な共有メモリの量に依存する。なぜなら、マルチプロセッサのレジスタと共有メモリは、ブロックのバッチで動くスレッド全てに分けられるからである。もし、少なくとも１つのブロックを実行するのにもレジスタや共有メモリが足りないのであれば、そのカーネルは実行できない。１つのマルチプロセッサは最大８ブロックを同時に実行できる。
- もし１つのwarp内でnon-atomicな命令が、グローバル化共有メモリの同じ場所に書き込みをした場合、直列化された書き込みの回数やどの順番で書かれたかはundefinedであるが、どれかの書き込みは成功することが保証されている。atomicな命令がグローバルメモリに同時に読出し、変更し、又は書きこんだ場合は、それぞれの操作は「起こり」、かつすべて直列化されるが、その順番はundefinedである。
Chapter 4 API †
![[Toho University]](image/CLogo_Kihon_process.jpg "[Toho University]")
![[YYLAB]](image/yamanouc2.png "[YYLAB]")
cuda_hardware_model.png 1518件
[詳細]
cuda_host_and_device.png 1533件
[詳細]
cuda_grid_and_thread.png 1535件
[詳細]
