11/10/22 11:46:32.47
ベンチマークで良い数値が出てもドライバがバグだらけなので実用的じゃないんです。
151:デフォルトの名無しさん
11/10/22 16:00:19.52
AMDとintelのドライバ、どっちが悪いかってくらいダメだからなぁ
152:デフォルトの名無しさん
11/10/22 16:09:57.98
いっそのことプンソにして作って貰った方がいいんじゃね
153:デフォルトの名無しさん
11/10/22 18:35:35.82
>>151
その2つ合わせたよりも高いOSクラッシュ率のNVも相当なもんだぞ
154:デフォルトの名無しさん
11/10/23 11:26:03.73
>>152
AMDのLinux向けドライバはもう大分前からオープンソースでやってるんじゃ?
155:デフォルトの名無しさん
11/10/23 12:44:21.72
両方あるよ。
AMDとNVIDIAの違いはオープンソースコミュニティに
ハードウェアの仕様をドライバ書けるレベルまで公開しているかどうか。
156:デフォルトの名無しさん
11/10/23 13:34:45.12
オープンな方のドライバはopenclをまだ実行できないんじゃないかな
AMDが配布してるクローズドな方は実行できるしもう研究に使ってるとこもあるみたい
157:デフォルトの名無しさん
11/10/26 04:22:15.81
ゲフォで一般向け、と言うかteslaじゃないのは計算誤りが含まれてるから選別したとか、
どこかの大学で言ってたと思うんだけど、その辺の事情はアムも同じなの?
158:デフォルトの名無しさん
11/10/26 05:21:34.17
ソレ言った長崎大の先生はそのあとHD5870でクラスタ組んで論文出してるけどそのへんの事情は言ってないね
まあ保証が欲しいならFirePro買ってくれって立場なのはAMDも変わらんだろうけど
159:デフォルトの名無しさん
11/11/05 13:53:48.44
こんな技術、BEEP音鳴らすブザーで音階を奏でる類の技術なわけで、
広く実用されることは永遠にないと思うな。
160:デフォルトの名無しさん
11/11/05 14:15:01.69
だいぶ違うと思うが
初期はともかく、現在のハードウェアはほぼ完全にHPCのトレンドに乗せてきてるし
161:デフォルトの名無しさん
11/11/06 02:36:21.19
だいぶ違うのは確かだが、広く実用化されるかというとどうだろう。
フィットする問題領域がいくつかあるし、将来もなくならないだろう。
そして多くの領域ではGPGPUが必要にならない。
ほんと下らない当たり前なことで、>159 は何を言いたかったんだろう。
162:デフォルトの名無しさん
11/11/08 10:32:58.02
需要がニッチ過ぎるしなぁ
CPUもコア数増やすしか無くなってきたし、一時的な技術なのは確か
今のところは、単純な計算を繰り返すだけならGPGPUのほうが優位ってだけで
163:デフォルトの名無しさん
11/11/08 10:40:38.94
そりゃまあ、一家に1台クレイ風ベクトルプロセッサ、みたいにはならなかった、
という意味では、たいていの技術は「広く実用されることは永遠にない」ものだろうが。
164:デフォルトの名無しさん
11/12/22 14:36:12.59
ラデの7000シリーズが出たな。
165:デフォルトの名無しさん
11/12/22 14:54:46.40
金のある研究員の人
ラデの人柱になってください
166:デフォルトの名無しさん
11/12/22 15:15:25.43
某東北の公立大の中の人がなるんじゃないか?
167:デフォルトの名無しさん
11/12/22 21:17:08.51
4Gamer.net ― AMD,新世代ハイエンドGPU「Radeon HD 7970」を発表―Southern Island世代のGPUアーキテクチャを整理する
URLリンク(www.4gamer.net)
>つまり,GNCでSPの演算機能に手は入っておらず,単精度の浮動小数点積和演算・
>積和算・乗算・加算と整数演算のみをサポートしたものだということだ。
>言い換えると,VLIW5アーキテクチャにおける“ビッグSP”のような,
>倍精度演算や超越関数演算に対応した特別機能ユニットは搭載されていない。
Graphics Core Nextって倍精度が激遅になったりしないんだろうか。
一応次のようなCTOの発言↓があるのでそれをカバーするしくみがあるのかもしれないけど。
>「依存関係にない複数の命令を1命令としてまとめて実行できるVLIW方式も,
>グラフィックス用途では十分に活用できるアーキテクチャだが,
>GCNは汎用コンピューティング用途などでも優れたパフォーマンスを発揮できるアーキテクチャだ」
168:デフォルトの名無しさん
11/12/22 21:38:45.64
自己解決しました。何にせよ新しいモノは楽しみですな。
【後藤弘茂のWeekly海外ニュース】 大きく進化を遂げた新世代GPU「Radeon HD 7900」
URLリンク(pc.watch.impress.co.jp)
>倍精度演算などは命令発行に16サイクルかかるため、Graphics Core Nextでは単精度と倍精度の
>ピークパフォーマンス比率は1対4となっている。925MHz動作なら947GFLOPSの倍精度性能となる。
>ここはトレードオフで、HPC(High Performance Computing)を重視するNVIDIAは、
>ダイ効率のトレードオフを払っても、単精度と倍精度の比率を1対2にしたが、AMDは1対4に抑えた。
169:デフォルトの名無しさん
11/12/22 22:53:29.18
つーか乗算の筆算考えてみれば
SPvsDPの計算負荷は1:4の方が自然なのは当たり前
NVIDIAのは倍精度重視というより、単精度軽視という方が近い
SP2倍結局計算できる回路があるのに使えなくしているだけなのだもの
まあ、レジスタ帯域とか考えれば1:2も分かるけど
これはレジスタの割り当て自由度を無駄に上げて
レジスタ帯域を上げる事を難しくしている
Fermiの構造的な問題だからな。
170:,,・´∀`・,,)っ-○○○
11/12/22 23:33:03.71
Tesla1のときは1:8だったけど理論値あたりの実効性能は90%に達してたんだよね
LSUやメモリなど足回りのパフォーマンスがDP演算器を支えるだけの余裕があったが
FermiでDPを増強したら今度は足回りが追いつかなくなったという皮肉
171:デフォルトの名無しさん
11/12/22 23:56:33.23
CUDAのアーキテクトがAMDに移籍したらしいがGCNはそれがモロに出たアーキテクチャだな
公式スライドそのままの画像をリークしてたVR-Zoneによると549ドルらしいし
HD6970、GT580より省電力とのことでかなり面白そうだ
172:デフォルトの名無しさん
11/12/23 01:10:26.44
確かに筆算の掛け算を考えると横に2倍×縦に2倍で4倍の計算量ってのが確認できますね。
仮数部が24bitと53bitで実際は4倍よりもうちょい要るはずだから、1:4でも少なからずリソースを割いているか。
64bit整数演算はまだかなぁ…使うアテないけど。
64bit整数が32bitと並ぶのはVRAMが4GB超えるのとセットなのかな。
173:デフォルトの名無しさん
11/12/23 03:05:28.01
>>170
行列積の話?
174:デフォルトの名無しさん
11/12/23 09:41:52.36
64bit整数演算はできないけど、53bitはできるって理解していいのでしょうか?
175:デフォルトの名無しさん
11/12/23 13:49:56.45
金のない俺としてはどのクラスまで倍精度をサポートするかが気になるな。
贅沢は言わんから7700シリーズまで頼む。
176:デフォルトの名無しさん
11/12/23 18:42:26.48
>>174
いいと思うけど、あくまで倍精度演算の速さでだよ。
さらに前後でdoubleに格納したり取り出したりするならその分もかかる。
177:デフォルトの名無しさん
11/12/28 02:26:18.12
おすすめのGPUを教えてください
178:デフォルトの名無しさん
11/12/28 09:43:57.74
最新の一番高いやつをその都度
179:デフォルトの名無しさん
11/12/28 11:54:58.22
URLリンク(www.freepatentsonline.com)
Multipurpose functional unit with double-precision and filtering operations
180:デフォルトの名無しさん
11/12/28 11:57:58.69
Warpのダイナミック再構成がつくって話だなkepler
181:デフォルトの名無しさん
11/12/31 14:45:52.09
AMD's got an ace up it's sleeve: Tahiti-ASIC probably has 36 CUs/2304 Shaders
URLリンク(www.gpu-tech.org)
182: 忍法帖【Lv=4,xxxP】
12/01/05 23:26:19.60
URLリンク(ascii.jp)
大原雄介の記事来た
URLリンク(news.mynavi.jp)
183:デフォルトの名無しさん
12/01/06 12:45:44.61
URLリンク(news.mynavi.jp)
めがっさ速い?
184:デフォルトの名無しさん
12/01/13 16:13:40.53
NVIDIAのステマ・広告攻勢がすごいから、
性能が同じくらい=AMDの圧勝ぐらいの意味だからな
そのあたりを気にしながら記事を読む必要がある
185:デフォルトの名無しさん
12/01/14 17:46:44.04
公式の7970用ドライバ来る前の記事だからな
186:デフォルトの名無しさん
12/01/15 12:54:20.11
nVidiaのステマ能力ははんぱない。
全世界のスパコンシェアを圧倒してしまった。
我々はこの独裁にどう対抗していけばよいのだろうか・・・
187: ◆QZaw55cn4c
12/01/15 13:08:06.32
>>186
ATI stream と nvidia CUDA の両刀使いが現れるまで待つしかない、と。
188:デフォルトの名無しさん
12/01/25 10:42:49.23
Revenge is Sweet: PowerVR Discrete GPGPU PCIe Card Coming Later in 2012
URLリンク(vr-zone.com)
レイトレを効率よく実行できるアーキテクチャなら使いでがありそうだな。
189:デフォルトの名無しさん
12/01/25 18:08:37.86
日本でのCUDA最先鋒な
東工大のAFDSでの発表が面白い。
URLリンク(developer.amd.com)
FFTをOpenCLで6970に移植したら、さくっとC2050や580の
最速実装を超えてしまったけど、でもだからと言って
AMDの方が良いとは言わないよとか。
190:デフォルトの名無しさん
12/01/25 18:22:40.59
TSUBAME2.0のためにTesla数千枚買っちゃったのに今更AMDが速いとか言ったら各方面から暗殺されかねんからな
191:デフォルトの名無しさん
12/01/25 20:10:51.84
超要約すると、
「なぜなら、このプログラムがあるのはNvidia賛のおかげだから(キリッ」
って事か?
192:デフォルトの名無しさん
12/01/26 11:48:32.10
「汎用のOpenCLよりnVIDIA特化したCUDAの方が速いよ」だと思う
当たり前のことだがw
せめてATIStreamと比較してくれ
193:デフォルトの名無しさん
12/01/26 13:04:08.77
なぜならECCが無いから問題外
194:デフォルトの名無しさん
12/01/26 15:03:56.23
ソフトウェアECCって処理によってがくっと性能下がるらしいからなぁ
195:デフォルトの名無しさん
12/01/26 17:07:39.79
たとえば?
196:デフォルトの名無しさん
12/01/26 17:49:06.78
>>192
「nVIDIA特化したCUDAでnVIDIAが最適化したCUFFTの方が
当然速いと思ったら、それよりは遅い自分たちがCUDAで書いた
コードを汎用のOpenCLに移植してAMDで動かした方が速かった」
という結果が出た上での結論だから。
197:デフォルトの名無しさん
12/01/26 19:55:34.70
>>195
話題に上がってる東工大の先生がGPUのソフトウェアECCについてまとめてたよ
あれ見るとハードでECC処理するTeslaのありがたみがよくわかる
198:デフォルトの名無しさん
12/01/26 20:13:42.50
ECCってハミング符号?
だったら、ソフトウェアで実装したら遅くなるのは道理だね
199:デフォルトの名無しさん
12/01/26 20:21:24.68
ECCに気を使うのも結構だけど、GPUの場合
そもそも計算間違う可能性についても
議論した方が良いと思うんだが。
長崎大でのGeforceのCUDA不適格率を考えてみても
Teslaがどの程度信頼性あるのか分からん。
200:デフォルトの名無しさん
12/01/26 21:10:41.07
統計的バラつきを消すために最初から複数回回す必要があるような
アルゴリズムを選ぶ方がいいのかもな。
201:デフォルトの名無しさん
12/03/10 17:43:48.01
今度のラデは800,700番台でも倍精度演算を実装してるみたいだ。
ちょっと勉強するのに買いやすくていいな。
202:デフォルトの名無しさん
12/03/11 05:18:50.99
おまえはもう達人だろ。
203:デフォルトの名無しさん
12/03/20 09:49:51.77
URLリンク(news.mynavi.jp)
77x0,78x0の倍精度はエミュレーションっぽいな。
やっぱり79x0買うしかないか。
204:デフォルトの名無しさん
12/03/20 13:12:54.98
GPUでIIRフィルタを効率よくやる方法ってあります?
205:デフォルトの名無しさん
12/03/21 00:40:57.22
IIR位何も考えずとも効率良くやれるだろ。
並列に計算できるデータが物凄く多ければな。
並列に計算できるデータが少なくて、時系列方向に長い場合は、
FIRなら時系列方向の並列化も可能だけど
IIRは無理じゃね。
206:デフォルトの名無しさん
12/03/21 10:31:16.75
"IIR filter vectorize"とか"~ SIMD"とかでググったらなんか出てくるし読めばいいんじゃねぇの?
207:デフォルトの名無しさん
12/03/22 01:33:34.92
>>203
一応単精度では6950ぶっちぎる場面もあるんでまあグラフィックフィルタとか単精度で良いソフトなら試す価値はあるかな
今までと違って動かないわけじゃないから倍精度の開発検証ぐらいはできるしその程度だと思えば…
208:デフォルトの名無しさん
12/03/22 23:22:18.46
>204
けち臭いこと言わずFFTしてフィルタリングしろ。
209:デフォルトの名無しさん
12/03/23 00:55:17.20
GTX680出たな
210:デフォルトの名無しさん
12/03/23 09:11:06.91
だが、APUあたりのレジスタやキャッシュが激減してるし、チップあたりの
実行可能なカーネル数も減ってる。倍精度もエミュレーションぽいし、GPGPU向きじゃないな。
211:デフォルトの名無しさん
12/03/23 09:11:36.03
APU→ALU
212:デフォルトの名無しさん
12/03/23 13:08:36.71
トレードオフしっかり設計に反映させないと難しい世代になってきたって事かねぇ
213:デフォルトの名無しさん
12/03/23 21:26:08.99
>210
具体的な数値って分かる?
214:デフォルトの名無しさん
12/03/23 21:29:40.24
pcヲチの後藤さんの記事のブロック図を見れ
215:デフォルトの名無しさん
12/05/07 12:34:41.04
Linux上でGPGPU使って遊びたいんだけど64bitの方がいいかな
もしそうならOS入れ直すんだけど
216:デフォルトの名無しさん
12/08/19 10:31:27.66
Ubuntu12.04 x64版上で、OpenCLを使って3x3行列の固有値&固有ベクトルを
単精度のJacobi法で計算させてみた。
CPU:intel i7-2600k (4.4GHz)
GPU:AMD Radeon-HD7970 (1.05GHz)
行列の個数5百万個(全て同じ行列を使用)
OpenCLソース内で同じ計算の繰り返しを10回
(->5千万個の3x3行列の固有値&固有ベクトルを求めたことに相当?)
OpenCLのソースコードはCPUとGPUで同じものを使用。
<結果>
CPUを使ったOpenCL(Jacobi法の反復回数は6回で収束)
カーネルの実行時間:47.68秒
GPUを使ったOpenCL(Jacobi法の反復回数はCPUと同じで6回で収束)
カーネルの実行時間:0.03389秒
GPUが1400倍も速いと言う結果になった。
OpenCLがAMD製なので、IntelのSSE、AVXなどへの最適化がうまく
行われていないのだろうか?
217:デフォルトの名無しさん
12/08/19 19:29:32.18
そもそもSSEやAVX使うようなコード書いているの?
218:デフォルトの名無しさん
12/08/19 19:37:00.00
float8とか明示的に256bit幅使うように指定しないと、AVX使わないのでは
219:216
12/08/19 20:08:03.83
float4とfloat4を引数に使ったmadは多用してるよ。
220:デフォルトの名無しさん
12/08/20 13:15:46.35
いくらなんでも1400倍はおかしくないか?
OpenCLのことはよくしらないけど、理論性能から考えてもそんなに
差が出るはずが無いと思うが・・・
140倍なら、CPUのコードがクソならあり得るが
221:デフォルトの名無しさん
12/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 しているならばなおさら。
222:デフォルトの名無しさん
12/08/20 19:26:12.93
倍精度だとSSEもAVXも同じなの?
223:デフォルトの名無しさん
12/08/21 10:06:59.44
x87比
224:216
12/08/22 23:02:09.92
同じ問題を APU E1-1200 のミニノートPCで解いてみた。
OSは、 >216 と同じくUbuntu12.04 x64版。
メモリ資源の制限から、行列の個数は1/5の100万個。
Jacobi法のOpenCLモジュール内で10回同じ計算をループで
回しているので、1000万個の3x3行列の固有値を求めたことになる。
当然ながら、このミニノートPCでは O.C. はしていない。
<結果>
反復回数(6回)、計算精度共にデスクトップマシンで計算結果と同じ。
CPU 7.542 sec
GPU 0.4235 sec
OpenCLイベントタイマーで計測した正味の計算時間比較では
GPUが18倍速いと言う結果になった。
これぐらいの比率なら正常?
225:216
12/08/22 23:10:51.04
ついでに >216 で使用したデスクトップマシン環境で同じ100万個の計算(1/5の量)を
再計測してみると
<結果>
CPU 9.311 sec
GPU 0.006517 sec やはり、1429倍となる。
ただし、皆さんお気付きのように100万個同士の計算時間比較を行うと
<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 の65倍となり、CPU側の計算時間が?
どちらも、C++で作成したOpenCLの環境セットアップ部を含めたJacobi法の
OpenCLコードまで全て同じソースコードを用いているが
CPU側のOpenCLイベントタイマー計測がどう見ても変なことが分かった。
デスクトップ側、ミニノートの両ドライバーファイルは12.8 catalyst + AMDAPP V2.7
で同じ。 またUbuntu付属のg++4.6.3+NetBeans上で実行モジュールを作成している。
226:216
12/08/22 23:26:18.10
>225
> CPU側のOpenCLイベントタイマー計測がどう見ても変なことが分かった。
上記行を削除します(書き間違えました)。
実際に腕時計で測っても、同等な8秒程度の時間でした。
OpenCLイベントタイマー計測が変なのではなく、実際に i7-2600k の方が若干遅かったです。
227:デフォルトの名無しさん
12/08/23 01:59:01.61
ループでもっとぶんまわせば?
計測短すぎてクロック上がってないんだろ。
228:デフォルトの名無しさん
12/08/23 02:51:56.85
そもそも3x3行列ってのが小さすぎ
229:デフォルトの名無しさん
12/08/23 10:37:48.70
たしかに3x3は小さいが、全部の固有値を馬鹿並列で求められるなら、むしろGPU向きな気もする
230:デフォルトの名無しさん
12/08/23 10:39:27.25
「100万個の固有値計算を馬鹿並列にできるなら」という意味ね。1つの行列の相異なる固有値を並列に求めるという意味では無く。俺日本語不自由すぎワロタ
231:デフォルトの名無しさん
12/08/24 01:59:31.99
AMDのOpenCLランタイムは、CPUが
ちょっと異様なぐらい遅いとか聞いたことがある気がする
232:デフォルトの名無しさん
12/08/24 12:23:43.91
それってSDKが出たばっかの頃の話じゃね
たしかデバッグか何かに絡んだ挙動だったと思ったが
233:デフォルトの名無しさん
12/08/28 11:59:01.94
7970もECC化できればなぁ
234:デフォルトの名無しさん
12/08/29 21:33:10.08
レイトレースをCPU、GPU側で同様のロジック作って走らせても1000倍くらい違ったし
GPUは並列度の高いコードだと思った以上に早くなる
CPUはそれだけ理論値からの落ち込みが激しいってことなんだろうけど
235:デフォルトの名無しさん
12/08/30 05:20:58.76
CPUの方のコードが最適化されてないんじゃないの?
最新のGPUでも単精度ピークが数Tflopsだから、
1000倍差だとCPUは数Gflops以下しか出てないことになる
236:デフォルトの名無しさん
12/08/30 08:19:16.75
x87比か
237:デフォルトの名無しさん
12/08/31 16:20:09.66
>>235
実際、素人が行列積書くと1-2GFlops しかでないし、まぁそんなもんだろうね
238:デフォルトの名無しさん
12/08/31 18:47:52.23
Intelがメニーコア「MIC」とAtom SoCの「Medfield」を発表
URLリンク(pc.watch.impress.co.jp)
GPU勢もうかうかしていられない?
239:デフォルトの名無しさん
12/09/09 12:44:54.16
GPUコアは、これも PowerVR なのね。人気者すぎ。
240:デフォルトの名無しさん
12/09/09 23:36:12.36
アーキテクチャがx86だってのは個人的にどうでもいいんだが、
メモリ階層がどうなるのか、メモリバンド幅がどれくらいでるのか、
CUDA5のdynamic parallelismに対抗できる機能があるのか、くらいが
勝負の分かれ目だな。
241:デフォルトの名無しさん
12/10/07 21:53:45.44
>>240
dynamic parallelismなんてわざわざご大層な名前を付けなくても
x86CPUなのだから同等以上のことが出来るに決まっているだろ。
242:デフォルトの名無しさん
12/10/07 22:01:03.70
対抗してCUDAのバイナリ仕様公開とかないかな
243:デフォルトの名無しさん
12/10/07 23:20:36.04
ナイコナの汎用性は別に誰も否定してなくね
問題はその効率だよな。やべーさすがIntel様だ、となるのかやっぱ汎用コア並べたらそんなもんだよね、で終わるのか
うちも研究枠で調達予定なので普通にwktkしてます
GPU性能を維持したままにじり寄ってるからこそ、dynamic parallelismなるご大層な名前…というか変態的な局所性能の誇張に縋ってるんだろう
244:デフォルトの名無しさん
12/10/08 00:16:04.33
また大人のおもちゃが税金でたくさん作られるね
245:デフォルトの名無しさん
12/10/08 15:19:35.15
えー?
コア一つ当たりの性能ではGPGPUを圧倒するんじゃないの?
問題は、それだけの性能を発揮するために必要なコアサイズが大きすぎることであって。
一つのダイに集積可能なコアの数が、GPGPUのプロセッサより一桁以上も少なくなるので、
コア性能では圧倒していても、総合性能では大幅に負けるに違いないと予想する。
>やっぱ汎用コア並べたらそんなもんだよね、で終わる
としか考えられないよなあ...
246:,,・´∀`・,,)っ-○○○
12/10/12 14:31:36.83
Coreの定義がIntelとGPUメーカーとで違うから当たり前だろ
1コアあたりのベクトルユニットが16SP/8DP積和だから
NVIDIAのCUDA Coreに換算すると16コア相当で
つまり50coreのPhiは800 CUDA coreに相当するんだけど?
247:デフォルトの名無しさん
12/10/12 15:05:22.84
あー、そか。たしかにそうだよな。
コアサイズが二桁ぐらい違うとかじゃないと性能で見劣りすることはないかもしれんのか。
まあ、コアの世代を古い方にして回路規模を小さくしちゃったりするとスループットが落ちて不利だがな、
動作クロックがGPUより上回ってる分と併せて考えると接戦になるのかな?
248:デフォルトの名無しさん
12/10/12 15:10:50.51
排他処理のルーチンはボトルネックとしか考えられないよなあ...
249:デフォルトの名無しさん
12/10/12 16:11:57.09
GPGPUはピークはともかく実効値がな・・・
先日出たTSUBAME2.0の1mメッシュ気流解析で実効15%ぐらいだっけ?
複雑な分岐や並列度低いものが入るとしぬし
ほかのGPUアクセラレータ積んだクラスタもそうだけどLINPACKばっか速くても仕方ない
250:デフォルトの名無しさん
12/10/12 16:39:36.75
その残りの85%は何してるの?
251:デフォルトの名無しさん
12/10/12 16:56:08.79
リラックス
252:デフォルトの名無しさん
12/10/12 18:14:08.29
良いこと考えた。
働いていない、85%を装置から除去すれば、消費電力も下がるんじゃね?
253:デフォルトの名無しさん
12/10/12 21:11:04.99
それこそSMTできるようにして同時実行させろって話じゃね
254:デフォルトの名無しさん
12/10/12 23:36:36.25
そしてまた効率低くてそぎ落としか
設計レベルでループしてんじゃねえよw
255:デフォルトの名無しさん
12/10/13 00:37:04.86
人間社会もGPGPU社会も一緒なんだな
256:デフォルトの名無しさん
12/10/13 00:49:49.97
分岐というか、単純に帯域の問題じゃね。
257:デフォルトの名無しさん
12/10/13 02:04:35.67
あのプレゼン見たけど、プログラミングに自信のある計算機寄りの人が、
てきとーに題材みつけてきて、計算機科学の研究として発表してる感じだった。
その結果が15%
258:デフォルトの名無しさん
12/10/13 02:37:46.42
計算目的よりも、tubameを生かすための作業やね。まあ察しはつくがw
259:デフォルトの名無しさん
12/10/13 03:38:41.47
>>252
蟻の巣からよく働く2割の蟻以外を取り除いても、
残った蟻の8割がなぜか怠け始めるんだそうな
260:,,・´∀`・,,)っ-○○○
12/10/13 04:49:03.94
>>250
働いてないんじゃなくて何かしらがボトルネックになってFMACの稼働率が抑えられてるのでは?
メモリ帯域が足りてる場合でも同時命令発行数の制約でload/store命令と積和命令を
同時に実行できない、とか
(これは京のVenusアーキテクチャにもある制約)
261:デフォルトの名無しさん
12/10/13 16:36:04.21
唐突ですまんけど、AMDはファイル名通りのこういう資料があるけど
URLリンク(developer.amd.com)
nVIDIAはこういうの無いの?
262:デフォルトの名無しさん
12/10/13 18:21:13.66
>>249
効率はFLOPSに対する効率であって、気象系のアプリは
バンド幅律速だからLINPACKと比べること自体ナンセンス。
ちなみに15%は気象系のステンシルアプリとしては高い方。
263:デフォルトの名無しさん
12/10/13 18:24:01.07
>>257
そらまぁ、青木研究室の存在意義が「GPGPUでできることを広げる」
だからね。仕方ない。
NECのSXが沈没の今、京かGPGPUしか無いわけで・・・
それともBlueGeneを買うか?
264:デフォルトの名無しさん
12/10/13 18:27:15.78
FLOPSに対する効率だから、アルゴリズム自体の並列化効率以上にはなりえないわけか。
気象系のアプリというものの並列化効率がどんくらいなのか知らないけど(なんとなく高そうではあるが)。
265:デフォルトの名無しさん
12/10/13 18:51:02.05
>>264
うーん、微妙に違うな。
今時のアーキテクチャだとメモリがすごく遅いってのは知ってるだろ?
ちなみにメモリの速度はバンド幅って言う。
それに対して演算速度はムーアの法則でどんどん上がってる(上がってた)
から、バランスが全然とれてないんだ。だから、普通のCPUはキャッシュ階層を
深くして(L1、L2、L3とか)なるべくデータの再利用をしてる。
つまり、データをメモリから持ってきたら、なるべくそれを
使い回して計算したいわけ。
で、ここで問題になるのが、浮動小数点演算(FLOP)と、データ転送量の比。
業界だと byte/flops とか言われる。言い方を変えると、どれくらいデータの
使い回しが効くかってこと。
続く
266:デフォルトの名無しさん
12/10/13 18:52:43.77
承前
アプリの要求byte/flops高いと、せかっくデータを持ってきても
たいした計算をせずに、すぐに次のデータが必要になっちゃう。
データの使い回しが効かないんだ。これはイマドキの計算機にとってはキツい。
で、気象とか流体とかはそういうアプリなわけ。
だから、基本的に流体とかの計算は演算は余ってて、バンド幅がボトルネックになる。
これはアプリの特性だから仕方ない。
気象・流体屋さんは、未だに地球シミュレーター信者みたいな人が多い。
ちなみに、気象・流体系アプリは並列度だけで言ったらはものすごくあるよ。
ちなみにスレ的に言うと、GPUはGDDR5っていう容量が少ない代わりに
バンド幅が出る贅沢なメモリを搭載して、かつ大量のスレッドを使うことで
レイテンシを隠蔽してる。これがGPGPUが気象・流体計算に使える理由。
267:264
12/10/13 19:29:29.68
なるほど、つまり気象系のアプリはコンピュートバウンドではなくメモリバウンドのものが多い、と。
AMDのAPUの使い道の話になったときに、DDR3という制約に悲観的になってる人が少なそうだったから
世の中にはメモリバウンドのものは少数派なのかなあとも思ってたけど普通にあるんですね。
268:デフォルトの名無しさん
12/10/13 19:43:21.60
バンド幅と言うか、物理原理的にDRAMのセルは応答が現代のCPUの演算速度に対してずっと遅い。
だから、メモリセルを並列にして同時に読み出し・書き込みするしかないのが根底にある。それをやれDDRxだの、
何ビット同時にアクセス=複数セルに同時アクセスでなんとかしてるね。
気象系や気体がどうのって言うよりも、規則正しく計算点を並べて偏微分方程式の数値近似式を一斉に走らせる
=SIMDだのシェーダだの言ってる演算回路がみんな一斉に同じ計算式をやる場合が最大性能が出る。
三次元の方がそりゃ、並列度は高いだろうね。境界条件とか面倒くさそうだが。
こういう根本的なところはもう全然進歩しないなw
269:,,・´∀`・,,)っ-○○○
12/10/13 20:27:18.96
bytes/flopsに関して言えば現状GPUはCPUより低いです。
ただbyte/sが数倍高くて、flopsが更に数倍高いというだけの話で。
Haswellあたりで、電力当たりFLOPS数はGPUとCPUは互角くらいになるのではという話も出ていますな。
270:デフォルトの名無しさん
12/10/13 20:32:24.78
TSVって、GPUみたいな爆熱チップでもできるの?
271:デフォルトの名無しさん
12/10/14 11:35:52.75
>>241
亀レスだけど、何のペナルティも無く普通のx86コアを50個
並べられるなら誰も苦労してないわけで・・・
272:デフォルトの名無しさん
12/10/14 12:44:22.63
DRAM内部のクロック
PC133 133MHz
DDR3-1600 200MHz
この間十数年・・・MRAMならもっと上がるかな?
273:デフォルトの名無しさん
12/10/20 09:38:39.11
GNU MPのGPGPU対応って、あんまりやってる人が居ないな。
274:,,・´∀`・,,)っ-○○○
12/10/20 09:42:31.34
依存関係があってなかなか並列化しにくいからね。
もちろん並列度の高いコンピュータに適した数式に置き換えればいいんだが
CPU側もAVX, FMAなどでGPGPUの電力効率の優位も揺らぎつつあるからね
275:デフォルトの名無しさん
12/11/17 23:26:42.82
C++AMP調べてるとこなんだけど
textureでピクセル間を補間する方法がわからない。
ひょっとして自前でやらないといけないの?
「C++ AMP API は、サンプリング テクスチャのフィルター処理機能を提供されません。」
とはこのことなのかな。
276:デフォルトの名無しさん
12/11/18 21:58:01.39
c++ampにfloat_3のmadやfmaが無いね。 short vector typeのmadやsin, cos などの超越関数も用意してもらいたい。 swizzlingはできるみたい。 OpenCLの方がコードをコンパクトに書けるかな。次期版期待してます。
277:デフォルトの名無しさん
13/02/17 08:49:52.18
何故かsea islandのISAが公開
URLリンク(developer.amd.com)
278:デフォルトの名無しさん
13/02/17 13:26:30.82
いつもに比べると異常に早いがやっぱsouthernislandのリネームなのかね
279:デフォルトの名無しさん
13/02/18 13:29:16.68
ディスクリートとしては出なくても
次世代APUのGCNはsea islandなのかね。
ISA見る限り、SIから確かに機能アップされている。
一番大きいのはFLAT メモリアクセス命令の追加か。
cudaのUVAみたいなもの
280:デフォルトの名無しさん
13/03/27 13:39:55.44
SL#(えすえるしゃーぷ)とは、GPUで実行されるプログラマブルシェーダーを、超高級言語である
C#で書けてしまうという夢のようなオープンソースのフレームワークである。
URLリンク(monobook.org)
281:デフォルトの名無しさん
13/03/30 11:07:23.56
>>280
まじかよ、ニンジャコンパイラが全部解決してくれるのか?
282:デフォルトの名無しさん
13/04/12 18:59:55.71
超高級言語w
283:デフォルトの名無しさん
13/04/14 14:30:38.59
抽象化の度合いが高いぐらいの意味合いでいいんじゃないの
実際問題プログラマが意識してGPUに仕事をさせるコードを書かずに勝手に解決してくれちゃう未来像は来るの?
284:デフォルトの名無しさん
13/04/14 18:29:31.68
たぶん、デザインパターンと同じくらいの敷居にはなるんじゃねーの?と言うか、デザインパターンとして確立されるんじゃね。
判らない人には判らないけど、自分の中に消化され始めると自然とそう言う書き方になるみたいな。
あとは、フレームワークでそう言う書き方を強制されて、なにも考えずにそう言うもんだと思うパターン。
285:デフォルトの名無しさん
13/04/15 02:40:47.34
コア数が100くらいのGPUがあるとして
整数大規模行列の
全要素の和を求めるのってGPGPUで100倍程度の高速化は可能ですよね?
286:デフォルトの名無しさん
13/04/15 07:36:47.18
1コアあたりのシングルスレッド能力がCPUと同程度あって、メモリ帯域幅がそれに見合うほど広ければな
287:デフォルトの名無しさん
13/04/22 21:00:09.19
>>283
実際、解きたい問題依存だけど、Hadoopのように、プログラミングモデルを固定して制限を付けて縛ることによって簡単な記述で自動的に高速実行ってのは比較的現実的な未来だと思うよ。
Hadoopってのも、書きたい処理をHadoopの"型"に押し込んで書ければ、残りの面倒な問題は全部Hadoopが面倒見てくれますって事だし。
288:デフォルトの名無しさん
13/09/21 12:06:20.79
このスレ立った頃のGPUってもうゴミだよな?
289:デフォルトの名無しさん
13/09/21 13:04:48.76
OpenCLで書いたプログラムを9600GTとQ9650で動かしたら同じくらいの速度だった。
今のCPUには絶対負ける。
290:デフォルトの名無しさん
13/10/02 19:53:46.66
なんか来てた。
URLリンク(www.phoronix.com)
291:デフォルトの名無しさん
13/10/20 14:42:26.77
GPGPUとMantleの関係を語ってくれよ
mantleってなんなんだ?
292:デフォルトの名無しさん
13/11/06 09:10:41.13
>>291
DirectXやOpenGLは長い歴史があって
色々なハードで動かす必要があるので
互換性維持のために滅茶苦茶厚いAPIコールを通さないといけないが
MantleはGCNに絞った最適化さえすればいいので
そういうものがないっつーこと
据え置きはAMD一色だから移植も相互に楽になるし、面倒な最適化も最小限で済む
というのが言い分
293:デフォルトの名無しさん
13/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をやってた層には移行先として十分かなと言った感じ
294:デフォルトの名無しさん
13/12/30 03:37:10.21
Mantleって互換性維持のため分厚くなり過ぎたDirectXに対して
低レベルなところを柔軟に叩けるようにするとかそういうものじゃなかったっけ?
Project Sumatraが楽しみだな
Java9ぐらいでparallelStreamにぺぺっと渡すと
OpenCLやってくれちゃうとか最高じゃん
逆に言うとそれぐらい猿でも簡単に扱えないとなっていうか
295:デフォルトの名無しさん
14/01/16 15:13:31.68
Mantle.Netはよ
296:デフォルトの名無しさん
14/01/16 19:52:34.70
kaveriのFLOPSのDP:SPが気になるんだが、誰か情報持ってないか?
297:デフォルトの名無しさん
14/01/16 20:00:22.38
OpenACCってどうよ
298:デフォルトの名無しさん
14/01/26 21:12:44.48
最近CUDAを始めたのですが、簡単な計算(大きな配列のベクタ足し算)を
CPUとGPUにやらせると明らかにGPUが遅いです。
具体的には、VC2012+CUDA5.5でコード(URLリンク(www1.axfc.net))を走らせると、
計算時間が次のようになりました(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:デフォルトの名無しさん
14/01/26 21:24:38.81
メモリ転送のオーバーヘッドがあるから、もっと大きな問題じゃないと効果は出ないよ。
300:デフォルトの名無しさん
14/01/26 21:54:26.48
>>299
それは知っているのですが、数値を大きくするとすぐにc00000fdでクラッシュするんですよ……。
<<<grid, block>>>もいろいろ弄っているのですが、どうにも効果が得られません。
今試してみたら、arraySizeの値で実行できるのは25600が最大みたいです。
301:デフォルトの名無しさん
14/01/26 22:08:07.85
GLSLからOpenCLへの移行を昨日から始めたけど
GLSLより書きやすいのはいいけど最適化を追い込まないととんでもなく遅くなるんだな
GLSLで複雑な汎用計算やらせるのは難解なパズルゲームみたいで嫌になってたけど
結局最適化の手間を考えたらどっちが楽ということはないんだね・・・
>>298みたいな単純な計算ならGLSLだとバグったような速度が簡単に出るから別世界感が凄い
302:298
14/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までは大丈夫ということなのでしょうか?
303:デフォルトの名無しさん
14/01/26 23:22:14.60
>>299
GPGPUはメモリ転送のオーバーヘッドがないHSA(Huma)だよな
PCではど重い処理でない限りAMDのHSAがGPGPU処理の主流になるだろうな
304:デフォルトの名無しさん
14/01/27 00:23:03.99
>>298
残念なお知らせ。
そのソースコードでは、GPUの演算時間ではなくGPUの呼び出し時間しか計測してないね。
「実際の演算時間」=「内部ブロック数」*(「内部ブロックの呼び出し時間」+「内部ブロックの演算時間」)だとすると、
「実際の演算時間」-「内部ブロックの演算時間」になっているはず。
ブロック数が充分大きければ誤差だけど、内部ブロック数が1のときは激速になってしまう。
まぁ、実際の運用ではCPUとGPUが並列に動作することを期待するからそれでもいいんだけどね。
いずれにしても、CPUぶん回すよりも手っ取り早いと思っていたら大間違いだよ。
それと、CUDAスレも宜しく。
305:298
14/01/27 00:47:11.47
>>304
>そのソースコードでは
え!? ……つまり、
普通にtimeGetTimeかQueryPerformanceCounterとかを使えってことなんですか?
それとも、測定する位置が間違っているということなんですか?
>CUDAスレも宜しく
分かりました。次回以降はそちらにレスすることにします。
306:デフォルトの名無しさん
14/01/27 08:23:49.16
>>304
何言ってんだ、こいつ?
307:デフォルトの名無しさん
14/01/27 21:34:08.90
>>298 arraySizeが大きいと、CPU版すらStackOverflowになるよ。
URLリンク(pastebin.com)
308:307
14/01/27 21:39:43.68
うっかり、166行目を「cudaStatus = cudaSetDevice(1);」にしちゃったので、適当に直しておいて。
309:デフォルトの名無しさん
14/01/27 23:30:12.43
ローカルメモリを使う場合って確保しようとした容量が大き過ぎると
グローバルのほうへ確保されてしまうんだよね?
AMDのGCNはどれくらいまでローカルメモリがあるのか分からないんだけど
試行錯誤して調べるしかないのか
310:298
14/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]らしいです。
311:デフォルトの名無しさん
14/01/28 01:09:12.72
>>307
冗長なOpenCLに比べてやっぱりCUDAはスマートでいいな
312:デフォルトの名無しさん
14/01/29 01:59:06.39
OpenCLのclEnqueueNDRangeKernelでカーネルを実行するときに
global_work_sizeとlocal_work_sizeに同じ値(256,256など)を入力すると
何もエラーは返されずにメモリの参照が壊れて?しまいclEnqueueReadBufferで
CPU側で読み取った値が全て0になってしまいます。
これは仕様なのでしょうか?
313:デフォルトの名無しさん
14/02/25 21:16:18.98
visual studio 2013でCUDAが使えないからC++AMPでやるお!
314:デフォルトの名無しさん
14/02/25 21:43:30.35
>>313
そのためだけにVS2012と2013使い分けてる俺……
315:デフォルトの名無しさん
14/04/04 10:44:13.17 YtPgho8U
openCL始めたお(・∀・)ノ
316:デフォルトの名無しさん
14/04/15 02:32:13.65 vGWbAtXL
(・∀・)ノ CPUの300倍くらいの性能が出たお!
比較したCPUはE2-2000っていうCPU+GPU=APUだけど全くGPUとしての機能をもってないのでガッカリしたお。
317:デフォルトの名無しさん
14/04/19 12:16:56.16 Firi/9oq
(・∀・)ノ ALU(IGP)のE2-2000はHD7770の1/50のパワーしかないが並列性はあるようだ。
318:デフォルトの名無しさん
14/04/22 04:44:14.02 aREYskwN
AIDA64に測定メニューあるよな