要出典なCUDAメモ

CUDAの挙動をよく知らないので, 現時点でぼくが理解している（真偽不明な）内容をメモ.
真偽はかなり怪しいです.

カーネルは非同期実行

カーネル関数は非同期なので処理はすぐ帰ってくる.
処理時間を計測するのにタイマーを使うときに注意が必要.
カーネル関数の呼び出しをタイマー処理ではさんで時間を計測する場合には,
こんな感じで呼び出すといいかも.

// タイマー処理
unsigned int timer;
CUT_SAFE_CALL(cutCreateTimer(&timer));
CUT_SAFE_CALL(cutResetTimer(timer));
CUT_SAFE_CALL(cutStartTimer(timer));

// カーネル関数の呼び出し
Kernel <<< blocks, threads >>> (a, b, c);

// カーネル関数の終了待ち
CUDA_SAFE_CALL(cudaThreadSynchronize());

// タイマーストップ
CUT_SAFE_CALL(cutStopTimer(hTimer));

カーネル関数内での分岐処理(if文)

カーネル関数内に分岐処理があり, ワープ内のスレッドで真偽が分かれる場合は,
ワープ内のスレッドは真と偽の両方の処理を（必要なくても）行う.

たとえば, 1次元で要素数Nの配列を処理する場合,

__device__ void
sample()
{
  if (cond) {
    // 処理A
  } else {
    // 処理B
  }
}

というような場合に, ワープ内のスレッドがAとBの処理を両方実行する場合,
ワープ内のスレッドはすべてAとBの処理を行う.

ワープ内のスレッドはすべて同じ命令を実行しなければならないため,
真と偽の処理を別々に実行することができないからだ.
(実行時, スレッドをワープに振り分けるとき, 分岐先によって自動で振り分けてくれる, なんてことはしてくれるのだろうか)

(Fermiでは32SP/SMなので多分ワープは1クロックで実行されるが, それ以前のアーキテクチャでは8SP/SMなのでワープは4クロックで実行される.
このスレッド分けのときに, 真と偽によって振り分けてくれるのだろうか(要出典))

カーネル関数内での繰り返し処理(for文)

ワープ内のスレッドで繰り返し回数が違う場合も無駄な処理が行われる.
ワープ内のスレッドの中で最も回数の多いスレッドの回数分だけ(必要がなくても)処理が実行されるからである.
(何もしないスレッドがあるのは当然効率が悪いので)

さいごに

なんどもいうけどしんぎは（ｒｙ