GPGPU#5
283 ：デフォルトの名無しさん：2013/04/14(日) 14:30:38.59
抽象化の度合いが高いぐらいの意味合いでいいんじゃないの
実際問題プログラマが意識してGPUに仕事をさせるコードを書かずに勝手に解決してくれちゃう未来像は来るの？
実際問題プログラマが意識してGPUに仕事をさせるコードを書かずに勝手に解決してくれちゃう未来像は来るの？
284 ：デフォルトの名無しさん：2013/04/14(日) 18:29:31.68
たぶん、デザインパターンと同じくらいの敷居にはなるんじゃねーの？と言うか、デザインパターンとして確立されるんじゃね。
判らない人には判らないけど、自分の中に消化され始めると自然とそう言う書き方になるみたいな。
あとは、フレームワークでそう言う書き方を強制されて、なにも考えずにそう言うもんだと思うパターン。
判らない人には判らないけど、自分の中に消化され始めると自然とそう言う書き方になるみたいな。
あとは、フレームワークでそう言う書き方を強制されて、なにも考えずにそう言うもんだと思うパターン。
285 ：デフォルトの名無しさん：2013/04/15(月) 02:40:47.34
コア数が100くらいのGPUがあるとして
整数大規模行列の
全要素の和を求めるのってGPGPUで100倍程度の高速化は可能ですよね？
整数大規模行列の
全要素の和を求めるのってGPGPUで100倍程度の高速化は可能ですよね？
286 ：デフォルトの名無しさん：2013/04/15(月) 07:36:47.18
1コアあたりのシングルスレッド能力がCPUと同程度あって、メモリ帯域幅がそれに見合うほど広ければな
287 ：デフォルトの名無しさん：2013/04/22(月) 21:00:09.19
>>283
実際、解きたい問題依存だけど、Hadoopのように、プログラミングモデルを固定して制限を付けて縛ることによって簡単な記述で自動的に高速実行ってのは比較的現実的な未来だと思うよ。
Hadoopってのも、書きたい処理をHadoopの"型"に押し込んで書ければ、残りの面倒な問題は全部Hadoopが面倒見てくれますって事だし。
実際、解きたい問題依存だけど、Hadoopのように、プログラミングモデルを固定して制限を付けて縛ることによって簡単な記述で自動的に高速実行ってのは比較的現実的な未来だと思うよ。
Hadoopってのも、書きたい処理をHadoopの"型"に押し込んで書ければ、残りの面倒な問題は全部Hadoopが面倒見てくれますって事だし。
288 ：デフォルトの名無しさん：2013/09/21(土) 12:06:20.79
このスレ立った頃のGPUってもうゴミだよな？
289 ：デフォルトの名無しさん：2013/09/21(土) 13:04:48.76
OpenCLで書いたプログラムを9600GTとQ9650で動かしたら同じくらいの速度だった。
今のCPUには絶対負ける。
今のCPUには絶対負ける。
290 ：デフォルトの名無しさん：2013/10/02(水) 19:53:46.66
291 ：デフォルトの名無しさん：2013/10/20(日) 14:42:26.77
GPGPUとMantleの関係を語ってくれよ
mantleってなんなんだ？
mantleってなんなんだ？
292 ：デフォルトの名無しさん：2013/11/06(水) 09:10:41.13
>>291
DirectXやOpenGLは長い歴史があって
色々なハードで動かす必要があるので
互換性維持のために滅茶苦茶厚いAPIコールを通さないといけないが
MantleはGCNに絞った最適化さえすればいいので
そういうものがないっつーこと
据え置きはAMD一色だから移植も相互に楽になるし、面倒な最適化も最小限で済む
というのが言い分
DirectXやOpenGLは長い歴史があって
色々なハードで動かす必要があるので
互換性維持のために滅茶苦茶厚いAPIコールを通さないといけないが
MantleはGCNに絞った最適化さえすればいいので
そういうものがないっつーこと
据え置きはAMD一色だから移植も相互に楽になるし、面倒な最適化も最小限で済む
というのが言い分
293 ：デフォルトの名無しさん：2013/12/30(月) 01:57:57.57
DirectXやOpenGLやOpenCLだと実行時に
「君（GPU）この機能使える？」
「よし、通れ！」
というプロセスを踏んでるからnVidia、AMD、Intel、PowerVRなど全部のハードウェアでそれなりに動くわけ
これに対してCUDAやMantleは他社対応を完全に無くすことで最初から最適化されたランタイムで動いてる
商用製品でCUDAがOpenCLより人気なのを見れば分かるように、
他のハードウェアで動かない以上のメリットがあるわけ
とくにゲーム機だと全てAMDで、今はAMDがARMと仲が良いからCUDA以上に広まりやすい環境にあるわけ
このスレ的にMantleは用途が違うだろうけど、HLSL/GLSL直書きでGPGPUをやってた層には移行先として十分かなと言った感じ
「君（GPU）この機能使える？」
「よし、通れ！」
というプロセスを踏んでるからnVidia、AMD、Intel、PowerVRなど全部のハードウェアでそれなりに動くわけ
これに対してCUDAやMantleは他社対応を完全に無くすことで最初から最適化されたランタイムで動いてる
商用製品でCUDAがOpenCLより人気なのを見れば分かるように、
他のハードウェアで動かない以上のメリットがあるわけ
とくにゲーム機だと全てAMDで、今はAMDがARMと仲が良いからCUDA以上に広まりやすい環境にあるわけ
このスレ的にMantleは用途が違うだろうけど、HLSL/GLSL直書きでGPGPUをやってた層には移行先として十分かなと言った感じ
294 ：デフォルトの名無しさん：2013/12/30(月) 03:37:10.21
Mantleって互換性維持のため分厚くなり過ぎたDirectXに対して
低レベルなところを柔軟に叩けるようにするとかそういうものじゃなかったっけ？
Project Sumatraが楽しみだな
Java9ぐらいでparallelStreamにぺぺっと渡すと
OpenCLやってくれちゃうとか最高じゃん
逆に言うとそれぐらい猿でも簡単に扱えないとなっていうか
低レベルなところを柔軟に叩けるようにするとかそういうものじゃなかったっけ？
Project Sumatraが楽しみだな
Java9ぐらいでparallelStreamにぺぺっと渡すと
OpenCLやってくれちゃうとか最高じゃん
逆に言うとそれぐらい猿でも簡単に扱えないとなっていうか
295 ：デフォルトの名無しさん：2014/01/16(木) 15:13:31.68
Mantle.Netはよ
296 ：デフォルトの名無しさん：2014/01/16(木) 19:52:34.70
kaveriのFLOPSのDP:SPが気になるんだが、誰か情報持ってないか？
297 ：デフォルトの名無しさん：2014/01/16(木) 20:00:22.38
OpenACCってどうよ
298 ：デフォルトの名無しさん：2014/01/26(日) 21:12:44.48
最近CUDAを始めたのですが、簡単な計算(大きな配列のベクタ足し算)を
CPUとGPUにやらせると明らかにGPUが遅いです。
具体的には、VC2012+CUDA5.5でコード(http://www1.axfc.net/u/3155195.dat)を走らせると、
計算時間が次のようになりました(Releaseビルド、x64モード)。
CPU→1.51892e-005[s]
GPU→3.7824e-005[s]
一応計算はできているのですが、どうも性能を引き出せていない気がします。
また、コード中でarraySizeを65536にすると実行できなくなるのは何故なのでしょう？
どの辺書き換えればいいのかを教えて下さいお願いします。
ちなみにGPUはGeForce 610M(理論値で141.7GFlops)、
CPUはCore i5-3210M(1コアしか使わない状態なので理論値20GFlops)です。
CPUとGPUにやらせると明らかにGPUが遅いです。
具体的には、VC2012+CUDA5.5でコード(http://www1.axfc.net/u/3155195.dat)を走らせると、
計算時間が次のようになりました(Releaseビルド、x64モード)。
CPU→1.51892e-005[s]
GPU→3.7824e-005[s]
一応計算はできているのですが、どうも性能を引き出せていない気がします。
また、コード中でarraySizeを65536にすると実行できなくなるのは何故なのでしょう？
どの辺書き換えればいいのかを教えて下さいお願いします。
ちなみにGPUはGeForce 610M(理論値で141.7GFlops)、
CPUはCore i5-3210M(1コアしか使わない状態なので理論値20GFlops)です。
299 ：デフォルトの名無しさん：2014/01/26(日) 21:24:38.81
メモリ転送のオーバーヘッドがあるから、もっと大きな問題じゃないと効果は出ないよ。
300 ：デフォルトの名無しさん：2014/01/26(日) 21:54:26.48
>>299
それは知っているのですが、数値を大きくするとすぐにc00000fdでクラッシュするんですよ……。
<<<grid, block>>>もいろいろ弄っているのですが、どうにも効果が得られません。
今試してみたら、arraySizeの値で実行できるのは25600が最大みたいです。
それは知っているのですが、数値を大きくするとすぐにc00000fdでクラッシュするんですよ……。
<<<grid, block>>>もいろいろ弄っているのですが、どうにも効果が得られません。
今試してみたら、arraySizeの値で実行できるのは25600が最大みたいです。
301 ：デフォルトの名無しさん：2014/01/26(日) 22:08:07.85
GLSLからOpenCLへの移行を昨日から始めたけど
GLSLより書きやすいのはいいけど最適化を追い込まないととんでもなく遅くなるんだな
GLSLで複雑な汎用計算やらせるのは難解なパズルゲームみたいで嫌になってたけど
結局最適化の手間を考えたらどっちが楽ということはないんだね・・・
>>298みたいな単純な計算ならGLSLだとバグったような速度が簡単に出るから別世界感が凄い
GLSLより書きやすいのはいいけど最適化を追い込まないととんでもなく遅くなるんだな
GLSLで複雑な汎用計算やらせるのは難解なパズルゲームみたいで嫌になってたけど
結局最適化の手間を考えたらどっちが楽ということはないんだね・・・
>>298みたいな単純な計算ならGLSLだとバグったような速度が簡単に出るから別世界感が凄い
302 ：298：2014/01/26(日) 23:14:26.27
>>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までは大丈夫ということなのでしょうか？
ソースを弄って足し算を各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までは大丈夫ということなのでしょうか？
303 ：デフォルトの名無しさん：2014/01/26(日) 23:22:14.60
304 ：デフォルトの名無しさん：2014/01/27(月) 00:23:03.99
>>298
残念なお知らせ。
そのソースコードでは、GPUの演算時間ではなくGPUの呼び出し時間しか計測してないね。
「実際の演算時間」=「内部ブロック数」*(「内部ブロックの呼び出し時間」+「内部ブロックの演算時間」)だとすると、
「実際の演算時間」-「内部ブロックの演算時間」になっているはず。
ブロック数が充分大きければ誤差だけど、内部ブロック数が1のときは激速になってしまう。
まぁ、実際の運用ではCPUとGPUが並列に動作することを期待するからそれでもいいんだけどね。
いずれにしても、CPUぶん回すよりも手っ取り早いと思っていたら大間違いだよ。
それと、CUDAスレも宜しく。
残念なお知らせ。
そのソースコードでは、GPUの演算時間ではなくGPUの呼び出し時間しか計測してないね。
「実際の演算時間」=「内部ブロック数」*(「内部ブロックの呼び出し時間」+「内部ブロックの演算時間」)だとすると、
「実際の演算時間」-「内部ブロックの演算時間」になっているはず。
ブロック数が充分大きければ誤差だけど、内部ブロック数が1のときは激速になってしまう。
まぁ、実際の運用ではCPUとGPUが並列に動作することを期待するからそれでもいいんだけどね。
いずれにしても、CPUぶん回すよりも手っ取り早いと思っていたら大間違いだよ。
それと、CUDAスレも宜しく。
305 ：298：2014/01/27(月) 00:47:11.47
>>304
>そのソースコードでは
え！？ ……つまり、
普通にtimeGetTimeかQueryPerformanceCounterとかを使えってことなんですか？
それとも、測定する位置が間違っているということなんですか？
>CUDAスレも宜しく
分かりました。次回以降はそちらにレスすることにします。
>そのソースコードでは
え！？ ……つまり、
普通にtimeGetTimeかQueryPerformanceCounterとかを使えってことなんですか？
それとも、測定する位置が間違っているということなんですか？
>CUDAスレも宜しく
分かりました。次回以降はそちらにレスすることにします。
306 ：デフォルトの名無しさん：2014/01/27(月) 08:23:49.16
>>304
何言ってんだ、こいつ？
何言ってんだ、こいつ？
307 ：デフォルトの名無しさん：2014/01/27(月) 21:34:08.90
うっかり、166行目を「cudaStatus = cudaSetDevice(1);」にしちゃったので、適当に直しておいて。
309 ：デフォルトの名無しさん：2014/01/27(月) 23:30:12.43
ローカルメモリを使う場合って確保しようとした容量が大き過ぎると
グローバルのほうへ確保されてしまうんだよね？
AMDのGCNはどれくらいまでローカルメモリがあるのか分からないんだけど
試行錯誤して調べるしかないのか
グローバルのほうへ確保されてしまうんだよね？
AMDのGCNはどれくらいまでローカルメモリがあるのか分からないんだけど
試行錯誤して調べるしかないのか
310 ：298：2014/01/27(月) 23:50:13.09
>>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]らしいです。
調査ありがとうございました。そうか、メモリのせいだったのか……
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]らしいです。
311 ：デフォルトの名無しさん：2014/01/28(火) 01:09:12.72
>>307
冗長なOpenCLに比べてやっぱりCUDAはスマートでいいな
冗長なOpenCLに比べてやっぱりCUDAはスマートでいいな
312 ：デフォルトの名無しさん：2014/01/29(水) 01:59:06.39
OpenCLのclEnqueueNDRangeKernelでカーネルを実行するときに
global_work_sizeとlocal_work_sizeに同じ値(256,256など)を入力すると
何もエラーは返されずにメモリの参照が壊れて？しまいclEnqueueReadBufferで
CPU側で読み取った値が全て0になってしまいます。
これは仕様なのでしょうか？
global_work_sizeとlocal_work_sizeに同じ値(256,256など)を入力すると
何もエラーは返されずにメモリの参照が壊れて？しまいclEnqueueReadBufferで
CPU側で読み取った値が全て0になってしまいます。
これは仕様なのでしょうか？
313 ：デフォルトの名無しさん：2014/02/25(火) 21:16:18.98
visual studio 2013でCUDAが使えないからC++AMPでやるお！
314 ：デフォルトの名無しさん：2014/02/25(火) 21:43:30.35
>>313
そのためだけにVS2012と2013使い分けてる俺……
そのためだけにVS2012と2013使い分けてる俺……
315 ：デフォルトの名無しさん：2014/04/04(金) 10:44:13.17 ID:YtPgho8U
openCL始めたお(・∀・)ﾉ
316 ：デフォルトの名無しさん：2014/04/15(火) 02:32:13.65 ID:vGWbAtXL
(・∀・)ﾉ CPUの300倍くらいの性能が出たお！
比較したCPUはE2-2000っていうCPU+GPU=APUだけど全くGPUとしての機能をもってないのでガッカリしたお。
比較したCPUはE2-2000っていうCPU+GPU=APUだけど全くGPUとしての機能をもってないのでガッカリしたお。
317 ：デフォルトの名無しさん：2014/04/19(土) 12:16:56.16 ID:Firi/9oq
(・∀・)ﾉ ALU(IGP)のE2-2000はHD7770の1/50のパワーしかないが並列性はあるようだ。
318 ：デフォルトの名無しさん：2014/04/22(火) 04:44:14.02 ID:aREYskwN
AIDA64に測定メニューあるよな
319 ：デフォルトの名無しさん：2014/08/29(金) 13:33:23.65 ID:P9znXDYB
AMDとMS，GPU演算用途向けのコンパイラ「C＋＋ AMP v1.2」を発表
http://www.4gamer.net/games/032/G003263/20140828031/
http://www.4gamer.net/games/032/G003263/20140828031/
320 ：デフォルトの名無しさん：2014/09/12(金) 04:54:39.38 ID:jvr90R5c
テキスト処理ってGPUで高速化できないものでしょうか
具体的には
Appache Solr
の検索処理が遅いのでなんとか高速化したいのですが
具体的には
Appache Solr
の検索処理が遅いのでなんとか高速化したいのですが
321 ：デフォルトの名無しさん：2014/09/12(金) 09:59:47.51 ID:cxN2yFh/
ボトルネックはメモリでしょう。
322 ：デフォルトの名無しさん：2014/09/16(火) 05:52:05.01 ID:padeH6x3
テキスト処理なんてわざわざGPUでやるよりSSE/AVXでやったほうが億倍マシ
323 ：デフォルトの名無しさん：2014/09/20(土) 00:49:25.20 ID:NyWaXORh
324 ：デフォルトの名無しさん：2014/09/24(水) 15:13:16.24 ID:ltG1hZ24
OpenCLでプログラム組んでみたけど、CPUとGPUメモリのやり取りがネックになっているのか、思ったよりスピードが出ない
他の人はGPU利用するにあたってメモリのやり取りとか何か工夫している？
他の人はGPU利用するにあたってメモリのやり取りとか何か工夫している？
325 ：デフォルトの名無しさん：2014/09/24(水) 23:08:38.12 ID:psEUFh+R
そりゃ工夫するだろう。
326 ：デフォルトの名無しさん：2014/09/25(木) 13:28:12.29 ID:F8MulcGG
ごめん、どんな工夫してるか聞いてみたかったんだ
327 ：デフォルトの名無しさん：2014/09/25(木) 16:05:41.16 ID:Coq6ADbv
基本はメモリとのやりとりを少なくするって話でしょ
それ以上の個別の工夫を簡単に説明するのは難しいよね
ケーススタディしたいのならそういう本なり文献なり漁るべき
それ以上の個別の工夫を簡単に説明するのは難しいよね
ケーススタディしたいのならそういう本なり文献なり漁るべき
328 ：デフォルトの名無しさん：2014/09/25(木) 18:24:41.04 ID:RDrb9uGa
OpenCVのOpenCLバインディングのコードを参考にしたらいいんじゃないのかな
329 ：デフォルトの名無しさん：2014/09/25(木) 21:14:08.96 ID:YRvO5dcq
>>324
kaveri使えよ
kaveri使えよ
330 ：デフォルトの名無しさん：2014/09/25(木) 22:16:45.04 ID:Vf7t0liy
OpenCLの1.1と1.2に後方互換性ありますか？
331 ：デフォルトの名無しさん：2014/09/27(土) 00:53:04.08 ID:SNKkkpyl
kaveriってOpenCL使うとき、コピーせずにポインタ参照で渡していいって解釈でいいの？
最近GPGPUをやりはじめたばかりだから、的外れなことかもしれんが。。
最近GPGPUをやりはじめたばかりだから、的外れなことかもしれんが。。