- 262 名前：255 mailto:sage [2008/03/16(日) 18:27:15 ]
- 例えばこんな感じ。
__global__ void func(float const * a1, float * a2, unsigned num)
{ // 実際には一時変数を使った方がいい希ガス
for (unsigned idx = blockIdx.x * blockDim.x + threadIdx.x; idx < num; idx += gridDim.x * blockDim.x) {
a2[idx] = a1[idx] * a1[idx];
}
num=ブロック数*スレッド数ならどのスレッドでもループは一回だけ回る。
num<ブロック数*スレッド数なら、暇なスレッドが発生する(から効率は宜しくない)。
num/(ブロック数*スレッド数)が11と12の中間なら一部のスレッドは12回回って残りは11回回る。
このループ回数が1や2じゃなければ、暇なスレッドの割合が相対的に少なくなる寸法。
# これが理由で、ブロック数は無闇に増やせばいいというわけでもないということになる。
あー説明が下手な私（：；
# サンプル作ったら誰か買わない?w
こういう回し方をすることによって、近くのスレッドが近くのメモリをアクセスする状態のままループが進行する。
つまり、ProgrammingGuideで言うところの"coalesced"。これをプログラミングガイドでは「結合した」と訳しているけど……
|
 |
