GPGPU#5

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/

＞998 ：デフォルトの名無しさん [↓] ：2010/08/15(日) 21:51:01
＞と言いつつAgeiaの中の人も今じゃAMDにいるからなぁ
＞とんだ詐欺師なのかねあの人

金です。

nvにとっちゃすでに用済みで、要らない子

専用設計とはいえPPUは58gflopsしかないんだが

基本的には、GPGPUが得意な処理を "適切なサイズ" に並列分割して
その分割された小包の集団をどかっとCUDAに押し込んでやると、分割が上手ければ
それなりに速く結果が出る。ただ、GPGPUで効率が出る並列化は簡単ではない。Larrabeeがこけたのもここ。

しかもC++のCUDA方言は不思議挙動だったりで、技術者がCUDAに習熟して十分な速度が
出せるようになるまでの時間を考えると、結構経費がかかる。だから、相当大きな話、というか
CUDAのX86@Intel CPUに対するワットパフォーマンス優位性が技術者の勉強代をカバーできる規模で無いと
わざわざわけわからん方言を勉強したくない。しかも、この方言は、いつまで有効かも怪しい。

だから、ほとんどの用途では、Nehalem-Ex とか、速いCPU乗せたマシンを増やした方が良い。
他のプログラムが、"全部" 速くなりますからねｗ

今後のCPUコアの高速化が鈍化するから
その対策として出てきたのがCPUのマルチコア化と
グラボのGPGPUとしての活用なわけで・・・

大部分の人には上位CPUなんて必要ないのと同様に
大部分のアプリにもGPGPUなんて必要ない。

６コアもGPGPUも本当に必要な人・アプリが使えばいいだけ

大部分って、静的WEBページを見るだけのユーザーのことか？ｗ
そんなもん無視でいいだろｗ

WEBブラウズだろうがオフィスアプリだろうが
音楽・動画再生だろうがゲームだろうが大部分のアプリには
高価な上位CPUも高速なGPGPUも必要じゃないだろ。

そこそこヘビーな自分でも4コア(疑似8コア)や
1TFLOPS以上のGPUをフル活用できるのは全PC作業の１割程度だし

LAMEとかiTunesとかで、GPGPUが効けばもっと広がると思うんだけど…
やる気無いですよねぇ

やる気程度で速くなってくれるなら今ごろみんな取り掛かってるだろうよ

LAME（音声の非可逆圧縮）程度じゃ処理が軽すぎるし
条件分岐も少なくないからCPUで計算したほうがいい。

映像編集ソフトですらエフェクト処理がメインでエンコードにはGPGPUが使えなかったりする。

iTunes（映像再生ソフト）にGPGPUとして使うなんて問題外。
大人しくOpenGLやDirect2DなんかでGPUとして活用すべき。
リアルタイムで映像にエフェクト処理を加えながら再生したいなら別だがiTunesの仕事じゃないｗ

ATI Stream使ってエンコードして負荷軽減してるソフトなかったけか？

>>11
PowerDirector？

エンコードに使うなら売りは速度ではなく品質にすべき。
データ転送がボトルネックなのだから
単位データあたりの演算量を増やさなきゃメリットが無い。

演算量が増えてもプログラムのフローが複雑になるようでは

>>14
どんだけ複雑になったって、大量に並列実行できればGPGPUにとってアドバンテージがある。
データに対して演算量が少なすぎると転送や処理待ちばかりになってパフォーマンスが上がらない。
だから問題は複雑性よりもデータの相互依存性とデータに対する演算量の少なさ。

複雑性ってなに？

文脈から鑑みるに、プログラムの複雑さじゃないの？

もっと端的に言ってしまえば分岐命令の数

この場合、相互依存性と複雑性は同義だと思うけどね。

>>18
この場合は違う。
>14 はそのつもりで発言しているのかも知れないが、>13 は違う。

そう言い切るのなら、どう違うかまでを説明せんといかんよ。

>>13って8x8DCTを4x4DCTにするみたいな話でしょ？
演算回数は増えるがGPUなら並列数を増やせる感じで

S|A What is AMD's Northern Islands?　A look at what is coming in October
http://www.semiaccurate.com/2010/09/06/what-amds-northern-islands/

ごめんSIって言ってたけど実はNIだったよ。えへ。
だから今度出るのはHD6000ファミリーはNIね。
32nmでNIテープアウトしてたけど40nmで出すよ。
コアは○○な感じで、アンコアは××な感じで強化してるよ。
なんでチップがEvergreenより10-15%大きくなるよ。
リリーススケジュールは10月12日にイベントで25日前後に店頭並ぶよ。
まずはAMDの穴の開いてる$175-250帯のHD6700から始めるよ。
次にHD6800、HD6900、年初にローエンド、28nmまでこのラインナップだよ。

HD6000出たら緑チームはHD5000よりコスト高いのに値下げしなくちゃだし、それでなくても冷め切ったセールスにもろ影響しちゃうよ。
だって、トップエンドは価格維持でHD5000は下がり始めるしね。
Nvidiaの夢と希望を打ち砕いちゃうね。
打つ手もないしね。
AMDはDX11のトーナメント1回戦をHD5000で勝利して、第2回戦もHD6000で勝利しちゃて、Nvidiaには財務的にもパフォーマンスでっかいマージンを取っちゃうよ。

28nmまではNvidiaにチャンスはないね。

余所に作らせたGPUを使ったプログラムが、CUDA部分でメモリリークくさいエラーを吐いてまともに動かないんですが、
窓から投げ捨てるべきでしょうか？

窓から投げるべき

証拠資料を作ろうとしても、「いつ止まるか」の再現性が微妙
やっぱり実績の無いハウスに委託したのが間違いだったか・・・

メモリの確保と解放を繰り返しているんじゃないかな。
弊社ではソースがあればデバッグも承りますw

ソースないっす・・・
その辺だけはしっかりしているという・・・

ていうか、ウチ（受け入れ側）のマネージャーが完全に「ドモホルンリンクル」で
どんなゴミを渡されても「努力あるのみ」とかの類の精神論を吐いて話にならないし




どっか、受託開発や納入後の展開方法についての客観的な評価をしてくれる
コンサルタントはないですかね・・・

CUDAでソースなし納品はありえんやろ
いつバージョンアップでバイナリが動かなくなってもおかしくないのに

>>28
コストの問題だろ？
ソースを要求すると価格が上がる

いや、将来動かない可能性が低くないのにコストカットされてもｗ

将来動かなくなる可能性が高いから値切るんだろうが

gpgpuを使用した場合、 CPUの性能はどの程度影響しますか？
teslaを用いた計算機を導入しようとしているのですが、i7-980xにするかi7-930にするか
迷っています。

CUDAやOpen CL以外のCPUコードの実行速度にモロに影響する。
他にもGPGPU用中間コードのコンパイルにも影響するが誤差範囲。

聞きたいのはCPUの性能によってGPUの性能が変わるかどうかじゃないの

初心者なんですけどフリーソフトでATI技術に対応してて
MP4に変換できるソフトってありますか？
あとRADEONのカードってエンコードなら値段と性能みてどれがコスパいいですか？

板違いです
ソフト板か自作板、DTV板へgo

板違いです。

ここは「ATI技術に対応しててMP4に変換できるソフト」を作る側の板です。

caymanは期待できそうだな。

GPGPU使って何かしたいけどこれっていう何かが見つからないのー
Actor とか Map Reduce とか上位層で駆逐されてしまうねん

俺はいっぱいアイデアあるけどな。

あら、気になるじゃない。聞きたいわ

突然申し訳ありません
cudaやってるんですけど・・・
カーネル関数起動させるところでエラーが出てしまいます
サンプルコードでアウトなんです
考えられる可能性を挙げていただきたいです
エロい人助けてください

ちなみに、
win7professional32
グラボ１：８６００gs（出力用）
グラボ２：４６０gtx(→cuda)
開発環境：visual studio 2008

質問あればできるものはすべて答えますんでよろしくお願いします

>>42
エラーメッセージぐらいのせろやカスが

密かにevergreenのISA仕様書が更新されているな。
ttp://developer.amd.com/gpu/ATIStreamSDK/assets/AMD_Evergreen-Family_Instruction_Set_Architecture.pdf
メモリアクセス周りの挙動について言及されているのがなかなか面白い。
コアレス化が余り重要じゃないという話がどういう意味か分かる。
要は、アーキテクチャ的に1スレッドが複数のメモリアクセス命令を同時発行可能で
1wavefront単位で発行された複数のメモリアクセス命令の間だけ
キャッシュ無しアクセスでもキャッシュが有効になっているから
複数のメモリアクセス命令間でコアレス化と同様の効果が得られるらしい。

CAL+ILの情報が少ないので、書き込みがあるだけで嬉しい。

>>43
スイマセン


CUDA error: Kernel execution failed


コンパイルはできていて、
他のマシンでは同じソースコードのプログラムは動かせます
原因は何なんでしょうか

>>46
カーネルの起動に失敗したというのだから、起動要件を満たしていないのだろう。
まさかとは思うが、CUDA用のドライバをインストールしていないというオチではあるまいな。

>>47
ドライバは入っています

それと、今日起動することに成功しました。
main()の変数宣言のすぐ後に、

cudaSetDevice(1);

を記述したら、それで通りました。
なぜ起動できなかったかは分かりません。
今可能性を探っているのですが、
タイムアウトが起こったのかもしれないと考えています。

なんだ、動いたなら後はcudaスレへ。

>>48
CUDAの命令はよく分からんけど、名前から察するに単純に処理するGPU指定してなかっただけじゃ。
8400GSだってCUDA対応なんだし。

>>48 >>50
デバイスを指定しなければデフォルトのデバイス0、8600GSで実行されるはずだけど。
Capability 2.0以降限定の関数を呼び出しているサンプルコードだったとか？
まあCUDAスレあるし、そっちでやるべきかな。

A Fast GEMM Implementation On a Cypress GPU
http://galaxy.u-aizu.ac.jp/trac/note/wiki/Fast_GEMM_Implementation_On_Cypress

Cayman GPUではスーパーファンクションユニットが削除されて５VLIWプロセッサーから
４VLIWプロセッサーになるとのことですが、現在のCALでサポートされているsin/cos等の
超越関数命令は、自分で多項式近似計算をしろと言うことなのでしょうか？

>>53
今までtレーンが担当してきた命令は
xyzwの複数レーンで1命令を実行するように変更されている。

で、超越関数は3SPを使った1命令で実行される。

>>54
53 です。　CORDICやチェビシェフ多項式をWeb上で漁っていたのですが
安心しました。

チェビシェフは自分で作るもんじゃろ

CORDICって条件分岐ばっかなのでGPGPUには不向きだという先入観があるんだけどどうなの？

それは全部ソフトウェアでやったら、の話だろ。

おやお久しぶり。
ソフトウェア無線専用ハード（微妙に矛盾？）でCORDICを使ってるとか話には聞いたことがありますな。

あと自分で実装してみて気づいたんだけど、分岐と言ってもある数を足すか引くかなので、
分岐しないようにしてビット操作に落とせるんですよね。

GPUは分岐が苦手とはいっても、単純なプレディケートに落とし込めるものならむしろ効率がいいくらいです。
GPUは同一ワープ内で命令ストリームを共有してますから同じ方向にしか進めない。
Cでいうif-elseは一見分岐だけど、GPUではプレディケート情報によって実行・不実行（あるいは結果に反映させない）を
選択する単一の流れに展開されています。

プレディケート自体はそんなに重たくないです。
むしろ分岐先が増えると増えた分だけ処理時間が増えるだけで。

A's Video Converterって、10/31付けで配布サイトが閉鎖されてるな
もう手に入らんの？(´・ω・｀)今日HD5670GETしたのに・・・

移転してるよ
http://bluesky23.yu-nagi.com/

>>62
うぉ！マジでありがとう
ググっても見つけられなかったんだ(T-T)ｳﾚｼｲ

プレディケートは現行Intel系GPUでは使い物になりません
GPGPU向けに機能追加されたSandy Bridge GPUコアの登場を待ちましょう。

そもそもGPGPUできないし

OpenCL対応するんでしょ＞＞Sandy Bridge GPUコア

L3使えるから規模の割には速いかもな

主にGPUコアで回してるかAVXつこてるかはインテル任せ
だんごに好かれてしまったからGPGPUも端ッパの技術バリエーションの一つに転落決定だな

AVXでGPUくらい速くできるなら寧ろ大歓迎だが。
ただしアセンブリ言語で書くのは嫌。

そんなもん規模しだいだろ

Sandy Bridgeの1EU＝4Way-FMACと仮定しても、まだCPU（AVX）のほうが速いですから

ION乗ってるノートでXP入れました！！！！
これでCUDAできますよね？
俺の夢かなえられますよね？

ひゃっはあああああああああああああああああああ

IONのCUDAベンチ＆レビューよろしく

GPU
グレートプログラマー初春

P=>パイパン

DSPやFPGA叩いて高速度・複雑なシステム作るよりは
CPU+GPU叩いて作ったほうがはるかにましだがなあ。生産効率が桁違いだわ。

e?

>>76
今までDSPやFPGAでやっていた事をCPU+GPUでできる用途ってどんなものがあるの？
考えたけど一つも思い浮かばなかった。

具体的なカテゴリは勘弁だが、サンプリングしたデータをフィルタで処理して画面に表示
みたいな処理では、GPUでの代替はかなり強烈だよ。
FPGAだとたかだか数百タップの複素FIRフィルタを40〜50MHzの動作速度でさばくのにも
現状だと5〜15万ぐらいのデバイスがいるし。

大量生産するものなら処理をチップ化して安くあげちゃうんだろうけど、
俺のとこみたいな数のでない無線通信製品だとGPGPUはかなり魅力的。
たぶん、画像検査装置みたいな分野でもGPGPUは強力だと思う。

DX11世代だと、本当に何でも出来そうだよな。

Linux版のAMD APP 2.4にCALのサンプルが付属していないのですが、
Windows版は付属していますか？

してません。
それどころかCALは(IL含め)2.5で死滅。
かわりにLLVM IR使え。
そんな感じです。

>>82
ありがとうございます。
そうですか。。
せっかくIL習得したところなんですが、困りましたね。

http://developer.amd.com/gpu/AMDAPPSDK/assets/AMD_APP_SDK_Release_Notes_Samples.pdf
ここにはCALのsampleはcalといディレクトリにあるとかいてますが、
旧バージョンから修正されてないだけでしょうか？

これからGPGPUを勉強する場合、どれを勉強しておくのが良いのでしょうか？
無難という意味では、OpenCLですか？

ソリューションは結局、問題や環境が決定するもの。

CUDAでいいんでない

>>85
ソフトをやるのかハードをやるのか？
ソフトの場合は上位をやるのか下位をやるのか？

５００万個の３×３行列の固有値を
（１）CPU Intel Q9450 (4 posix threads) GCC 4.4.3　（最適化無し、１２５万個／スレッド）
（２）ATI HD4870 ＋ OpenCL (AMD APP SDK 2.4) （最適化無し）

で計算させてみた（行列は正定値実対称の素直な行列）。

ハードはCPU、GPUともに定格で使用。OSはUbuntu x64 10.4 LTS , AMDドライバはCatalyst 11.5
GCC4.4.3とOpenCLで使用したソースコードは略同じものを使用（相違点は OpenCL側コードに__global 指定が付いた程度）
時間測定はC言語側の計算ルーチン呼び出し元でgettimeofday()を使用してマイクロ秒単位で測定。

(1) Q9450 4スレッド
------------------------------------------------------------------
Posix thread creation : 0.000124(sec)
Thread #1 Exection time: 5.992(sec)
Thread #2 Exection time: 6.08448(sec)
Thread #3 Exection time: 5.9132(sec)
Thread #4 Exection time: 5.91843(sec)

Total Exection time : 6.08452(sec)　<ーースレッド中の最大値 + α

(2) HD4870 (800スレッド？）
------------------------------------------------------------------
GPU Kernel Compile : 1.6e-05(sec)
GPU Kernel Build 　 : 5.02891(sec)
GPU Kernel Creation : 6e-06(sec)
GPU Kernel Set Args. : 2e-06(sec)
GPU Kernel Execution : 4e-05(sec) --- clEnqueueNDRangeKernel ()を挟むgettimeofday()の時間差
Memory mapping(READ MODE) : 5.38053(sec)
＜この間でデータ読み出し＞
Memory UnMapping(from READ MODE) : 0.020504(sec)

OpenCLソースのビルド＆結果データの読み出しまで含めるとGPUが1.7倍遅いが計算実行時間の単純比較だと

6.08452 / 4.0E-5 = 1.5E5 = 15万倍速い！　と言う結果になりました。 いくら何でも１５万倍は速すぎのような気が・・・・(^_^;;)

以上。

>90 (自己レス)

ごめんなさい。　ミススペルしてました。

(1) Q9450 4スレッド
------------------------------------------------------------------
Posix thread creation : 0.000124(sec)
Thread #1 Execution time: 5.992(sec)
Thread #2 Execution time: 6.08448(sec)
Thread #3 Execution time: 5.9132(sec)
Thread #4 Execution time: 5.91843(sec)

Total Execution time : 6.08452(sec)　<ーースレッド中の最大値 + α

GPGPUについては詳しくないんだけど、
(sizeof float)*3*3*5000000≒180[MB]
これがシステムメモリとVRAM間で往復するから360[MB]
所要時間が2[ms]だから、1[s]に180[GB]も動いてることになる
何か変だ

ところで、結果は一致してるのか？ｗ

89　です。

使った行列は
2.000000E+00, 1.000000E+00, -1.000000E+00
1.000000E+00, 3.000000E+00, 2.000000E+00
-1.000000E+00, 2.000000E+00, 4.000000E+00

ただしデータは対称性の為、(2.000000E+00, 3.000000E+00, 4.000000E+00, 1.000000E+00, 2.000000E+00, -1.000000E+00)
の６成分のみで、 システムメモリ〜VRAM間の転送量は
３×３行列　　　　　sizeof(cl_float)*6*5000000 = 114MB
固有値 sizeof(cl_float)*3*5000000 = 57MB
固有ベクトル　　　　sizeof(cl_float)*9*5000000 = 171MB
反復解法の収束回数　sizeof(cl_int )*5000000 = 19MB

CPUでの解は
Eigen value e1 = 2.855207E-01, Eigen vector1 = ( 6.345872E-01, -5.961551E-01, 4.918314E-01)
Eigen value e2 = 3.143277E+00, Eigen vector2 = ( 7.708385E-01, 5.341269E-01, -3.471549E-01)
Eigen value e3 = 5.571202E+00, Eigen vector3 = ( 5.574225E-02, -5.994226E-01, -7.984895E-01)
<e1,e2> = 5.960464E-08　←　固有ベクトル間の内積での直交性チェック
<e2,e3> = 0.000000E+00
<e3,e1> = 0.000000E+00

GPUでの解は
Eigen value e1 = 2.855215E-01, Eigen vector1 = ( 6.345873E-01, -5.961550E-01, 4.918314E-01)
Eigen value e2 = 3.143277E+00, Eigen vector2 = ( 7.708384E-01, 5.341271E-01, -3.471551E-01)
Eigen value e3 = 5.571201E+00, Eigen vector3 = ( 5.574221E-02, -5.994227E-01, -7.984894E-01)
<e1,e2> = -4.470348E-08
<e2,e3> = -5.960464E-08
<e3,e1> = 0.000000E+00

で略一致してます。

89 です。

同じ問題を maxima の eigens_by_jacobi で解くと
(%i1) A:matrix([2,1,-1],[1,3,2],[-1,2,4]);
(%i2) eigens_by_jacobi(A);
(%o2) [[0.28552125561224, 3.143277321839643, 5.571201422548121],
[ 0.63458730239817 0.77083835004074 - 0.055742207899264　]
[ 　　　　　　　　　　　　　]
[ - 0.59615502320039 0.53412697029887 0.59942269552653　　]]
[ 　　　　　　　　　　　　　]
[ 0.49183141821965 - 0.347155034107 0.79848934767235 　]

（こちらは、固有ベクトルの成分が縦方向に並んでいます）

GPUのほうは最適化の有無でガラっと変わるんでそこんとこどうなのよ

89です。

rtn = clBuildProgram ( pgm,　the_number_of_devices,　devices,　"-cl-opt-disable",　NULL,　NULL );
でBUILDしています。

KHRONOSのPDFマニュアル p115に

-cl-opt-disable
　　　This option disables all optimizations. The default is optimizations are enabled.

と記述があります。

またkernel 実行は
rtn = clEnqueueNDRangeKernel ( CommandQueue,　
kernel,
1,
NULL,
&pe_size, // 5000000
&group_size, // 64
0,
NULL,
NULL // No triger event will be used.
);

今気がついたのですが、ｐ１３２に

clEnqueueNDRangeKernel returns CL_SUCCESS if the kernel execution was successfully　〜〜　queued　〜〜.
Otherwise, it returns one of the following errors:

とありました。　"Execution Time" と上で書いた時間は実行キューへの登録時間でした。
お騒がせしてすみませんでした。

89 です。以下の方法で、Kernel実行時間とメモリマッピング時間の計測が可能であることが分かりましたので再計測してみました。

　　cl_event event;

rtn = clEnqueueNDRangeKernel ( CommandQueue, kernel, 1, NULL,　&pe_size　/* 5000000 */, &group_size /* 64 */,
0,　NULL,
&event 　<- イベント追加
);
if( event ){
(void)clWaitForEvents( 1, &event );
(void)clGetEventProfilingInfo ( event, CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &device_time_counter[0], NULL );
(void)clGetEventProfilingInfo ( event, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &device_time_counter[1], NULL );
(void)clReleaseEvent( event );
}

実行時間 device_time_counter[1] - device_time_counter[0] (nsec);

GPU Kernel Compile : 1.5e-05(sec)
GPU Kernel Build : 5.02459(sec)
GPU Kernel Creation : 6e-06(sec)
GPU Kernel Set Args. : 1e-06(sec)
*GPU Kernel Execution : 0.114895(sec)
*C[114MB] memory mapping(READ MODE): 0.0358828(SEC) 3177.01(MB/sec)
*E[ 57MB] memory mapping(READ MODE): 0.0179288(SEC) 3179.24(MB/sec)
*V[171MB] memory mapping(READ MODE): 0.0537894(SEC) 3179.07(MB/sec)
*iter[19MB] memory mapping(READ MODE): 0.00600078(SEC) 3166.26(MB/sec)

*はOpenCLのプロファイリング機能で測定した時間。　それ以外はgettimeofday()を使用して呼び出し元から測定した時間。

結局　6.08452 / 0.114895 = 52.96倍　　　　次期 HD7000 が楽しみになってきました (^_^)。　

visual studio pro, radeon 6000 台　で ati stream ないし　open cl
使って、　並列FPU高速化を確認するだけ、ってどのくらい大変ですか？
前提としてCは出来ます

GPUすごいなあ・・・。もうお前イラネで、失業の危機を感じるわ…。

>>100
解決させる問題にもよるだろうけれど、本質的には４台でも6000台でも一緒。
それが並列化ってもんだろう。

GPUをCPUのように扱えるFusion System ArchitectureをAMDが発表
http://pc.watch.impress.co.jp/docs/column/kaigai/20110616_453498.html

.NETで使えるように・・・したのがWPFだっけか

>>103
全然違う
WPFはGDIの代わりにDirectX使ってるだけ(Vista/7のみ)

FSAに先駆けてVisual Studio2012でついにAMDとnVidiaのGPUがC++プログラミングに完全対応

Rubyバカにしてる子ってさ
変数に＄ついてる言語触ってるって事だよね

いちいちSHIFT＋４キーおして ＄ 打ちまくってる感触はどう？

【レス抽出】
対象スレ：GPGPU#5
キーワード：ruby
検索方法：マルチワード(OR)

105 名前：天使 ◆uL5esZLBSE [sage] 投稿日：2011/07/01(金) 18:26:05.76
Rubyバカにしてる子ってさ
変数に＄ついてる言語触ってるって事だよね

いちいちSHIFT＋４キーおして ＄ 打ちまくってる感触はどう？

GPGPUプログラムしている人、どこののGPU使っている人が多いの？
Intel、nVidia、ATi、どれ？

速度求めるならATIただしライブラリとドライバが糞

intelってできるんだ？

intelからOpenCLのSDKは出てるようだが

それCPU用だから。でもインテルのだから速いんだろうなきっと。他力本願

来年のIvyBridgeからだろうなあ

OpenCV2.3がCUDA対応になったけど、どこまで対応してるんだろうな

ってすと

>>111
SandyBridgeのCore i5にインテルとアムドのOpenCL入れてるが、アムドの方が速いっす
インテルの方はAVX対応とか言ってるくせに効果なし

CUDAとかそういうGPGPU向けに作られた言語でなくて、プログラマブルシェーダでできるGPGPUってどんなのあるかな
そういう初期のGPGPU用の参考サイトある?

Cgとかじゃなくて？

>>117
純粋に古の技術を学びたいです

DXCSがまさにそれだから、むしろ最新の技術なんじゃないか？

そうなのか
レンダリング結果のピクセルカラーから値を読み取るやつをやりたかったんだけど、DCCSってのはそうなの？

DXCSは先祖がえり

わざわざ先祖返りしたってことは、
書くのは大変だけど速さはGPGPU用よりも速いってことかな

ちょっと調べてみるわサンクス

おや、こんなスレなんてあったんですね

JOCLって

とりあえず安定ならCUDAが一番なのかな？

安定＝他人のソースをコピペできる
ならそうかな

不安定だと他人のソースのコピペもできないのか！

性能ならATi
汎用性ならCUDA
でおｋ？

>>128
こうじゃね？

性能：ATI Stream
汎用：OpenCL
実用：CUDA

DirectComputeは？
あとATI Streamは悲しいほどに資料がみつからないんだけどそんなに高性能だったの？

ATI Streamが高性能というより
AMDGPUの演算性能自体がNVIDIAより2-3倍高い。
NVIDIAが力を入れているC2090のDP性能でも
6970と理論値で互角、実効値では後塵を拝している。

以前、ATIのIL(アセンブリ)で組んだことあるけど、
チップセット内蔵GPUしか持ってなかったから糞遅かった。
ちゃんとしたGPUで動かすと速いのかな。

>>131
そうだったんだ。OpenCLやDirectComputeでの比較がないか探してみよっと。

ATI Streamは本当に資料が無いよな・・・
CUDAの本は何冊か出てるのに

OpenCLで最新ATIの性能をフルに引き出せる？

前に例の長崎大のGPUスパコンの人がOpenCLでCypressのDGEMMベンチマークやってたよ

もしかして俺がCAL+ILの日本語本を書いたら大もうけできるのだろうか。
でも需要無いか。

次期東大スパコンにも使われる予定のIntel MICのほうが良いかも。
LarrabeeはGPUではなくなってしまったからGP「GPU」ではないかもしれないが。

>>136
それよりはOpenCLで書いて各ハード向けへの最適化手法を本にしたほうが儲かると思うよ

NVIDIAのGPUでOpenCLやろうとすると徹底的に最適化しなきゃお話にならないあたりでどうにも

NVIDIA＜CUDA使え

って意味だな

CUDAはC#版もいちおー出てるのが大きいな

それ言ったらOpenCLもJava版出てるよ

両方いろいろ出てるよな。Python版とか。C/C++しか使わんけど

pythonからpycuda経由で使ってみているけど、結構便利。
細かいことやろうとすると、結局python内にC(CUDA)のコード埋め込む事になるけど。

VisualStudio11のDPでC++AMPって使るんだねコンパイルしただけだから
実際どうだかとかわからないけど

ATI Streamは資料も無いしラッパーも無いし、何も無いのが痛い

>>146
え？
ttp://developer.amd.com/sdks/AMDAPPSDK/documentation/Pages/default.aspx
じゃだめなの？

英語読めない人にも配慮してもらわないとな

>>131
2〜3倍ってそんなに違うのか。ATIはいいもの作ってもマーケッティング面が弱いのかな

ベンチマークで良い数値が出てもドライバがバグだらけなので実用的じゃないんです。

AMDとintelのドライバ、どっちが悪いかってくらいダメだからなぁ

いっそのことプンソにして作って貰った方がいいんじゃね

>>151
その2つ合わせたよりも高いOSクラッシュ率のNVも相当なもんだぞ

>>152
AMDのLinux向けドライバはもう大分前からオープンソースでやってるんじゃ？

両方あるよ。
AMDとNVIDIAの違いはオープンソースコミュニティに
ハードウェアの仕様をドライバ書けるレベルまで公開しているかどうか。

オープンな方のドライバはopenclをまだ実行できないんじゃないかな
AMDが配布してるクローズドな方は実行できるしもう研究に使ってるとこもあるみたい

ゲフォで一般向け、と言うかteslaじゃないのは計算誤りが含まれてるから選別したとか、
どこかの大学で言ってたと思うんだけど、その辺の事情はアムも同じなの？

ソレ言った長崎大の先生はそのあとHD5870でクラスタ組んで論文出してるけどそのへんの事情は言ってないね
まあ保証が欲しいならFirePro買ってくれって立場なのはAMDも変わらんだろうけど

こんな技術、BEEP音鳴らすブザーで音階を奏でる類の技術なわけで、
広く実用されることは永遠にないと思うな。

だいぶ違うと思うが

初期はともかく、現在のハードウェアはほぼ完全にHPCのトレンドに乗せてきてるし

だいぶ違うのは確かだが、広く実用化されるかというとどうだろう。
フィットする問題領域がいくつかあるし、将来もなくならないだろう。
そして多くの領域ではGPGPUが必要にならない。
ほんと下らない当たり前なことで、>159 は何を言いたかったんだろう。

需要がニッチ過ぎるしなぁ
CPUもコア数増やすしか無くなってきたし、一時的な技術なのは確か
今のところは、単純な計算を繰り返すだけならGPGPUのほうが優位ってだけで

そりゃまあ、一家に1台クレイ風ベクトルプロセッサ、みたいにはならなかった、
という意味では、たいていの技術は「広く実用されることは永遠にない」ものだろうが。

ラデの7000シリーズが出たな。

金のある研究員の人
ラデの人柱になってください

某東北の公立大の中の人がなるんじゃないか？

4Gamer.net ― AMD，新世代ハイエンドGPU「Radeon HD 7970」を発表――Southern Island世代のGPUアーキテクチャを整理する
http://www.4gamer.net/games/135/G013536/20111221078/

＞つまり，GNCでSPの演算機能に手は入っておらず，単精度の浮動小数点積和演算・
＞積和算・乗算・加算と整数演算のみをサポートしたものだということだ。
＞言い換えると，VLIW5アーキテクチャにおける“ビッグSP”のような，
＞倍精度演算や超越関数演算に対応した特別機能ユニットは搭載されていない。

Graphics Core Nextって倍精度が激遅になったりしないんだろうか。
一応次のようなCTOの発言↓があるのでそれをカバーするしくみがあるのかもしれないけど。

＞「依存関係にない複数の命令を1命令としてまとめて実行できるVLIW方式も，
＞グラフィックス用途では十分に活用できるアーキテクチャだが，
＞GCNは汎用コンピューティング用途などでも優れたパフォーマンスを発揮できるアーキテクチャだ」

自己解決しました。何にせよ新しいモノは楽しみですな。

【後藤弘茂のWeekly海外ニュース】 大きく進化を遂げた新世代GPU「Radeon HD 7900」
http://pc.watch.impress.co.jp/docs/column/kaigai/20111222_501138.html

＞倍精度演算などは命令発行に16サイクルかかるため、Graphics Core Nextでは単精度と倍精度の
＞ピークパフォーマンス比率は1対4となっている。925MHz動作なら947GFLOPSの倍精度性能となる。
＞ここはトレードオフで、HPC(High Performance Computing)を重視するNVIDIAは、
＞ダイ効率のトレードオフを払っても、単精度と倍精度の比率を1対2にしたが、AMDは1対4に抑えた。

つーか乗算の筆算考えてみれば
SPvsDPの計算負荷は1:4の方が自然なのは当たり前
NVIDIAのは倍精度重視というより、単精度軽視という方が近い
SP2倍結局計算できる回路があるのに使えなくしているだけなのだもの

まあ、レジスタ帯域とか考えれば1:2も分かるけど
これはレジスタの割り当て自由度を無駄に上げて
レジスタ帯域を上げる事を難しくしている
Fermiの構造的な問題だからな。

Tesla1のときは1:8だったけど理論値あたりの実効性能は90％に達してたんだよね
LSUやメモリなど足回りのパフォーマンスがDP演算器を支えるだけの余裕があったが

FermiでDPを増強したら今度は足回りが追いつかなくなったという皮肉

CUDAのアーキテクトがAMDに移籍したらしいがGCNはそれがモロに出たアーキテクチャだな
公式スライドそのままの画像をリークしてたVR-Zoneによると549ドルらしいし
HD6970、GT580より省電力とのことでかなり面白そうだ

確かに筆算の掛け算を考えると横に2倍×縦に2倍で4倍の計算量ってのが確認できますね。
仮数部が24bitと53bitで実際は4倍よりもうちょい要るはずだから、1:4でも少なからずリソースを割いているか。

64bit整数演算はまだかなぁ…使うアテないけど。
64bit整数が32bitと並ぶのはVRAMが4GB超えるのとセットなのかな。

>>170
行列積の話？

64bit整数演算はできないけど、53bitはできるって理解していいのでしょうか？

金のない俺としてはどのクラスまで倍精度をサポートするかが気になるな。

贅沢は言わんから7700シリーズまで頼む。

>>174
いいと思うけど、あくまで倍精度演算の速さでだよ。
さらに前後でdoubleに格納したり取り出したりするならその分もかかる。

おすすめのGPUを教えてください

最新の一番高いやつをその都度

http://www.freepatentsonline.com/8051123.html
Multipurpose functional unit with double-precision and filtering operations

Warpのダイナミック再構成がつくって話だなkepler

AMD's got an ace up it's sleeve: Tahiti-ASIC probably has 36 CUs/2304 Shaders
http://www.gpu-tech.org/content.php/181-AMD-s-got-an-ace-up-it-s-sleeve-Tahiti-ASIC-probably-has-36-CUs-2304-Shaders

http://ascii.jp/elem/000/000/659/659849/index-3.html


大原雄介の記事来た
http://news.mynavi.jp/special/2012/tahiti/index.html

http://news.mynavi.jp/special/2012/tahiti/005.html
めがっさ速い？

NVIDIAのステマ・広告攻勢がすごいから、
性能が同じくらい＝AMDの圧勝ぐらいの意味だからな
そのあたりを気にしながら記事を読む必要がある

公式の7970用ドライバ来る前の記事だからな

nVidiaのステマ能力ははんぱない。
全世界のスパコンシェアを圧倒してしまった。
我々はこの独裁にどう対抗していけばよいのだろうか・・・

>>186
ATI stream と nvidia CUDA の両刀使いが現れるまで待つしかない、と。

Revenge is Sweet: PowerVR Discrete GPGPU PCIe Card Coming Later in 2012
ttp://vr-zone.com/articles/revenge-is-sweet-powervr-discrete-gpgpu-pcie-card-coming-later-in-2012/14609.html

レイトレを効率よく実行できるアーキテクチャなら使いでがありそうだな。

日本でのCUDA最先鋒な
東工大のAFDSでの発表が面白い。
http://developer.amd.com/afds/assets/presentations/2913_3_final.pdf
FFTをOpenCLで6970に移植したら、さくっとC2050や580の
最速実装を超えてしまったけど、でもだからと言って
AMDの方が良いとは言わないよとか。

TSUBAME2.0のためにTesla数千枚買っちゃったのに今更AMDが速いとか言ったら各方面から暗殺されかねんからな

超要約すると、

「なぜなら、このプログラムがあるのはNvidia賛のおかげだから（ｷﾘｯ」

って事か？

「汎用のOpenCLよりnVIDIA特化したCUDAの方が速いよ」だと思う
当たり前のことだがｗ
せめてATIStreamと比較してくれ

なぜならECCが無いから問題外

ソフトウェアECCって処理によってがくっと性能下がるらしいからなぁ

たとえば？

>>192
「nVIDIA特化したCUDAでnVIDIAが最適化したCUFFTの方が
当然速いと思ったら、それよりは遅い自分たちがCUDAで書いた
コードを汎用のOpenCLに移植してAMDで動かした方が速かった」
という結果が出た上での結論だから。

>>195
話題に上がってる東工大の先生がGPUのソフトウェアECCについてまとめてたよ
あれ見るとハードでECC処理するTeslaのありがたみがよくわかる

ECCってハミング符号？
だったら、ソフトウェアで実装したら遅くなるのは道理だね

ECCに気を使うのも結構だけど、GPUの場合
そもそも計算間違う可能性についても
議論した方が良いと思うんだが。

長崎大でのGeforceのCUDA不適格率を考えてみても
Teslaがどの程度信頼性あるのか分からん。

統計的バラつきを消すために最初から複数回回す必要があるような
アルゴリズムを選ぶ方がいいのかもな。

今度のラデは800,700番台でも倍精度演算を実装してるみたいだ。
ちょっと勉強するのに買いやすくていいな。

おまえはもう達人だろ。

http://news.mynavi.jp/articles/2012/03/19/radeon/004.html

77x0,78x0の倍精度はエミュレーションっぽいな。
やっぱり79x0買うしかないか。

ＧＰＵでIIRフィルタを効率よくやる方法ってあります？

IIR位何も考えずとも効率良くやれるだろ。
並列に計算できるデータが物凄く多ければな。

並列に計算できるデータが少なくて、時系列方向に長い場合は、
FIRなら時系列方向の並列化も可能だけど
IIRは無理じゃね。

"IIR filter vectorize"とか"~ SIMD"とかでググったらなんか出てくるし読めばいいんじゃねぇの？

>>203
一応単精度では6950ぶっちぎる場面もあるんでまあグラフィックフィルタとか単精度で良いソフトなら試す価値はあるかな
今までと違って動かないわけじゃないから倍精度の開発検証ぐらいはできるしその程度だと思えば…

>204
けち臭いこと言わずFFTしてフィルタリングしろ。

GTX680出たな

だが、APUあたりのレジスタやキャッシュが激減してるし、チップあたりの
実行可能なカーネル数も減ってる。倍精度もエミュレーションぽいし、GPGPU向きじゃないな。

APU→ALU

トレードオフしっかり設計に反映させないと難しい世代になってきたって事かねぇ

>210
具体的な数値って分かる？

pcヲチの後藤さんの記事のブロック図を見れ

Linux上でGPGPU使って遊びたいんだけど64bitの方がいいかな
もしそうならOS入れ直すんだけど

Ubuntu12.04 x64版上で、OpenCLを使って３ｘ３行列の固有値＆固有ベクトルを
単精度のJacobi法で計算させてみた。

CPU：intel 　 i7-2600k (4.4GHz)
GPU:AMD Radeon-HD7970 (1.05GHz)

行列の個数５百万個（全て同じ行列を使用）
OpenCLソース内で同じ計算の繰り返しを１０回
（->５千万個の３ｘ３行列の固有値＆固有ベクトルを求めたことに相当？）

OpenCLのソースコードはCPUとGPUで同じものを使用。

＜結果＞
CPUを使ったOpenCL（Jacobi法の反復回数は６回で収束）
カーネルの実行時間：４７．６８秒

GPUを使ったOpenCL（Jacobi法の反復回数はCPUと同じで６回で収束）
カーネルの実行時間：0.03389秒

GPUが１４００倍も速いと言う結果になった。

OpenCLがAMD製なので、IntelのSSE、AVXなどへの最適化がうまく
行われていないのだろうか？

そもそもSSEやAVX使うようなコード書いているの？

float8とか明示的に256bit幅使うように指定しないと、AVX使わないのでは

float4とfloat4を引数に使ったmadは多用してるよ。

いくらなんでも1400倍はおかしくないか？
OpenCLのことはよくしらないけど、理論性能から考えてもそんなに
差が出るはずが無いと思うが･･･
140倍なら、ＣＰＵのコードがクソならあり得るが

i7-2600k (AVX) が単精度 220 GFlops、倍精度 110 GFlops
同 (SSE) が単精度 54 GFlops、倍精度 110 GFlops
Radeon HD 7970 が単精度 3.8 TFlops、倍精度 950 GFlops

SSE でも100 倍以内に収まらないとちと CPU 遅すぎ。
OC しているならばなおさら。

倍精度だとSSEもAVXも同じなの？

x87比

同じ問題を APU E1-1200 のミニノートPCで解いてみた。
OSは、 >216 と同じくUbuntu12.04 x64版。

メモリ資源の制限から、行列の個数は１/５の１００万個。
Jacobi法のOpenCLモジュール内で１０回同じ計算をループで
回しているので、１０００万個の３ｘ３行列の固有値を求めたことになる。
当然ながら、このミニノートPCでは O.C. はしていない。

＜結果＞

反復回数（６回）、計算精度共にデスクトップマシンで計算結果と同じ。

CPU　 7.542 sec
GPU 0.4235 sec

OpenCLイベントタイマーで計測した正味の計算時間比較では
GPUが１８倍速いと言う結果になった。
これぐらいの比率なら正常？

ついでに >２１６ で使用したデスクトップマシン環境で同じ１００万個の計算（１/５の量）を
再計測してみると
＜結果＞
CPU　 9.311 　sec
GPU 0.006517 sec　　やはり、１４２９倍となる。

ただし、皆さんお気付きのように１００万個同士の計算時間比較を行うと

＜CPU同士の比較＞
APU　E1-1200 : 7.522 sec
i7-2600k (4.4GHz) : 9.311 sec

Woo！　i7-2600k(4.4GHz)が E1-1200 の　0.8倍　?????

＜GPU同士の比較＞
APU　E1-1200 : 0.4235 sec
HD7970 (1.05GHz) : 0.006517 sec

HD7970 (1.05GHz) が APU E1-1200 の６５倍となり、CPU側の計算時間が？

どちらも、C++で作成したOpenCLの環境セットアップ部を含めたJacobi法の
OpenCLコードまで全て同じソースコードを用いているが
CPU側のOpenCLイベントタイマー計測がどう見ても変なことが分かった。
デスクトップ側、ミニノートの両ドライバーファイルは12.8 catalyst + AMDAPP V2.7
で同じ。 またUbuntu付属のg++4.6.3+NetBeans上で実行モジュールを作成している。

>225

> CPU側のOpenCLイベントタイマー計測がどう見ても変なことが分かった。

上記行を削除します（書き間違えました）。

実際に腕時計で測っても、同等な８秒程度の時間でした。
OpenCLイベントタイマー計測が変なのではなく、実際に i7-2600k の方が若干遅かったです。

ループでもっとぶんまわせば？
計測短すぎてクロック上がってないんだろ。

そもそも3x3行列ってのが小さすぎ

たしかに3x3は小さいが、全部の固有値を馬鹿並列で求められるなら、むしろＧＰＵ向きな気もする

「100万個の固有値計算を馬鹿並列にできるなら」という意味ね。1つの行列の相異なる固有値を並列に求めるという意味では無く。俺日本語不自由すぎﾜﾛﾀ

AMDのOpenCLランタイムは、CPUが
ちょっと異様なぐらい遅いとか聞いたことがある気がする

それってSDKが出たばっかの頃の話じゃね
たしかデバッグか何かに絡んだ挙動だったと思ったが

7970もECC化できればなぁ

レイトレースをCPU、GPU側で同様のロジック作って走らせても1000倍くらい違ったし
GPUは並列度の高いコードだと思った以上に早くなる
CPUはそれだけ理論値からの落ち込みが激しいってことなんだろうけど

CPUの方のコードが最適化されてないんじゃないの？
最新のGPUでも単精度ピークが数Tflopsだから、
1000倍差だとCPUは数Gflops以下しか出てないことになる

x87比か

>>235
実際、素人が行列積書くと1−2GFlops しかでないし、まぁそんなもんだろうね

Intelがメニーコア「MIC」とAtom SoCの「Medfield」を発表
ttp://pc.watch.impress.co.jp/docs/column/kaigai/20120831_556528.html

GPU勢もうかうかしていられない？

GPUコアは、これも PowerVR なのね。人気者すぎ。

アーキテクチャがx86だってのは個人的にどうでもいいんだが、
メモリ階層がどうなるのか、メモリバンド幅がどれくらいでるのか、
CUDA5のdynamic parallelismに対抗できる機能があるのか、くらいが
勝負の分かれ目だな。

>>240
dynamic parallelismなんてわざわざご大層な名前を付けなくても
x86CPUなのだから同等以上のことが出来るに決まっているだろ。

対抗してCUDAのバイナリ仕様公開とかないかな

ナイコナの汎用性は別に誰も否定してなくね
問題はその効率だよな。やべーさすがIntel様だ、となるのかやっぱ汎用コア並べたらそんなもんだよね、で終わるのか
うちも研究枠で調達予定なので普通にwktkしてます
GPU性能を維持したままにじり寄ってるからこそ、dynamic parallelismなるご大層な名前…というか変態的な局所性能の誇張に縋ってるんだろう

また大人のおもちゃが税金でたくさん作られるね

えー?
コア一つ当たりの性能ではGPGPUを圧倒するんじゃないの?
問題は、それだけの性能を発揮するために必要なコアサイズが大きすぎることであって。

一つのダイに集積可能なコアの数が、GPGPUのプロセッサより一桁以上も少なくなるので、
コア性能では圧倒していても、総合性能では大幅に負けるに違いないと予想する。


＞やっぱ汎用コア並べたらそんなもんだよね、で終わる

としか考えられないよなあ...

Coreの定義がIntelとGPUメーカーとで違うから当たり前だろ
1コアあたりのベクトルユニットが16SP/8DP積和だから
NVIDIAのCUDA Coreに換算すると16コア相当で
つまり50coreのPhiは800 CUDA coreに相当するんだけど？

あー、そか。たしかにそうだよな。
コアサイズが二桁ぐらい違うとかじゃないと性能で見劣りすることはないかもしれんのか。

まあ、コアの世代を古い方にして回路規模を小さくしちゃったりするとスループットが落ちて不利だがな、
動作クロックがGPUより上回ってる分と併せて考えると接戦になるのかな?

排他処理のルーチンはボトルネックとしか考えられないよなあ...

GPGPUはピークはともかく実効値がな･･･
先日出たTSUBAME2.0の1mメッシュ気流解析で実効15%ぐらいだっけ？
複雑な分岐や並列度低いものが入るとしぬし
ほかのGPUアクセラレータ積んだクラスタもそうだけどLINPACKばっか速くても仕方ない

その残りの85%は何してるの？

リラックス

良いこと考えた。
働いていない、85%を装置から除去すれば、消費電力も下がるんじゃね？

それこそSMTできるようにして同時実行させろって話じゃね

そしてまた効率低くてそぎ落としか
設計レベルでループしてんじゃねえよw

人間社会もGPGPU社会も一緒なんだな

分岐というか、単純に帯域の問題じゃね。

あのプレゼン見たけど、プログラミングに自信のある計算機寄りの人が、
てきとーに題材みつけてきて、計算機科学の研究として発表してる感じだった。

その結果が15%

計算目的よりも、tubameを生かすための作業やね。まあ察しはつくがw

>>252
蟻の巣からよく働く2割の蟻以外を取り除いても、
残った蟻の8割がなぜか怠け始めるんだそうな

>>250
働いてないんじゃなくて何かしらがボトルネックになってFMACの稼働率が抑えられてるのでは？
メモリ帯域が足りてる場合でも同時命令発行数の制約でload/store命令と積和命令を
同時に実行できない、とか
（これは京のVenusアーキテクチャにもある制約）

唐突ですまんけど、AMDはファイル名通りのこういう資料があるけど
ttp://developer.amd.com/gpu_assets/R700-Family_Instruction_Set_Architecture.pdf

nVIDIAはこういうの無いの？

>>249
効率はFLOPSに対する効率であって、気象系のアプリは
バンド幅律速だからLINPACKと比べること自体ナンセンス。
ちなみに15％は気象系のステンシルアプリとしては高い方。

>>257
そらまぁ、青木研究室の存在意義が「GPGPUでできることを広げる」
だからね。仕方ない。
NECのSXが沈没の今、京かGPGPUしか無いわけで･･･
それともBlueGeneを買うか？

FLOPSに対する効率だから、アルゴリズム自体の並列化効率以上にはなりえないわけか。
気象系のアプリというものの並列化効率がどんくらいなのか知らないけど（なんとなく高そうではあるが）。

>>264
うーん、微妙に違うな。

今時のアーキテクチャだとメモリがすごく遅いってのは知ってるだろ？
ちなみにメモリの速度はバンド幅って言う。
それに対して演算速度はムーアの法則でどんどん上がってる（上がってた）
から、バランスが全然とれてないんだ。だから、普通のCPUはキャッシュ階層を
深くして（L1、L2、L3とか）なるべくデータの再利用をしてる。

つまり、データをメモリから持ってきたら、なるべくそれを
使い回して計算したいわけ。

で、ここで問題になるのが、浮動小数点演算（FLOP）と、データ転送量の比。
業界だと byte/flops とか言われる。言い方を変えると、どれくらいデータの
使い回しが効くかってこと。
続く

承前
アプリの要求byte/flops高いと、せかっくデータを持ってきても
たいした計算をせずに、すぐに次のデータが必要になっちゃう。
データの使い回しが効かないんだ。これはイマドキの計算機にとってはキツい。
で、気象とか流体とかはそういうアプリなわけ。

だから、基本的に流体とかの計算は演算は余ってて、バンド幅がボトルネックになる。
これはアプリの特性だから仕方ない。

気象･流体屋さんは、未だに地球シミュレーター信者みたいな人が多い。
ちなみに、気象・流体系アプリは並列度だけで言ったらはものすごくあるよ。

ちなみにスレ的に言うと、GPUはGDDR５っていう容量が少ない代わりに
バンド幅が出る贅沢なメモリを搭載して、かつ大量のスレッドを使うことで
レイテンシを隠蔽してる。これがGPGPUが気象・流体計算に使える理由。

なるほど、つまり気象系のアプリはコンピュートバウンドではなくメモリバウンドのものが多い、と。
AMDのAPUの使い道の話になったときに、DDR3という制約に悲観的になってる人が少なそうだったから
世の中にはメモリバウンドのものは少数派なのかなあとも思ってたけど普通にあるんですね。

バンド幅と言うか、物理原理的にDRAMのセルは応答が現代のCPUの演算速度に対してずっと遅い。
だから、メモリセルを並列にして同時に読み出し・書き込みするしかないのが根底にある。それをやれDDRxだの、
何ビット同時にアクセス＝複数セルに同時アクセスでなんとかしてるね。

気象系や気体がどうのって言うよりも、規則正しく計算点を並べて偏微分方程式の数値近似式を一斉に走らせる
＝SIMDだのシェーダだの言ってる演算回路がみんな一斉に同じ計算式をやる場合が最大性能が出る。
三次元の方がそりゃ、並列度は高いだろうね。境界条件とか面倒くさそうだが。

こういう根本的なところはもう全然進歩しないなw

bytes/flopsに関して言えば現状GPUはCPUより低いです。
ただbyte/sが数倍高くて、flopsが更に数倍高いというだけの話で。
Haswellあたりで、電力当たりFLOPS数はGPUとCPUは互角くらいになるのではという話も出ていますな。

TSVって、GPUみたいな爆熱チップでもできるの？

>>241
亀レスだけど、何のペナルティも無く普通のx86コアを50個
並べられるなら誰も苦労してないわけで･･･

DRAM内部のクロック
PC133　133MHz
DDR3-1600　200MHz
この間十数年・・・MRAMならもっと上がるかな？

GNU MPのGPGPU対応って、あんまりやってる人が居ないな。

依存関係があってなかなか並列化しにくいからね。
もちろん並列度の高いコンピュータに適した数式に置き換えればいいんだが
CPU側もAVX, FMAなどでGPGPUの電力効率の優位も揺らぎつつあるからね

C++AMP調べてるとこなんだけど
textureでピクセル間を補間する方法がわからない。
ひょっとして自前でやらないといけないの？
「C++ AMP API は、サンプリング テクスチャのフィルター処理機能を提供されません。」
とはこのことなのかな。

c++ampにfloat_3のmadやfmaが無いね。 short vector typeのmadやsin, cos などの超越関数も用意してもらいたい。 swizzlingはできるみたい。 OpenCLの方がコードをコンパクトに書けるかな。次期版期待してます。

何故かsea islandのISAが公開
ttp://developer.amd.com/tools/heterogeneous-computing/amd-accelerated-parallel-processing-app-sdk/documentation/

いつもに比べると異常に早いがやっぱsouthernislandのリネームなのかね

ディスクリートとしては出なくても
次世代APUのGCNはsea islandなのかね。

ISA見る限り、SIから確かに機能アップされている。
一番大きいのはFLAT メモリアクセス命令の追加か。
cudaのUVAみたいなもの

SL#（えすえるしゃーぷ）とは、GPUで実行されるプログラマブルシェーダーを、超高級言語である
C#で書けてしまうという夢のようなオープンソースのフレームワークである。
http://monobook.org/wiki/SL_Sharp

>>280
まじかよ、ニンジャコンパイラが全部解決してくれるのか？

超高級言語ｗ

抽象化の度合いが高いぐらいの意味合いでいいんじゃないの
実際問題プログラマが意識してGPUに仕事をさせるコードを書かずに勝手に解決してくれちゃう未来像は来るの？

たぶん、デザインパターンと同じくらいの敷居にはなるんじゃねーの？と言うか、デザインパターンとして確立されるんじゃね。

判らない人には判らないけど、自分の中に消化され始めると自然とそう言う書き方になるみたいな。
あとは、フレームワークでそう言う書き方を強制されて、なにも考えずにそう言うもんだと思うパターン。

コア数が100くらいのGPUがあるとして
整数大規模行列の
全要素の和を求めるのってGPGPUで100倍程度の高速化は可能ですよね？

1コアあたりのシングルスレッド能力がCPUと同程度あって、メモリ帯域幅がそれに見合うほど広ければな

>>283
実際、解きたい問題依存だけど、Hadoopのように、プログラミングモデルを固定して制限を付けて縛ることによって簡単な記述で自動的に高速実行ってのは比較的現実的な未来だと思うよ。
Hadoopってのも、書きたい処理をHadoopの"型"に押し込んで書ければ、残りの面倒な問題は全部Hadoopが面倒見てくれますって事だし。

このスレ立った頃のGPUってもうゴミだよな？

OpenCLで書いたプログラムを9600GTとQ9650で動かしたら同じくらいの速度だった。
今のCPUには絶対負ける。

なんか来てた。
ttp://www.phoronix.com/scan.php?page=news_item&px=MTQ3NDU

GPGPUとMantleの関係を語ってくれよ
mantleってなんなんだ？

>>291
DirectXやOpenGLは長い歴史があって
色々なハードで動かす必要があるので
互換性維持のために滅茶苦茶厚いAPIコールを通さないといけないが
MantleはGCNに絞った最適化さえすればいいので
そういうものがないっつーこと
据え置きはAMD一色だから移植も相互に楽になるし、面倒な最適化も最小限で済む
というのが言い分

DirectXやOpenGLやOpenCLだと実行時に
「君（GPU）この機能使える？」
「よし、通れ！」
というプロセスを踏んでるからnVidia、AMD、Intel、PowerVRなど全部のハードウェアでそれなりに動くわけ
これに対してCUDAやMantleは他社対応を完全に無くすことで最初から最適化されたランタイムで動いてる

商用製品でCUDAがOpenCLより人気なのを見れば分かるように、
他のハードウェアで動かない以上のメリットがあるわけ
とくにゲーム機だと全てAMDで、今はAMDがARMと仲が良いからCUDA以上に広まりやすい環境にあるわけ

このスレ的にMantleは用途が違うだろうけど、HLSL/GLSL直書きでGPGPUをやってた層には移行先として十分かなと言った感じ

Mantleって互換性維持のため分厚くなり過ぎたDirectXに対して
低レベルなところを柔軟に叩けるようにするとかそういうものじゃなかったっけ？

Project Sumatraが楽しみだな
Java9ぐらいでparallelStreamにぺぺっと渡すと
OpenCLやってくれちゃうとか最高じゃん
逆に言うとそれぐらい猿でも簡単に扱えないとなっていうか

Mantle.Netはよ

kaveriのFLOPSのDP:SPが気になるんだが、誰か情報持ってないか？

OpenACCってどうよ

最近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)です。

メモリ転送のオーバーヘッドがあるから、もっと大きな問題じゃないと効果は出ないよ。

>>299
それは知っているのですが、数値を大きくするとすぐにc00000fdでクラッシュするんですよ……。
<<<grid, block>>>もいろいろ弄っているのですが、どうにも効果が得られません。
今試してみたら、arraySizeの値で実行できるのは25600が最大みたいです。

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スレも宜しく
分かりました。次回以降はそちらにレスすることにします。

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

>>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使い分けてる俺……

openCL始めたお(・∀・)ﾉ

(・∀・)ﾉ CPUの300倍くらいの性能が出たお！
比較したCPUはE2-2000っていうCPU+GPU=APUだけど全くGPUとしての機能をもってないのでガッカリしたお。

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

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

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

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

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

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

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

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

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

そりゃ工夫するだろう。

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

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

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

>>324
kaveri使えよ

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

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

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