この記事には誤りがあります。それについて最後のほうに追記があります。

例のGPUレイキャストのカーネルを弄ってたら、Geforce 9400M で100万レイの処理に1.7秒ぐらい掛かってたのが0.9秒ぐらいに短縮した。

レイごとにBVHをトラバースするカーネルで、末端ノードが見つかったときにグローバルメモリに対してatomicAddをしてる部分があったけど、それが遅かったみたい。

共有メモリに1warp分、つまり32個の要素を持つintの配列を確保し、atomicAddをしてた部分で各スレッドで自分が該当する場所に1をセットし、最後に0番目のスレッドで集計して1warpで1回だけatomicAddを実行するようにしたら、2倍近く高速化した。
基本的な高速化手法だけどここまで効果があるとは思ってなかったので後回しにしてた。
交差判定の方のカーネルはそのままなので、トラバースのカーネルだけでいうと2倍以上高速化してるかもしれない。

変更した部分は以下のとおり。

// カーネルの外でマクロ定義
#define WARP_SIZE 32
#define WARP_INDEX (threadIdx.x % WARP_SIZE)

…省略…

// カーネルのはじめのほう
__shared__ int staged[WARP_SIZE];
staged[WARP_INDEX] = 0;

…省略…

// 直接グローバルメモリにatomicAddしてた部分
staged[WARP_INDEX] = 1;

…省略…

// カーネルの最後のほう
__syncthreads();
…省略…
if (WARP_INDEX == 0) {
    int sum = 0;
    for (int i = 0; i < WARP_SIZE; i++) sum += staged&#91;i&#93;;
    if (sum > 0) atomicAdd(stagedCount, sum);
}

Geforce 9400MがCompute Capability 1.1止まりだから少し回りくどい書き方になってるけど、1.2以上なら共有メモリに対して直接atomicAddが呼べるのでもっとスッキリ書けると思う。（厳密には同じ処理内容にはならないけど）

GTX 460で試してみたら、デバイスメモリの帯域幅が9400Mと比べて広いからか、グローバルメモリに対するatomicAddが効率化されてるからか、0.06秒が0.055秒になる程度だった。

いまのアルゴリズム的には、交差判定が必要なときにstagedCountが全体として1以上になりさえすればよく、正確な値である必要はないので、上のコードはもっと効率化出来る。
例えば__shared__ int staged[WARP_SIZE];の代わりに、__shared__ bool staged;として（配列じゃなくする）、各スレッドからatomicAddの代わりにstagedにtrueを書き込み、最後に0番目のスレッドでstagedがtrueの場合に実際にatomicAddするようにすればいい。
（そうすると変数名としてstagedCountという名前はふさわしくなくなってしまうが、代わりのいい名前が思いつかない）

// カーネルの外でマクロ定義
#define WARP_SIZE 32
#define WARP_INDEX (threadIdx.x % WARP_SIZE)

…省略…

// カーネルのはじめのほう
__shared__ bool staged;
staged = false;

…省略…

// 直接グローバルメモリにatomicAddしてた部分
staged = true;

…省略…

// カーネルの最後のほう
__syncthreads();
…省略…
if (WARP_INDEX == 0) {
    if (staged) atomicAdd(stagedCount, 1);
}

ただ、この方法では共有メモリの同じ場所に複数のスレッドから同時に書き込まれる可能性があり、まぁ同じ値を書き込むんなら大丈夫だと思うけど、少々不安な方法でもある。
実際、9400M（Mac）とGTX 460（Ubuntu）で全体的なstagedCountの値に違いが出た。
最終的な出力画像に違いは無かったけど、配列を使う場合と比べてほとんど処理時間に差が出なかった。
ただ、共有メモリの節約にはなると思うので、ここまでやるのも無意味ではないと思う。

// マクロ不要

…省略…

// カーネルのはじめのほう
__shared__ bool staged;
staged = false;

…省略…

// 直接グローバルメモリにatomicAddしてた部分
staged = true;

…省略…

// カーネルの最後のほう
__syncthreads();
…省略…
if (threadIdx.x == 0) {
    if (staged) atomicAdd(stagedCount, 1);
}