忍者ブログ
CUDA+GPGPU、C++、C#などのプログラムについての備忘録がわり
[1] [2] [3] [4
Posted by - 2019.10.17,Thu
×

[PR]上記の広告は3ヶ月以上新規記事投稿のないブログに表示されています。新しい記事を書く事で広告が消えます。

Posted by サンマヤ - 2008.08.04,Mon
VistaマシンにCUDAを導入ってのを考えたんですが、
VistaSP1にするとCUDAのドライバが入れられないっぽい。
あるいはGeforce8800Mだからか。
やれやれ。

VistaでCUDAはしばらく無理そうです。
プリインストールでSP1だからどうにもならない。
PR
Posted by サンマヤ - 2008.08.01,Fri
CUDA Zoneのページが新しくなっていました。
一覧性が悪くなって使いづらくなったような・・・

CUDA2.0がβ2にバージョンアップしていました。
Vistaマシンを手に入れたので、そちらにインストールして使ってみます。
動く状態になったら、レポートします。
Posted by サンマヤ - 2008.07.25,Fri

結局、1ブロックあたりのスレッド数を128にしても、512にしても、
扱える個数が変化するだけでグラフの形には影響がないので、
今のままで分析を進めることにします。
ただし、1ブロック当りのスレッド数を増やすと、アルゴリズムにもよりますが、
要素数が2^16を越えたあたりでパフォーマンスが落ちます。
これの理由はアルゴリズムの内容によりますが、Kernel6でいえば、
ブロック数が64を超えると、1つのブロックを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個なら同期不要というのはパフォーマンス向上にはよいのかも。

Posted by サンマヤ - 2008.07.22,Tue
Webページで書きかけのReductionのテスト結果を考察中、
Reductionがうまく動かないときにいろいろいじったことによって
1ブロックあたりのスレッド数を128にしてしまったようです。
これにより、結果に影響が出ているので、
近日中にテストしなおしてUPします。。。
Posted by サンマヤ - 2008.07.18,Fri
共有メモリのテスト結果をホームページに反映させました。
そのほか、いくつかの表現、記述を更新しています。

やっつけで書いたので、記述がいい加減ですが、
ソースコードなどもこちらに載せておきます。

ページはこちら
カレンダー
09 2019/10 11
S M T W T F S
1 2 3 4 5
6 7 8 9 10 11 12
13 14 15 16 17 18 19
20 21 22 23 24 25 26
27 28 29 30 31
フリーエリア
最新コメント
[11/19 矢野 忠]
[02/25 山本義和]
[07/08 hirota]
[07/06 hirota]
[02/05 矢野 忠]
最新トラックバック
プロフィール
HN:
サンマヤ
性別:
非公開
職業:
プログラマ
趣味:
ゲーム
バーコード
ブログ内検索
カウンター
忍者アナライズ
Template by mavericyard*
Powered by "Samurai Factory"
忍者ブログ [PR]