CUDAを使ったGPUプログラミング超入門 #GPU

CUDAを使ったプログラミングに触れる機会があるため、下記、ざっと学んだことを記します。

細かいところは端折って、ざっとCUDAを使ったGPUプログラミングがどういったものを理解します。

GPUとは

Graphics Processing Unitの略で、描画処理用のCPUのこと
AMD社とNVidia社が有名
GPGPUはGeneral Purpose Graphics Processing Unitの略で、GPUを画像処理ではなく他の用途（ex. 暗号解読、音声処理）に使用するときに使用される。

GPUの特徴

CPUより並列処理の長けており演算性能が高い

種類	処理速度
Corei7(4コア)	100G FLOPS
GeForce GTX 285	1T FLOPS

ゼロから始めるGPU Computing

http://www.gdep.jp/page/view/248

一般的に数倍～100倍以上の速度と言われている

補足：なぜ演算性能が高いのか？

GPUでは、CPUに搭載されている分岐予測・投機実行 アウトオブオーダーが簡素化されており、その分、コア数を増やすことができる。

分岐予測・投機実行とは？
条件分岐した際の処理を前もって実行しておくこと。過去の実行履歴から実行される可能性の高い命令を実行しておく。
[イメージ]
----------------------------------------------------
  arr = [1, 3, 5, 7, 8];
  for (elm as arr){  
    if(elm % 2 == 1){
      echo "odd"; //奇数が続くためCPUは事前にこの処理を予測し実行しておく
    }else{
      echo "even"; //偶数が来た時にはoddの結果を破棄してこちらを実行する
    }
  }
----------------------------------------------------
こうすることで命令処理全体の実行時間を短縮できる。
アウトオブオーダーとは？
プログラムに記述された命令の順番に関係なく処理が行われる仕組みのこと。
処理に必要なデータが揃った命令から実行するようにする。
[イメージ]
----------------------------------------------------
x = 1 + 1;
y = x + 1; //→yはxに依存しているためxの計算の後に処理が行われないといけない
z = 1 + 1; //→zは独立して計算できるため、yの処理でxのメモリを並行して行うことができる
----------------------------------------------------
こうすることで命令処理全体の実行時間を短縮できる。

CUDAとは

NVIDIAが提供するGPU向けのC/C++言語の統合開発環境
SPMD方式をとる

SPMD方式・・・同一のプログラムがで並列に実行される仕組みのこと。　
（↔︎MPMD・・・複数のプログラムが協調的に動作する並列処理のこと）

http://h50146.www5.hp.com/solutions/hpc/stc/soft/pdfs/mpi_training.pdf

CUDAの基礎的な用語

用語	意味
ホスト	呼び出し側のこと
デバイス	GPU側のこと
カーネル	デバイス上で実行される処理のこと。ホスト上のソースコードに定義する
スレッド	カーネルを実行する最小単位のプロセスのこと。各々のスレッドは3次元で管理される
ブロック	スレッドのかたまり。
グリッド	ブロックの塊
SM(ストリーミングマルチプロセッサ)	ストリームを管理するプロセッサ
ストリーム	実際にCUDAコアが逐次的に処理実行するための処理の流れ
CUDAコア	実際に処理を行う演算装置
ウォープ	スレッドのかたまり（スレッド単位で見た時、CUDAが処理で切る最大並列数はCUDAコアの数分だけで、その数をウォープという）

CUDAのアーキテクチャ

http://www.nvidia.co.jp/docs/IO/51174/NVIDIA_CUDA_Programming_Guide_1.1_JPN.pdf

補足：

上図はThreadも2次元で管理されているが、実際は3次元である
ブロックは、65535*65535*1の3次元で管理されている
スレッドは5*5*5の3次元で管理されている
最大スレッド実行数は65535 * 65535 * 512

CUDAの仕組み

#include <stdio.h>

__device__ void device_strcpy(char *dst, const char *src) {
    while (*dst++ = *src++);
}

__global__ void kernel(char *A) {
    device_strcpy(A, "Hello, World!");
}

int main() {
   char *d_hello;
   char hello[32];

   cudaMalloc((void**)&d_hello, 32);

   kernel<<<1,1>>>(d_hello);

   cudaMemcpy(hello, d_hello, 32, cudaMemcpyDeviceToHost);

   cudaFree(d_hello);

   puts(hello);
}

基本的な流れとして、

デバイス用のソースを定義する  
↓
デバイスとホストの各々でデータのやり取りのためにポインターを定義する
↓
カーネルを呼ぶ
↓
先に定義したポインターを介してデータのやり取りを行う

だけである。

なお、kernel<<<1,1>>>(d_hello);では、カーネルを実行するときに使用するリソースの数（ブロック数、スレッド数）を定義している。

分かりやすいアニメーション

CUDAを使った開発時の注意点

デバイスはホストメモリにはアクセスできない
カーネルからの戻り値の型はvoid型（※）
可変引数は使用不可
再帰処理は不可
static変数は使えない
（※）cudaMemCpyなどのCUDAのAPIに関しては、cudaError_t型の構造体が返ってくる

CUDAを使ったGPUプログラミング超入門