X

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にとっちゃすでに用済みで、要らない子
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が効けばもっと広がると思うんだけど…
やる気無いですよねぇ
0010デフォルトの名無しさん
垢版 |
2010/08/18(水) 23:32:04
LAME（音声の非可逆圧縮）程度じゃ処理が軽すぎるし
条件分岐も少なくないからCPUで計算したほうがいい。

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

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

もっと端的に言ってしまえば分岐命令の数
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部分でメモリリークくさいエラーを吐いてまともに動かないんですが、
窓から投げ捨てるべきでしょうか？
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でソースなし納品はありえんやろ
いつバージョンアップでバイナリが動かなくなってもおかしくないのに
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用中間コードのコンパイルにも影響するが誤差範囲。
0035デフォルトの名無しさん
垢版 |
2010/10/25(月) 12:02:17
初心者なんですけどフリーソフトでATI技術に対応してて
MP4に変換できるソフトってありますか？
あとRADEONのカードってエンコードなら値段と性能みてどれがコスパいいですか？
0037デフォルトの名無しさん
垢版 |
2010/10/26(火) 01:43:50
板違いです。

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

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

質問あればできるものはすべて答えますんでよろしくお願いします
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単位で発行された複数のメモリアクセス命令の間だけ
キャッシュ無しアクセスでもキャッシュが有効になっているから
複数のメモリアクセス命令間でコアレス化と同様の効果が得られるらしい。
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);

を記述したら、それで通りました。
なぜ起動できなかったかは分かりません。
今可能性を探っているのですが、
タイムアウトが起こったのかもしれないと考えています。
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スレあるし、そっちでやるべきかな。
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命令で実行される。
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したのに・・・
0064デフォルトの名無しさん
垢版 |
2010/12/16(木) 04:47:34
プレディケートは現行Intel系GPUでは使い物になりません
GPGPU向けに機能追加されたSandy Bridge GPUコアの登場を待ちましょう。
0068デフォルトの名無しさん
垢版 |
2010/12/17(金) 06:45:10
主にGPUコアで回してるかAVXつこてるかはインテル任せ
だんごに好かれてしまったからGPGPUも端ッパの技術バリエーションの一つに転落決定だな
0071,,・´∀｀・,,）っ-○○○
垢版 |
2010/12/18(土) 00:01:16
Sandy Bridgeの1EU＝4Way-FMACと仮定しても、まだCPU（AVX）のほうが速いですから
0072デフォルトの名無しさん
垢版 |
2010/12/24(金) 16:15:10
ION乗ってるノートでXP入れました！！！！
これでCUDAできますよね？
俺の夢かなえられますよね？

ひゃっはあああああああああああああああああああ
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は強力だと思う。
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習得したところなんですが、困りましたね。
0085デフォルトの名無しさん
垢版 |
2011/05/27(金) 14:48:43.51
これからGPGPUを勉強する場合、どれを勉強しておくのが良いのでしょうか？
無難という意味では、OpenCLですか？
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]も動いてることになる
何か変だ
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 　]

（こちらは、固有ベクトルの成分が縦方向に並んでいます）
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は出来ます
0102デフォルトの名無しさん
垢版 |
2011/06/16(木) 08:09:43.34
>>100
解決させる問題にもよるだろうけれど、本質的には４台でも6000台でも一緒。
それが並列化ってもんだろう。
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、どれ？
0109デフォルトの名無しさん
垢版 |
2011/07/05(火) 20:26:53.66
intelってできるんだ？
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用の参考サイトある?
0120デフォルトの名無しさん
垢版 |
2011/08/12(金) 20:34:42.35
そうなのか
レンダリング結果のピクセルカラーから値を読み取るやつをやりたかったんだけど、DCCSってのはそうなの？
0122デフォルトの名無しさん
垢版 |
2011/08/12(金) 21:07:24.66
わざわざ先祖返りしたってことは、
書くのは大変だけど速さはGPGPU用よりも速いってことかな

ちょっと調べてみるわサンクス
0124デフォルトの名無しさん
垢版 |
2011/09/27(火) 00:26:31.09
JOCLって
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で動かすと速いのかな。
0134デフォルトの名無しさん
垢版 |
2011/10/03(月) 09:33:34.11
ATI Streamは本当に資料が無いよな・・・
CUDAの本は何冊か出てるのに

OpenCLで最新ATIの性能をフルに引き出せる？
0136デフォルトの名無しさん
垢版 |
2011/10/03(月) 23:52:36.43
もしかして俺がCAL+ILの日本語本を書いたら大もうけできるのだろうか。
でも需要無いか。
0137,,・´∀｀・,,）っ-○○○
垢版 |
2011/10/04(火) 01:30:15.92
次期東大スパコンにも使われる予定のIntel MICのほうが良いかも。
LarrabeeはGPUではなくなってしまったからGP「GPU」ではないかもしれないが。
0139デフォルトの名無しさん
垢版 |
2011/10/04(火) 02:11:32.30
NVIDIAのGPUでOpenCLやろうとすると徹底的に最適化しなきゃお話にならないあたりでどうにも
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って使るんだねコンパイルしただけだから
実際どうだかとかわからないけど
0150デフォルトの名無しさん
垢版 |
2011/10/22(土) 11:46:32.47
ベンチマークで良い数値が出てもドライバがバグだらけなので実用的じゃないんです。
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台クレイ風ベクトルプロセッサ、みたいにはならなかった、
という意味では、たいていの技術は「広く実用されることは永遠にない」ものだろうが。
0167デフォルトの名無しさん
垢版 |
2011/12/22(木) 21:17:08.51
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は汎用コンピューティング用途などでも優れたパフォーマンスを発揮できるアーキテクチャだ」
0168デフォルトの名無しさん
垢版 |
2011/12/22(木) 21:38:45.64
自己解決しました。何にせよ新しいモノは楽しみですな。

【後藤弘茂の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に抑えた。
0169デフォルトの名無しさん
垢版 |
2011/12/22(木) 22:53:29.18
つーか乗算の筆算考えてみれば
SPvsDPの計算負荷は1:4の方が自然なのは当たり前
NVIDIAのは倍精度重視というより、単精度軽視という方が近い
SP2倍結局計算できる回路があるのに使えなくしているだけなのだもの

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


0170,,・´∀｀・,,）っ-○○○
垢版 |
2011/12/22(木) 23:33:03.71
Tesla1のときは1:8だったけど理論値あたりの実効性能は90％に達してたんだよね
LSUやメモリなど足回りのパフォーマンスがDP演算器を支えるだけの余裕があったが

FermiでDPを増強したら今度は足回りが追いつかなくなったという皮肉
0171デフォルトの名無しさん
垢版 |
2011/12/22(木) 23:56:33.23
CUDAのアーキテクトがAMDに移籍したらしいがGCNはそれがモロに出たアーキテクチャだな
公式スライドそのままの画像をリークしてたVR-Zoneによると549ドルらしいし
HD6970、GT580より省電力とのことでかなり面白そうだ
0172デフォルトの名無しさん
垢版 |
2011/12/23(金) 01:10:26.44
確かに筆算の掛け算を考えると横に2倍×縦に2倍で4倍の計算量ってのが確認できますね。
仮数部が24bitと53bitで実際は4倍よりもうちょい要るはずだから、1:4でも少なからずリソースを割いているか。

64bit整数演算はまだかなぁ…使うアテないけど。
64bit整数が32bitと並ぶのはVRAMが4GB超えるのとセットなのかな。
0175デフォルトの名無しさん
垢版 |
2011/12/23(金) 13:49:56.45
金のない俺としてはどのクラスまで倍精度をサポートするかが気になるな。

贅沢は言わんから7700シリーズまで頼む。
0176デフォルトの名無しさん
垢版 |
2011/12/23(金) 18:42:26.48
>>174
いいと思うけど、あくまで倍精度演算の速さでだよ。
さらに前後でdoubleに格納したり取り出したりするならその分もかかる。
0184デフォルトの名無しさん
垢版 |
2012/01/13(金) 16:13:40.53
NVIDIAのステマ・広告攻勢がすごいから、
性能が同じくらい＝AMDの圧勝ぐらいの意味だからな
そのあたりを気にしながら記事を読む必要がある
0186デフォルトの名無しさん
垢版 |
2012/01/15(日) 12:54:20.11
nVidiaのステマ能力ははんぱない。
全世界のスパコンシェアを圧倒してしまった。
我々はこの独裁にどう対抗していけばよいのだろうか・・・
0187 ◆QZaw55cn4c
垢版 |
2012/01/15(日) 13:08:06.32
>>186
ATI stream と nvidia CUDA の両刀使いが現れるまで待つしかない、と。
0188デフォルトの名無しさん
垢版 |
2012/01/25(水) 10:42:49.23
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

レイトレを効率よく実行できるアーキテクチャなら使いでがありそうだな。
0190デフォルトの名無しさん
垢版 |
2012/01/25(水) 18:22:40.59
TSUBAME2.0のためにTesla数千枚買っちゃったのに今更AMDが速いとか言ったら各方面から暗殺されかねんからな
0191デフォルトの名無しさん
垢版 |
2012/01/25(水) 20:10:51.84
超要約すると、

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

って事か？
0192デフォルトの名無しさん
垢版 |
2012/01/26(木) 11:48:32.10
「汎用のOpenCLよりnVIDIA特化したCUDAの方が速いよ」だと思う
当たり前のことだがｗ
せめてATIStreamと比較してくれ
0196デフォルトの名無しさん
垢版 |
2012/01/26(木) 17:49:06.78
>>192
「nVIDIA特化したCUDAでnVIDIAが最適化したCUFFTの方が
当然速いと思ったら、それよりは遅い自分たちがCUDAで書いた
コードを汎用のOpenCLに移植してAMDで動かした方が速かった」
という結果が出た上での結論だから。
0197デフォルトの名無しさん
垢版 |
2012/01/26(木) 19:55:34.70
>>195
話題に上がってる東工大の先生がGPUのソフトウェアECCについてまとめてたよ
あれ見るとハードでECC処理するTeslaのありがたみがよくわかる
0199デフォルトの名無しさん
垢版 |
2012/01/26(木) 20:21:24.68
ECCに気を使うのも結構だけど、GPUの場合
そもそも計算間違う可能性についても
議論した方が良いと思うんだが。

長崎大でのGeforceのCUDA不適格率を考えてみても
Teslaがどの程度信頼性あるのか分からん。
0200デフォルトの名無しさん
垢版 |
2012/01/26(木) 21:10:41.07
統計的バラつきを消すために最初から複数回回す必要があるような
アルゴリズムを選ぶ方がいいのかもな。
0201デフォルトの名無しさん
垢版 |
2012/03/10(土) 17:43:48.01
今度のラデは800,700番台でも倍精度演算を実装してるみたいだ。
ちょっと勉強するのに買いやすくていいな。
0205デフォルトの名無しさん
垢版 |
2012/03/21(水) 00:40:57.22
IIR位何も考えずとも効率良くやれるだろ。
並列に計算できるデータが物凄く多ければな。

並列に計算できるデータが少なくて、時系列方向に長い場合は、
FIRなら時系列方向の並列化も可能だけど
IIRは無理じゃね。
0206デフォルトの名無しさん
垢版 |
2012/03/21(水) 10:31:16.75
"IIR filter vectorize"とか"~ SIMD"とかでググったらなんか出てくるし読めばいいんじゃねぇの？
0207デフォルトの名無しさん
垢版 |
2012/03/22(木) 01:33:34.92
>>203
一応単精度では6950ぶっちぎる場面もあるんでまあグラフィックフィルタとか単精度で良いソフトなら試す価値はあるかな
今までと違って動かないわけじゃないから倍精度の開発検証ぐらいはできるしその程度だと思えば…
0210デフォルトの名無しさん
垢版 |
2012/03/23(金) 09:11:06.91
だが、APUあたりのレジスタやキャッシュが激減してるし、チップあたりの
実行可能なカーネル数も減ってる。倍精度もエミュレーションぽいし、GPGPU向きじゃないな。

0215デフォルトの名無しさん
垢版 |
2012/05/07(月) 12:34:41.04
Linux上でGPGPU使って遊びたいんだけど64bitの方がいいかな
もしそうならOS入れ直すんだけど
0216デフォルトの名無しさん
垢版 |
2012/08/19(日) 10:31:27.66
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などへの最適化がうまく
行われていないのだろうか？
0219216
垢版 |
2012/08/19(日) 20:08:03.83
float4とfloat4を引数に使ったmadは多用してるよ。

0220デフォルトの名無しさん
垢版 |
2012/08/20(月) 13:15:46.35
いくらなんでも1400倍はおかしくないか？
OpenCLのことはよくしらないけど、理論性能から考えてもそんなに
差が出るはずが無いと思うが･･･
140倍なら、ＣＰＵのコードがクソならあり得るが
0221デフォルトの名無しさん
垢版 |
2012/08/20(月) 18:30:40.57
i7-2600k (AVX) が単精度 220 GFlops、倍精度 110 GFlops
同 (SSE) が単精度 54 GFlops、倍精度 110 GFlops
Radeon HD 7970 が単精度 3.8 TFlops、倍精度 950 GFlops

SSE でも100 倍以内に収まらないとちと CPU 遅すぎ。
OC しているならばなおさら。
0224216
垢版 |
2012/08/22(水) 23:02:09.92
同じ問題を APU E1-1200 のミニノートPCで解いてみた。
OSは、 >216 と同じくUbuntu12.04 x64版。

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

＜結果＞

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

CPU　 7.542 sec
GPU 0.4235 sec

OpenCLイベントタイマーで計測した正味の計算時間比較では
GPUが１８倍速いと言う結果になった。
これぐらいの比率なら正常？
0225216
垢版 |
2012/08/22(水) 23:10:51.04
ついでに >２１６ で使用したデスクトップマシン環境で同じ１００万個の計算（１/５の量）を
再計測してみると
＜結果＞
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上で実行モジュールを作成している。
0226216
垢版 |
2012/08/22(水) 23:26:18.10
>225

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

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

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

0229デフォルトの名無しさん
垢版 |
2012/08/23(木) 10:37:48.70
たしかに3x3は小さいが、全部の固有値を馬鹿並列で求められるなら、むしろＧＰＵ向きな気もする
0230デフォルトの名無しさん
垢版 |
2012/08/23(木) 10:39:27.25
「100万個の固有値計算を馬鹿並列にできるなら」という意味ね。1つの行列の相異なる固有値を並列に求めるという意味では無く。俺日本語不自由すぎﾜﾛﾀ
0231デフォルトの名無しさん
垢版 |
2012/08/24(金) 01:59:31.99
AMDのOpenCLランタイムは、CPUが
ちょっと異様なぐらい遅いとか聞いたことがある気がする
0232デフォルトの名無しさん
垢版 |
2012/08/24(金) 12:23:43.91
それってSDKが出たばっかの頃の話じゃね
たしかデバッグか何かに絡んだ挙動だったと思ったが
0234デフォルトの名無しさん
垢版 |
2012/08/29(水) 21:33:10.08
レイトレースをCPU、GPU側で同様のロジック作って走らせても1000倍くらい違ったし
GPUは並列度の高いコードだと思った以上に早くなる
CPUはそれだけ理論値からの落ち込みが激しいってことなんだろうけど
0235デフォルトの名無しさん
垢版 |
2012/08/30(木) 05:20:58.76
CPUの方のコードが最適化されてないんじゃないの？
最新のGPUでも単精度ピークが数Tflopsだから、
1000倍差だとCPUは数Gflops以下しか出てないことになる
0238デフォルトの名無しさん
垢版 |
2012/08/31(金) 18:47:52.23
Intelがメニーコア「MIC」とAtom SoCの「Medfield」を発表
ttp://pc.watch.impress.co.jp/docs/column/kaigai/20120831_556528.html

GPU勢もうかうかしていられない？
0240デフォルトの名無しさん
垢版 |
2012/09/09(日) 23:36:12.36
アーキテクチャがx86だってのは個人的にどうでもいいんだが、
メモリ階層がどうなるのか、メモリバンド幅がどれくらいでるのか、
CUDA5のdynamic parallelismに対抗できる機能があるのか、くらいが
勝負の分かれ目だな。
0241デフォルトの名無しさん
垢版 |
2012/10/07(日) 21:53:45.44
>>240
dynamic parallelismなんてわざわざご大層な名前を付けなくても
x86CPUなのだから同等以上のことが出来るに決まっているだろ。
0243デフォルトの名無しさん
垢版 |
2012/10/07(日) 23:20:36.04
ナイコナの汎用性は別に誰も否定してなくね
問題はその効率だよな。やべーさすがIntel様だ、となるのかやっぱ汎用コア並べたらそんなもんだよね、で終わるのか
うちも研究枠で調達予定なので普通にwktkしてます
GPU性能を維持したままにじり寄ってるからこそ、dynamic parallelismなるご大層な名前…というか変態的な局所性能の誇張に縋ってるんだろう
0245デフォルトの名無しさん
垢版 |
2012/10/08(月) 15:19:35.15
えー?
コア一つ当たりの性能ではGPGPUを圧倒するんじゃないの?
問題は、それだけの性能を発揮するために必要なコアサイズが大きすぎることであって。

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


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

としか考えられないよなあ...
0246,,・´∀｀・,,）っ-○○○
垢版 |
2012/10/12(金) 14:31:36.83
Coreの定義がIntelとGPUメーカーとで違うから当たり前だろ
1コアあたりのベクトルユニットが16SP/8DP積和だから
NVIDIAのCUDA Coreに換算すると16コア相当で
つまり50coreのPhiは800 CUDA coreに相当するんだけど？
0247デフォルトの名無しさん
垢版 |
2012/10/12(金) 15:05:22.84
あー、そか。たしかにそうだよな。
コアサイズが二桁ぐらい違うとかじゃないと性能で見劣りすることはないかもしれんのか。

まあ、コアの世代を古い方にして回路規模を小さくしちゃったりするとスループットが落ちて不利だがな、
動作クロックがGPUより上回ってる分と併せて考えると接戦になるのかな?
0248デフォルトの名無しさん
垢版 |
2012/10/12(金) 15:10:50.51
排他処理のルーチンはボトルネックとしか考えられないよなあ...
0249デフォルトの名無しさん
垢版 |
2012/10/12(金) 16:11:57.09
GPGPUはピークはともかく実効値がな･･･
先日出たTSUBAME2.0の1mメッシュ気流解析で実効15%ぐらいだっけ？
複雑な分岐や並列度低いものが入るとしぬし
ほかのGPUアクセラレータ積んだクラスタもそうだけどLINPACKばっか速くても仕方ない
0252デフォルトの名無しさん
垢版 |
2012/10/12(金) 18:14:08.29
良いこと考えた。
働いていない、85%を装置から除去すれば、消費電力も下がるんじゃね？
0257デフォルトの名無しさん
垢版 |
2012/10/13(土) 02:04:35.67
あのプレゼン見たけど、プログラミングに自信のある計算機寄りの人が、
てきとーに題材みつけてきて、計算機科学の研究として発表してる感じだった。

その結果が15%
0260,,・´∀｀・,,）っ-○○○
垢版 |
2012/10/13(土) 04:49:03.94
>>250
働いてないんじゃなくて何かしらがボトルネックになってFMACの稼働率が抑えられてるのでは？
メモリ帯域が足りてる場合でも同時命令発行数の制約でload/store命令と積和命令を
同時に実行できない、とか
（これは京のVenusアーキテクチャにもある制約）
0261デフォルトの名無しさん
垢版 |
2012/10/13(土) 16:36:04.21
唐突ですまんけど、AMDはファイル名通りのこういう資料があるけど
ttp://developer.amd.com/gpu_assets/R700-Family_Instruction_Set_Architecture.pdf

nVIDIAはこういうの無いの？
0262デフォルトの名無しさん
垢版 |
2012/10/13(土) 18:21:13.66
>>249
効率はFLOPSに対する効率であって、気象系のアプリは
バンド幅律速だからLINPACKと比べること自体ナンセンス。
ちなみに15％は気象系のステンシルアプリとしては高い方。
0263デフォルトの名無しさん
垢版 |
2012/10/13(土) 18:24:01.07
>>257
そらまぁ、青木研究室の存在意義が「GPGPUでできることを広げる」
だからね。仕方ない。
NECのSXが沈没の今、京かGPGPUしか無いわけで･･･
それともBlueGeneを買うか？
0264デフォルトの名無しさん
垢版 |
2012/10/13(土) 18:27:15.78
FLOPSに対する効率だから、アルゴリズム自体の並列化効率以上にはなりえないわけか。
気象系のアプリというものの並列化効率がどんくらいなのか知らないけど（なんとなく高そうではあるが）。
0265デフォルトの名無しさん
垢版 |
2012/10/13(土) 18:51:02.05
>>264
うーん、微妙に違うな。

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

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

で、ここで問題になるのが、浮動小数点演算（FLOP）と、データ転送量の比。
業界だと byte/flops とか言われる。言い方を変えると、どれくらいデータの
使い回しが効くかってこと。
続く
0266デフォルトの名無しさん
垢版 |
2012/10/13(土) 18:52:43.77
承前
アプリの要求byte/flops高いと、せかっくデータを持ってきても
たいした計算をせずに、すぐに次のデータが必要になっちゃう。
データの使い回しが効かないんだ。これはイマドキの計算機にとってはキツい。
で、気象とか流体とかはそういうアプリなわけ。

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

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

ちなみにスレ的に言うと、GPUはGDDR５っていう容量が少ない代わりに
バンド幅が出る贅沢なメモリを搭載して、かつ大量のスレッドを使うことで
レイテンシを隠蔽してる。これがGPGPUが気象・流体計算に使える理由。
0267264
垢版 |
2012/10/13(土) 19:29:29.68
なるほど、つまり気象系のアプリはコンピュートバウンドではなくメモリバウンドのものが多い、と。
AMDのAPUの使い道の話になったときに、DDR3という制約に悲観的になってる人が少なそうだったから
世の中にはメモリバウンドのものは少数派なのかなあとも思ってたけど普通にあるんですね。
0268デフォルトの名無しさん
垢版 |
2012/10/13(土) 19:43:21.60
バンド幅と言うか、物理原理的にDRAMのセルは応答が現代のCPUの演算速度に対してずっと遅い。
だから、メモリセルを並列にして同時に読み出し・書き込みするしかないのが根底にある。それをやれDDRxだの、
何ビット同時にアクセス＝複数セルに同時アクセスでなんとかしてるね。

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

こういう根本的なところはもう全然進歩しないなw
0269,,・´∀｀・,,）っ-○○○
垢版 |
2012/10/13(土) 20:27:18.96
bytes/flopsに関して言えば現状GPUはCPUより低いです。
ただbyte/sが数倍高くて、flopsが更に数倍高いというだけの話で。
Haswellあたりで、電力当たりFLOPS数はGPUとCPUは互角くらいになるのではという話も出ていますな。
0271デフォルトの名無しさん
垢版 |
2012/10/14(日) 11:35:52.75
>>241
亀レスだけど、何のペナルティも無く普通のx86コアを50個
並べられるなら誰も苦労してないわけで･･･
0272デフォルトの名無しさん
垢版 |
2012/10/14(日) 12:44:22.63
DRAM内部のクロック
PC133　133MHz
DDR3-1600　200MHz
この間十数年・・・MRAMならもっと上がるかな？
0273デフォルトの名無しさん
垢版 |
2012/10/20(土) 09:38:39.11
GNU MPのGPGPU対応って、あんまりやってる人が居ないな。
0274,,・´∀｀・,,）っ-○○○
垢版 |
2012/10/20(土) 09:42:31.34
依存関係があってなかなか並列化しにくいからね。
もちろん並列度の高いコンピュータに適した数式に置き換えればいいんだが
CPU側もAVX, FMAなどでGPGPUの電力効率の優位も揺らぎつつあるからね
0275デフォルトの名無しさん
垢版 |
2012/11/17(土) 23:26:42.82
C++AMP調べてるとこなんだけど
textureでピクセル間を補間する方法がわからない。
ひょっとして自前でやらないといけないの？
「C++ AMP API は、サンプリング テクスチャのフィルター処理機能を提供されません。」
とはこのことなのかな。
0276デフォルトの名無しさん
垢版 |
2012/11/18(日) 21:58:01.39
c++ampにfloat_3のmadやfmaが無いね。 short vector typeのmadやsin, cos などの超越関数も用意してもらいたい。 swizzlingはできるみたい。 OpenCLの方がコードをコンパクトに書けるかな。次期版期待してます。
0277デフォルトの名無しさん
垢版 |
2013/02/17(日) 08:49:52.18
何故かsea islandのISAが公開
ttp://developer.amd.com/tools/heterogeneous-computing/amd-accelerated-parallel-processing-app-sdk/documentation/
0279デフォルトの名無しさん
垢版 |
2013/02/18(月) 13:29:16.68
ディスクリートとしては出なくても
次世代APUのGCNはsea islandなのかね。

ISA見る限り、SIから確かに機能アップされている。
一番大きいのはFLAT メモリアクセス命令の追加か。
cudaのUVAみたいなもの
0283デフォルトの名無しさん
垢版 |
2013/04/14(日) 14:30:38.59
抽象化の度合いが高いぐらいの意味合いでいいんじゃないの
実際問題プログラマが意識してGPUに仕事をさせるコードを書かずに勝手に解決してくれちゃう未来像は来るの？
0284デフォルトの名無しさん
垢版 |
2013/04/14(日) 18:29:31.68
たぶん、デザインパターンと同じくらいの敷居にはなるんじゃねーの？と言うか、デザインパターンとして確立されるんじゃね。

判らない人には判らないけど、自分の中に消化され始めると自然とそう言う書き方になるみたいな。
あとは、フレームワークでそう言う書き方を強制されて、なにも考えずにそう言うもんだと思うパターン。
0285デフォルトの名無しさん
垢版 |
2013/04/15(月) 02:40:47.34
コア数が100くらいのGPUがあるとして
整数大規模行列の
全要素の和を求めるのってGPGPUで100倍程度の高速化は可能ですよね？
0286デフォルトの名無しさん
垢版 |
2013/04/15(月) 07:36:47.18
1コアあたりのシングルスレッド能力がCPUと同程度あって、メモリ帯域幅がそれに見合うほど広ければな
0287デフォルトの名無しさん
垢版 |
2013/04/22(月) 21:00:09.19
>>283
実際、解きたい問題依存だけど、Hadoopのように、プログラミングモデルを固定して制限を付けて縛ることによって簡単な記述で自動的に高速実行ってのは比較的現実的な未来だと思うよ。
Hadoopってのも、書きたい処理をHadoopの"型"に押し込んで書ければ、残りの面倒な問題は全部Hadoopが面倒見てくれますって事だし。
0289デフォルトの名無しさん
垢版 |
2013/09/21(土) 13:04:48.76
OpenCLで書いたプログラムを9600GTとQ9650で動かしたら同じくらいの速度だった。
今のCPUには絶対負ける。
0292デフォルトの名無しさん
垢版 |
2013/11/06(水) 09:10:41.13
>>291
DirectXやOpenGLは長い歴史があって
色々なハードで動かす必要があるので
互換性維持のために滅茶苦茶厚いAPIコールを通さないといけないが
MantleはGCNに絞った最適化さえすればいいので
そういうものがないっつーこと
据え置きはAMD一色だから移植も相互に楽になるし、面倒な最適化も最小限で済む
というのが言い分
0293デフォルトの名無しさん
垢版 |
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をやってた層には移行先として十分かなと言った感じ
0294デフォルトの名無しさん
垢版 |
2013/12/30(月) 03:37:10.21
Mantleって互換性維持のため分厚くなり過ぎたDirectXに対して
低レベルなところを柔軟に叩けるようにするとかそういうものじゃなかったっけ？

Project Sumatraが楽しみだな
Java9ぐらいでparallelStreamにぺぺっと渡すと
OpenCLやってくれちゃうとか最高じゃん
逆に言うとそれぐらい猿でも簡単に扱えないとなっていうか
0298デフォルトの名無しさん
垢版 |
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)です。
0299デフォルトの名無しさん
垢版 |
2014/01/26(日) 21:24:38.81
メモリ転送のオーバーヘッドがあるから、もっと大きな問題じゃないと効果は出ないよ。
0300デフォルトの名無しさん
垢版 |
2014/01/26(日) 21:54:26.48
>>299
それは知っているのですが、数値を大きくするとすぐにc00000fdでクラッシュするんですよ……。
<<<grid, block>>>もいろいろ弄っているのですが、どうにも効果が得られません。
今試してみたら、arraySizeの値で実行できるのは25600が最大みたいです。
0301デフォルトの名無しさん
垢版 |
2014/01/26(日) 22:08:07.85
GLSLからOpenCLへの移行を昨日から始めたけど
GLSLより書きやすいのはいいけど最適化を追い込まないととんでもなく遅くなるんだな
GLSLで複雑な汎用計算やらせるのは難解なパズルゲームみたいで嫌になってたけど
結局最適化の手間を考えたらどっちが楽ということはないんだね・・・

>>298みたいな単純な計算ならGLSLだとバグったような速度が簡単に出るから別世界感が凄い
0302298
垢版 |
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までは大丈夫ということなのでしょうか？
0303デフォルトの名無しさん
垢版 |
2014/01/26(日) 23:22:14.60
>>299
GPGPUはメモリ転送のオーバーヘッドがないHSA(Huma)だよな
PCではど重い処理でない限りAMDのHSAがGPGPU処理の主流になるだろうな
0304デフォルトの名無しさん
垢版 |
2014/01/27(月) 00:23:03.99
>>298
残念なお知らせ。
そのソースコードでは、GPUの演算時間ではなくGPUの呼び出し時間しか計測してないね。
「実際の演算時間」=「内部ブロック数」*(「内部ブロックの呼び出し時間」+「内部ブロックの演算時間」)だとすると、
「実際の演算時間」-「内部ブロックの演算時間」になっているはず。
ブロック数が充分大きければ誤差だけど、内部ブロック数が1のときは激速になってしまう。

まぁ、実際の運用ではCPUとGPUが並列に動作することを期待するからそれでもいいんだけどね。
いずれにしても、CPUぶん回すよりも手っ取り早いと思っていたら大間違いだよ。
それと、CUDAスレも宜しく。
0305298
垢版 |
2014/01/27(月) 00:47:11.47
>>304
>そのソースコードでは
え！？　……つまり、
普通にtimeGetTimeかQueryPerformanceCounterとかを使えってことなんですか？
それとも、測定する位置が間違っているということなんですか？
>CUDAスレも宜しく
分かりました。次回以降はそちらにレスすることにします。
0308307
垢版 |
2014/01/27(月) 21:39:43.68
うっかり、166行目を「cudaStatus = cudaSetDevice(1);」にしちゃったので、適当に直しておいて。
0309デフォルトの名無しさん
垢版 |
2014/01/27(月) 23:30:12.43
ローカルメモリを使う場合って確保しようとした容量が大き過ぎると
グローバルのほうへ確保されてしまうんだよね？
AMDのGCNはどれくらいまでローカルメモリがあるのか分からないんだけど
試行錯誤して調べるしかないのか
0310298
垢版 |
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]らしいです。
0312デフォルトの名無しさん
垢版 |
2014/01/29(水) 01:59:06.39
OpenCLのclEnqueueNDRangeKernelでカーネルを実行するときに
global_work_sizeとlocal_work_sizeに同じ値(256,256など)を入力すると
何もエラーは返されずにメモリの参照が壊れて？しまいclEnqueueReadBufferで
CPU側で読み取った値が全て0になってしまいます。

これは仕様なのでしょうか？
0314デフォルトの名無しさん
垢版 |
2014/02/25(火) 21:43:30.35
>>313
そのためだけにVS2012と2013使い分けてる俺……
0316デフォルトの名無しさん
垢版 |
2014/04/15(火) 02:32:13.65ID:vGWbAtXL
(・∀・)ﾉ CPUの300倍くらいの性能が出たお！
比較したCPUはE2-2000っていうCPU+GPU=APUだけど全くGPUとしての機能をもってないのでガッカリしたお。
0317デフォルトの名無しさん
垢版 |
2014/04/19(土) 12:16:56.16ID:Firi/9oq
(・∀・)ﾉ ALU(IGP)のE2-2000はHD7770の1/50のパワーしかないが並列性はあるようだ。
0320デフォルトの名無しさん
垢版 |
2014/09/12(金) 04:54:39.38ID:jvr90R5c
テキスト処理ってGPUで高速化できないものでしょうか
具体的には
Appache Solr
の検索処理が遅いのでなんとか高速化したいのですが
0324デフォルトの名無しさん
垢版 |
2014/09/24(水) 15:13:16.24ID:ltG1hZ24
OpenCLでプログラム組んでみたけど、CPUとGPUメモリのやり取りがネックになっているのか、思ったよりスピードが出ない

他の人はGPU利用するにあたってメモリのやり取りとか何か工夫している？
0327デフォルトの名無しさん
垢版 |
2014/09/25(木) 16:05:41.16ID:Coq6ADbv
基本はメモリとのやりとりを少なくするって話でしょ
それ以上の個別の工夫を簡単に説明するのは難しいよね
ケーススタディしたいのならそういう本なり文献なり漁るべき
0331デフォルトの名無しさん
垢版 |
2014/09/27(土) 00:53:04.08ID:SNKkkpyl
>>329
買えたら買ってるよ
メモリの転送の処理が要らなくなったらと思うと幸せな気分になれるよ
0332デフォルトの名無しさん
垢版 |
2014/10/14(火) 21:40:56.17ID:noiOU3fL
kaveriってOpenCL使うとき、コピーせずにポインタ参照で渡していいって解釈でいいの？
最近GPGPUをやりはじめたばかりだから、的外れなことかもしれんが。。
0336デフォルトの名無しさん
垢版 |
2015/05/10(日) 00:43:24.29ID:60tvXotD
vexclを少し使ってみたけど便利だ
あとはC++AMPみたいにradeonのドライバーの
バージョン上がると使えなくなったりしないなら
0337デフォルトの名無しさん
垢版 |
2015/06/23(火) 13:14:06.30ID:AOM31ZzX
GPUの行列演算ライブラリってないですか？
具体的には特異値分解できるのを探してます
0338デフォルトの名無しさん
垢版 |
2015/06/23(火) 13:41:37.08ID:DUXK3D31
>>337
機械学習スレで書いてた人かな？
ちゃんと調べてないけどMAGMAなら入ってるかも
http://icl.cs.utk.edu/magma/overview/index.html
{sdcz}gesvd はサポートしてるって書いてある

ただし、GPUカーネル内から直接呼びたい場合は使えないらしい
CPUからカーネル呼び出しする必要がある
（SC14時点の資料）
0339デフォルトの名無しさん
垢版 |
2015/06/26(金) 11:47:09.05ID:JVzNXP51
>>338
ありがとうございます。
0340デフォルトの名無しさん
垢版 |
2015/07/08(水) 11:17:46.41ID:i7xBLVJ6
最大固有値
最大固有ベクトル
だけを求めたい場合って、
0341デフォルトの名無しさん
垢版 |
2015/07/08(水) 11:18:15.73ID:i7xBLVJ6
最大固有値
最大固有ベクトル
だけを求めたい場合って、べき乗法が最速でしょうか？
0343デフォルトの名無しさん
垢版 |
2015/10/11(日) 13:46:03.27ID:9Az+Dnte
VS2015のc++amp仕様が変わった？
ＣＰＵで実行するrestrict(cpu)のマイクロソフトのサンプルコードがコンパイルエラーになる。
0345125
垢版 |
2016/01/30(土) 08:09:23.91ID:gCqMUv9A
マイクロソフトの開発ブログで

＞ Is C＋＋AMP dead ？

との質問にレスが無い。

モスさんどこ行った？
0346デフォルトの名無しさん
垢版 |
2016/06/15(水) 14:41:37.67ID:d2Xou3GL
test
0347わたしはぐぷぐぷ派です
垢版 |
2016/12/28(水) 13:24:22.43ID:6d1C8mET
は？ごぽごぽに決まってるだろ？って言われた．．．
頭ごなしに言ってくる人って何なんでしょうね
0348デフォルトの名無しさん
垢版 |
2018/05/23(水) 23:05:05.10ID:Au5e7VGg
僕の知り合いの知り合いができたパソコン一台でお金持ちになれるやり方
役に立つかもしれません
グーグルで検索するといいかも『ネットで稼ぐ方法　モニアレフヌノ』

MYAIQ
0349デフォルトの名無しさん
垢版 |
2018/07/04(水) 23:04:08.53ID:gFgZc5FG
2XM
■ このスレッドは過去ログ倉庫に格納されています
ニューススポーツなんでも実況