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

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

第6回 CUDAプログラミングモデル①

第6回 「CUDAプログラミングモデル①」 

<< 第5回   |   目次に戻る   |   第7回 >>  

 今まで色々とCUDAを用いたGPUコンピューティングについてお話してきましたが、いまいちCUDAってまだよく分からないという方もいるかもしれません。

 今回は、CUDAがどうGPUと絡んでくるのか、C言語とどう違うのかなど、CUDAプログラミングの概念についてなるべく分かりやすく要点を押さえて解説していきたいと思います。このCUDAプログラミングモデルこそGPUコンピューティングのキモになってきますので、しっかりと勉強するため次回と合わせて2回に分けてお送りします。

それではまずはもう一度CUDAについての話から始めましょう。

6.1 CUDAとは?

 NVIDIAはGPUシリーズを開発・展開していく中で、2006年に発表されたGeForce 8800では、それまでハードの部分でグラフィックス専用の構造だったところを汎用的なプログラムを動かせるような構造にすることで、より自由度を増加させてグラフィックス専用の機能も同時に向上させました。そして、その汎用的なプログラムが動かせるという構造に対して、グラフィックス向けだけでなくより汎用性の高いプログラミングをしてみてはどうかという考えから、CUDAを発表しました。

 つまり、CUDA(Compute Unified Device Architecture)とは、「NVIDIAGPU製品において汎用的なプログラムを動かすためのプラットフォーム・統合開発環境」です。

CUDAの特徴を挙げると、

  • 容易な習得…C言語ベースの拡張言語。CGの知識が不要。専用の関数・変数修飾子とドライバ&API。
  • 自由度の高いメモリアクセス…スレッドは必要なだけのメモリを読み書き。特性の異なる複数種類のメモリ使用可能。
  • 高い並列性…数千~数万単位の軽量スレッドが並列に動作。スレッドスケジューラによる動的スケジューリング。
  • オープン性…クロスプラットフォーム(Windows、Mac OS X、Linux)。様々な言語からの呼び出し。既存のソフトウェアに対するプラグインの提供
  • ロイヤリティフリー…作成したソースコードおよびバイナリの取り扱いが自由。

などがあります。

 

6.2 ホストとデバイス

 GPUコンピューティングではCPUとGPUの両方を用いて計算します。そこで、CPUやメインメモリ側を「ホスト」、GPUやビデオメモリ側を「デバイス」と呼び、区別しています。また、CUDAにおいてCPUに実行させるプログラムを「ホストプログラム(または単にプログラム)」、GPUに実行させるプログラムを「カーネル(またはカーネルプログラム」と呼びます。ホストプログラムで扱うデータはメインメモリに、カーネルで扱うデータはデバイス側のメモリに保存され、やり取りされます。

 CUDAを動作させるためのGPUは複数取り付けることができ、ホストプログラムをマルチスレッドのプログラムにすることにより、同時にそれら複数のデバイスを利用することができます。

プログラムの一連の流れとしては以下のようになります。

  1. プログラムを起動
  2. カーネルがデバイスにロードされる
  3. ホスト側でデータを作成し、値をデバイスに渡す
  4. ホストがデバイスに対し、カーネルを起動するよう命令(このことをキックと呼びます)
  5. デバイスがカーネルを実行し、処理開始
  6. 計算結果をホスト側に渡す
  7. ホストがもらったデータを処理

 ここで大事なのはホスト)⇔デバイス間のメモリ転送速度を考慮することです。前回GPUのメモリ転送速度についてお話しましたが、メインメモリとデバイス側メモリの転送速度はGPU内のものより一段と遅くなります。そのため、CUDAで並列処理させることで短縮できる時間が、メモリ転送にかかる時間よりも大きくなければCUDAを利用する意味がなくなってしまいます。トータルで考えたとき、実際の計算時間はこうした伝送速度も絡んでくることを念頭に置いておいてください。

 

6.3 nvcc(CUDAコードのコンパイル)

 nvccとはCUDAソフトウェアに含まれるコンパイラのことです。各OSのSDKに含まれており、コマンドとして直接利用できますし、WindowsではまたVisual Studioから呼び出して使うことも可能です。

 処理フローは以下のようになります。

  1. CUDAソースコードをnvccにより、ホスト(CPU)用コードとカーネル(GPU)用コードに分割
  2. ホスト用コードをCコンパイラに渡し、カーネル用コードはnvcc自身でPTXコードに変換
  3. PTXコードはさらにGPUハードウェア上でコンパイルされながら実行される
出典:「CUDAプログラミングの基本 パートⅠ」 by NVIDIA

 

6.4 スレッド・ブロック・グリッド

 CUDAプログラムにおいて、このスレッド・ブロック(スレッドブロック)・グリッドという言葉とその概念は非常に重要です。この先幾度も出てくる単語ですので、しっかりその概念を覚えていってください。

では一つずつ見ていきます。

○スレッド
 カーネルを動作させたときの多数のプログラムの最小単位を指します。これはCPUにおいても使われる単語ですが、異なるのはCPUではそのコア数とほぼ同数のスレッドが動作するのに対し、GPUではコアに対し数千~数万と圧倒的な数のスレッドが並列に動作することにより、その高い性能を引き出しているという点です。
 
一歩踏み込んだ話をしますと、スレッド自体はホスト側から起動されて各スレッドプロセッサ(CUDAコア)で同じ処理が行われるわけですが、その実行タイミングがそれぞれ異なります。これはカーネルの呼び出しタイミングが同時ではなく1つあたり1クロックずつずれるためであり、例えば32スレッドを同時に動かすと最初と最後のスレッドで32クロックのずれが生じます。このとき、最初のスレッドが終わった時点で最後のスレッドの処理が終わってないことも考えられます。このことを非同期の動作と言います。
 CUDAプログラミングにおいては多数の処理が並列に行われるという点が重要ですが、その処理は全て非同期であるために、処理が同時に働くという前提で考えてしまうと計算結果の値を再利用する場合などに困ってしまいます。それを解決するために__syncthreads()という関数がありますが、この話はまた後ですることにします。
 
 
○ブロック(スレッドブロック)
 スレッドをまとめたもので、1つのブロック当たり最大512スレッドが格納します。ブロックというだけあってx方向、y方向、z方向に8スレッドずつ、つまり「1ブロックに8×8×8スレッド」と3次元的表現で管理することができます。また、「1ブロックに512スレッド」「1ブロックに16×16スレッド」と1次元的、2次元的にまとめることも可能です。
 
 
○グリッド(カーネルグリッド)
 ブロックをさらにまとめたものがグリッドと呼ばれます。これもブロック同様、xyzの3次元で表されますが、現在z方向のブロック数は1でなければならないため、実質2次元で管理されます。1グリッド当たり、x方向あるいはy方向に配置できる最大ブロック数は65535個で、それを超えて配置させようとするとエラーを起こしてしまいます。
 
ちなみに、さらに上位の概念というのが「デバイス」となりますが、デバイスとグリッドを分ける定義が明確でないため、同等に考えて良いそうです。そのため、複数のCUDA対応GPUが複数搭載されている場合、グリッドもその数だけ切り替わります。

 

なぜこのような階層的かつ次元配列的な概念が必要かというと、CUDAプログラムではGPUのCUDAコアに対し最高で「65535×65535×512」個のスレッドの実行を命令することができますが、このような多数のスレッドに対して1つの整理番で管理するのは賢明ではないためです。

以下のものは、これら実行モデルをソフト側とハード側で対応させた図です。(クリックで拡大します)

出典:「CUDAプログラミングの基本 パートⅠ」 by NVIDIA

 

 

6.5 カーネル関数

 カーネルとはGPUで動くプログラムを指すことは前述しました。CUDAではどれくらいのスレッド数で「カーネル関数」を実行させるか指定することができます。SPMD(Single Program Multiple Data)と呼ばれるプログラミングモデルが採用されているため、実行するプログラムは全てのスレッドで同じです。

どうやってスレッド数を指定するのか、C言語に対する拡張部分はどこかといったことを試しにコードを見ていきましょう。後のコラムで詳しく説明しますが、とりあえずCUDAの基本的なコードとしてどのようなものがあるか程度で理解してもらえれば構いません。

なお、ここに書いてあるコードは便宜上のため、省略している部分や矛盾している部分があります。

__global__ void
kernelFunction(int* inA, int* inB, int* inC)
{
    int x = threadIdx.x;
    int y = threadIdx.y;
    int z = threadIdx.z;
    :
    :
}

int main()
{
    kernelFunction<<<4,512>>>(A,B);
    :
    dim3 grid(10, 10);
    dim3 block(8, 8, 8);
    :
    :
}

 見慣れない部分として__global__(globalの左右はアンダーバー2つ)、threadIdx.x、kernelFunction<<<1,512>>>、dim3といったところが挙げられるのではないでしょうか。演習してみないと理解しにくい部分もありますが、とりあえずここで一通り説明しておきます。

__global__
この関数(ここではkernelFunction)がGPUで動作する関数であるということを宣言しています。
 
threadIdx.x
この変数はカーネルが動作しているスレッドのインデックス番号になります。ここではmain()関数で呼ばれたカーネルが実行されるとき、スレッドにインデックス番号が付き、情報が格納されるのがthreadIdxとなります。また、今回続いて「threadIdx.y」、「threadIdx.z」という変数も書きましたが、これはブロック内のスレッドが3次元で構成されている場合のインデックス番号となります。
 
blockIdx
上のサンプルコードには出てきていませんが、これはthreadIdxの上位の変数であり、グリッド上で動作しているブロックのインデックス番号を指します。ブロックの配列は実質2次元で構成されるため、「blockIdx.x」、「blockIdx.y」にインデックス番号が代入されます。
 
カーネル関数名<<< >>>
main関数の中にあるkernelFunctionの右に<<< >>>というものがありますね。これは<<<グリッド中のブロック数と、ブロック中のスレッド数>>>を指しており、つまりここで4つのブロックを、それぞれ最大512スレッドが起動する、ということを表しています。
 
dim3
これはdim3宣言と呼ばれ、CUDAのブロックやスレッドは3次元構造で管理することができると述べましたが、その数を指定するために3次元ベクトルの整数型の変数の宣言として使われます。
ここでは「dim3 grid(10, 10)」は「グリッドの中に10×10のブロックを起動します(gridは変数名)」、「dim3 block(8, 8, 8)」は「ブロックの中に8×8×8のスレッドを起動します(blockは変数名)」という意味で使っています。よって例えば、上のkernelFunction<<<4,512 >>>をkernelFunction<<<grid,block>>>とすれば、100個のブロック数に、最大512個ずつスレッドを起動させることになります。

 

6.6 ウォープ(Warp)

ここでは青木先生著「はじめてのCUDAプログラミング」という本に習い、意味合いの差別化を図るためにワープではなくウォープと発音しています。

先ほど、すべてのスレッドは並列に動作すると述べましたが、厳密には32スレッドごとに動作します。この数字、どこかで見たことありませんか?そう、前回のコラムでSM当たり32個のCUDAコアが入っているということを学びましたが、これと関わってくるのです。つまり、32個のCUDAコアが1クロックサイクルで1スレッド動くため32スレッドごとが一つの単位となり、これをウォープと呼んでいます。

ここで重要なのが、「スレッド処理数は32の倍数が望ましい」ということです。例えば、20スレッドだと残り12個のCUDAコアが余ってしまいますし、42スレッドだと32スレッドを一度計算した後10スレッドを2サイクル目で処理を行い、そのとき残りの22スレッドは余ることとなりGPUの処理に無駄ができてしまいます。

 

 

ここまで理解して頂けたでしょうか?

やっとCUDAプログラミングの話に入ることができましたが、今回では話しきれないため残りは次回にします。

次回は、メモリアクセスや同期・非同期処理などについて詳しく説明していきます。
 

<謝辞>
本コンテンツにおいて、NVIDIA社より資料の引用を快諾頂きました。ここに記して謝意を表します。

 

 

<< 第5回   |   目次に戻る   |   第7回 >>


 

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