GPGPU#5
■ このスレッドは過去ログ倉庫に格納されています
GLSLからOpenCLへの移行を昨日から始めたけど
GLSLより書きやすいのはいいけど最適化を追い込まないととんでもなく遅くなるんだな
GLSLで複雑な汎用計算やらせるのは難解なパズルゲームみたいで嫌になってたけど
結局最適化の手間を考えたらどっちが楽ということはないんだね・・・
>>298みたいな単純な計算ならGLSLだとバグったような速度が簡単に出るから別世界感が凄い >>300からの続きですが、arraySizeをあまり大きくできないので、
ソースを弄って足し算を各100万回行うように改造しました。結果、
Releaseビルド、x64モードで
CPU→16.5872[s]
GPU→5.77132[s]
となりました。ここからFlopsを出してみると、>>298では
CPUが1078.66MFLOPS、GPUが433.164MFLOPSだったのが、
今回はCPUが1975.5MFLOPS、GPUが5677.73MFLOPSとなりました。
理論値からは明らかに小さいですが、少なくともGPUはより活用できているように感じます。
……結局arraySizeを大きくできない問題は解決していません。
ただ、float・int型にしてみると倍(51200)まで設定出来ました。
つまり、流し込むデータは200KBまでは大丈夫ということなのでしょうか？ >>299
GPGPUはメモリ転送のオーバーヘッドがないHSA(Huma)だよな
PCではど重い処理でない限りAMDのHSAがGPGPU処理の主流になるだろうな >>298
残念なお知らせ。
そのソースコードでは、GPUの演算時間ではなくGPUの呼び出し時間しか計測してないね。
「実際の演算時間」=「内部ブロック数」*(「内部ブロックの呼び出し時間」+「内部ブロックの演算時間」)だとすると、
「実際の演算時間」-「内部ブロックの演算時間」になっているはず。
ブロック数が充分大きければ誤差だけど、内部ブロック数が1のときは激速になってしまう。
まぁ、実際の運用ではCPUとGPUが並列に動作することを期待するからそれでもいいんだけどね。
いずれにしても、CPUぶん回すよりも手っ取り早いと思っていたら大間違いだよ。
それと、CUDAスレも宜しく。 >>304
>そのソースコードでは
え！？ ……つまり、
普通にtimeGetTimeかQueryPerformanceCounterとかを使えってことなんですか？
それとも、測定する位置が間違っているということなんですか？
>CUDAスレも宜しく
分かりました。次回以降はそちらにレスすることにします。 >>298 arraySizeが大きいと、CPU版すらStackOverflowになるよ。
http://pastebin.com/Av3YzTGs うっかり、166行目を「cudaStatus = cudaSetDevice(1);」にしちゃったので、適当に直しておいて。 ローカルメモリを使う場合って確保しようとした容量が大き過ぎると
グローバルのほうへ確保されてしまうんだよね？
AMDのGCNはどれくらいまでローカルメモリがあるのか分からないんだけど
試行錯誤して調べるしかないのか >>307-308
調査ありがとうございました。そうか、メモリのせいだったのか……
gridsizeの65536制限は知っていたのですが、block・gridでの
分割方法がイマイチよく分かっていなかったので、実コードで
示してくださって助かります。こちらの環境でテストしてみると、
Releaseビルド、x64モードで
> CPU計算時間:0.060652126[s] -> 276.614[MFLOPS]
> size: 16777216
> size_x,y: 262144,64
> blockSize: 256,1
> gridSize: 1024,64
> GPU計算時間:0.034433924[s] -> 487.229[MFLOPS]
> 最大絶対誤差:0.0000000000000000
となりました。>>298より微妙に速くなった程度ですが、
負荷が軽すぎるせいだということは>>302で確認しています。
ちなみにCUDA-Z でこちらのグラボを計測すると、スレッドの次元が1024x1024x64、
グリッドの次元が65535x65535x65535、演算性能は
int32=47.1[Giop/s]・float=94.0[Gflop/s]・double=11.8[Gflop/s]らしいです。 >>307
冗長なOpenCLに比べてやっぱりCUDAはスマートでいいな OpenCLのclEnqueueNDRangeKernelでカーネルを実行するときに
global_work_sizeとlocal_work_sizeに同じ値(256,256など)を入力すると
何もエラーは返されずにメモリの参照が壊れて？しまいclEnqueueReadBufferで
CPU側で読み取った値が全て0になってしまいます。
これは仕様なのでしょうか？ visual studio 2013でCUDAが使えないからC++AMPでやるお！ >>313
そのためだけにVS2012と2013使い分けてる俺…… (・∀・)ﾉ CPUの300倍くらいの性能が出たお！
比較したCPUはE2-2000っていうCPU+GPU=APUだけど全くGPUとしての機能をもってないのでガッカリしたお。 (・∀・)ﾉ ALU(IGP)のE2-2000はHD7770の1/50のパワーしかないが並列性はあるようだ。 テキスト処理ってGPUで高速化できないものでしょうか
具体的には
Appache Solr
の検索処理が遅いのでなんとか高速化したいのですが テキスト処理なんてわざわざGPUでやるよりSSE/AVXでやったほうが億倍マシ OpenCLでプログラム組んでみたけど、CPUとGPUメモリのやり取りがネックになっているのか、思ったよりスピードが出ない
他の人はGPU利用するにあたってメモリのやり取りとか何か工夫している？ 基本はメモリとのやりとりを少なくするって話でしょ
それ以上の個別の工夫を簡単に説明するのは難しいよね
ケーススタディしたいのならそういう本なり文献なり漁るべき OpenCVのOpenCLバインディングのコードを参考にしたらいいんじゃないのかな OpenCLの1.1と1.2に後方互換性ありますか？ >>329
買えたら買ってるよ
メモリの転送の処理が要らなくなったらと思うと幸せな気分になれるよ kaveriってOpenCL使うとき、コピーせずにポインタ参照で渡していいって解釈でいいの？
最近GPGPUをやりはじめたばかりだから、的外れなことかもしれんが。。 Boost.ComputeあったらC++ AMPいらなくない？ vexclを少し使ってみたけど便利だ
あとはC++AMPみたいにradeonのドライバーの
バージョン上がると使えなくなったりしないなら GPUの行列演算ライブラリってないですか？
具体的には特異値分解できるのを探してます >>337
機械学習スレで書いてた人かな？
ちゃんと調べてないけどMAGMAなら入ってるかも
http://icl.cs.utk.edu/magma/overview/index.html
{sdcz}gesvd はサポートしてるって書いてある
ただし、GPUカーネル内から直接呼びたい場合は使えないらしい
CPUからカーネル呼び出しする必要がある
（SC14時点の資料） 最大固有値
最大固有ベクトル
だけを求めたい場合って、 最大固有値
最大固有ベクトル
だけを求めたい場合って、べき乗法が最速でしょうか？ VS2015のc++amp仕様が変わった？
ＣＰＵで実行するrestrict(cpu)のマイクロソフトのサンプルコードがコンパイルエラーになる。 マイクロソフトの開発ブログで
＞ Is C＋＋AMP dead ？
との質問にレスが無い。
モスさんどこ行った？ は？ごぽごぽに決まってるだろ？って言われた．．．
頭ごなしに言ってくる人って何なんでしょうね 僕の知り合いの知り合いができたパソコン一台でお金持ちになれるやり方
役に立つかもしれません
グーグルで検索するといいかも『ネットで稼ぐ方法 モニアレフヌノ』
MYAIQ ■ このスレッドは過去ログ倉庫に格納されています
