スポンサーサイト

上記の広告は1ヶ月以上更新のないブログに表示されています。
新しい記事を書く事で広告が消せます。

CUDAでのGPUによる総和について

GPUで総和(リダクション)をする機会があって調べてみると

色々と情報があるもので

すんなりいけると思った割には失敗したので

自分みたいに苦労する人が少なるなるように

紹介したいと思います
CUDAを提供しているnVidiaから

詳しい文献が出ています

こちら

基本的には載っているソースをそのまま使えば動作しますが

何点か気をつけなければいけないことが有りました

(普通ならわかって当然のことかもしれませんが...)

まず、#4 First Add During Load の部分です

この部分では、シェアードメモリに格納する処理のお陰で

実行するスレッド数を半分にすることが出来ます

よって、#1~#3までに設定していたスレッド数を半分にしないといけません

要注意です

次は#5 Unroll the Last Warp

載っているソースのままでも問題ないはずですが

私の場合は上手く動きませんでした

いろいろ調べてみましたら、シェアードメモリの宣言の部分に

extern __shared__ int sdata[];



volatile extern __shared__ int sdata[];


この修飾子「volatile」をつけることで

最適化を防ぎ、そのまんま動作するというものらしいですが

実際、よくはわかっていません

しかし、これをつけることでちゃんと動作したので

動かないという方はつけてみるのはどうでしょうか?

最後に#6 Completely Unrolled

この部分で再びスレッド数を元に戻し、逆にtemplateに渡す値を半分にすると

動くっていう

ソース元がおかしいのか、私がおかしいのか

どちらを信じるか信じないかはあなた次第です

Comment

Comment Form
公開設定

Trackback


→ この記事にトラックバックする(FC2ブログユーザー)
上記広告は1ヶ月以上更新のないブログに表示されています。新しい記事を書くことで広告を消せます。