CUDA
リファレンス読みながら組んでいたら気がついたらRuntime APIを混ぜてしまった。NVIDIA CUDA Library: Module IndexたぶんDriver APIが低レベルでRuntime APIが高レベルで、混ぜて使っちゃいけないはず。Driver APIは名前がcuMemAlloc()とかcuで始まってRunt…
int __all(int predicate);int __any(int predicate); 忘れがちだけど、cudaのwarp vote関数は便利。 場合によってはWarp内のちょっとしたやり取りにshared memory経由や外部変数のatomic関数による更新の代わりに使える。機能は名前のまんまなので覚えやす…
【イベントレポート】【SIGGRAPH Asia 2009レポート】 東工大、スクウェアエニックスがCUDA実装事例を紹介 - PC Watch12/16に参加したカンファレンスの記事。 一番の見所はCUDAの最適化方法についての解説だったけど。 抽選であたるGPUに外れたのは残念だっ…
SIGGRAPH ASIA 2009の横で行われているNVIDIA主催のGPU Computing Master Classに参加した。 領収書を貰ったからたぶん会社からお金が出るだろう。 受付でCUDA本を貰った。（既に持ってるから誰かに上げる予定）http://www.siggraph.org/asia2009/jp/gpu/公…
gccのバージョンの関係でSDKのビルドが通らない・・。 OSの選定も含めて完全に使えるのはもう少し先
ついにGTX 295 + CUDA 2.3環境ができた。まあ、自分は横で見てただけだけど。 Linuxでの環境構築は大変だ。これじゃあWindowsよりもユーザー数が増えることはないな。DeviceQueryを行うとGPUが２個ついているように見える。 GTX 295での最適化はどんな落とし…
CUDAに合うソートって何だろう。 最近10種類以上のソートを試したけど、いまだに満足できるものに出会えていない。 (bubbleから並列bitonicまで試した) もちろん、ソート対象の要素数に依存するんだろうけど。CUDAでは再帰が使えないので、Quick Sortとかは…
2009/12/10 GPUコンピューティングの現状とスーパーコンピューティングの未来2chで見つけた。京都大学の天体系の人が書いたスライド。 わかりやすくて面白かったのでボリュームがあるが一気に読んでしまった。 いろいろ考えさせられる。
今日は最後の課題に取りかかった。 内容は128K(131072)個のfloat値の合計CUDAを利用するためのオーバーヘッドだけでCPU(i7)で行った場合の時間を超えていたので、残念な感じだった。 データ量を増やしてみたらCPUよりも速くなったけど・・。レポートも書いて…
自分の組んだCUDAのKernelがたまに正しく動かず、かなり時間をかけて調査をしていたが原因が分からなかった。 同じデータに対して、同じ処理を行っているのに時々結果がランダムにおかしくなる。 （ランダムな結果を出力するKernelを作る方が、同じ結果を出…
今日はcudaMemcpyToSymbol()とcudaMemcpyFromSymbol()の２つの関数でひどい罠にはまった。 cudaError_t cudaMemcpyToSymbol (const char* symbol, const void * src, size_t count, size_t offset, enum cudaMemcpyKind kind) cudaError_t cudaMemcpyFromSym…
2048 x 2048のCUDAを使った行列の転置も完成。 最適化していないCPUでの実行の10倍ぐらいは速くなった。この処理の場合はHostとDevice間の転送がボトルネックになってくるから、その部分をどうするかが鍵になるな。 と言っても、選択肢はそれ程ないけど。
いろいろあって、次の３つの関数をCUDAで最適化しないといけない。 512x512の大きさの正方行列同士の乗算 2048x2048の正方行列の転置 128k個のfloatの値の総和 とりあえず乗算は終わった。 時間は毎回のぶれが大きいが速いときは3msぐらい。 最適化せずにcpu…
パフォーマンスの点では値段の高いTeslaよりも安いGTX285やGTX290の方が高い。 おまけにGTXは画像の出力機能まである。 Teslaの値段が高いのは信頼性が高いからと言われるがどのくらい違うのか。そのヒントになりそうなのが次の論文（右側のdownloadの所から…
昨日の説明は文章だけでわかりにくかったから、矩形領域をglobal memoryからshared memoryへ読み込むコードと結果の追加。 uint4で読むときには一列をhalf warp(16thread)の１回の読み込みで済ましている。(今回は横幅の最大を64要素に制限しているので) int…
はじめてのCUDAプログラミング―驚異の開発環境[GPU+CUDA]を使いこなす! (I・O BOOKS)作者: 青木尊之,額田彰,第二I O編集部出版社/メーカー: 工学社発売日: 2009/11メディア: 単行本購入: 11人 クリック: 310回この商品を含むブログ (29件) を見る11/20にCUDA…
複数のthreadを利用して画像のような矩形領域をglobal memoryからshared memoryへ読み込む場合、しかも全てのthreadがint型で読んでも１回じゃ読み切れない場合にはint型で複数回読むのとuint4型で読むのとどちらが良い速いのか。（uint4でなら１回で読み込…
blockの数を増やしてみるとうまくいくことがあるらしい。
今日はCUDAについて、良い情報がいろいろ得られた。 しかし、どれも会社の別の人が編み出した事ばかりなのでどこまで書いて良いか微妙。 なので、ほどほどな感じで書く予定。CUDAのプログラムで初期化（最初に呼んだCUDA API）が1秒近くかかるのは改善可能。…
社内の持ち回りの勉強会で、今日の当番の人が忙しくて辞退したので、代わりに担当してみた。 CUDAでの最適化と画像に関するアルゴリズムの話をした。 自分で話しておいて何だが、アルゴリズムに関する数学的な質問は答えられないので一緒の仕事をしている頭…
11章の粒子計算や12章の偏微分を最適化しながら実効性能を乗せてくれていたのは良かった。 （自分には計算式に、ついていけない部分もあったけど）付録Aに載っていたマニュアルへのパスの通し方は早速.bashrcに追加しておいた。 export MANPATH=$MANPATH:/us…
CUDA本はまだ読み途中。性能を引き出す上で6章に書いてあった、Occupancyを高めるよりもマルチプロセッサあたりのアクティブなブロック数の方が大事というのはいいヒントになりそうだ。ついついBlockあたりのThread数を上げようとしてしまうが、同時におよそ…
__device__をつけてバッファを用意する。 HOST側のプログラムからは次の関数を利用してアドレスを取得する。 cudaError_t cudaGetSymbolAddress (void **devPtr, const char * symbol); こんな感じになるかな __device__ int d_buff[256]; int main(void) { …
秋葉ヨドバシで最後の一冊を手に入れた。 あまり仕入れていないのか、それとも凄い売れているのか？まだ5章までしか読んでいないけど、4章に書いてあった-deviceemuと-D__DEVICE_EMULATION__を 同時に指定するのは良い方法かも。 Kernel中にprintf()とかasse…
doubleを利用するならcudaドライバとtool kitは最新の2.3を使った方が良いかもしれない。 doubleを利用して-arch=sm_13を指定したときに2.2でのみ発生する問題に遭遇した。 （もう少し切り分けのための検証が必要そうだけど）でも使っているFedora 9には2.3…
CUDAでは最適化のテクニックとしてKernel関数をC++のtemplateを利用して量産して、適切な関数を呼び出すというのが結構使える。 たとえばKernel関数内部の最内のループ回数が1〜10の間で可変のときに、 ループする回数ごとに10種類の関数を用意して利用する。…
いよいよ明日発売か。 今から読むのが楽しみだ。今日フラゲした人の話ではかなりいいらしい。 （本を発売日前に手に入れることもフラゲっていうんだ。ゲームソフトに対してのみ使うと思っていた）はじめてのCUDAプログラミング―驚異の開発環境[GPU+CUDA]を使…
CUDAのプログラムではBlock内のThread数は多くすると必ずしも速くなるとは限らない。 Occupancyとかの関係でThread数を変えていくと突然速くなるポイントがあったりする。Thread数を減らしても同時にSMに入るBlock数が増えてOccupancyが上がれば速度が上がる…
cudaMalloc()を一切使わないで、代わりに__device__変数を利用してみた。 結果としては、アプリケーション全体では大した差は出なかった。 もともとcudaMalloc()を一回しか呼ばないようにしてたしなぁ。たくさんcudaMalloc()呼んでるプログラムなら、もう少…
内輪の勉強会に行ってきた。 10時から15時ぐらいまで。 残念ながら終わっても飲みは無かった。とりあえず、今日の話を聞いて試してみようと思ったのは次の事。 __device__をつけて変数を用意すればcudaMalloc()無しでdevice用のメモリが確保できるらしいので…
