• トップ
  • G-DEPについて
  • ご購入ガイド
  • サポート
  • お問い合わせ

G-DEPトップ  >  第7回 CUDAプログラミングモデル②

第7回 CUDAプログラミングモデル②

第7回 「CUDAプログラミングモデル②」 

<< 第6回   |   目次に戻る  |   第8回 >>  

 第7回はCUDAプログラミングの続きということで、実際のCUDAのC言語拡張について、関数識別子、メモリへのアクセス、同期処理などについての説明を混ぜながら解説していきます。

 

7.1 識別子/Qualifier

 CUDAではホスト側とデバイス側で分けて処理が行われ、デバイス側だけで処理が終わるということはありません。そこで、関数がホストとデバイスのどちらで実行させるものなのか、またメモリ上のどこに変数や配列を置くのか、といったことを指定することが出来る識別子があります。

 7.1.1 ホスト、デバイスで実行する関数の指定

 前回6.2で、CUDAでは動作するハードウェアをCPUとGPUで分けて考え、それぞれをホストとデバイスと呼ぶ、という説明をしました。また、そのデバイス側で実行されるプログラムの事を「カーネル」と呼びます。

 カーネルはC関数に以下の制約がついた特徴を持っています。これらの制約に関して分からない場合は、今は無理して覚える必要はありません。

  • ホストメモリはアクセスできない
  • 戻り値 型はvoid
  • 可変引数("vaargs")は不可
  • 再起処理は出来ない
  • 静的変数は使えない

 C言語ベースであるCUDAプログラミングにおいて処理の単位としては関数になりますが、それらを実行する上で特別な識別子を付けてあげることで、ホストで動く関数あるいはデバイスで動く関数という様に区別することが出来ます。これらの識別子として「__global__」、「__device__」、「__host__」と3つのものがあります。いずれも、「_(アンダーバー)」×2で挟んでいます。


__global__

例)

__global__ void

func<<<…>>>() {…}

ホスト側からのみ呼び出す事が出来、デバイスで実行させるカーネルの指定に使います。必ず戻り値はvoidです。


 __device__

例)

__device__ void

func<<<…>>>() {…}

 デバイス側でのみ呼び出す事が出来、デバイス側で実行されます。ホストコードからは呼び出せません。


__host__

例)

__host__ void

func<<<…>>>() {…}

 
 ホスト側でのみ呼び出す事が出来、ホスト側で実行されます。通常の関数と同じ扱いになります。
__device__と同時に利用することで、ホスト・デバイスの双方で利用できる関数を作成できます。しかし、__global__とは同時に利用することが出来ません。
 
 
 7.1.2 メモリの指定

 CUDAのメモリモデルに合わせた以下のような識別子があります。

__device__
 グローバルメモリ領域に確保されます。全てのスレッドからアクセスすることが出来、ホスト側からは読み書きが可能です。
 
__constant__
 コンスタントメモリ領域に確保されます。__device__同様、全てのスレッドから読み出せます。
 
__shared__
 シェアードメモリ領域に確保されます。スレッドの実行中はブロック単位で確保され、ブロック内のスレッドから読み書きが可能となります。

 

 

7.2 CUDAにおける特別な変数と型

 前回のコラムでも述べましたが、CUDAではC言語を拡張した機能があり、独自の型とその宣言方法があります。

kernelFunc<<<Grid_dim, Block_dim, Sm, Stream>>>(a, b, c);

例のようにカーネルプログラムに対し<<<...>>>で囲まれた指定をすることで様々な設定が行えます。(Grid_dim, Block_dim, Sm, Nstrmはあくまで変数名なので、変更できます。)

それぞれの引数の意味は

  • Grid_dim…dim3型の変数で、グリッド中のブロック数を指定します。
  • Block_dim…dim3型の変数で、ブロック中のスレッド数を指定します。
  • Sm…size_t型の変数で、ブロックあたりに割り当てるシェアードメモリ容量をバイトで指定します。
  • Stream…ストリームの番号を示します。詳細は7.4で述べます。

これらのうち、SmとNstrmの部分に関しては省略が可能です。

 7.2.1 dim3

 ブロックとスレッドは3次元で管理されるため、その数を指定するためには3次元要素が必要となってきます。そこで、3次元ベクトルの整数(int)型の宣言「dim3」を利用します。

例1)

dim3 A(8, 8, 1);

dim3 B = dim3(8, 8, 1);

dim3 C = make_dim3(8, 8, 1);

 異なった値の代入の仕方を示しましたが、いずれも動作は同じです。また、ここでは3つのデータを指定していますが、dim3宣言では指定されない要素には1が入るという機能があります。

例2)

dim3 gd(10, 10);

dim3 bd(16, 16, 16);

kernelFunc<<<gd, bd>>>(x, y, z);

 ここでgdは「グリッド中のブロック数」を示しており、(10, 10, 1)と同じ意味です。グリッドは現状の仕様では2次元配列で管理され、3つめ(Z方向)の要素が不要なため、このように省略することが出来ます。また必要に応じて1次元配列で使用したい場合も同じように省略できます。

 ちなみに上の例2の意味は、kernelFuncという関数が「グリッドあたり10×10のブロック数、ブロックあたり16×16×16のスレッド数」を確保して実行されるという設定を示しています。

 7.2.2 ベクタ型

 CUDAにおいてデータをまとめて利用するためにベクタ型と呼ばれる特別な型が利用できます。
 
次元数
1
2
3
4
char
char1
char2
char3
char4
unsigned char
uchar1
uchar2
uchar3
uchar4
short
short1
short2
short3
short4
unsigned short
ushort1
ushort2
ushort3
ushort4
int
int1
int2
int3
int4
unsigned int
uint1
uint2
uint3
uint4
long
long1
long2
long3
long4
unsigned long
ulong1
ulong2
ulong3
ulong4
longlong
longlong1
longlong2
-
-
float
float1
float2
float3
float4
double
double1
double2
-
-
 
例えば、unsigned char3の変数を宣言したいとすると
uchar3 A = make_uchar3(A, B, C);
のように、「make_型名()」という関数でデータを初期化します。引数はベクタ型の元になっている型となります。

 

7.3 同期処理

 CPUコンピューティングではCUDAコアの数よりも多くのスレッドが動きます。呼び出された関数は並列に動作しても、それは非同期になっており実行の順序がずれています。動作が非同期だとあるスレッドの計算が終わっても別のスレッドの処理が終わっていないということがあるわけです。その実行されるスレッドそれぞれが完全に独立であれば良いのですが、例えばN番目のスレッドで計算された結果をN+1番目のスレッドで使いたいといった場合などでは動作が非同期のままだと上手く計算ができません。

 そこで、CUDAでは「__syncthreads()」という同一ブロック内のスレッドにのみ作用する関数を用意しています。この関数によりある種のバリア(壁)を生成することでブロック内の全てのスレッド処理がそのバリアに到達するまで通過できないようにし、全処理が到達すると再度処理を開始する、という働きを持っています。そうすることで同期をとることができます。

 ただ、この関数は全てのスレッドがバリアまで到達するまで処理を待つ時間がかかってしまうため、多用しすぎるとかえって計算効率を下げてしまいます。必要に応じて、スレッドのどの位置で同期を取るべきかを考えてプログラミングをする必要があります。

 また、例えば異なる性能のGPUデバイスが複数個あったとして、全体で同期を取ってしまうと遅いGPUに対して早いGPUがそれに合わせる形になってしまい、それこそ計算効率を悪くしてしまいますが、同一ブロック内のスレッドにのみ適用されることでそれぞれのデバイスの性能を引き出すことができます。仮に異なるブロック間で同期を取りたい場合は、それぞれのブロックを生成したカーネルを終了させて新たにカーネルを生成することで実現できます。

 

7.4 メモリアクセス

 CPUとGPUは別々にメモリを持っていますが、デバイスメモリはホストコードで管理します。具体的にはデバイスメモリの領域確保および解放、デバイスとホスト間のデータコピーです。

 7.4.1 GPUメモリのアロケーション/領域解放

  • cudaMalloc(void ** pointer, size_t nbytes) …メモリ領域の確保
  • cudaMemset(void * pointer, int value, size_t count) …確保したメモリ領域に値をセット
  • cudaFree(void* pointer) …メモリの解放

 cudaMalloc()の第1引数は割り当てるメモリ領域の先頭アドレスへのポインタで、第2引数は割り当てるメモリ要領をバイト単位で指定します。cudaFree()の引数はメモリ領域の先頭アドレスです。

 7.4.2 データの転送(メモリコピー)

  • cudaMemcpy(void *dst, void *src, size_t nbytes, enum cudaMemcpyKind direction)

 dstは転送先アドレス、srcは転送元アドレス、第3引数は転送する容量(バイト)、第4が転送形式(dstとsrcの場所)を指定します。

転送形式は以下の4つがあります。

  • cudaMemcpyHostToHost ホスト→ホスト(CPU内の移動)
  • cudaMemcpyHostToDevice ホスト→デバイス(CPU→GPU)
  • cudaMemcpyDeviceToHost デバイス→ホスト(GPU→CPU)
  • cudaMemcpyDeviceToDevice デバイス→デバイス(GPU内の移動)

NVIDIAの資料にデータ移動の例がありましたので、参考までに載せておきます。GIFアニメーションになっています。

引用元:「CUDA Programming Basics PartⅠ」 by NVIDIA

 

 今回はまだ基礎的な部分ではありますが、CUDAの拡張部分について説明しました。もっと高度な処理やCUDAで利用できるCUBLASなどといった演算関数については追々紹介していきます。

おそらく実際にコードを見ながらやったほうが身につくのも早いと思いますので、次回は演習を行います。

 

<< 第6回   |   目次に戻る   |   第8回 >>


 

(執筆 G-DEP Associate Research Engineer 東京大学大学院工学系研究科 岡安優)