VistaSP1にするとCUDAのドライバが入れられないっぽい。
あるいはGeforce8800Mだからか。
やれやれ。
VistaでCUDAはしばらく無理そうです。
プリインストールでSP1だからどうにもならない。
まず、runtimeのビルド時に、C++＞コード生成を、
/MT --> /Mtd
とデバッグ版を使うように変更して、リビルド。
その上で、付属のfiltering.slnをビルド
まだ、名前解決ができないといわれる。
そこで、リンカのオプションで、追加のライブラリに
user32.lib gdi32.lib
を追加。
これでリンクもできるようになります。
しかし、この二つって元からリンクされるものじゃないのかなあ。。。
さて、リンクもうまくいく、CPUモードでも走るんですが、
肝心のGPUモード(BRT_RUNTIME=ogl or dx9)
にするとstd::vector関係で落ちます。
なんだこりゃ。。。
あきらめて、以前動かなかったBrookbench（発行元がリンクきれてます。。。）をビルド。
実行してみると、見事動きました。
これで、BrookGPUのテストもできるようになりますね。
もう少しいじってみてからHPにも反映させたいと思います。
ちなみに、同じGPUモードでも、
oglとdx9では、OGLのほうがずっと速い。
何十倍と速度差が出るのはなんでなんでしょうね。
CUDAはGeForceのごく限られたバージョンしか使えず、
現在の職場環境においても、一般的な環境においても、
かなり制約が高いという問題があります。
そこで、もう少し一般的な環境で動く言語として
BrookGPUを検討していたのですが、
先日購入したGPUGem2にBrookの環境が入っていたので
使ってみたのですが、やはりビルドできず。
LIBC.LIBが見つからないだの、名前が解決できないだの、いろいろ出ます。。。
もう少しで使えそうなんですがね。。。
一覧性が悪くなって使いづらくなったような・・・
CUDA2.0がβ2にバージョンアップしていました。
Vistaマシンを手に入れたので、そちらにインストールして使ってみます。
動く状態になったら、レポートします。
結局、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個なら同期不要というのはパフォーマンス向上にはよいのかも。
Powered by "Samurai Factory"