GPGPU#5

1 ：デフォルトの名無しさん：2010/08/15(日) 21:47:50 .net

GPGPUについて語りましょう

前スレ
GPGPU#4
http://hibari.2ch.net/test/read.cgi/tech/1255256230/l50

関連スレ
OpenCLプログラミング#1
http://hibari.2ch.net/test/read.cgi/tech/1228891105/l50
【GPGPU】くだすれCUDAスレ pert3【NVIDIA】
http://hibari.2ch.net/test/read.cgi/tech/1271587710/l50

参考リンク
総本山? gpgpu.org
http://www.gpgpu.org/
OpenCL
http://www.khronos.org/opencl/
NVIDIA CUDA
http://developer.nvidia.com/object/cuda.html
ATI Stream
http://developer.amd.com/gpu/ATIStreamSDK/Pages/default.aspx
GPUをCPU的に活用するGPGPUの可能性
http://pcweb.mycom.co.jp/articles/2005/09/06/siggraph2/

302 ：298：2014/01/26(日) 23:14:26.27 .net

>>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までは大丈夫ということなのでしょうか？

303 ：デフォルトの名無しさん：2014/01/26(日) 23:22:14.60 .net

>>299
GPGPUはメモリ転送のオーバーヘッドがないHSA(Huma)だよな
PCではど重い処理でない限りAMDのHSAがGPGPU処理の主流になるだろうな

304 ：デフォルトの名無しさん：2014/01/27(月) 00:23:03.99 .net

>>298
残念なお知らせ。
そのソースコードでは、GPUの演算時間ではなくGPUの呼び出し時間しか計測してないね。
「実際の演算時間」=「内部ブロック数」*(「内部ブロックの呼び出し時間」+「内部ブロックの演算時間」)だとすると、
「実際の演算時間」-「内部ブロックの演算時間」になっているはず。
ブロック数が充分大きければ誤差だけど、内部ブロック数が1のときは激速になってしまう。

まぁ、実際の運用ではCPUとGPUが並列に動作することを期待するからそれでもいいんだけどね。
いずれにしても、CPUぶん回すよりも手っ取り早いと思っていたら大間違いだよ。
それと、CUDAスレも宜しく。

305 ：298：2014/01/27(月) 00:47:11.47 .net

>>304
>そのソースコードでは
え！？　……つまり、
普通にtimeGetTimeかQueryPerformanceCounterとかを使えってことなんですか？
それとも、測定する位置が間違っているということなんですか？
>CUDAスレも宜しく
分かりました。次回以降はそちらにレスすることにします。

306 ：デフォルトの名無しさん：2014/01/27(月) 08:23:49.16 .net

>>304
何言ってんだ、こいつ？

307 ：デフォルトの名無しさん：2014/01/27(月) 21:34:08.90 .net

>>298 arraySizeが大きいと、CPU版すらStackOverflowになるよ。
http://pastebin.com/Av3YzTGs

308 ：307：2014/01/27(月) 21:39:43.68 .net

うっかり、166行目を「cudaStatus = cudaSetDevice(1);」にしちゃったので、適当に直しておいて。

309 ：デフォルトの名無しさん：2014/01/27(月) 23:30:12.43 .net

ローカルメモリを使う場合って確保しようとした容量が大き過ぎると
グローバルのほうへ確保されてしまうんだよね？
AMDのGCNはどれくらいまでローカルメモリがあるのか分からないんだけど
試行錯誤して調べるしかないのか

310 ：298：2014/01/27(月) 23:50:13.09 .net

>>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]らしいです。

311 ：デフォルトの名無しさん：2014/01/28(火) 01:09:12.72 .net

>>307
冗長なOpenCLに比べてやっぱりCUDAはスマートでいいな

312 ：デフォルトの名無しさん：2014/01/29(水) 01:59:06.39 .net

OpenCLのclEnqueueNDRangeKernelでカーネルを実行するときに
global_work_sizeとlocal_work_sizeに同じ値(256,256など)を入力すると
何もエラーは返されずにメモリの参照が壊れて？しまいclEnqueueReadBufferで
CPU側で読み取った値が全て0になってしまいます。

これは仕様なのでしょうか？

313 ：デフォルトの名無しさん：2014/02/25(火) 21:16:18.98 .net

visual studio 2013でCUDAが使えないからC++AMPでやるお！

314 ：デフォルトの名無しさん：2014/02/25(火) 21:43:30.35 .net

>>313
そのためだけにVS2012と2013使い分けてる俺……

315 ：デフォルトの名無しさん：2014/04/04(金) 10:44:13.17 ID:YtPgho8U.net

openCL始めたお(・∀・)ﾉ

316 ：デフォルトの名無しさん：2014/04/19(土) 12:16:56.16 ID:Firi/9oq.net

(・∀・)ﾉ ALU(IGP)のE2-2000はHD7770の1/50のパワーしかないが並列性はあるようだ。

317 ：デフォルトの名無しさん：2014/04/22(火) 04:44:14.02 ID:aREYskwN.net

AIDA64に測定メニューあるよな

318 ：デフォルトの名無しさん：2014/08/29(金) 13:33:23.65 ID:P9znXDYB.net

AMDとMS，GPU演算用途向けのコンパイラ「C＋＋ AMP v1.2」を発表
http://www.4gamer.net/games/032/G003263/20140828031/

319 ：デフォルトの名無しさん：2014/09/12(金) 04:54:39.38 ID:jvr90R5c.net

テキスト処理ってGPUで高速化できないものでしょうか
具体的には
Appache Solr
の検索処理が遅いのでなんとか高速化したいのですが

320 ：デフォルトの名無しさん：2014/09/12(金) 09:59:47.51 ID:cxN2yFh/.net

ボトルネックはメモリでしょう。

321 ：デフォルトの名無しさん：2014/09/16(火) 05:52:05.01 ID:padeH6x3.net

テキスト処理なんてわざわざGPUでやるよりSSE/AVXでやったほうが億倍マシ

322 ：デフォルトの名無しさん：2014/09/20(土) 00:49:25.20 ID:NyWaXORh.net

>>318
Linuxで動かすの以外と簡単みたいです
http://d.hatena.ne.jp/niitsuma/20080102/1411114218

323 ：デフォルトの名無しさん：2014/09/24(水) 15:13:16.24 ID:ltG1hZ24.net

OpenCLでプログラム組んでみたけど、CPUとGPUメモリのやり取りがネックになっているのか、思ったよりスピードが出ない

他の人はGPU利用するにあたってメモリのやり取りとか何か工夫している？

324 ：デフォルトの名無しさん：2014/09/24(水) 23:08:38.12 ID:psEUFh+R.net

そりゃ工夫するだろう。

325 ：デフォルトの名無しさん：2014/09/25(木) 13:28:12.29 ID:F8MulcGG.net

ごめん、どんな工夫してるか聞いてみたかったんだ

326 ：デフォルトの名無しさん：2014/09/25(木) 16:05:41.16 ID:Coq6ADbv.net

基本はメモリとのやりとりを少なくするって話でしょ
それ以上の個別の工夫を簡単に説明するのは難しいよね
ケーススタディしたいのならそういう本なり文献なり漁るべき

327 ：デフォルトの名無しさん：2014/09/25(木) 18:24:41.04 ID:RDrb9uGa.net

OpenCVのOpenCLバインディングのコードを参考にしたらいいんじゃないのかな

328 ：デフォルトの名無しさん：2014/09/25(木) 21:14:08.96 ID:YRvO5dcq.net

>>323
kaveri使えよ

329 ：デフォルトの名無しさん：2014/09/25(木) 22:16:45.04 ID:Vf7t0liy.net

OpenCLの1.1と1.2に後方互換性ありますか？

330 ：デフォルトの名無しさん：2014/09/27(土) 00:53:04.08 ID:SNKkkpyl.net

>>328
買えたら買ってるよ
メモリの転送の処理が要らなくなったらと思うと幸せな気分になれるよ

331 ：デフォルトの名無しさん：2014/10/14(火) 21:40:56.17 ID:noiOU3fL.net

kaveriってOpenCL使うとき、コピーせずにポインタ参照で渡していいって解釈でいいの？
最近GPGPUをやりはじめたばかりだから、的外れなことかもしれんが。。

332 ：デフォルトの名無しさん：2014/12/05(金) 22:27:50.90 ID:ri2FK3vXR

kaveriのアドレス共有はOpenCL2.0かHSAじゃないと対応できないから、今はメモリコピーするしかない

333 ：デフォルトの名無しさん：2015/04/12(日) 00:04:42.92 ID:g4+PudFo.net

Boost.ComputeあったらC++ AMPいらなくない？

334 ：デフォルトの名無しさん：2015/04/12(日) 19:07:28.43 ID:b726GPIq.net

どうだろう？

335 ：デフォルトの名無しさん：2015/04/24(金) 08:27:23.88 ID:A3qraRkp.net

http://stackoverflow.com/questions/20154179/differences-between-vexcl-thrust-and-boost-compute
一番高速なのはどれ？

336 ：デフォルトの名無しさん：2015/05/10(日) 00:43:24.29 ID:60tvXotD.net

vexclを少し使ってみたけど便利だ
あとはC++AMPみたいにradeonのドライバーの
バージョン上がると使えなくなったりしないなら

337 ：デフォルトの名無しさん：2015/06/23(火) 13:14:06.30 ID:AOM31ZzX.net

GPUの行列演算ライブラリってないですか？
具体的には特異値分解できるのを探してます

338 ：デフォルトの名無しさん：2015/06/23(火) 13:41:37.08 ID:DUXK3D31.net

>>337
機械学習スレで書いてた人かな？
ちゃんと調べてないけどMAGMAなら入ってるかも
http://icl.cs.utk.edu/magma/overview/index.html
{sdcz}gesvd はサポートしてるって書いてある

ただし、GPUカーネル内から直接呼びたい場合は使えないらしい
CPUからカーネル呼び出しする必要がある
（SC14時点の資料）

339 ：デフォルトの名無しさん：2015/06/26(金) 11:47:09.05 ID:JVzNXP51.net

>>338
ありがとうございます。

340 ：デフォルトの名無しさん：2015/07/08(水) 11:17:46.41 ID:i7xBLVJ6.net

最大固有値
最大固有ベクトル
だけを求めたい場合って、

341 ：デフォルトの名無しさん：2015/07/08(水) 11:18:15.73 ID:i7xBLVJ6.net

最大固有値
最大固有ベクトル
だけを求めたい場合って、べき乗法が最速でしょうか？

342 ：デフォルトの名無しさん：2015/07/10(金) 17:33:26.71 ID:QJI1WR+Q.net

万病に効く薬はないんやで

343 ：デフォルトの名無しさん：2015/10/11(日) 13:46:03.27 ID:9Az+Dnte.net

VS2015のc++amp仕様が変わった？
ＣＰＵで実行するrestrict(cpu)のマイクロソフトのサンプルコードがコンパイルエラーになる。

344 ：デフォルトの名無しさん：2016/01/29(金) 12:38:56.49 ID:VedX2j8l.net

>>343
C++AMPは終わりやな

345 ：125：2016/01/30(土) 08:09:23.91 ID:gCqMUv9A.net

マイクロソフトの開発ブログで

＞ Is C＋＋AMP dead ？

との質問にレスが無い。

モスさんどこ行った？

346 ：デフォルトの名無しさん：2016/06/15(水) 14:41:37.67 ID:d2Xou3GL.net

test

347 ：わたしはぐぷぐぷ派です：2016/12/28(水) 13:24:22.43 ID:6d1C8mET.net

は？ごぽごぽに決まってるだろ？って言われた．．．
頭ごなしに言ってくる人って何なんでしょうね

348 ：デフォルトの名無しさん：2018/05/23(水) 23:05:05.10 ID:Au5e7VGg.net

僕の知り合いの知り合いができたパソコン一台でお金持ちになれるやり方
役に立つかもしれません
グーグルで検索するといいかも『ネットで稼ぐ方法　モニアレフヌノ』

MYAIQ

349 ：デフォルトの名無しさん：2018/07/04(水) 23:04:08.53 ID:gFgZc5FG.net

2XM

350 ：デフォルトの名無しさん：2018/07/06(金) 12:37:07.14 ID:uTPDH9XV.net

MYAIQ

351 ：デフォルトの名無しさん：2019/09/05(木) 13:16:11.62 ID:va7N80X/.net

CUDAスレはあるのにROCmスレがない