CUDAをちょっといじってみたので、そこで疑問に思ったことや解決したことを適当に書いてみる。
CUDAはNVIDIAが開発した、GPUを汎用用途で使用するためのC言語拡張ライブラリ。
CELLのSPUみたいな感じでGPUを使うのを簡単に行うためのものです。
現在のバージョンは2.0。ここから落とせます。
で、templateサンプルを利用してちょっとお勉強。
まず、grid, block, thread の概念。
GPUでは同じ計算を大量に行う場合に強く、シェーダとかはまさにそれなわけで。
例えば、ピクセルシェーダの計算は1ピクセルずつ行われるわけではなく、複数のピクセルを並列計算します。
この1ピクセルの計算（みたいなもの）を thread といいます。
thread は理由はよくわからないのですが、3次元配列的に作成することが可能です。例えば、100個のスレッドを立てたい場合、100*1*1と1次元配列にしてもいいし、50*2*1と2次元配列にしてもいいし、5*5*4と3次元配列にしてもいいわけです。
で、このように3次元配列的に作成された thread 群を block といいます。
ここでまた、block は3次元配列的に作成可能です。で、作成された block 群を grid と呼びます。
こんな面倒な2段構えになっている理由は簡単で、thread は立てられる数が限られているからです。限界はGPUによって異なるかと思います。
512個の thread しか立てられないなら 1000個のデータは1blockでは扱いきれません。この場合、最低2blockが必要です。ちなみに、block には理論上限界がないようです。
で、実際の計算処理（つまり thread の処理）を行う kernel で thread や block のインデックスを取得するには以下の変数にアクセスします。
threadIdx : thread のインデックス
blockDim : block 中の thread の幅
blockIdx : block のインデックス
gridDim : grid 中の block の幅
これらはすべて3次元なので、.x, .y, .z で各次元の値にアクセスできます。
与えられたグローバル変数の特定要素にアクセスしたい場合はこれらを利用してインデックスを計算します。
たとえば、block が 100*1*1 thread の場合と、50*2*1 thread の場合では1次元配列データへのアクセスは以下のようになります。
100*1*1 : index = threadIdx.x
50*2*1 : index = threadIdx.x + threadIdx.y * blockDim.x
block も増えてくるとかなり面倒なのでインデックスの生成には気をつけてください。
