トップページ⇒tech
163コメント49KB

GPGPU#5

■ このスレッドは過去ログ倉庫に格納されています
0001デフォルトの名無しさん2010/08/15(日) 21:47:50
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/
0002デフォルトの名無しさん2010/08/15(日) 21:55:10
＞998 ：デフォルトの名無しさん [↓] ：2010/08/15(日) 21:51:01
＞と言いつつAgeiaの中の人も今じゃAMDにいるからなぁ
＞とんだ詐欺師なのかねあの人

金です。

nvにとっちゃすでに用済みで、要らない子
0003デフォルトの名無しさん2010/08/15(日) 21:56:26
専用設計とはいえPPUは58gflopsしかないんだが
0004デフォルトの名無しさん2010/08/15(日) 23:08:03
基本的には、GPGPUが得意な処理を "適切なサイズ" に並列分割して
その分割された小包の集団をどかっとCUDAに押し込んでやると、分割が上手ければ
それなりに速く結果が出る。ただ、GPGPUで効率が出る並列化は簡単ではない。Larrabeeがこけたのもここ。

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

だから、ほとんどの用途では、Nehalem-Ex とか、速いCPU乗せたマシンを増やした方が良い。
他のプログラムが、"全部" 速くなりますからねｗ
0005デフォルトの名無しさん2010/08/16(月) 00:00:44
今後のCPUコアの高速化が鈍化するから
その対策として出てきたのがCPUのマルチコア化と
グラボのGPGPUとしての活用なわけで・・・

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

６コアもGPGPUも本当に必要な人・アプリが使えばいいだけ
0006デフォルトの名無しさん2010/08/16(月) 11:53:00
大部分って、静的WEBページを見るだけのユーザーのことか？ｗ
そんなもん無視でいいだろｗ
0007デフォルトの名無しさん2010/08/16(月) 23:36:42
WEBブラウズだろうがオフィスアプリだろうが
音楽・動画再生だろうがゲームだろうが大部分のアプリには
高価な上位CPUも高速なGPGPUも必要じゃないだろ。

そこそこヘビーな自分でも4コア(疑似8コア)や
1TFLOPS以上のGPUをフル活用できるのは全PC作業の１割程度だし
0008デフォルトの名無しさん2010/08/18(水) 18:34:14
LAMEとかiTunesとかで、GPGPUが効けばもっと広がると思うんだけど…
やる気無いですよねぇ
0009デフォルトの名無しさん2010/08/18(水) 20:00:08
やる気程度で速くなってくれるなら今ごろみんな取り掛かってるだろうよ
0010デフォルトの名無しさん2010/08/18(水) 23:32:04
LAME（音声の非可逆圧縮）程度じゃ処理が軽すぎるし
条件分岐も少なくないからCPUで計算したほうがいい。

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

iTunes（映像再生ソフト）にGPGPUとして使うなんて問題外。
大人しくOpenGLやDirect2DなんかでGPUとして活用すべき。
リアルタイムで映像にエフェクト処理を加えながら再生したいなら別だがiTunesの仕事じゃないｗ
0011デフォルトの名無しさん2010/08/19(木) 01:06:20
ATI Stream使ってエンコードして負荷軽減してるソフトなかったけか？
0012デフォルトの名無しさん2010/08/19(木) 07:21:50
>>11
PowerDirector？
0013デフォルトの名無しさん2010/08/19(木) 07:22:32
エンコードに使うなら売りは速度ではなく品質にすべき。
データ転送がボトルネックなのだから
単位データあたりの演算量を増やさなきゃメリットが無い。
0014デフォルトの名無しさん2010/08/19(木) 11:52:21
演算量が増えてもプログラムのフローが複雑になるようでは
0015デフォルトの名無しさん2010/08/19(木) 14:52:07
>>14
どんだけ複雑になったって、大量に並列実行できればGPGPUにとってアドバンテージがある。
データに対して演算量が少なすぎると転送や処理待ちばかりになってパフォーマンスが上がらない。
だから問題は複雑性よりもデータの相互依存性とデータに対する演算量の少なさ。
0016デフォルトの名無しさん2010/08/22(日) 07:57:15
複雑性ってなに？
0017デフォルトの名無しさん2010/08/22(日) 13:20:37
文脈から鑑みるに、プログラムの複雑さじゃないの？

もっと端的に言ってしまえば分岐命令の数
0018デフォルトの名無しさん2010/08/22(日) 13:29:10
この場合、相互依存性と複雑性は同義だと思うけどね。
0019デフォルトの名無しさん2010/08/22(日) 15:16:18
>>18
この場合は違う。
>14 はそのつもりで発言しているのかも知れないが、>13 は違う。
0020デフォルトの名無しさん2010/08/22(日) 18:36:09
そう言い切るのなら、どう違うかまでを説明せんといかんよ。
0021デフォルトの名無しさん2010/09/01(水) 10:24:46
>>13って8x8DCTを4x4DCTにするみたいな話でしょ？
演算回数は増えるがGPUなら並列数を増やせる感じで
0022デフォルトの名無しさん2010/09/07(火) 22:59:56
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にチャンスはないね。
0023デフォルトの名無しさん2010/09/22(水) 20:39:53
余所に作らせたGPUを使ったプログラムが、CUDA部分でメモリリークくさいエラーを吐いてまともに動かないんですが、
窓から投げ捨てるべきでしょうか？
0024デフォルトの名無しさん2010/09/22(水) 20:52:15
窓から投げるべき
0025デフォルトの名無しさん2010/09/22(水) 22:09:50
証拠資料を作ろうとしても、「いつ止まるか」の再現性が微妙
やっぱり実績の無いハウスに委託したのが間違いだったか・・・
0026デフォルトの名無しさん2010/09/23(木) 02:16:54
メモリの確保と解放を繰り返しているんじゃないかな。
弊社ではソースがあればデバッグも承りますw
0027デフォルトの名無しさん2010/09/24(金) 18:46:38
ソースないっす・・・
その辺だけはしっかりしているという・・・

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




どっか、受託開発や納入後の展開方法についての客観的な評価をしてくれる
コンサルタントはないですかね・・・
0028デフォルトの名無しさん2010/09/24(金) 21:25:44
CUDAでソースなし納品はありえんやろ
いつバージョンアップでバイナリが動かなくなってもおかしくないのに
0029デフォルトの名無しさん2010/09/24(金) 21:40:12
>>28
コストの問題だろ？
ソースを要求すると価格が上がる
0030デフォルトの名無しさん2010/10/23(土) 14:49:24
いや、将来動かない可能性が低くないのにコストカットされてもｗ
0031デフォルトの名無しさん2010/10/23(土) 14:50:21
将来動かなくなる可能性が高いから値切るんだろうが
0032デフォルトの名無しさん2010/10/24(日) 23:06:42
gpgpuを使用した場合、 CPUの性能はどの程度影響しますか？
teslaを用いた計算機を導入しようとしているのですが、i7-980xにするかi7-930にするか
迷っています。
0033デフォルトの名無しさん2010/10/25(月) 00:15:47
CUDAやOpen CL以外のCPUコードの実行速度にモロに影響する。
他にもGPGPU用中間コードのコンパイルにも影響するが誤差範囲。
0034デフォルトの名無しさん2010/10/25(月) 00:53:13
聞きたいのはCPUの性能によってGPUの性能が変わるかどうかじゃないの
0035デフォルトの名無しさん2010/10/25(月) 12:02:17
初心者なんですけどフリーソフトでATI技術に対応してて
MP4に変換できるソフトってありますか？
あとRADEONのカードってエンコードなら値段と性能みてどれがコスパいいですか？
0036デフォルトの名無しさん2010/10/25(月) 13:26:20
板違いです
ソフト板か自作板、DTV板へgo
0037デフォルトの名無しさん2010/10/26(火) 01:43:50
板違いです。

ここは「ATI技術に対応しててMP4に変換できるソフト」を作る側の板です。
0038デフォルトの名無しさん2010/11/02(火) 10:30:33
caymanは期待できそうだな。
0039デフォルトの名無しさん2010/11/23(火) 13:37:02
GPGPU使って何かしたいけどこれっていう何かが見つからないのー
Actor とか Map Reduce とか上位層で駆逐されてしまうねん
0040デフォルトの名無しさん2010/11/23(火) 14:15:02
俺はいっぱいアイデアあるけどな。
0041デフォルトの名無しさん2010/11/23(火) 17:23:59
あら、気になるじゃない。聞きたいわ
0042デフォルトの名無しさん2010/11/25(木) 17:47:25
突然申し訳ありません
cudaやってるんですけど・・・
カーネル関数起動させるところでエラーが出てしまいます
サンプルコードでアウトなんです
考えられる可能性を挙げていただきたいです
エロい人助けてください

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

質問あればできるものはすべて答えますんでよろしくお願いします
0043デフォルトの名無しさん2010/11/25(木) 23:02:46
>>42
エラーメッセージぐらいのせろやカスが
0044デフォルトの名無しさん2010/11/26(金) 00:11:02
密かにevergreenのISA仕様書が更新されているな。
ttp://developer.amd.com/gpu/ATIStreamSDK/assets/AMD_Evergreen-Family_Instruction_Set_Architecture.pdf
メモリアクセス周りの挙動について言及されているのがなかなか面白い。
コアレス化が余り重要じゃないという話がどういう意味か分かる。
要は、アーキテクチャ的に1スレッドが複数のメモリアクセス命令を同時発行可能で
1wavefront単位で発行された複数のメモリアクセス命令の間だけ
キャッシュ無しアクセスでもキャッシュが有効になっているから
複数のメモリアクセス命令間でコアレス化と同様の効果が得られるらしい。
0045デフォルトの名無しさん2010/11/26(金) 00:25:48
CAL+ILの情報が少ないので、書き込みがあるだけで嬉しい。
0046デフォルトの名無しさん2010/11/27(土) 03:42:20
>>43
スイマセン


CUDA error: Kernel execution failed


コンパイルはできていて、
他のマシンでは同じソースコードのプログラムは動かせます
原因は何なんでしょうか
0047デフォルトの名無しさん2010/11/27(土) 08:37:10
>>46
カーネルの起動に失敗したというのだから、起動要件を満たしていないのだろう。
まさかとは思うが、CUDA用のドライバをインストールしていないというオチではあるまいな。
0048デフォルトの名無しさん2010/11/28(日) 01:15:20
>>47
ドライバは入っています

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

cudaSetDevice(1);

を記述したら、それで通りました。
なぜ起動できなかったかは分かりません。
今可能性を探っているのですが、
タイムアウトが起こったのかもしれないと考えています。
0049デフォルトの名無しさん2010/11/28(日) 01:31:27
なんだ、動いたなら後はcudaスレへ。
0050デフォルトの名無しさん2010/12/01(水) 20:09:43
>>48
CUDAの命令はよく分からんけど、名前から察するに単純に処理するGPU指定してなかっただけじゃ。
8400GSだってCUDA対応なんだし。
0051デフォルトの名無しさん2010/12/01(水) 21:28:19
>>48 >>50
デバイスを指定しなければデフォルトのデバイス0、8600GSで実行されるはずだけど。
Capability 2.0以降限定の関数を呼び出しているサンプルコードだったとか？
まあCUDAスレあるし、そっちでやるべきかな。
0052デフォルトの名無しさん2010/12/11(土) 01:06:52
A Fast GEMM Implementation On a Cypress GPU
http://galaxy.u-aizu.ac.jp/trac/note/wiki/Fast_GEMM_Implementation_On_Cypress
0053デフォルトの名無しさん2010/12/12(日) 09:17:49
Cayman GPUではスーパーファンクションユニットが削除されて５VLIWプロセッサーから
４VLIWプロセッサーになるとのことですが、現在のCALでサポートされているsin/cos等の
超越関数命令は、自分で多項式近似計算をしろと言うことなのでしょうか？
0054デフォルトの名無しさん2010/12/12(日) 11:57:59
>>53
今までtレーンが担当してきた命令は
xyzwの複数レーンで1命令を実行するように変更されている。

で、超越関数は3SPを使った1命令で実行される。
0055デフォルトの名無しさん2010/12/12(日) 12:37:42
>>54
53 です。　CORDICやチェビシェフ多項式をWeb上で漁っていたのですが
安心しました。
0056デフォルトの名無しさん2010/12/12(日) 12:40:19
チェビシェフは自分で作るもんじゃろ
0057デフォルトの名無しさん2010/12/12(日) 14:37:54
CORDICって条件分岐ばっかなのでGPGPUには不向きだという先入観があるんだけどどうなの？
0058,,・´∀｀・,,）っ-○○○ 2010/12/12(日) 16:41:21
それは全部ソフトウェアでやったら、の話だろ。
0059デフォルトの名無しさん2010/12/12(日) 17:01:09
おやお久しぶり。
ソフトウェア無線専用ハード（微妙に矛盾？）でCORDICを使ってるとか話には聞いたことがありますな。

あと自分で実装してみて気づいたんだけど、分岐と言ってもある数を足すか引くかなので、
分岐しないようにしてビット操作に落とせるんですよね。
0060,,・´∀｀・,,）っ-○○○ 2010/12/12(日) 17:42:39
GPUは分岐が苦手とはいっても、単純なプレディケートに落とし込めるものならむしろ効率がいいくらいです。
GPUは同一ワープ内で命令ストリームを共有してますから同じ方向にしか進めない。
Cでいうif-elseは一見分岐だけど、GPUではプレディケート情報によって実行・不実行（あるいは結果に反映させない）を
選択する単一の流れに展開されています。

プレディケート自体はそんなに重たくないです。
むしろ分岐先が増えると増えた分だけ処理時間が増えるだけで。
0061デフォルトの名無しさん2010/12/12(日) 19:38:44
A's Video Converterって、10/31付けで配布サイトが閉鎖されてるな
もう手に入らんの？(´・ω・｀)今日HD5670GETしたのに・・・
0062デフォルトの名無しさん2010/12/13(月) 12:44:47
移転してるよ
http://bluesky23.yu-nagi.com/
0063デフォルトの名無しさん2010/12/13(月) 22:14:55
>>62
うぉ！マジでありがとう
ググっても見つけられなかったんだ(T-T)ｳﾚｼｲ
0064デフォルトの名無しさん2010/12/16(木) 04:47:34
プレディケートは現行Intel系GPUでは使い物になりません
GPGPU向けに機能追加されたSandy Bridge GPUコアの登場を待ちましょう。
0065デフォルトの名無しさん2010/12/17(金) 05:55:20
そもそもGPGPUできないし
0066デフォルトの名無しさん2010/12/17(金) 06:04:34
OpenCL対応するんでしょ＞＞Sandy Bridge GPUコア
0067デフォルトの名無しさん2010/12/17(金) 06:36:57
L3使えるから規模の割には速いかもな
0068デフォルトの名無しさん2010/12/17(金) 06:45:10
主にGPUコアで回してるかAVXつこてるかはインテル任せ
だんごに好かれてしまったからGPGPUも端ッパの技術バリエーションの一つに転落決定だな
0069デフォルトの名無しさん2010/12/17(金) 07:00:14
AVXでGPUくらい速くできるなら寧ろ大歓迎だが。
ただしアセンブリ言語で書くのは嫌。
0070デフォルトの名無しさん2010/12/17(金) 21:25:49
そんなもん規模しだいだろ
0071,,・´∀｀・,,）っ-○○○ 2010/12/18(土) 00:01:16
Sandy Bridgeの1EU＝4Way-FMACと仮定しても、まだCPU（AVX）のほうが速いですから
0072デフォルトの名無しさん2010/12/24(金) 16:15:10
ION乗ってるノートでXP入れました！！！！
これでCUDAできますよね？
俺の夢かなえられますよね？

ひゃっはあああああああああああああああああああ
0073デフォルトの名無しさん2010/12/28(火) 07:59:13
IONのCUDAベンチ＆レビューよろしく
0074デフォルトの名無しさん2010/12/28(火) 16:38:30
GPU
グレートプログラマー初春
0075デフォルトの名無しさん2010/12/31(金) 20:09:33
P=>パイパン
0076 [―{}@{}@{}-] デフォルトの名無しさん2011/03/29(火) 02:35:48.68
DSPやFPGA叩いて高速度・複雑なシステム作るよりは
CPU+GPU叩いて作ったほうがはるかにましだがなあ。生産効率が桁違いだわ。
0077デフォルトの名無しさん2011/03/31(木) 21:45:22.56
e?
0078デフォルトの名無しさん2011/03/31(木) 22:10:51.72
>>76
今までDSPやFPGAでやっていた事をCPU+GPUでできる用途ってどんなものがあるの？
考えたけど一つも思い浮かばなかった。
0079デフォルトの名無しさん2011/03/31(木) 23:32:44.83
具体的なカテゴリは勘弁だが、サンプリングしたデータをフィルタで処理して画面に表示
みたいな処理では、GPUでの代替はかなり強烈だよ。
FPGAだとたかだか数百タップの複素FIRフィルタを40〜50MHzの動作速度でさばくのにも
現状だと5〜15万ぐらいのデバイスがいるし。

大量生産するものなら処理をチップ化して安くあげちゃうんだろうけど、
俺のとこみたいな数のでない無線通信製品だとGPGPUはかなり魅力的。
たぶん、画像検査装置みたいな分野でもGPGPUは強力だと思う。
0080デフォルトの名無しさん2011/04/03(日) 16:35:24.45
DX11世代だと、本当に何でも出来そうだよな。
0081デフォルトの名無しさん2011/05/04(水) 19:46:46.10
Linux版のAMD APP 2.4にCALのサンプルが付属していないのですが、
Windows版は付属していますか？
0082デフォルトの名無しさん2011/05/05(木) 22:26:57.83
してません。
それどころかCALは(IL含め)2.5で死滅。
かわりにLLVM IR使え。
そんな感じです。
0083デフォルトの名無しさん2011/05/06(金) 02:02:50.78
>>82
ありがとうございます。
そうですか。。
せっかくIL習得したところなんですが、困りましたね。
0084デフォルトの名無しさん2011/05/06(金) 23:57:54.77
http://developer.amd.com/gpu/AMDAPPSDK/assets/AMD_APP_SDK_Release_Notes_Samples.pdf
ここにはCALのsampleはcalといディレクトリにあるとかいてますが、
旧バージョンから修正されてないだけでしょうか？
0085デフォルトの名無しさん2011/05/27(金) 14:48:43.51
これからGPGPUを勉強する場合、どれを勉強しておくのが良いのでしょうか？
無難という意味では、OpenCLですか？
0086デフォルトの名無しさん2011/05/27(金) 18:28:55.31
ソリューションは結局、問題や環境が決定するもの。
0087デフォルトの名無しさん2011/05/27(金) 22:37:24.01
CUDAでいいんでない
0088デフォルトの名無しさん2011/05/29(日) 19:31:28.97
>>85
ソフトをやるのかハードをやるのか？
ソフトの場合は上位をやるのか下位をやるのか？
0089デフォルトの名無しさん 2011/05/30(月) 21:43:43.33
５００万個の３×３行列の固有値を
（１）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()を使用してマイクロ秒単位で測定。
0090デフォルトの名無しさん2011/05/30(月) 21:44:13.63
(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)　<ーースレッド中の最大値 + α
0091デフォルトの名無しさん 2011/05/30(月) 21:44:46.50
(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万倍速い！　と言う結果になりました。 いくら何でも１５万倍は速すぎのような気が・・・・(^_^;;)

以上。
0092デフォルトの名無しさん 2011/05/30(月) 21:56:52.70
>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)　<ーースレッド中の最大値 + α

0093デフォルトの名無しさん2011/05/31(火) 00:50:53.32
GPGPUについては詳しくないんだけど、
(sizeof float)*3*3*5000000≒180[MB]
これがシステムメモリとVRAM間で往復するから360[MB]
所要時間が2[ms]だから、1[s]に180[GB]も動いてることになる
何か変だ
0094デフォルトの名無しさん2011/05/31(火) 03:26:44.39
ところで、結果は一致してるのか？ｗ
0095デフォルトの名無しさん2011/05/31(火) 04:53:50.24
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

で略一致してます。
0096デフォルトの名無しさん2011/05/31(火) 05:05:03.85

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 　]

（こちらは、固有ベクトルの成分が縦方向に並んでいます）
0097デフォルトの名無しさん2011/05/31(火) 07:21:15.15
GPUのほうは最適化の有無でガラっと変わるんでそこんとこどうなのよ
0098デフォルトの名無しさん 2011/05/31(火) 19:27:14.54

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" と上で書いた時間は実行キューへの登録時間でした。
お騒がせしてすみませんでした。
0099デフォルトの名無しさん2011/05/31(火) 22:53:09.21
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 が楽しみになってきました (^_^)。　
0100デフォルトの名無しさん2011/06/15(水) 12:15:39.69
visual studio pro, radeon 6000 台　で ati stream ないし　open cl
使って、　並列FPU高速化を確認するだけ、ってどのくらい大変ですか？
前提としてCは出来ます
0101デフォルトの名無しさん2011/06/16(木) 00:28:27.42
GPUすごいなあ・・・。もうお前イラネで、失業の危機を感じるわ…。
0102デフォルトの名無しさん2011/06/16(木) 08:09:43.34
>>100
解決させる問題にもよるだろうけれど、本質的には４台でも6000台でも一緒。
それが並列化ってもんだろう。
0103デフォルトの名無しさん2011/06/16(木) 13:12:43.15
GPUをCPUのように扱えるFusion System ArchitectureをAMDが発表
http://pc.watch.impress.co.jp/docs/column/kaigai/20110616_453498.html

.NETで使えるように・・・したのがWPFだっけか
0104デフォルトの名無しさん2011/06/16(木) 23:17:05.13
>>103
全然違う
WPFはGDIの代わりにDirectX使ってるだけ(Vista/7のみ)

FSAに先駆けてVisual Studio2012でついにAMDとnVidiaのGPUがC++プログラミングに完全対応
0105天使 ◆uL5esZLBSE 2011/07/01(金) 18:26:05.76
Rubyバカにしてる子ってさ
変数に＄ついてる言語触ってるって事だよね

いちいちSHIFT＋４キーおして ＄ 打ちまくってる感触はどう？
0106デフォルトの名無しさん2011/07/01(金) 18:56:59.71
【レス抽出】
対象スレ：GPGPU#5
キーワード：ruby
検索方法：マルチワード(OR)

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

いちいちSHIFT＋４キーおして ＄ 打ちまくってる感触はどう？
0107デフォルトの名無しさん2011/07/05(火) 19:30:20.01
GPGPUプログラムしている人、どこののGPU使っている人が多いの？
Intel、nVidia、ATi、どれ？
0108デフォルトの名無しさん2011/07/05(火) 20:26:42.62
速度求めるならATIただしライブラリとドライバが糞
0109デフォルトの名無しさん2011/07/05(火) 20:26:53.66
intelってできるんだ？
0110デフォルトの名無しさん2011/07/06(水) 03:52:32.44
intelからOpenCLのSDKは出てるようだが
0111デフォルトの名無しさん2011/07/06(水) 07:54:48.31
それCPU用だから。でもインテルのだから速いんだろうなきっと。他力本願
0112デフォルトの名無しさん2011/07/06(水) 08:10:04.21
来年のIvyBridgeからだろうなあ
0113デフォルトの名無しさん2011/07/06(水) 09:51:30.14
OpenCV2.3がCUDA対応になったけど、どこまで対応してるんだろうな
0114ふぁふぁ ◆mPVQxchC5E 2011/07/21(木) 11:41:22.23
ってすと
0115デフォルトの名無しさん2011/07/27(水) 23:23:53.51
>>111
SandyBridgeのCore i5にインテルとアムドのOpenCL入れてるが、アムドの方が速いっす
インテルの方はAVX対応とか言ってるくせに効果なし
0116デフォルトの名無しさん2011/08/05(金) 21:37:53.19
CUDAとかそういうGPGPU向けに作られた言語でなくて、プログラマブルシェーダでできるGPGPUってどんなのあるかな
そういう初期のGPGPU用の参考サイトある?
0117デフォルトの名無しさん2011/08/12(金) 15:29:18.34
Cgとかじゃなくて？
0118デフォルトの名無しさん2011/08/12(金) 20:15:53.69
>>117
純粋に古の技術を学びたいです
0119デフォルトの名無しさん2011/08/12(金) 20:23:35.76
DXCSがまさにそれだから、むしろ最新の技術なんじゃないか？
0120デフォルトの名無しさん2011/08/12(金) 20:34:42.35
そうなのか
レンダリング結果のピクセルカラーから値を読み取るやつをやりたかったんだけど、DCCSってのはそうなの？
0121デフォルトの名無しさん2011/08/12(金) 21:02:24.14
DXCSは先祖がえり
0122デフォルトの名無しさん2011/08/12(金) 21:07:24.66
わざわざ先祖返りしたってことは、
書くのは大変だけど速さはGPGPU用よりも速いってことかな

ちょっと調べてみるわサンクス
0123デフォルトの名無しさん2011/08/27(土) 04:28:09.76
おや、こんなスレなんてあったんですね
0124デフォルトの名無しさん2011/09/27(火) 00:26:31.09
JOCLって
0125デフォルトの名無しさん2011/09/27(火) 09:43:30.24
とりあえず安定ならCUDAが一番なのかな？
0126デフォルトの名無しさん2011/09/28(水) 00:15:54.14
安定＝他人のソースをコピペできる
ならそうかな
0127デフォルトの名無しさん2011/09/28(水) 02:40:38.94
不安定だと他人のソースのコピペもできないのか！
0128デフォルトの名無しさん2011/09/28(水) 16:27:35.03
性能ならATi
汎用性ならCUDA
でおｋ？
0129デフォルトの名無しさん2011/09/28(水) 17:17:55.93
>>128
こうじゃね？

性能：ATI Stream
汎用：OpenCL
実用：CUDA
0130デフォルトの名無しさん2011/10/02(日) 09:07:43.00
DirectComputeは？
あとATI Streamは悲しいほどに資料がみつからないんだけどそんなに高性能だったの？
0131デフォルトの名無しさん2011/10/02(日) 11:11:37.87
ATI Streamが高性能というより
AMDGPUの演算性能自体がNVIDIAより2-3倍高い。
NVIDIAが力を入れているC2090のDP性能でも
6970と理論値で互角、実効値では後塵を拝している。
0132デフォルトの名無しさん2011/10/02(日) 13:07:08.51
以前、ATIのIL(アセンブリ)で組んだことあるけど、
チップセット内蔵GPUしか持ってなかったから糞遅かった。
ちゃんとしたGPUで動かすと速いのかな。
0133デフォルトの名無しさん2011/10/02(日) 19:07:21.59
>>131
そうだったんだ。OpenCLやDirectComputeでの比較がないか探してみよっと。
0134デフォルトの名無しさん2011/10/03(月) 09:33:34.11
ATI Streamは本当に資料が無いよな・・・
CUDAの本は何冊か出てるのに

OpenCLで最新ATIの性能をフルに引き出せる？
0135デフォルトの名無しさん2011/10/03(月) 15:09:16.19
前に例の長崎大のGPUスパコンの人がOpenCLでCypressのDGEMMベンチマークやってたよ
0136デフォルトの名無しさん2011/10/03(月) 23:52:36.43
もしかして俺がCAL+ILの日本語本を書いたら大もうけできるのだろうか。
でも需要無いか。
0137,,・´∀｀・,,）っ-○○○2011/10/04(火) 01:30:15.92
次期東大スパコンにも使われる予定のIntel MICのほうが良いかも。
LarrabeeはGPUではなくなってしまったからGP「GPU」ではないかもしれないが。
0138デフォルトの名無しさん2011/10/04(火) 01:34:43.66
>>136
それよりはOpenCLで書いて各ハード向けへの最適化手法を本にしたほうが儲かると思うよ
0139デフォルトの名無しさん2011/10/04(火) 02:11:32.30
NVIDIAのGPUでOpenCLやろうとすると徹底的に最適化しなきゃお話にならないあたりでどうにも
0140デフォルトの名無しさん2011/10/04(火) 09:36:23.59
NVIDIA＜CUDA使え

って意味だな
0141デフォルトの名無しさん2011/10/06(木) 09:36:21.40
CUDAはC#版もいちおー出てるのが大きいな
0142デフォルトの名無しさん2011/10/14(金) 15:14:25.33
それ言ったらOpenCLもJava版出てるよ
0143デフォルトの名無しさん2011/10/14(金) 19:41:56.68
両方いろいろ出てるよな。Python版とか。C/C++しか使わんけど
0144デフォルトの名無しさん2011/10/15(土) 11:09:55.49
pythonからpycuda経由で使ってみているけど、結構便利。
細かいことやろうとすると、結局python内にC(CUDA)のコード埋め込む事になるけど。
0145デフォルトの名無しさん2011/10/15(土) 15:25:56.50
VisualStudio11のDPでC++AMPって使るんだねコンパイルしただけだから
実際どうだかとかわからないけど
0146デフォルトの名無しさん2011/10/17(月) 09:42:39.10
ATI Streamは資料も無いしラッパーも無いし、何も無いのが痛い
0147デフォルトの名無しさん2011/10/17(月) 09:54:44.42
>>146
え？
ttp://developer.amd.com/sdks/AMDAPPSDK/documentation/Pages/default.aspx
じゃだめなの？
0148デフォルトの名無しさん2011/10/18(火) 19:29:48.59
英語読めない人にも配慮してもらわないとな
0149デフォルトの名無しさん2011/10/22(土) 11:35:46.31
>>131
2〜3倍ってそんなに違うのか。ATIはいいもの作ってもマーケッティング面が弱いのかな
0150デフォルトの名無しさん2011/10/22(土) 11:46:32.47
ベンチマークで良い数値が出てもドライバがバグだらけなので実用的じゃないんです。
0151デフォルトの名無しさん2011/10/22(土) 16:00:19.52
AMDとintelのドライバ、どっちが悪いかってくらいダメだからなぁ
0152デフォルトの名無しさん2011/10/22(土) 16:09:57.98
いっそのことプンソにして作って貰った方がいいんじゃね
0153デフォルトの名無しさん2011/10/22(土) 18:35:35.82
>>151
その2つ合わせたよりも高いOSクラッシュ率のNVも相当なもんだぞ
0154デフォルトの名無しさん2011/10/23(日) 11:26:03.73
>>152
AMDのLinux向けドライバはもう大分前からオープンソースでやってるんじゃ？
0155デフォルトの名無しさん2011/10/23(日) 12:44:21.72
両方あるよ。
AMDとNVIDIAの違いはオープンソースコミュニティに
ハードウェアの仕様をドライバ書けるレベルまで公開しているかどうか。
0156デフォルトの名無しさん2011/10/23(日) 13:34:45.12
オープンな方のドライバはopenclをまだ実行できないんじゃないかな
AMDが配布してるクローズドな方は実行できるしもう研究に使ってるとこもあるみたい
0157デフォルトの名無しさん2011/10/26(水) 04:22:15.81
ゲフォで一般向け、と言うかteslaじゃないのは計算誤りが含まれてるから選別したとか、
どこかの大学で言ってたと思うんだけど、その辺の事情はアムも同じなの？
0158デフォルトの名無しさん2011/10/26(水) 05:21:34.17
ソレ言った長崎大の先生はそのあとHD5870でクラスタ組んで論文出してるけどそのへんの事情は言ってないね
まあ保証が欲しいならFirePro買ってくれって立場なのはAMDも変わらんだろうけど
0159デフォルトの名無しさん2011/11/05(土) 13:53:48.44
こんな技術、BEEP音鳴らすブザーで音階を奏でる類の技術なわけで、
広く実用されることは永遠にないと思うな。
0160デフォルトの名無しさん2011/11/05(土) 14:15:01.69
だいぶ違うと思うが

初期はともかく、現在のハードウェアはほぼ完全にHPCのトレンドに乗せてきてるし
0161デフォルトの名無しさん2011/11/06(日) 02:36:21.19
だいぶ違うのは確かだが、広く実用化されるかというとどうだろう。
フィットする問題領域がいくつかあるし、将来もなくならないだろう。
そして多くの領域ではGPGPUが必要にならない。
ほんと下らない当たり前なことで、>159 は何を言いたかったんだろう。
0162デフォルトの名無しさん2011/11/08(火) 10:32:58.02
需要がニッチ過ぎるしなぁ
CPUもコア数増やすしか無くなってきたし、一時的な技術なのは確か
今のところは、単純な計算を繰り返すだけならGPGPUのほうが優位ってだけで
0163デフォルトの名無しさん2011/11/08(火) 10:40:38.94
そりゃまあ、一家に1台クレイ風ベクトルプロセッサ、みたいにはならなかった、
という意味では、たいていの技術は「広く実用されることは永遠にない」ものだろうが。
■ このスレッドは過去ログ倉庫に格納されています