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/
GPGPU#5
■ このスレッドは過去ログ倉庫に格納されています
1デフォルトの名無しさん
2010/08/15(日) 21:47:50253デフォルトの名無しさん
2012/10/12(金) 21:11:04.99 それこそSMTできるようにして同時実行させろって話じゃね
254デフォルトの名無しさん
2012/10/12(金) 23:36:36.25 そしてまた効率低くてそぎ落としか
設計レベルでループしてんじゃねえよw
設計レベルでループしてんじゃねえよw
255デフォルトの名無しさん
2012/10/13(土) 00:37:04.86 人間社会もGPGPU社会も一緒なんだな
256デフォルトの名無しさん
2012/10/13(土) 00:49:49.97 分岐というか、単純に帯域の問題じゃね。
257デフォルトの名無しさん
2012/10/13(土) 02:04:35.67 あのプレゼン見たけど、プログラミングに自信のある計算機寄りの人が、
てきとーに題材みつけてきて、計算機科学の研究として発表してる感じだった。
その結果が15%
てきとーに題材みつけてきて、計算機科学の研究として発表してる感じだった。
その結果が15%
258デフォルトの名無しさん
2012/10/13(土) 02:37:46.42 計算目的よりも、tubameを生かすための作業やね。まあ察しはつくがw
259デフォルトの名無しさん
2012/10/13(土) 03:38:41.472012/10/13(土) 04:49:03.94
>>250
働いてないんじゃなくて何かしらがボトルネックになってFMACの稼働率が抑えられてるのでは?
メモリ帯域が足りてる場合でも同時命令発行数の制約でload/store命令と積和命令を
同時に実行できない、とか
(これは京のVenusアーキテクチャにもある制約)
働いてないんじゃなくて何かしらがボトルネックになってFMACの稼働率が抑えられてるのでは?
メモリ帯域が足りてる場合でも同時命令発行数の制約でload/store命令と積和命令を
同時に実行できない、とか
(これは京のVenusアーキテクチャにもある制約)
261デフォルトの名無しさん
2012/10/13(土) 16:36:04.21 唐突ですまんけど、AMDはファイル名通りのこういう資料があるけど
ttp://developer.amd.com/gpu_assets/R700-Family_Instruction_Set_Architecture.pdf
nVIDIAはこういうの無いの?
ttp://developer.amd.com/gpu_assets/R700-Family_Instruction_Set_Architecture.pdf
nVIDIAはこういうの無いの?
262デフォルトの名無しさん
2012/10/13(土) 18:21:13.66263デフォルトの名無しさん
2012/10/13(土) 18:24:01.07264デフォルトの名無しさん
2012/10/13(土) 18:27:15.78 FLOPSに対する効率だから、アルゴリズム自体の並列化効率以上にはなりえないわけか。
気象系のアプリというものの並列化効率がどんくらいなのか知らないけど(なんとなく高そうではあるが)。
気象系のアプリというものの並列化効率がどんくらいなのか知らないけど(なんとなく高そうではあるが)。
265デフォルトの名無しさん
2012/10/13(土) 18:51:02.05 >>264
うーん、微妙に違うな。
今時のアーキテクチャだとメモリがすごく遅いってのは知ってるだろ?
ちなみにメモリの速度はバンド幅って言う。
それに対して演算速度はムーアの法則でどんどん上がってる(上がってた)
から、バランスが全然とれてないんだ。だから、普通のCPUはキャッシュ階層を
深くして(L1、L2、L3とか)なるべくデータの再利用をしてる。
つまり、データをメモリから持ってきたら、なるべくそれを
使い回して計算したいわけ。
で、ここで問題になるのが、浮動小数点演算(FLOP)と、データ転送量の比。
業界だと byte/flops とか言われる。言い方を変えると、どれくらいデータの
使い回しが効くかってこと。
続く
うーん、微妙に違うな。
今時のアーキテクチャだとメモリがすごく遅いってのは知ってるだろ?
ちなみにメモリの速度はバンド幅って言う。
それに対して演算速度はムーアの法則でどんどん上がってる(上がってた)
から、バランスが全然とれてないんだ。だから、普通のCPUはキャッシュ階層を
深くして(L1、L2、L3とか)なるべくデータの再利用をしてる。
つまり、データをメモリから持ってきたら、なるべくそれを
使い回して計算したいわけ。
で、ここで問題になるのが、浮動小数点演算(FLOP)と、データ転送量の比。
業界だと byte/flops とか言われる。言い方を変えると、どれくらいデータの
使い回しが効くかってこと。
続く
266デフォルトの名無しさん
2012/10/13(土) 18:52:43.77 承前
アプリの要求byte/flops高いと、せかっくデータを持ってきても
たいした計算をせずに、すぐに次のデータが必要になっちゃう。
データの使い回しが効かないんだ。これはイマドキの計算機にとってはキツい。
で、気象とか流体とかはそういうアプリなわけ。
だから、基本的に流体とかの計算は演算は余ってて、バンド幅がボトルネックになる。
これはアプリの特性だから仕方ない。
気象・流体屋さんは、未だに地球シミュレーター信者みたいな人が多い。
ちなみに、気象・流体系アプリは並列度だけで言ったらはものすごくあるよ。
ちなみにスレ的に言うと、GPUはGDDR5っていう容量が少ない代わりに
バンド幅が出る贅沢なメモリを搭載して、かつ大量のスレッドを使うことで
レイテンシを隠蔽してる。これがGPGPUが気象・流体計算に使える理由。
アプリの要求byte/flops高いと、せかっくデータを持ってきても
たいした計算をせずに、すぐに次のデータが必要になっちゃう。
データの使い回しが効かないんだ。これはイマドキの計算機にとってはキツい。
で、気象とか流体とかはそういうアプリなわけ。
だから、基本的に流体とかの計算は演算は余ってて、バンド幅がボトルネックになる。
これはアプリの特性だから仕方ない。
気象・流体屋さんは、未だに地球シミュレーター信者みたいな人が多い。
ちなみに、気象・流体系アプリは並列度だけで言ったらはものすごくあるよ。
ちなみにスレ的に言うと、GPUはGDDR5っていう容量が少ない代わりに
バンド幅が出る贅沢なメモリを搭載して、かつ大量のスレッドを使うことで
レイテンシを隠蔽してる。これがGPGPUが気象・流体計算に使える理由。
267264
2012/10/13(土) 19:29:29.68 なるほど、つまり気象系のアプリはコンピュートバウンドではなくメモリバウンドのものが多い、と。
AMDのAPUの使い道の話になったときに、DDR3という制約に悲観的になってる人が少なそうだったから
世の中にはメモリバウンドのものは少数派なのかなあとも思ってたけど普通にあるんですね。
AMDのAPUの使い道の話になったときに、DDR3という制約に悲観的になってる人が少なそうだったから
世の中にはメモリバウンドのものは少数派なのかなあとも思ってたけど普通にあるんですね。
268デフォルトの名無しさん
2012/10/13(土) 19:43:21.60 バンド幅と言うか、物理原理的にDRAMのセルは応答が現代のCPUの演算速度に対してずっと遅い。
だから、メモリセルを並列にして同時に読み出し・書き込みするしかないのが根底にある。それをやれDDRxだの、
何ビット同時にアクセス=複数セルに同時アクセスでなんとかしてるね。
気象系や気体がどうのって言うよりも、規則正しく計算点を並べて偏微分方程式の数値近似式を一斉に走らせる
=SIMDだのシェーダだの言ってる演算回路がみんな一斉に同じ計算式をやる場合が最大性能が出る。
三次元の方がそりゃ、並列度は高いだろうね。境界条件とか面倒くさそうだが。
こういう根本的なところはもう全然進歩しないなw
だから、メモリセルを並列にして同時に読み出し・書き込みするしかないのが根底にある。それをやれDDRxだの、
何ビット同時にアクセス=複数セルに同時アクセスでなんとかしてるね。
気象系や気体がどうのって言うよりも、規則正しく計算点を並べて偏微分方程式の数値近似式を一斉に走らせる
=SIMDだのシェーダだの言ってる演算回路がみんな一斉に同じ計算式をやる場合が最大性能が出る。
三次元の方がそりゃ、並列度は高いだろうね。境界条件とか面倒くさそうだが。
こういう根本的なところはもう全然進歩しないなw
2012/10/13(土) 20:27:18.96
bytes/flopsに関して言えば現状GPUはCPUより低いです。
ただbyte/sが数倍高くて、flopsが更に数倍高いというだけの話で。
Haswellあたりで、電力当たりFLOPS数はGPUとCPUは互角くらいになるのではという話も出ていますな。
ただbyte/sが数倍高くて、flopsが更に数倍高いというだけの話で。
Haswellあたりで、電力当たりFLOPS数はGPUとCPUは互角くらいになるのではという話も出ていますな。
270デフォルトの名無しさん
2012/10/13(土) 20:32:24.78 TSVって、GPUみたいな爆熱チップでもできるの?
271デフォルトの名無しさん
2012/10/14(日) 11:35:52.75272デフォルトの名無しさん
2012/10/14(日) 12:44:22.63 DRAM内部のクロック
PC133 133MHz
DDR3-1600 200MHz
この間十数年・・・MRAMならもっと上がるかな?
PC133 133MHz
DDR3-1600 200MHz
この間十数年・・・MRAMならもっと上がるかな?
273デフォルトの名無しさん
2012/10/20(土) 09:38:39.11 GNU MPのGPGPU対応って、あんまりやってる人が居ないな。
2012/10/20(土) 09:42:31.34
依存関係があってなかなか並列化しにくいからね。
もちろん並列度の高いコンピュータに適した数式に置き換えればいいんだが
CPU側もAVX, FMAなどでGPGPUの電力効率の優位も揺らぎつつあるからね
もちろん並列度の高いコンピュータに適した数式に置き換えればいいんだが
CPU側もAVX, FMAなどでGPGPUの電力効率の優位も揺らぎつつあるからね
275デフォルトの名無しさん
2012/11/17(土) 23:26:42.82 C++AMP調べてるとこなんだけど
textureでピクセル間を補間する方法がわからない。
ひょっとして自前でやらないといけないの?
「C++ AMP API は、サンプリング テクスチャのフィルター処理機能を提供されません。」
とはこのことなのかな。
textureでピクセル間を補間する方法がわからない。
ひょっとして自前でやらないといけないの?
「C++ AMP API は、サンプリング テクスチャのフィルター処理機能を提供されません。」
とはこのことなのかな。
276デフォルトの名無しさん
2012/11/18(日) 21:58:01.39 c++ampにfloat_3のmadやfmaが無いね。 short vector typeのmadやsin, cos などの超越関数も用意してもらいたい。 swizzlingはできるみたい。 OpenCLの方がコードをコンパクトに書けるかな。次期版期待してます。
277デフォルトの名無しさん
2013/02/17(日) 08:49:52.18 何故かsea islandのISAが公開
ttp://developer.amd.com/tools/heterogeneous-computing/amd-accelerated-parallel-processing-app-sdk/documentation/
ttp://developer.amd.com/tools/heterogeneous-computing/amd-accelerated-parallel-processing-app-sdk/documentation/
278デフォルトの名無しさん
2013/02/17(日) 13:26:30.82 いつもに比べると異常に早いがやっぱsouthernislandのリネームなのかね
279デフォルトの名無しさん
2013/02/18(月) 13:29:16.68 ディスクリートとしては出なくても
次世代APUのGCNはsea islandなのかね。
ISA見る限り、SIから確かに機能アップされている。
一番大きいのはFLAT メモリアクセス命令の追加か。
cudaのUVAみたいなもの
次世代APUのGCNはsea islandなのかね。
ISA見る限り、SIから確かに機能アップされている。
一番大きいのはFLAT メモリアクセス命令の追加か。
cudaのUVAみたいなもの
280デフォルトの名無しさん
2013/03/27(水) 13:39:55.44 SL#(えすえるしゃーぷ)とは、GPUで実行されるプログラマブルシェーダーを、超高級言語である
C#で書けてしまうという夢のようなオープンソースのフレームワークである。
http://monobook.org/wiki/SL_Sharp
C#で書けてしまうという夢のようなオープンソースのフレームワークである。
http://monobook.org/wiki/SL_Sharp
281デフォルトの名無しさん
2013/03/30(土) 11:07:23.56 >>280
まじかよ、ニンジャコンパイラが全部解決してくれるのか?
まじかよ、ニンジャコンパイラが全部解決してくれるのか?
282デフォルトの名無しさん
2013/04/12(金) 18:59:55.71 超高級言語w
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 なんか来てた。
ttp://www.phoronix.com/scan.php?page=news_item&px=MTQ3NDU
ttp://www.phoronix.com/scan.php?page=news_item&px=MTQ3NDU
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だとバグったような速度が簡単に出るから別世界感が凄い
302298
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.60304デフォルトの名無しさん
2014/01/27(月) 00:23:03.99 >>298
残念なお知らせ。
そのソースコードでは、GPUの演算時間ではなくGPUの呼び出し時間しか計測してないね。
「実際の演算時間」=「内部ブロック数」*(「内部ブロックの呼び出し時間」+「内部ブロックの演算時間」)だとすると、
「実際の演算時間」-「内部ブロックの演算時間」になっているはず。
ブロック数が充分大きければ誤差だけど、内部ブロック数が1のときは激速になってしまう。
まぁ、実際の運用ではCPUとGPUが並列に動作することを期待するからそれでもいいんだけどね。
いずれにしても、CPUぶん回すよりも手っ取り早いと思っていたら大間違いだよ。
それと、CUDAスレも宜しく。
残念なお知らせ。
そのソースコードでは、GPUの演算時間ではなくGPUの呼び出し時間しか計測してないね。
「実際の演算時間」=「内部ブロック数」*(「内部ブロックの呼び出し時間」+「内部ブロックの演算時間」)だとすると、
「実際の演算時間」-「内部ブロックの演算時間」になっているはず。
ブロック数が充分大きければ誤差だけど、内部ブロック数が1のときは激速になってしまう。
まぁ、実際の運用ではCPUとGPUが並列に動作することを期待するからそれでもいいんだけどね。
いずれにしても、CPUぶん回すよりも手っ取り早いと思っていたら大間違いだよ。
それと、CUDAスレも宜しく。
305298
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.90308307
2014/01/27(月) 21:39:43.68 うっかり、166行目を「cudaStatus = cudaSetDevice(1);」にしちゃったので、適当に直しておいて。
309デフォルトの名無しさん
2014/01/27(月) 23:30:12.43 ローカルメモリを使う場合って確保しようとした容量が大き過ぎると
グローバルのほうへ確保されてしまうんだよね?
AMDのGCNはどれくらいまでローカルメモリがあるのか分からないんだけど
試行錯誤して調べるしかないのか
グローバルのほうへ確保されてしまうんだよね?
AMDのGCNはどれくらいまでローカルメモリがあるのか分からないんだけど
試行錯誤して調べるしかないのか
310298
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.17ID:YtPgho8U openCL始めたお(・∀・)ノ
316デフォルトの名無しさん
2014/04/15(火) 02:32:13.65ID:vGWbAtXL (・∀・)ノ CPUの300倍くらいの性能が出たお!
比較したCPUはE2-2000っていうCPU+GPU=APUだけど全くGPUとしての機能をもってないのでガッカリしたお。
比較したCPUはE2-2000っていうCPU+GPU=APUだけど全くGPUとしての機能をもってないのでガッカリしたお。
317デフォルトの名無しさん
2014/04/19(土) 12:16:56.16ID:Firi/9oq (・∀・)ノ ALU(IGP)のE2-2000はHD7770の1/50のパワーしかないが並列性はあるようだ。
318デフォルトの名無しさん
2014/04/22(火) 04:44:14.02ID:aREYskwN AIDA64に測定メニューあるよな
319デフォルトの名無しさん
2014/08/29(金) 13:33:23.65ID: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.38ID:jvr90R5c テキスト処理ってGPUで高速化できないものでしょうか
具体的には
Appache Solr
の検索処理が遅いのでなんとか高速化したいのですが
具体的には
Appache Solr
の検索処理が遅いのでなんとか高速化したいのですが
321デフォルトの名無しさん
2014/09/12(金) 09:59:47.51ID:cxN2yFh/ ボトルネックはメモリでしょう。
322デフォルトの名無しさん
2014/09/16(火) 05:52:05.01ID:padeH6x3 テキスト処理なんてわざわざGPUでやるよりSSE/AVXでやったほうが億倍マシ
323デフォルトの名無しさん
2014/09/20(土) 00:49:25.20ID:NyWaXORh324デフォルトの名無しさん
2014/09/24(水) 15:13:16.24ID:ltG1hZ24 OpenCLでプログラム組んでみたけど、CPUとGPUメモリのやり取りがネックになっているのか、思ったよりスピードが出ない
他の人はGPU利用するにあたってメモリのやり取りとか何か工夫している?
他の人はGPU利用するにあたってメモリのやり取りとか何か工夫している?
325デフォルトの名無しさん
2014/09/24(水) 23:08:38.12ID:psEUFh+R そりゃ工夫するだろう。
326デフォルトの名無しさん
2014/09/25(木) 13:28:12.29ID:F8MulcGG ごめん、どんな工夫してるか聞いてみたかったんだ
327デフォルトの名無しさん
2014/09/25(木) 16:05:41.16ID:Coq6ADbv 基本はメモリとのやりとりを少なくするって話でしょ
それ以上の個別の工夫を簡単に説明するのは難しいよね
ケーススタディしたいのならそういう本なり文献なり漁るべき
それ以上の個別の工夫を簡単に説明するのは難しいよね
ケーススタディしたいのならそういう本なり文献なり漁るべき
328デフォルトの名無しさん
2014/09/25(木) 18:24:41.04ID:RDrb9uGa OpenCVのOpenCLバインディングのコードを参考にしたらいいんじゃないのかな
329デフォルトの名無しさん
2014/09/25(木) 21:14:08.96ID:YRvO5dcq >>324
kaveri使えよ
kaveri使えよ
330デフォルトの名無しさん
2014/09/25(木) 22:16:45.04ID:Vf7t0liy OpenCLの1.1と1.2に後方互換性ありますか?
331デフォルトの名無しさん
2014/09/27(土) 00:53:04.08ID:SNKkkpyl332デフォルトの名無しさん
2014/10/14(火) 21:40:56.17ID:noiOU3fL kaveriってOpenCL使うとき、コピーせずにポインタ参照で渡していいって解釈でいいの?
最近GPGPUをやりはじめたばかりだから、的外れなことかもしれんが。。
最近GPGPUをやりはじめたばかりだから、的外れなことかもしれんが。。
333デフォルトの名無しさん
2015/04/12(日) 00:04:42.92ID:g4+PudFo Boost.ComputeあったらC++ AMPいらなくない?
334デフォルトの名無しさん
2015/04/12(日) 19:07:28.43ID:b726GPIq どうだろう?
335デフォルトの名無しさん
2015/04/24(金) 08:27:23.88ID:A3qraRkp336デフォルトの名無しさん
2015/05/10(日) 00:43:24.29ID:60tvXotD vexclを少し使ってみたけど便利だ
あとはC++AMPみたいにradeonのドライバーの
バージョン上がると使えなくなったりしないなら
あとはC++AMPみたいにradeonのドライバーの
バージョン上がると使えなくなったりしないなら
337デフォルトの名無しさん
2015/06/23(火) 13:14:06.30ID:AOM31ZzX GPUの行列演算ライブラリってないですか?
具体的には特異値分解できるのを探してます
具体的には特異値分解できるのを探してます
338デフォルトの名無しさん
2015/06/23(火) 13:41:37.08ID:DUXK3D31 >>337
機械学習スレで書いてた人かな?
ちゃんと調べてないけどMAGMAなら入ってるかも
http://icl.cs.utk.edu/magma/overview/index.html
{sdcz}gesvd はサポートしてるって書いてある
ただし、GPUカーネル内から直接呼びたい場合は使えないらしい
CPUからカーネル呼び出しする必要がある
(SC14時点の資料)
機械学習スレで書いてた人かな?
ちゃんと調べてないけどMAGMAなら入ってるかも
http://icl.cs.utk.edu/magma/overview/index.html
{sdcz}gesvd はサポートしてるって書いてある
ただし、GPUカーネル内から直接呼びたい場合は使えないらしい
CPUからカーネル呼び出しする必要がある
(SC14時点の資料)
339デフォルトの名無しさん
2015/06/26(金) 11:47:09.05ID:JVzNXP51 >>338
ありがとうございます。
ありがとうございます。
340デフォルトの名無しさん
2015/07/08(水) 11:17:46.41ID:i7xBLVJ6 最大固有値
最大固有ベクトル
だけを求めたい場合って、
最大固有ベクトル
だけを求めたい場合って、
341デフォルトの名無しさん
2015/07/08(水) 11:18:15.73ID:i7xBLVJ6 最大固有値
最大固有ベクトル
だけを求めたい場合って、べき乗法が最速でしょうか?
最大固有ベクトル
だけを求めたい場合って、べき乗法が最速でしょうか?
342デフォルトの名無しさん
2015/07/10(金) 17:33:26.71ID:QJI1WR+Q 万病に効く薬はないんやで
343デフォルトの名無しさん
2015/10/11(日) 13:46:03.27ID:9Az+Dnte VS2015のc++amp仕様が変わった?
CPUで実行するrestrict(cpu)のマイクロソフトのサンプルコードがコンパイルエラーになる。
CPUで実行するrestrict(cpu)のマイクロソフトのサンプルコードがコンパイルエラーになる。
344デフォルトの名無しさん
2016/01/29(金) 12:38:56.49ID:VedX2j8l >>343
C++AMPは終わりやな
C++AMPは終わりやな
345125
2016/01/30(土) 08:09:23.91ID:gCqMUv9A マイクロソフトの開発ブログで
> Is C++AMP dead ?
との質問にレスが無い。
モスさんどこ行った?
> Is C++AMP dead ?
との質問にレスが無い。
モスさんどこ行った?
346デフォルトの名無しさん
2016/06/15(水) 14:41:37.67ID:d2Xou3GL test
347わたしはぐぷぐぷ派です
2016/12/28(水) 13:24:22.43ID:6d1C8mET は?ごぽごぽに決まってるだろ?って言われた...
頭ごなしに言ってくる人って何なんでしょうね
頭ごなしに言ってくる人って何なんでしょうね
348デフォルトの名無しさん
2018/05/23(水) 23:05:05.10ID:Au5e7VGg 僕の知り合いの知り合いができたパソコン一台でお金持ちになれるやり方
役に立つかもしれません
グーグルで検索するといいかも『ネットで稼ぐ方法 モニアレフヌノ』
MYAIQ
役に立つかもしれません
グーグルで検索するといいかも『ネットで稼ぐ方法 モニアレフヌノ』
MYAIQ
349デフォルトの名無しさん
2018/07/04(水) 23:04:08.53ID:gFgZc5FG 2XM
350デフォルトの名無しさん
2018/07/06(金) 12:37:07.14ID:uTPDH9XV MYAIQ
351デフォルトの名無しさん
2019/09/05(木) 13:16:11.62ID:va7N80X/ CUDAスレはあるのにROCmスレがない
■ このスレッドは過去ログ倉庫に格納されています
ニュース
- 「おこめ券は米以外の食品も買える。効果的な活用を」 地元で農水相 [山形県] [少考さん★]
- 高市首相の答弁書に「台湾有事答えない」と明記 存立危機発言当時 ★11 [蚤の市★]
- 【速報】「女芸人No.1決定戦 THE W」9代目女王にニッチェ! 7年ぶり3度目で悲願の優勝 [牛丼★]
- 【芸能】『女芸人No.1決定戦THE W』 粗品が最後にバッサリ「優勝賞金1000万円にしてはレベル低い大会」 [冬月記者★]
- 【沖縄】開業4ヵ月でこれは…“国民の税金”投入の『ジャングリア沖縄』で見た衝撃的な光景と、モチベーションが低い一部スタッフの現状 [ぐれ★]
- 今年の流行語大賞 『働いて働いて働いてまいります』が受賞で不快感… 過労自殺の遺族らが会見「家族にむち打つような行為だ」 [冬月記者★]
- __イスラエル紙、ベネズエラ政権交代をトランプに促したのはイスラエル、影響力の大きさを示唆 [827565401]
- __ブルガリア、Z世代の抗議が増税予算と汚職政治への怒りへ、政権が崩壊、若者を無視する政治への警告 [827565401]
- クズ「勉強頑張らなかった奴は一生DQNと一緒に肉体労働しろ」☚勉強頑張れるのも環境と巡り合わせなんだが? [783475554]
- 俺の口癖が「へー」「そう」「どうも」なんだが
- キャッシュレスに対応してない店、手数料が問題ならその分値上げすればいいじゃない、現金の管理や手数料、両替もただじゃない [943688309]
- 【正論】検察「山上よ、どんな事情があろうと暴力が許されない」 [442080748]
