CUDA by Example

CUDA by Example 汎用GPUプログラミング入門

自宅マシンにはGeForce GT520が刺さっていて、CUDAが使えるのだが、ディスプレイ出力と同時に使用すると性能が悪かったりエラーになったり色々問題があるようだ。

ざっと目を通したまとめ。

Hostから呼ばれ、GPUで動作する関数には__global__をつける。
GPUから呼ばれ、GPUで動作する関数には__device__をつける。
__global__の関数は関数ポインタが取得でき、関数テンプレートも使えるが、__device__関数はポインタがとれない。
Linux kernel的なポリモーフィズムはGPUでは無理。
→追記　4.0からC++サポートでvirtual __device__な関数も使えるらしい。

cudaGetDeviceProperties()でGPUの情報を取得できる。

Device 0: "GeForce GT 520"
  CUDA Driver Version / Runtime Version          4.0 / 4.0
  CUDA Capability Major/Minor version number:    2.1
  Total amount of global memory:                 1023 MBytes (1072889856 bytes)
  ( 1) Multiprocessors x (48) CUDA Cores/MP:     48 CUDA Cores
  GPU Clock Speed:                               1.62 GHz
  Memory Clock rate:                             535.00 Mhz
  Memory Bus Width:                              64-bit
  L2 Cache Size:                                 65536 bytes
... 

threadIdx.x、blockIdx.xなどで自分のブロック、スレッドを知ることができ、blockDim.xなどを使って自分のスレッドの担当部分を判断する。

int tid = threadIdx.x + blockIdx.x * blockDim.x;
c[tid] = a[tid] + b[tid];
...

__shared__メモリはブロック内で共有。
__constant__メモリはHost側からのみ、書き込み可。

cudaEventRecord()、cudaEventElapsedTime()などでGPUの処理時間を計測できる。

ストリームでメモリ転送・カーネル関数などを非同期化できる。

複数GPUがある環境ではcudaSetDevice()で切り替え。