CUDA技術を利用したGPUコンピューティングの実際(後編) ―― FFTを利用した光波の伝播(フレネル回折)をGPUで高速計算
● ホスト側の処理の流れを見る
まず,リスト1の(1)はGPUの初期化を行う部分です.初期化には,ホストに接続されているCUDA対応のGPUの個数を調べたり,認識できたGPUの特性(ブロック当たりのレジスタ数や搭載メモリ容量など)を取得したりするステップを踏む必要があります.この手順を一まとめにしたマクロCUT_DEVICE_INITがヘッダ・ファイルcutil.hに用意されています.このマクロを使用することで,手順を簡略化できます(ただし,実際にはこのマクロを省いても動作する).
次に(2)の部分で,GPU上のグローバル・メモリに行列データを格納する領域を確保します.GPUはこの確保されたメモリ中のデータに対して演算を行うことになります.GPU上のグローバル・メモリを確保するにはCUDAのAPIであるcudaMallocを使用します(図1).このAPIを使用することで1次元から多次元までのメモリ領域を確保できます.ほかにもメモリ確保用のAPIが用意されていますが,本稿ではこのAPIのみを使用します.
図1 cudaMallocの書式
GPU上のグローバル・メモリを確保する際に利用する.
このプログラムでは,WIDTH×WIDTHサイズの行列Aと行列BをGPU上で乗算し,その結果をGPU上のメモリCという領域に格納するため,float型でWIDTH×WIDTHサイズの配列をそれぞれd_a,d_b,d_cに確保します.
(3)の部分で,行列Aと行列Bのデータをホスト上に用意します.このデータはホスト側の配列h_a,h_bに格納します.なお,h_cはGPUが計算した乗算結果をホストが回収するための配列になっています.
(4)の部分で,ホスト側で用意したデータを,先ほどGPU上に確保したメモリ領域(グローバル・メモリ)へコピーします.コピーの作業にはCUDAのAPIであるcudaMemcpy(図2)を使用します.ここでは,行列Aのデータh_a(ホスト側)をd_a(GPU側)へ,行列Bのデータh_b(ホスト側)をd_b(GPU側)へ,float型でWIDTH×WIDTH個コピーします.
図2 cudaMemcpyの書式
ホスト側で用意したデータをGPU上に確保したメモリ領域(グローバル・メモリ)へコピーする際に利用する.
(5)の部分で,実際にGPUで行列の乗算を行わせるカーネルを起動します.カーネルを起動する際に,そのカーネルをどの程度のブロック数とスレッド数で実行するかを指定しなければなりません.カーネル起動の書式を図3に示します.
図3 カーネル起動の書式
GPUで行列の乗算を行わせるカーネルを起動する際に利用する.カーネルをどの程度のブロック数とスレッド数で実行するかを指定する.