CUDAを用いてGPUを用いた並列処理を行うとき、GPU側では多数の「スレッド」が生成され、各スレッドがその処理をこなすことで並列化が実現されます。
例えば処理を1024並列で走らせる場合、スレッドが1024個必要となります。
CUDAプログラミングにおいては、この1024並列で処理を実行したい場合に、1024並列で処理を実行したいということをプログラムの中で記述する必要があります。
既にCUDAのプログラミングの経験がある方はご存じの方が多いと思いますが、CUDAのスレッドの構成は階層構造になっており少々複雑で、設定にあたっては単に並列数を指定するだけではなく、「グリッド」「ブロック」「スレッド」というものについて理解をする必要があります。
今回は、CUDAの並列処理を理解する上で欠かせない「グリッド」「ブロック」「スレッド」への理解を深めることを目的に、これらの概要を説明します。
目次
CUDAのスレッドの構成
まずは以下の図をベースに説明します。
CUDAのスレッドの構造は以下の階層構造で抽象化されています。
- グリッド
- ブロック
- スレッド
まず、GPU側で動作する関数が実行されると、多数のスレッドにより並列動作がされるわけですが、このスレッドを所定の数でグループにしたものがブロックです。そして、全てのブロックをひとまとめにしたものがグリッドです。
すなわち、階層構造としては「1個のグリッドの中にブロックがあり」「ブロックの中にスレッドがある」です。
このように、CUDAのプログラミングではブロック32個、（1ブロックあたりのスレッド）32個のようにブロックとスレッドの数を決めて処理します。この例だと32×32＝1024個の並列処理が動きます。
以下のように例えるとわかりやすいでしょう。
- グリッド：学校
- ブロック：クラス
- スレッド：生徒
ある学校に、全学年合わせて30クラスあり、各クラスが50人の生徒から成るとします。合計の生徒数（スレッド数）は1500であり、クラスの数（ブロック数）は30です。
このように、ある集団が階層構造になっているのは、管理などの観点でメリットがあるのは現実世界とも同じと言えるかもしれません。
何故階層構造が必要か
何故、階層構造が必要かということですが、もちろん理由があります。その理由を説明していきましょう。
同一ブロックのスレッドのみがデータを共有できるシェアードメモリの存在
もう一度最初の図に戻ります。この図に登場する「Global Memory」はGPU側でデータを保存しておくときに使う、最もポピュラーなメモリです。
CPU側からもアクセスができるため、CPUからGlobal Memoryにデータを移して、GPU側でGlobal Memory上のデータを処理する場合など、とにかくよく使います。
さて、このグローバルメモリは、GPU側から見た場合には全てのスレッドがアクセスすることができます。
学校の例えで言えば、全ての学生がアクセスできるオープンなSNSだと思ってください。違うクラスの人同士でも、共通に情報をアップロードしたり、閲覧したりすることができるわけです。
さて、CUDAにはグローバルメモリ以外にもメモリが用意されています。グローバルメモリは一般的にメモリへのアクセスが遅く、小さいデータに限定して、より高速にアクセスができる「シェアードメモリ」というメモリが用意されていますが、シェアードメモリは同一のブロックに属するスレッド間でしかデータのやり取りができません。
上の図のように、シェアードメモリはブロックごとに用意され、ブロック内の各スレッドのみがアクセスできます。
すなわち、クラスメイトしかアクセスできない共有空間が存在するわけです。
このように、シェアードメモリも考慮したプログラミングをする場合、スレッドをどのようにブロックに分けるかというのは重要な問題となってくるわけです。
スレッド間の同期
GPU側で、各スレッドの処理の同期を取りたいケースがあります。
デバイス側で同期を取る関数「__syncthreads()」を記述すると、各スレッドは「__syncthreads()」のところで他の全てのスレッドが追いつくまで待ち、再び処理を開始します。
さて、このスレッド間の同期ですが、同じブロックの中のスレッドでしか同期が取れません。
これも学校で例えると、クラス単位で同期は取るけれども、学校や学年全体で同期を取ったりはしないということですね。
まとめ
CUDAのスレッドの階層構造について説明しました。CUDAのプログラミングでは、「ブロック」と「スレッド」について理解して、プログラミングを行う必要があります。