Reduction最適化の分析 1ブロックあたりのスレッド数について

結局、1ブロックあたりのスレッド数を128にしても、512にしても、
扱える個数が変化するだけでグラフの形には影響がないので、
今のままで分析を進めることにします。
ただし、1ブロック当りのスレッド数を増やすと、アルゴリズムにもよりますが、
要素数が2^16を越えたあたりでパフォーマンスが落ちます。
これの理由はアルゴリズムの内容によりますが、Kernel6でいえば、
ブロック数が64を超えると、１つのブロックを2つ以上のワープで処理する必要があり、
そうなると並列リダクション部分で__syncthreads()を使う必要でてくるためだと思います。
これはソースコードを見れば分かりますので、SDKのサンプルを参照してください。

さて以下ではブロック数128の場合で考えます。
1スレッド当り1個の要素をコピーする部分を持つKernel2までは、
ブロックの最大数65535（2^16-1)から、
128x32768=2^22=約400万要素数の最大値となっています。
Kernel3-Kernel5は、最初の部分で1スレッドが2つの要素の和を計算するところから
始めていますので、要素数も倍まで扱えるようになっています。
Kernel6では直列リダクションと併用しますので、要素数に限界はありません。

以上の考察の結果、最終的なKernel6でいえば、
・ブロック当りのスレッド数を増やせば、直列リダクションのループ回数が減る。
直列の段階の最後に1回だけスレッド同期が入る。
・しかし、その後の並列の段階でスレッド同期が必要になる。
512で3回、256で2回、128で1回。同期のコストは高い。
・このカーネルでは要素数の制限はない。
ということを踏まえると、スレッド同期の回数が少ない128にしたほうが、
パフォーマンスが高くなるということになります。
ただし、要素数が増えていくと（テストでは2^24個あたりから）、
直列のループ回数が増えていくのでスレッド数を多くしたほうが速くなります。
といってもわずかな差でしかありませんが。

というわけで、1ブロックあたりのスレッド数は、共有メモリの使用サイズだけでなく、
アルゴリズム上の問題（スレッド同期の必要性）にもよるということがわかりました。

アルゴリズムのチューニングは奥が深いですね。
この辺はハードの特性によるところが大きいですが。
一般にスレッド同期というのはコストが高そうです。
32個なら同期不要というのはパフォーマンス向上にはよいのかも。