10/04/18 19:48:30
このスレッドは、他のスレッドでは書き込めない超低レベル、
もしくは質問者自身何が何だが分からない質問を勇気を持って書き込むスレッドです。
CUDA使いが優しくコメントを返しますが、
お礼はCUDAの布教と初心者の救済をお願いします。
CUDA・HomePage
URLﾘﾝｸ(www.nvidia.com)
関連スレ
GPUで汎用コンピューティングを行うスレ
ｽﾚﾘﾝｸ(tech板)
GPGPU#3
ｽﾚﾘﾝｸ(tech板)
前スレ
【GPGPU】くだすれCUDAスレ【NVIDIA】
ｽﾚﾘﾝｸ(tech板)
【GPGPU】くだすれCUDAスレ pert2【NVIDIA】
ｽﾚﾘﾝｸ(tech板)
2:デフォルトの名無しさん
10/04/18 19:49:11
関連サイト
CUDA Zone
URLﾘﾝｸ(www.nvidia.co.jp)
フォーラム
URLﾘﾝｸ(forum.nvidia.co.jp)
CUDAに触れてみる
URLﾘﾝｸ(chihara.naist.jp)
cudaさわってみた
URLﾘﾝｸ(gpgpu.jp)
CUDA のインストール
URLﾘﾝｸ(blog.goo.ne.jp)
NVIDIAの「GeForce 8800 GT(G92)」と次に控える64-bit GPUアーキテクチャ
URLﾘﾝｸ(pc.watch.impress.co.jp)
CUDAを使う
URLﾘﾝｸ(tech.ckme.co.jp)
NVIDIA CUDAを弄ってみた その2
URLﾘﾝｸ(dvd-r.sblo.jp)
CUDAベンチ
URLﾘﾝｸ(wataco.air-nifty.com)
KNOPPIX for CUDA
URLﾘﾝｸ(www.yasuoka.mech.keio.ac.jp)
3:デフォルトの名無しさん
10/04/18 21:42:11
いちおつ
4:デフォルトの名無しさん
10/04/19 00:54:58
一乙、前スレの、魔二来vsTeslaのコストパフォーマンス比較面白いな
もっとやってくれ
5:デフォルトの名無しさん
10/04/19 05:45:00
>>4
もう結論でたよ
理論ピーク性能はTeslaの勝ち
大部分のプログラムでの実行性能はマニクールの勝ち
プログラミングのしやすさではマニクールが圧勝
6:デフォルトの名無しさん
10/04/19 07:05:25
すげーな
48コア使いこなせるとかｗ
7:デフォルトの名無しさん
10/04/19 07:09:48
で、比較したんだろ988さんよ
988 ：デフォルトの名無しさん [] ：2010/04/18(日) 17:40:17
teslaよりマニクールのほうが倍精度で性能を比較すれば
ずっとリーズナブルだっていう事実
995 ：デフォルトの名無しさん [↓] ：2010/04/18(日) 18:06:57
6168
744*4=2976
1.9*12*4*2*2=364.8
6176se
1386*4=5544
2.3*12*4*2*2=441.6
c2050
2499
1.215*448=544.3
8:デフォルトの名無しさん
10/04/19 07:38:14
>>7
これって、いまいち理解出来ないのだが、誰か説明して。
9:デフォルトの名無しさん
10/04/19 09:51:51
>>7
将来が危ぶまれるNVIDIA専用のCUDAと心中するつもりなら
Teslaでもいいかもね
10:デフォルトの名無しさん
10/04/19 12:33:45
>>8
たぶん
6168と6176seはマニクールの製品名でc2050はtesla
その下の２つは値段（ドル）と性能（GFLOPS)
値段はメーカーから直接購入した時の価格
CPUは
（総費用）＝（価格）X（コア数）
（GFLOPS）＝（クロック数）X（コア数）X（CPU数）X（SIMD演算器数）X（１SIMD演算器当たりの倍精度演算数）
GPUは
（総費用）＝（価格）
（GFLOPS) ＝（クロック数）X（コア数）
11:デフォルトの名無しさん
10/04/19 12:35:22
>>10
間違えた
×（総費用）＝（価格）X（コア数）
○（総費用）＝（価格）X（CPU数）
12:デフォルトの名無しさん
10/04/19 14:12:24
OpenMPとかえらい簡単に並列化できるしな。
あと何覚えればいいですか？
13:デフォルトの名無しさん
10/04/19 15:20:32
普通のスレッド
MPI
14:デフォルトの名無しさん
10/04/19 16:22:11
>>10
OpteronってSSEとかが1コアあたり2個ある訳？
15:デフォルトの名無しさん
10/04/19 16:24:58
OpenMPもCUDA並に最適化が面倒。
その代わりMPIの知識があれば以外と簡単。
16:デフォルトの名無しさん
10/04/19 18:03:05
で、比較したのまだぁ？
17:デフォルトの名無しさん
10/04/19 18:05:27
そもそもC2050が出てないのに
比較なんてできるかよ
18:デフォルトの名無しさん
10/04/19 18:22:54
そもそもスレチ
19:デフォルトの名無しさん
10/04/19 21:40:15
秋葉にオープンしたfaithにCUDAのブースができたね
実体はunitcomの宣伝かな
20:デフォルトの名無しさん
10/04/19 22:03:30
GTX480でデモしてくれるならすぐにでも行くが
21:デフォルトの名無しさん
10/04/19 22:04:09
長崎大の部品を入れたとか書いてたね。CUDA話のできる店員さんいるのかしら。
22:デフォルトの名無しさん
10/04/20 07:05:14
チップセット・・・完全撤退
GPU・・・AMDの大攻勢の前に大苦戦
Tesla・・・CPUの高性能化・マルチコア化の前で微妙に
NVIDIAさんこれからどうするつもりなんでしょうね
23:デフォルトの名無しさん
10/04/20 07:18:02
自作板に帰れ
24:デフォルトの名無しさん
10/04/20 07:29:44
SSEとかAVXとか覚えるよりも、
CUDAのお作法覚えるほうが簡単な気がしない？おれだけ？
25:デフォルトの名無しさん
10/04/20 08:09:46
>>24
でも将来性が最悪だね
遅かれ早かれOenpCLに移行するだろうし
26:デフォルトの名無しさん
10/04/20 08:53:34
CUDAとSIMDでは考え方が違うような気が。
個人的にはCUDAは割と大ざっぱに並べる。
SSEとかは細かなパズル。
自分でも何を言っているかわからんが、そういう思いでいる。
27:デフォルトの名無しさん
10/04/20 08:57:23
前スレで480特攻すると書いた者だが
倍精度1/4と聞いてキャンセルしてしまった…
ゲーマーだったら失敗覚悟で注文したんだけど
28:デフォルトの名無しさん
10/04/20 09:09:12
ZOTACまだリリース前になってますね。キャンセルできてよかったですね。
ケース待ちだった方は購入できたのでしょうか。
29:デフォルトの名無しさん
10/04/20 09:20:53
マザーボード破損の書き込みもあるし手を出さなくて正解だろうね
30:デフォルトの名無しさん
10/04/21 18:23:51
カーネル中でcoutとかは使えないの？
どこが悪いのか調べるときはたいていcoutで変数を書き出していたのだけど、
CUDAの場合はどうすればいいの？
31:デフォルトの名無しさん
10/04/22 02:17:10
cudaMemcpy()でCPUに送る。
32:デフォルトの名無しさん
10/04/22 05:13:57
カーネルの途中で表示したいんだけど
33:デフォルトの名無しさん
10/04/22 08:07:33
どんな手法にしろCPU介さなかったら表示はできんよ
34:デフォルトの名無しさん
10/04/22 08:32:27
URLﾘﾝｸ(gpu.fixstars.com)を使う
自分はコーディングレビューしかしないけど
35:デフォルトの名無しさん
10/04/22 09:45:25
結果とは別にデバイス側に確保しておいて、
あとで見るのはだめなの？
36:デフォルトの名無しさん
10/04/22 10:12:38
>>32
・deviceemuでなんとか頑張る
・可能な限り詳細な出力をメモリに吐き出しておいて、ホストに戻ってから(転送して)解析する
それ以前に、そのレベルだとcudaの細かいノウハウを掴んでないだろ。
然るべき業者に委託したほうがよくないか?
37:デフォルトの名無しさん
10/04/22 20:00:42
なぜ仕事前提？
38:デフォルトの名無しさん
10/04/22 22:26:01
>>30
とゆうか、そうゆう発想する事態C++もCUDAも全然理解できてないよねｗ
39:デフォルトの名無しさん
10/04/23 01:36:42
CPUで考えると、「レジスタの計算の途中をCoutで見ることは出来ないか？」
と聞いているようなものだと思うのだが。
それを自前でやるにはレジスタの内容を一度メモリにコピーしてcoutの標準出力で表示させることになる。
それと一緒だと思うよ。
40:デフォルトの名無しさん
10/04/23 02:53:45
device emuで解決ジャン
41:デフォルトの名無しさん
10/04/23 03:31:06
待てればね。
42:デフォルトの名無しさん
10/04/23 03:47:59
エミュレーションモードってなくなるんだろ
43:デフォルトの名無しさん
10/04/23 05:53:13
すでに実機があるからなあ。
44:デフォルトの名無しさん
10/04/23 06:50:27
ダンプ用の領域を用意しておいて実行中に書き込んで後で見る
しかないんじゃね？
45:デフォルトの名無しさん
10/04/23 07:05:41
ubuntu 10.04でcudaが使えるようになるのはいつごろですか
46:デフォルトの名無しさん
10/04/23 08:28:37
GPUを搭載していないノートPCでエミュレーションモードを使ってたのに
なくなるってマジ？
47:デフォルトの名無しさん
10/04/23 11:36:52
マジ
48:デフォルトの名無しさん
10/04/23 11:42:20
リモートデスクトップでCUDA使えればいいのになー。
49:デフォルトの名無しさん
10/04/25 16:17:32
リナカフェいったひといないすか
50:デフォルトの名無しさん
10/04/25 16:22:48
サトリナのカフェなんてあるの？
51:デフォルトの名無しさん
10/04/26 10:35:43
このスレもしぼんだなあ
GTX480が倍精度外れた（らしい）のが効いているんだろうな
52:デフォルトの名無しさん
10/04/26 11:17:15
倍精度演算ユニットは少なくなっているだけで無くなったわけではないのだが、
そもそも480が入手できないことの方が原因ではないかと。
53:デフォルトの名無しさん
10/04/26 15:43:31
結局口だけの貧乏人しか居なかったって事だなｗ
54:デフォルトの名無しさん
10/04/26 17:33:22
おまえもな
55:デフォルトの名無しさん
10/04/26 19:14:13
URLﾘﾝｸ(www.amazon.co.jp)
これIPAの岡ちゃんが書いたの？
いつの間にかネットマイスターってところに転勤したのかな
56:デフォルトの名無しさん
10/04/26 19:24:07
もう入門書もでてるから、新刊はいらないな。
需要はあるのだろうか･･･。
57:デフォルトの名無しさん
10/04/26 19:28:27
つーか、ネットマイスター代表だってよ。
同名をtwitterで検索したらURLﾘﾝｸ(twitter.com) が引っ掛かったけど……
58:デフォルトの名無しさん
10/04/27 01:07:50
流石にロリコンばらされたら職場には居にくいか……
アプリも割り放題だったしな。
たくましく生きているようで何よりww
59:デフォルトの名無しさん
10/04/27 01:20:01
この本買ったけど、ちょっと入門過ぎる気がするなあ。
「高速」という言葉にてっきり「最適化」ということを連想してしまった。
環境設定とかで半分以上割いているよ。
まあ、CPUコードと両方書いてあって、Webからダウンロード出来るから、
最初にでた本よりかはいいかもしれない。
60:デフォルトの名無しさん
10/04/27 08:02:14
自分で買うほどじゃないかな
図書館に購入希望だしとこう
61:デフォルトの名無しさん
10/04/27 21:19:54
本にするほどの蓄積と技術の咀嚼があるとは思えないけどどうなんでしょ。
あるとすればまだ論文レベルなんじゃないか？
逆に、まだ広まってない今が新しいことするチャンスかっ！
62:デフォルトの名無しさん
10/04/27 21:27:03
プログラムの入門書なんてこんなもんだろ
63:デフォルトの名無しさん
10/04/28 04:21:24
GF100が失敗、GF104はキャッシュ削るみたいでGPGPUは絶望的の現在、今後広まる可能性が残っているかどうか…
64:デフォルトの名無しさん
10/04/28 05:19:00
は？
65:デフォルトの名無しさん
10/04/28 05:45:52
GPGPUに残されているのは個人では動画エンコードくらいか。
それなりの規模になると、マルチプロセッサのシステムを組んで、
CPUのコードを走らせた方が開発工数から考えると速いからなあ。
66:デフォルトの名無しさん
10/04/28 06:14:57
で、おまえは何しにここに来てる訳？
67:デフォルトの名無しさん
10/04/28 06:17:08
ばかだな
ただのAMD信者様だろｗ
68:デフォルトの名無しさん
10/04/28 09:19:52
GPGPUに将来性はないだろうね。
牧野もそういっているし。
69:デフォルトの名無しさん
10/04/28 09:29:19
コスト面がちょっと微妙かなあ
C10シリーズのようにアカデミックで10万！とかなら個人で使ってみようという流れになるだろうに
GTX480は倍精度1/4にしました！C20シリーズを売るためです！と公言されちゃったし
70:デフォルトの名無しさん
10/04/28 09:52:12
>>69
コスト面つか買う気失せることすんなよなー て感じっすよ
71:デフォルトの名無しさん
10/04/28 10:00:12
､--‐冖'⌒￣￣`ｰ-､
／⌒` 三ミヽｰ-ﾍ,_
__,{ ;;,, ミミ ｉ ´Z,
ゝ ''〃//,,, ,,..｀ミﾐ､_ノリ}j; f彡
_) 〃///, ,;彡'rffｯ､ｨ彡'ﾉ从iﾉ彡
>';;,, ノ丿川j !川｜; :.`7ﾗ公 '>了
_く彡川fﾞノ'ノﾉ ノ_ノノﾉイシﾉ｜ }.: '〈八ﾐ､､;.)
ヽ.:.:.:.:.:.;=､彡／‐-ニ''_ｰ＜、{_,ﾉ -一ヾ`~;.;.;）
く .:.:.:.:.:!ﾊ.Yイ ぇ'无ﾃ,`ヽ｝｝}ィt于 ｀|ィ"~
):.:.:.:.:|.Y }: :! ｀二´/' ; |丶ﾆ ﾉﾉ
） :.: ﾄ､リ: :!ヾ:、 丶 ; | ﾞ ｲ:} 逆に考えるんだ
｛ .:.: ｌ ｛: : } ｀ ,.__(__,} /ﾉ
ヽ ! `'ﾞ! ,.,,.`三'ﾞ､,_ /´ 「単精度なら４倍」と
,/´{ ﾐ l /ﾞ,:-…-～､ ） |
,r{ ＼ ﾐ ＼ `' '≡≡' " ﾉ 考えるんだ
__ノ ヽ ＼ ヽ＼ 彡 ,ｲ_
＼ ＼ ヽ 丶. ノ!|ヽ｀ヽ、
＼ ＼ヽ ｀¨¨¨¨´/ |ｌ ト､ ｀'ｰ-､__
＼ ｀'ｰ-､ /／ /:.:.} ｀'ｰ､_
｀､＼ ／⌒ヽ /!:.:.|
｀､ ＼ /ヽLｆ___ハ/ {
′ / ! ヽ
72:デフォルトの名無しさん
10/04/28 10:01:53
長崎大のセンセが8800GTで激安スパコン組んだときも、
そのうち２、３割がGPGPU用としてはぶっ壊れていたじゃん。
今度のは、そのままだと半分以上壊れたものを出荷してるから、まずいってんで、
止めてるんでないの。
73:デフォルトの名無しさん
10/04/28 12:24:26
牧野もそういっているし。
牧野もそういっているし。
牧野もそういっているし。
牧野もそういっているし。
牧野もそういっているし。
牧野乙
74:デフォルトの名無しさん
10/04/28 12:29:43
そもそも、+α要素であるアクセラレータが
終わるとか終わらないとか
おばかさん
75:デフォルトの名無しさん
10/04/29 01:52:54
金にならんかったら終わるだろ。
76:デフォルトの名無しさん
10/04/29 09:55:55
GPGPUはマルチスレッドな浮動小数点演算に特化したものだからな、使えるシーンはかなり限られてくる
77:デフォルトの名無しさん
10/04/29 10:26:40
forループの一部だけを実行する方法を教えてください
78:デフォルトの名無しさん
10/04/29 12:24:55
>>55
やっぱりこういう問い合わせが多いんだろうな
著者紹介にIPAに勤務したことはないって書いてあったよ
79:デフォルトの名無しさん
10/04/29 12:29:11
>>59
俺も買った。
全く同じ感想。
青木先生の本が理解できる人は全く買う必要なし。
>>55
本の経歴の部分に、「IPAに勤務したことはない。」と不自然に書いてあったよw
80:デフォルトの名無しさん
10/04/29 13:44:48
なんだ同名で迷惑している人なのか
81:デフォルトの名無しさん
10/04/29 15:21:40
んなわけない。
IPAを除名されたか、退職するときにそれに準じる誓約書を書かされたのだろう。
82:デフォルトの名無しさん
10/04/30 02:18:54
証拠は？
83:デフォルトの名無しさん
10/04/30 05:35:43
>>79
今それを見てワロタ。
「IPAに勤務したことはない。」
で株式会社ネットマイスターをググってみたけど、
すぐには見つからなかった。
でもこういった本を書く人が流出なんてするのだろうか？
84:デフォルトの名無しさん
10/04/30 08:01:21
GF100の板に出ていてみたいだけど、
GTX480でCUDA動作に不具合があるとか？
ドライバ改善待ち、みたい。
GTX480でのベンチマーク結果（単精度、倍精度問わず）が出ない理由はそれだったりするのかも？
また、GTX460が6月に出るとか。
お値段は480/470よりも安いんだろうけど、倍精度しょぼくなった480よりもさらにしょぼいんだろうなあ
いつかは出るであろう、Quadroにかすかに期待するしかないのか…
85:デフォルトの名無しさん
10/04/30 09:46:41
>>84
＞いつかは出るであろう、Quadroにかすかに期待するしかないのか…
お値段は同程度のGFの5～10倍でも?w
5月中にはGTX480が手に入る筈なんで、そしたら各種ベンチを取る予定なんだけどなんだけど。
処で、暫く更新していない間にnvccが新しくなっていてptx出力の段階では64bit実数処理のコードが出力されてて笑った。
-ptxをつけないと64bitコードは処理できないみたいなメッセージが出るけれど、ptxを眺めて最適化しているから
今後は実数定数の扱いにも気を遣わないとならないようだ。
それで思ったんだけど、ptx出力からメモリアクセスの回数や割り算の回数を数えるツールというか、
ptx出力を解析するツールを作ったら需要あるかな。私自身は自分で眺めたりgrepで数えたりしているんだけど。
# 要は、iccのレポート出力みたいな機能がnvccにあればいいんだけどね。
86:デフォルトの名無しさん
10/05/01 15:36:56
>>85
もちろん重要ありますよ。
87:デフォルトの名無しさん
10/05/01 16:09:33
【Fermi】GeForceGTX4xx総合Part31【GF100】
ｽﾚﾘﾝｸ(jisaku板)
の548以降でGTX480を含むGPUのベンチマークが出ている
LU分解のようです
倍精度に制限かかっているのはどうやら本当みたいだ
88:デフォルトの名無しさん
10/05/01 18:13:38
具体的にどう制限かかっているんだろ？
それによってCUDAのコードの書き方が変わってくると思う
89:デフォルトの名無しさん
10/05/01 22:13:39
>87
いや倍精度どころか単精度でも全然性能でてないわけだが
90:デフォルトの名無しさん
10/05/01 23:17:23
なんかそのベンチ遅くね
285の単精度で16Gflopsとか書いてあるけど
91:デフォルトの名無しさん
10/05/01 23:25:48
GTX285と比べる限り、GTX480の倍精度に期待できそうな気配はないなあ
1/4は本当くさいな
92:デフォルトの名無しさん
10/05/02 16:32:28
cublasのgemmを使ったベンチをだれかが提供すべきだな
もちろん、俺はやる気はないが
93:デフォルトの名無しさん
10/05/03 10:55:57
gemmもやって欲しいけど、gemmはピーク性能しか出てこないから俺はLU分解やFFTの方が多少は信用している。
94:デフォルトの名無しさん
10/05/03 12:02:20
というか一番多用しそうなFFTでどれだけ実効性能出るのかが気になる
95:デフォルトの名無しさん
10/05/03 16:43:35
>>94
例えばどんなサイズが？
96:デフォルトの名無しさん
10/05/04 18:08:02
URLﾘﾝｸ(journal.mycom.co.jp)
ここのコードについて質問なんだけど
これはマルチスレッドで実行する関数の内部だよね？
ということはAsやBsのデバイスメモリからシェアードメモリへの転送や、
最後のホストの変数Cへのデータ転送は無駄に全スレッドがやるの？
それとも同じデータの転送については１スレッドがやってる？
97:デフォルトの名無しさん
10/05/04 19:47:43
> これはマルチスレッドで実行する関数の内部だよね？
そのとおり。
> AsやBsのデバイスメモリからシェアードメモリへの
全スレッドで一要素ずつ分担して転送してる。無駄にはやってない。
Cのシェアードメモリからデバイスメモリへの転送も同じ。
シェアードメモリ内の As と Bs は各スレッドあたり BLOCK_SIZE 回使用する。
> 最後のホストの変数Cへのデータ転送は
今回のコードには入ってない。ホスト側のコードで、ブロック転送を指示する。
98:デフォルトの名無しさん
10/05/04 20:53:45
>>97
ｔｈｘ
99:デフォルトの名無しさん
10/05/07 12:16:01
GCC 4.5.0で使えるようにする設定を教えてください
100:デフォルトの名無しさん
10/05/07 18:46:18
for (int i=1;i<n-1;i++)
{
a[i]=b[i]+b[i+1]+b[i-1]
}
をGPUで並列実行するコードにしたいのですが、どのように書けばよいのでしょうか？
101:デフォルトの名無しさん
10/05/07 21:33:16
三つのループにばらしてBLASを３回使えばいいけど、
もっと速い方法があるかな。
102:デフォルトの名無しさん
10/05/08 01:11:10
>>100
for (int i=0; i<n; i++){
if (i == 0 || i == n){ break;}
a[i] = b[i] + b[i+1] + b[i-1];
}
のループを並列化すると考えるといいかと。
とりあえず1つのStreaming Multiprocessor(CUDA Core)を利用する場合は
dim3 dimBlock(ARRAY_SIZE_N);
dim3 dimGrid(1);
kernel_1<<<dimGrid, dimBlock>>>(a_dev, b_dev);
__global__ void kernel_1(float *a, float *b)
{
int tx = threadIdx.x;
if (tx == 0){ return;}
if (tx == ARRAY_SIZE_N){ return;}
a[tx] = b[tx] + b[tx+1] + b[tx-1];
}
こんな感じかな。
103:デフォルトの名無しさん
10/05/08 01:11:54
続き・・・
これだけではSM1つしか使わないし、ARRAY_SIZE_Nが大きい場合はSMあたりのスレッド数上限を超えたりする。
なので実際には更に複数ブロックに分割して、複数のSMで実行されるようにしないと動作しなかったり非常に低い性能しか得られない。
G80やG92では256スレッド/ブロック程度が良いことが多いのかな。
でも変数の数によってはレジスタが足りなくなったりするだろうし、しっかり計算して決めるかパラメータを変えながら試行錯誤することになると思う。
お勧めサイトはイリノイ大学の
URLﾘﾝｸ(courses.ece.illinois.edu)
国内ではGeForceのアーキテクチャについて解説されているページも結構有用というか
ハードのことをある程度知っていないとかなり厳しいかと。
104:デフォルトの名無しさん
10/05/08 02:45:33
ついでに言うと、>102ではカーネルの処理時間が短過ぎて呼び出しコストに見合わない。
ARRAY_SIZEをブロックとスレッドに分割したら更に残りをループにするようにしないと効率が悪そう。
それから、bのアクセスも競合するから待ちが入ることになりそう。
こっちについてはループを3回に分けた方がいいかもしれないし、強行した方がいいかもしれないし、
profileと睨めっこする必要があるかもね。
ハードの知識がなくてもptx出力を読めてprofileの意味を理解できればいいのだけれど、
それって結局ハードの知識を得るのと同じことだしね。
105:デフォルトの名無しさん
10/05/08 03:14:44
結局nがいくつかによるだけじゃない。
たぶん100はCUDAのくの時もまだ理解していないと思う。
106:デフォルトの名無しさん
10/05/08 03:16:30
>>104
あの程度だとコンテキストスイッチやホスト・デバイス間の転送等の
呼び出しコストに見合わないね。
学習用か極端にシンプルにしたコードだと思ってあまり気にしなかったけど。
bのアクセスの競合はブロックに分割してシェアードメモリを使うなら
バンクコンフリクトも起こらず、大丈夫かと思ったけど甘いかな？
ptx読めてprofileの意味を理解するにはアーキテクチャの理解が必要というか
もうその段階はとっくに超えているような気もする・・・
107:デフォルトの名無しさん
10/05/08 03:18:30
突っ込みどころが多すぎるんだが…
面倒なので2点。
>Streaming Multiprocessor(CUDA Core)
違う。CUDA Coreに対応するのはSMではなくSP。
>a[tx] = b[tx] + b[tx+1] + b[tx-1];
前提としているGPUがわからんが、根本的にコアレスメモリアクセスになっていない。
その後のレスにあるG80やG92のCC1.2未満では問題外。
CC1.2ではハーフワープあたり無駄なメモリロードが２回走る。
CC2.0ではキャッシュが効くのでARRAY_SIZE_Nのサイズしだいでは問題にならない可能性がある。
108:デフォルトの名無しさん
10/05/08 03:45:08
>>107
厳しい突っ込みどうも。
CUDA Coreと呼ばれるようになったのはSMじゃなくてSPだったのですね。
図を見れば一目瞭然ですね・・・うろ覚えで書いてすみませんでした。
コアレスアクセスはまだよく理解していないので勉強してきます・・・
109:デフォルトの名無しさん
10/05/08 04:15:14
だから、16人17脚でパン喰い競争やるようなもんだよw
110:デフォルトの名無しさん
10/05/08 05:18:45
この表現はわかりやすいね。
111:デフォルトの名無しさん
10/05/08 06:18:59
一人重い荷物を持たされると、全体の速度が遅くなる、的なｗ
112:デフォルトの名無しさん
10/05/08 08:19:25
青木本で最も評価するのがあの手書きの絵だww
113:デフォルトの名無しさん
10/05/08 09:18:32
青木本にこのスレから引用したと見られる節が幾つかある件w
>>109
>107の例だと、自分のレーンのパンだけじゃなくて左右のパンも食べないといけないんだな。
そりゃ遅くなるわさ。
114:デフォルトの名無しさん
10/05/08 11:04:49
コアレスのコアってなんのコアのことなの?
115:デフォルトの名無しさん
10/05/08 11:07:02
coalescedなので、
合体したという意味。
合体してメモリにアクセスするって意味だ。
116:デフォルトの名無しさん
10/05/08 12:01:16
正直その並びなら、texture使えばキャッシュ使えるじゃん
ただみんな、言うように並列度がどれくらい稼げるか(nに依存)で変わるから実際やってみないとわからないんじゃない？？
117:デフォルトの名無しさん
10/05/08 12:24:17
↓新明解(旧版)で合体の項を引くなよ
118:デフォルトの名無しさん
10/05/08 13:18:40
メモリにアクセスするときに、各スレッドが足並みそろえて
＋１ずつのアドレスを見に行かないといかんのです
すると32ワードなりまとめてとってくる。 とってくる間は
別のスレッドが実行されてる。
119:デフォルトの名無しさん
10/05/08 14:52:01
>>114
私は「連動した」メモリアクセスと説明している。要は16人17脚なんだけどねw
120:デフォルトの名無しさん
10/05/08 22:51:22
おまえら頭いいんだろうな
121:デフォルトの名無しさん
10/05/08 22:57:18
どうせ底辺大学生が研究で勉強してるだけでしょ？
122:デフォルトの名無しさん
10/05/08 23:11:19
と学歴昆布が申しております
123:デフォルトの名無しさん
10/05/09 15:32:57
でもGT200以降だと16人17脚でパン喰い競争って微妙に語弊あるよね。
16人17脚でパン喰い競争じゃない使い方したことないけど。
124:デフォルトの名無しさん
10/05/09 16:44:58
>>102のコードは
あるスレッドの遅れが、他のスレッドの足を引っ張る結果になるの？
コード内で参照する領域に依存関係がないのに
他のスレッドに影響するの？
なんだか納得出来ないなあ
125:デフォルトの名無しさん
10/05/09 17:42:10
>>102
if使ってる時点でありえないかと。。。
適当に(関数名とか間違ってるかも)書くけど、
int tx = threadIdx.x;
a[tx] = text1D(tx-1) + text1D(tx) + text1D(tx+1);
で、thread数とblock数を最適化しておけば、おｋでは？？
出力もコアレスだしさ
126:デフォルトの名無しさん
10/05/09 20:50:46
>>125
そうでもないだろ。
どっちみち、nの数によっては条件分岐が必要なんだし。
まぁ、a[i] = b[i] + b[i+1] + b[i-1];と書くよりはa[i + 1] = b[i] + b[i+1] + b[i + 2];として条件分岐を一つ減らすだろうけど。
いずれにしても、既出のようにcoalescedにはならないわけだが。
>>124
SIMDって理解できている?
127:デフォルトの名無しさん
10/05/09 20:59:34
>>107
>a[tx] = b[tx] + b[tx+1] + b[tx-1];
前提としているGPUがわからんが、根本的にコアレスメモリアクセスになっていない。
本当にわからんなあ。例えばコアレスアクセスになるにはどう書けばいいの？
つーかCUDAって根本的に間違ってるよなぁ。
C言語で書けることが利点？なのに、ベタにCで書くと
コアレスじゃないとか何とかで遅いというｗ
じゃあ、遅くないfloatN_tでも用意してそれ使わせればいいのに、と考えると
HLSLとかの方がいいんじゃねえの？って結論になるｗ
128:デフォルトの名無しさん
10/05/09 21:11:07
だから、16人17脚だってばw
先ず、simdだから16人(スレッド)は同時に同じことしかできない。
但し、「何もしない」という行動は有り得る。
次に、メモリアクセスはcoalescedであるためには16レーンにリニアに並んだメモリが割り当てられないといけない。
パン喰いのアナロジーならレーン1からレーン16までに1番から16番のパンが並んでいる状態だな。
件のコードは、例えば4番のスレッドが4レーンを走っていて4番と3番と5番のパンを必要としているんだ。
最初はスタッフが1～16のパンを順に置いておいてくれるからすぐ喰える(読み込める)けれど、
次に3番のパンを喰うにはスタッフが1～16レーンに0番から15番までのパンを並べてくれるのを待たないといけない。
ついでに言えば、このとき1番のスレッドは0番のパンなぞ喰いたくはないから「何もしない」をしていることが要求されるな。
>126の遣り方ならそのための条件判断を端折れるということだ。
更についでに言えば、実は条件分岐も意外に遅くはない場合もある。この辺りはptx出力を眺めていると判ってくるところだね。
129:デフォルトの名無しさん
10/05/09 21:18:40
揚げ足をとるようだけど、右足を出した隣の人は左足を最初に出してる。
__syncthreadsしなければ隣のひとたちを見ないでばりばり走るし、
過度な比喩は逆に混乱させるだけだと思う。
>>100みたいなのは計算式を丸ごと変えないといけないな。
プログラムより、数学的な知識が必要だと思う。
130:デフォルトの名無しさん
10/05/09 21:19:54
>>126
いやｗ
nの数の条件分岐とか必要ないでしょｗｗ
Host側がコール時に調整すればいいだけの話
CUDAの場合、無駄な命令入れるよりかは、少し出力メモリ先がはみ出ててもコアレスならおｋ
131:デフォルトの名無しさん
10/05/09 21:23:42
>>130
nが巨大なら、ループする必要があるでしょ。ループ継続条件は通常、nとの比較で行なわれるわけで。
132:デフォルトの名無しさん
10/05/09 21:24:41
>>128
キャッシュラインのことなのかレジスタなのかわからんが
bって、必ず特定のバッファに結び付けられてしまうの？
for (int i=1;i<n-1;i++)
{
registor r0 = b[i]
registor r1 = b[i+1]
registor r2 = b[i-1]
a[i]= r0 + r1 + r2
}
こうすりゃ別に、ストールはしないと思うが。
bは必ずr0にロードされるって制約があるなら別だが。
133:デフォルトの名無しさん
10/05/09 21:32:13
>>131
だったら数回カーネルをコールすればいいのでは？？
この場合カーネル内でループしちゃうと、キャッシュ外れる避けられない。キャッシュミスは大きいよ。
134:デフォルトの名無しさん
10/05/09 21:39:07
>>100 は一次元差分法からの連想だろうけど、
差分法なら青木本に２次元があるから参考にしたら。
自分は３次元の差分法を作ったけど信じられない速度が出た。
まあ、比較したCPU版が遅いんだろうけどｗ
135:デフォルトの名無しさん
10/05/09 22:06:07
>>132
CC1.1までを前提にするけど、コアレスであるためには
スレッドがアクセスするメモリがハーフワープ(16個)連続して配置されている必要がある。
SMがロードする際16個まとめてグローバルメモリをロードして各スレッドに割り当てるから。
だからハーフワープあたりのグローバルメモリのアクセスは1回になる。
じゃあコアレスじゃないとどうなるかというと、
キャッシュなんてないないから各SPは個別に毎回グローバルメモリに読みにいく。
だから遅い。
136:デフォルトの名無しさん
10/05/09 22:20:52
>>135
「bへのアクセス3つは16個連続していること」って用件は満たしてる。
多分アライメントも16個境界になってないと駄目ということか。
確かにそうすると、 b[i+1]とb[i-1]はアライメントに合ってないから
都度ロードが発生する訳だな。
理解しました。どうもありがとう。
137:デフォルトの名無しさん
10/05/10 00:33:06
CUDA Programing Guide 2.2を見てコアレスメモリアクセスをある程度理解したつもりですが、
Compute Capability 1.0/1.1では開始アドレスが64の倍数で、half-warp内のメモリアクセスする全てのスレッドが
開始アドレス + (スレッドID x 32bit)のアドレスにアクセスする場合に限って
トランザクションが1回にまとめられるということでしょうか？
CC 1.2以降はこのあたりの制限がかなり緩くなっているようではありますけど。
>>125
warpサイズを考慮せずに分岐のあるコードを書くと
両方の処理を行うことになって極端に効率が落ちると思いますが
分岐先で何もせずにreturnする場合でもかなり落ちるのでしょうか？
138:デフォルトの名無しさん
10/05/10 00:51:19
>>127
CUDAというかアーキテクチャの制限ではないでしょうかね？
そのままグローバルメモリにアクセスするのではなく、シェアードメモリを利用するという思想で作られていると思っています。
>>130
出力メモリ先がはみ出るのはまずいかと思いますが、分岐が複雑になる場合は
分岐を行わずに配列に0を入れておくとかの方が良さそうですね。
>>134
あの本はそういったことも載っているのですか。
やはり一度は読むべきなのでしょうかね・・・
139:デフォルトの名無しさん
10/05/10 02:36:11
>出力メモリ先がはみ出るのはまずいかと思いますが、
メモリを考慮してグリッドとブロックの数を決めるより
はじめにグリッドとブロックの数を決めてそれに合わせてメモリを確保したほうが良い。
ってことが言いたいんじゃないか。
140:デフォルトの名無しさん
10/05/10 07:11:13
>>137
分岐の種類によって違う結果になる。
・ワープ中幾つかのスレッドだけが条件に該当する場合
条件分岐のコスト(のみ)が余計に掛かる。
・ワープ中全てのスレッドが条件に該当する場合
ワープによっては条件分岐のコストが嵩むが
ワープによっては条件分岐以降の処理をパスするので早く開放される。
尚、if文の中身が複雑でループ内にあるがループには依存しない場合、
条件判断処理をループ外で行ない結果のBoolean値だけをレジスタに保存しておく最適化は行なわれる。
条件分岐は怖くありませーん。下手に複雑な演算するよりお徳でーす。
141:デフォルトの名無しさん
10/05/10 11:26:17
>>102
シェアードメモリは使ったほうがいいと思いますよ
142:デフォルトの名無しさん
10/05/10 11:47:48
うーん、CUDAの初歩も知らない奴に取り敢えず雛形を提示したら、
ちょっと齧った連中が寄って集ってけちつけるという構図になったなw
合間に慣れた人が適切な指摘をしてくれているからいいけどねぇ。
つーか、CUDAがそれだけ「注目」されていると言うことなんだろうけど。
143:デフォルトの名無しさん
10/05/10 13:31:41
やだ・・・なにこの失敗したCell並に難解なプログラミング・・・
144:デフォルトの名無しさん
10/05/10 14:51:44
大丈夫、Cellのspuみたいにインストラクションを一個ずつあれこれしなければならないほど複雑じゃない。
145:デフォルトの名無しさん
10/05/10 17:57:02
GTX4*シリーズ買う？
それとも買った？
146:デフォルトの名無しさん
10/05/10 23:14:36
前スレから
303 名前：デフォルトの名無しさん[sage] 投稿日：2009/12/05(土) 19:32:04
倍精度性能はGT200世代はおまけで１個演算機がついてるだけだから期待するだけ無駄だよ
Fermi世代から各コアに演算機付けるって言ってるけど、
一般向けには倍精度削ってGF100の名前で出すとか言ってるからどうなるか分からん
>一般向けには倍精度削ってGF100の名前で出すとか言ってるからどうなるか分からん
GTX480で倍精度1/4とかなった（ようだけど）
去年の12月の時点でこんな話が出ていたんだね
で、GTX4*買うかはやっぱり倍精度が気になるわけで、様子見の状態
147:デフォルトの名無しさん
10/05/12 09:25:19
C20シリーズは何とか予算用意して買うか、どこかの計算機センタに導入されるのを待つ形になる
それの準備的なものとして、GTX4シリーズを買うと思う
148:デフォルトの名無しさん
10/05/14 10:29:43
完全にHPCユーザーしかGPGPUに興味を持たなくなったな
149:デフォルトの名無しさん
10/05/14 11:10:38
画像処理とか（？）する人はどう思っているんでしょう
GPGPUについては
150:デフォルトの名無しさん
10/05/14 12:19:47
CG系研究室の俺が。
GPGPU(=CUDA)は今や当たり前の技術です。
なんか論文で新しい手法を提案するとき、実行時間の項にCPUでの計算時間だけ書いていたら、
「GPUでやったらどうなのよ？」と査読者から普通にコメントつけられるレベルです。
GPUに載せづらい既存の手法を、ちょっと改良してGPU適合なアルゴリズムにしました、的な論文が
3年ぐらい前から大量に出ています。
151:デフォルトの名無しさん
10/05/14 13:02:33
>>150
> CG系研究室の俺が。
CGの人ってコンピュータビジョンにもアンテナ伸ばしてるものなんですか？
> 「GPUでやったらどうなのよ？」と査読者から普通にコメントつけられるレベルです。
＊ただしリアルタイム応用に限る。
> 3年ぐらい前から大量に出ています。
出過ぎたせいで最近は実装してみましたレベルの論文はむしろ通りにくいですよ。
152:デフォルトの名無しさん
10/05/14 18:40:04
ATI stream とCUDA とどっちがよく使われているんでしょうかね
153:デフォルトの名無しさん
10/05/14 19:19:01
「実装してみました」系はCUDAだろうね。
ピークパフォーマンスたたき出すならATIだろJS
154:150
10/05/14 20:07:19
>>151
>CGの人ってコンピュータビジョンにもアンテナ伸ばしてるものなんですか？
ちょっとだけね。 CVのいくつかのツールは、こっちでも役に立つから。
>出過ぎたせいで最近は実装してみましたレベルの論文はむしろ通りにくいですよ。
あるあるwwww
>>152-153
俺の狭い視界だけで言うと、ATI streamは使っている人を見たことがない。
俺は昔試しに使ってみたけど、ちょっと複雑なことをやろうとするとすぐに
Direct3D8でやったようなシェーダアセンブリを書かなくちゃならなくなって辟易した。
今のところGPGPUではCUDA一択だと思います。
研究者の皆さんもツールの使い方にそれほど時間を掛けられるわけでもないので、
ほぼC/C++と同様に書けるCUDAが主流になるのは仕方ないと思います。
155:デフォルトの名無しさん
10/05/14 20:19:13
ATIを利用した論文を目にしたことがないな、確かに。
156:デフォルトの名無しさん
10/05/14 20:23:35
米NVIDIAは13日(現地時間)、2011年第1四半期(2010年5月2日締め)の決算を発表した。
これによると売上高は10億180万ドルと、前年同期比で51%の増収。
純利益は前年同期の2億130万ドルの赤字から、1億3,760万ドルの黒字へと回復した。1株当たり利益は0.23ドル。
同四半期中、同社はハイエンドGPU「GeForce GTX 400」シリーズを出荷。
また、ハイエンドサーバー向けのTeslaが過去最高の売り上げを記録するなど、上位製品が好調だった。
また、同社はリリース中で、Microsoftが先だって米国で発売したスマートフォン「KIN」がTegraを搭載していることを明らかにした。
157:デフォルトの名無しさん
10/05/15 01:12:46
>>154
GPGPUとしてはCUDAが主流になると思う。
だけどこれまではほんとに酷かったなあ。
ちょっとでもコンピュータサイエンスの知見を持っていたら恥ずかしくて発表出来ないレベルが多かった。
今までは流行でみんながCUDA、CUDAといっていたけど、多くの人はあくまでも手段の一つであって、
目的ではない。目的にするならコンピュータサイエンスの分野でやればよいわけで。
それにCUDAを使う場合って大抵最適化しているから、CPUとの比較も余り意味がないし。
これからはそういうのはなくなっていくだろうね。
158:& ◆6kLsuKGM9vPI
10/05/15 02:51:39
480と285の倍精度計算のベンチマークの比較はどこかにありませんか?
159:デフォルトの名無しさん
10/05/15 14:32:42
俺が勘違いしていないかの確認なんだが
a[tx] = b[tx] + b[tx+1] + b[tx-1];
はハーフワープ当たりの読み出しが3回でいいんだよな?
16*3回じゃないよな?
それより俺は音声処理とかのIIRっぽい場合で、bがaに依存するプログラムをどう書けばいいのか知りたい。
aを更新する処理で更新前のaと比較、分岐していたりすると俺の頭脳では依存関係が解消出来ない。
世の中の音声処理はSIMDとかSMPに対応しているんじゃないかと思うんだが、どうやっているんだろう。
160:デフォルトの名無しさん
10/05/15 15:40:25
>>157
どういうこと？
161:デフォルトの名無しさん
10/05/15 16:14:37
>>159
うまく行くケース(GPUのバージョン依存だっけ？？)なら３回でいけるだろうけど、
最悪の場合
b[tx]:コアレス
b[tx+1]:シリアライズ
b[tx-1]:シリアライズ
ってなるから、1+16+16 = 33 じゃね？？
162:デフォルトの名無しさん
10/05/15 16:39:09
>>161
全部のSPがb[tx]を読む→コアレス
全部のSPがb[tx+1]を読む→コアレス
全部のSPがb[tx-1]を読む→コアレス
とはならないの？
163:デフォルトの名無しさん
10/05/15 17:00:15
>>161
そういう場合、
スレッド i-1 はtx-1をグローバルからSharedに読む
スレッド i はtxをグローバルからSharedに読む
スレッド i+1 はtx+1をグローバルからSharedに読む
：
でひと固まりもってきて、
あとはShared上でCoaleshed? ｷﾆｼﾅｲ! で動けばいいんじゃね
164:デフォルトの名無しさん
10/05/15 20:35:59
>>162
だってアラインがあわないじゃん。
>>163
shared memory使うなら、
b[tx]:コアレス
b[tx+1]:コアレスx1, シリアライズx1
b[tx-1]:コアレスx1, シリアライズx1
かなー
GPUバージョンだっけ？これが2.0とかなら、たしか上記全てコアレスになると思うけど
165:デフォルトの名無しさん
10/05/16 02:31:46
>>162-163
「t*x」 になっているんでこのままだとtステップになっちゃうわけだよ
つまりもともとの配列の並べ方から考え直さないといかんよな。
x-1, x, x+1, ,,, になったら163みたいに、
・スレッド協調してSharedに持ってくる
・Shared上では好き勝手に読み書きしてよろし
・syncthreadする
・スレッド協調してGlobalに書き戻す
というやりかたが基本かな
166:デフォルトの名無しさん
10/05/16 09:56:22
>>165
>・Shared上では好き勝手に読み書きしてよろし
今扱ってる問題だと場合によっては、bank conconflictがあったりする。。。
167:デフォルトの名無しさん
10/05/16 09:57:09
>>166
まちがえたｗ
>>165
>・Shared上では好き勝手に読み書きしてよろし
今扱ってる問題だと大丈夫だけど、場合によってはbank conconflictがあったりする。。。
168:デフォルトの名無しさん
10/05/16 10:06:01
これだと全部コアレスになる？
const int tx = threadIdx.x;
const unsigned int unit_size = 16;
const unsigned int unit_min = 0;
const unsigned int unit_max = unit_size - 1;
const unsigned int txd = tx / unit_size;
const unsigned int txm = tx % unit_size;
const value_t * bc = b + ( (txd+0) * unit_size );
const value_t * bn = b + ( (txd-1) * unit_size );
const value_t * bp = b + ( (txd+1) * unit_size );
a[tx] = *( bc + txm );
if( txm > unit_min )
a[tx] += *( bc + txm - 1 );
else
a[tx] += *( bn + unit_max );
if( txm < unit_max )
a[tx] += *( bc + txm + 1 );
else
a[tx] += *( bp + unit_min );
169:デフォルトの名無しさん
10/05/16 11:06:33
テクスチャキャッシュを使うのが面倒なくていいと思うのは俺だけか?
170:デフォルトの名無しさん
10/05/16 11:33:22
>>169
少し前にテクスチャ使ったコードで例を示したけど、誰も反応しなかったぞ
なんで？？
171:デフォルトの名無しさん
10/05/16 14:23:32
どのレス？
172:デフォルトの名無しさん
10/05/16 18:59:40
>>171
>>125
173:デフォルトの名無しさん
10/05/16 19:31:54
OpenCLでなくCUDAを学ぶ価値ってあるの？
174:デフォルトの名無しさん
10/05/16 19:34:06
OpenCLってなにに使われてるの
175:デフォルトの名無しさん
10/05/16 19:38:18
OpenCLは今のところRadeon使う人用かと。
176:デフォルトの名無しさん
10/05/16 20:33:59
DirectComputeはどんな感じ？
177:デフォルトの名無しさん
10/05/16 21:33:22
ローカルメモリに対するテクスチャメモリのメリットがいまいち分からない
そもそもローカルメモリを使いこなせないから、テクスチャメモリまで手が回らない
178:デフォルトの名無しさん
10/05/16 23:49:15
>>175
え？
179:デフォルトの名無しさん
10/05/17 00:41:56
>>177
補間できるとか
180:デフォルトの名無しさん
10/05/17 20:52:47
これってcudaMallocHostってsizeof(構造体)でぴったりとその構造体の容量分だけちゃんと確保できてるんですか？
なんかsizeof使ってぴったり構造体分確保したら条件によってセグメントエラーが出るんですが（printfを適当なとこに放りこんだらエラーが出ずに通ったりします）
2倍の容量を確保したら問題なく通ります
181:デフォルトの名無しさん
10/05/17 23:08:13
>>177
ローカルメモリがLocal Memoryのことをいってるなら、役割がまったく違うからメリットもなにもない
Global Memoryのことをいってるなら、キャッシュに加えて補間、正規化アドレス、変換とかだな
キャッシュが有効に機能しているかは、プログラムによっては疑わしいんだけどね
あとコーディングがちょっとめんどくさい
182:デフォルトの名無しさん
10/05/18 09:35:45
>>180
構造体に問題があると思われます。
各要素がアラインメントの条件を満たすようにダミーの要素を追加して調整することをお勧めします。
くらいしかアドバイスしようがないな。
183:デフォルトの名無しさん
10/05/18 10:19:15
printf挟むと動くと言ってるからどっかでスタック壊してんじゃないかと予想
184:デフォルトの名無しさん
10/05/19 14:37:47
まぁ、晒せ、と。
185:デフォルトの名無しさん
10/05/20 11:59:21
ω
186:デフォルトの名無しさん
10/05/21 09:13:45
研究室にGTX 480届いた．
smokeParticlesで170fpsくらい．
倍精度FFTとかはこれからやる．
187:デフォルトの名無しさん
10/05/21 10:23:00
FFTやってみた．
1024×1024の二次元FFTの順変換&逆変換を1024回ループで回した．
・単精度
285 → 1.428[s]
480 → 0.736[s]
・倍精度
285 → 7.917[s]
480 → 3.829[s]
コアの数の比を考えると速度は変わらない気がする．
188:デフォルトの名無しさん
10/05/21 11:33:50
nbody --benchmarkよろ
189:デフォルトの名無しさん
10/05/21 11:53:41
nbody --benchmarkがよく分からんから全部貼る
C:\Documents and Settings\***>nbody --benchmark
Run "nbody -benchmark [-n=<numBodies>]" to measure perfomance.
-fullscreen (run n-body simulation in fullscreen mode)
-fp64 (use double precision floating point values for simulation)
> Windowed mode
> Simulation data stored in video memory
> Single precision floating point simulation
> Compute 2.0 CUDA device: [GeForce GTX 480]
15360 bodies, total time for 10 iterations: 73.829 ms
= 31.956 billion interactions per second
= 639.126 single-precision GFLOP/s at 20 flops per interaction
これでいい？
190:デフォルトの名無しさん
10/05/21 13:05:20
>>187
多分CUFFT3.0を使った複素２次元FFTだと思うのだが、
こちらの計測データとかなり違うような・・・。
こちらでは480の単精度で1回(片方向)あたり0.44ミリ秒、2*1024回分で0.91秒くらい。
ちなみに頑張ると倍精度は2.6秒くらいでいける。
191:187
10/05/21 13:14:35
CUFFTに慣れてないから何か不備があるのかも．
↓ソース
#include <cutil.h>
#include <cufft.h>
const int ElementNum = 1024;
int main(int argc, char **argv)
{
cudaSetDevice(0);
int i, j, k;
cufftDoubleComplex *HostData, *DeviceData;
int nbyte = sizeof(cufftDoubleComplex) * ElementNum * ElementNum;
cudaHostAlloc((void **)&HostData, nbyte, cudaHostAllocPortable);
cudaMalloc((void **)&DeviceData, nbyte);
192:187
10/05/21 13:16:38
for(i = 0; i < ElementNum; ++i)
{
for(j = 0; j < ElementNum; ++j)
{
if(i > (ElementNum/2-128) && i < (ElementNum/2+128))
{
if(j > (ElementNum/2-128) && j < (ElementNum/2+128))
{
HostData[i+ElementNum*j].x = 1.;
HostData[i+ElementNum*j].y = 0;
}
}
else
{
HostData[i+ElementNum*j].x = -1.;
HostData[i+ElementNum*j].y = 0;
}
}
}
cudaMemcpy(DeviceData, HostData, nbyte, cudaMemcpyHostToDevice);
cufftHandle plan;
cufftPlan2d(&plan, ElementNum, ElementNum, CUFFT_Z2Z);
printf("\nStart Calculation\n");
193:187
10/05/21 13:18:09
unsigned int timer;
CUT_SAFE_CALL(cutCreateTimer(&timer));
CUT_SAFE_CALL(cutResetTimer(timer));
CUT_SAFE_CALL(cutStartTimer(timer));
for(k = 0; k < 1024; k++)
{
cufftExecZ2Z(plan, DeviceData, DeviceData, CUFFT_FORWARD);
cufftExecZ2Z(plan, DeviceData, DeviceData, CUFFT_INVERSE);
}
CUT_SAFE_CALL(cutStopTimer(timer));
printf("\nEnd Calculation\n");
printf("\nProcessing time : %f [ms]\n", cutGetTimerValue(timer));
CUT_SAFE_CALL(cutDeleteTimer(timer));
cudaMemcpy(HostData, DeviceData, nbyte, cudaMemcpyDeviceToHost);
cufftDestroy(plan);
cudaFree(DeviceData);
cudaFreeHost(HostData);
CUT_EXIT(argc, argv);
return EXIT_SUCCESS;
}
194:187
10/05/21 13:21:21
nvcc cufft.cu -o cufft_double -arch=sm_20 -lcudart,cutil32,cufft -Xcompiler /O2
↑のソースをこれでコンパイルした．
195:デフォルトの名無しさん
10/05/21 13:47:54
>>189
おお、ありがとうございます。
GTX260でも370GFOPS位出るのであんまり速くないですね・・・（´・ω・｀）
196:デフォルトの名無しさん
10/05/21 13:54:24
>>187
>CUT_SAFE_CALL(cutStopTimer(timer));
これの前にcudaThreadSynchronize();を入れないといけない予感。
k の反復回数を512にした時のほぼ倍の時間になるかどうかをチェックしてみて。
197:187
10/05/21 15:00:01
同期をちゃんととってみた．
・単精度
285 → 1.710[s]
480 → 0.875[s]
・倍精度
285 → 8.438[s]
480 → 4.076[s]
>>190
倍精度を頑張るってどう頑張るんですか？
198:デフォルトの名無しさん
10/05/21 15:28:32
質問させてください
CPU...：Intel Core i7-860(2.80GHz)
メモリ........：4G
M/B .......：ASUS P7P55D Delux
VGA .....：FX５８０
電源..........：850W (80Plus)
の機材で現在ps4とpremireで地方の小さなブライダルの写真や映像をやっているのですが
友人からFX５８０を譲っていただけるとのことで２枚差しで使ってみようかと思うのですが
SLIは全く知識がなく、２枚差しの効果があるのかがわかりません。
どこで聞いたらよいのかわからずスレチかもしれませんが先輩方、アドバイスお願いします。
199:デフォルトの名無しさん
10/05/21 20:45:35
SLIというかマルチGPUで使う場合はOpenMPかなんかで複数スレッドを作って
それぞれののGPUを使えばいい、というかサンプルがSDKにあったはず。
ゲームみたいにSLI組んでシングルGPUに見せるのはまだ出来なかったはず
200:デフォルトの名無しさん
10/05/21 21:36:18
>>199
親切にありがとうございます。
まさにゲームのようなSLIの考えでいましたが、そういった使い方はできないということですので
自宅用に使用するか等、使い方を考えてみます。本当にありがとうございました。
201:デフォルトの名無しさん
10/05/21 22:56:10
>>197
もちろんコードを自分で書く。
202:デフォルトの名無しさん
10/05/26 16:39:37
GPU TeslaC1060x2 / OS CentOS5.4 x64
という環境で、
CUDA3.0 / R-2.11.0 / cula_1.3a-linux64 / gputools_0.21
を入れてみました。
自分でソースを書いたり、culaのexampleなどは動くのですが、
R上でgputoolsを使おうとするとエラーになってしまいます。
> gpuCor(A, B, method="pearson")
以下にエラー gpuCor(A, B, method = "pearson") : invalid argument
追加情報: 警告メッセージ：
In gpuCor(A, B, method = "pearson") : PMCC function : malloc and memcpy
> gpuCor(A, B, method="kendall")
以下にエラー gpuCor(A, B, method = "kendall") : invalid device function
追加情報: 警告メッセージ：
In gpuCor(A, B, method = "kendall") : executing gpu kernel
こんな感じです。chooseGpu()なんかは動くのですが。
どうすれば良いのか、ヒントでも良いので教えてください。
203:202
10/05/26 17:15:17
と聞いたばかりですみません。
解決しました。
ちゃんとgputoolsのinstall.txtを呼んでいなかったので
パラメータを間違えていて、--enable-emulationというのを
つけてしまっていました。
お騒がせしました。
204:デフォルトの名無しさん
10/05/26 23:26:49
http://倍精度、FP64演算については、取り組み方をGeForce GTX 2x0の時から変わっている。
GeForce GTX 2x0では1SM(8SP)あたり、1基の専用FP64スカラ演算器を有していたが、GeForce GTX 4x0でこれを削除したのだ。
しかし、GeForce GTX 4x0ではSP内のFP32スカラ演算器で2サイクルをかけてFP64演算を行うようにしている。
専用演算器はなくなったが、増加したSP群のおかげでピーク時のFP64演算性能は先代からちゃんと向上することにはなる。
いまさらだけど1/4にはならないって
205:デフォルトの名無しさん
10/05/26 23:38:38
>>204
言われてる1/4ってのは当初期待されていた性能に対して1/4ってことじゃなかったの？
いくらなんでもTesla世代と比較してないでしょ
206:デフォルトの名無しさん
10/05/27 03:41:07
当初は短精度に対して理論値で1/2の性能を出せるって話だったのだけど
Geforceでは削られて短精度に対して1/4しかでないことになっている
207:デフォルトの名無しさん
10/05/27 05:53:57
1/8じゃなかったの？
208:デフォルトの名無しさん
10/05/27 06:06:21
皆GeForceでは1/8 (Fermi版Teslaシリーズのさらに1/4)という意味で話していたけど
一部の人は誤解していたみたい。
209:デフォルトの名無しさん
10/05/27 20:51:44
どういうことだってばよ
210:デフォルトの名無しさん
10/05/28 00:28:01
けっきょくどういうことなんだってばよ！
211:デフォルトの名無しさん
10/05/28 08:37:10
こまけえことはいいんだよ！
>>205
俺の理解はこれなんだが
GTX480も在庫が増えて値段も下がってきたね
俺、5万円切ったら本気出すんだ…
212:デフォルトの名無しさん
10/05/28 17:09:48
Teslaじゃコードを書き換える手間を考えたら、たいしてコストパフォーマンスは良くないな。
213:デフォルトの名無しさん
10/05/28 21:44:15
？
214:デフォルトの名無しさん
10/05/29 05:36:25
>>100の様なコードだったら手間はかからんよ。
215:デフォルトの名無しさん
10/05/29 12:56:20
TeslaがGT世代の事じゃなくてGPGPU専用カードの事を言ってるなら
Teslaの利点はバグのあるコードを走らせてもOS巻き込んで死んだりしない点が一番じゃね？
開発機はTeslaが楽だわ
216:デフォルトの名無しさん
10/06/02 22:09:06
ドライバAPIでパラメータ取得したらこんな感じになりました。
device_name: GeForce GTX 480
major = 2, minor = 0
max_threads_per_block:1024
max_block_dim_x:1024
max_block_dim_y:1024
max_block_dim_z:64
max_grid_dim_x:65535
max_grid_dim_y:65535
max_grid_dim_z:64
max_shared_memory_per_block:49152
total_constant_memory:65536
warp_size:32
max_pitch:2147483647
max_registers_per_block:32768
clock_rate:810000
texture_alignment:512
gpu_overlap:1
multiprocessor_count:15
kernel_exec_timeout:0
integrated:0
can_map_host_memory:1
conpute_mode:0
concurrent_kernels:1
ecc_enabled:0
Total memory size: 1576468480
driver_version: 3000
217:デフォルトの名無しさん
10/06/03 06:46:24
2次元配列を使うと、遅くなりますが、
2次元ブロックも同様の問題をかかえていますか?
218:デフォルトの名無しさん
10/06/04 17:03:58
なんのこっちゃ
219:デフォルトの名無しさん
10/06/05 01:05:17
2次元配列なんてCUDAにはありません。
意図的に1次元に並べてください。
220:デフォルトの名無しさん
10/06/05 20:17:51
GTX480でおすすめのボードってありますか？
もしくはダメなボードはありますか？
221:デフォルトの名無しさん
10/06/05 21:43:43
こちらも480入手。取り敢えず動くのは確認。細かいテストはこれから。
>>220
今のところどこも経験値低いだろうからリファレンスボードのままだろうし、
ELSAが無難じゃね? 高いかもしれないけど。
ヒートパイプの所為で更に場所を喰うので要注意。
222:デフォルトの名無しさん
10/06/05 21:52:51
リファレンスだからドスパラのでいいんじゃないの
シール張るだけらしいから
223:220
10/06/07 11:06:57
>>221-222
ありがとうございます。
elsaのGTX480のページに対応OSがwindowsしか書いてないんですが、
Linuxでは動かないのでしょうか？
GeforceのCUDA用ドライバをいれたら普通に動くんでしょうか？
224:221
10/06/07 11:31:39
>>223
ドライバのリビジョンは今手元にないから判らないけど、CUDA2.3当時のドライバでCUDA2.3が動いている。
CUDA3に関しては未だテストしていない。
225:デフォルトの名無しさん
10/06/14 13:16:15
最近CUDAを勉強しはじめたんですが、質問させてください。
入力画像をグレースケール化して出力する、というプログラムを作ったところOpenCVだけで作ったプログラムに比べて、
プログラム全体の処理時間は速くなりませんでした。どうやらプログラムで最初に実行されるcudaMallocで100ms近く掛かっているようです。
次の行にも同じくcudaMallocがあるんですが、これは1msほどで完了します。
これは改善できませんか？
OSはUbuntu 9.04、GPUは9500GT、cudaのバージョンは2.3です。
226:デフォルトの名無しさん
10/06/14 13:20:32
できます。
一回目のcudaMalloc()の前にcudaThreadSynchronize()でもしておけば、一回目のcudaMalloc()は早くなります。
尤も、cudaThreadSynchronize()で時間が(恐らく100ms程度)掛かりますが。
227:デフォルトの名無しさん
10/06/14 13:29:32
>226
ありがとうございます。
早速試したところ、そのとおりの結果になりました。。。
これは諦めるしかないんですかね。
228:デフォルトの名無しさん
10/06/14 13:29:33
なんかCUDAは最初に呼ぶときに時間がかかるとかリファレンスかなんかに書いてあったような・・・
斜め読みしたけどみつけられんかった
229:デフォルトの名無しさん
10/06/14 13:41:43
>228
そんな仕様だったんですね、知らなかったです。
書き換えたもとの処理がたいして時間が掛からない場合、CUDAの利点を活かせないんですね。
230:226
10/06/14 13:45:06
>>227
cudaは元々単発プロセス向きじゃないので、諦めて常駐型のプロセスか処理時間の掛かるプロセスに使いましょう。
具体的に何をやっているかは判らないけど、cudaのAPIを呼ぶ初回は必ず時間が掛かるのよね。
例えばmain()の先頭ででも一回syncするだけでいいから、私の用途では問題にならないんだけど。
231:デフォルトの名無しさん
10/06/14 13:48:59
>>229
一回だけだとそうかも知れないけど、バッチやリアルタイムで複数ファイルを処理するとかなら
効いてくると思います。
232:デフォルトの名無しさん
10/06/16 10:00:49
研究にCUDAを使ってみたいと思っている学生です。＞＜
プログラマブルシェーダで物理演算をしてシミュレーションをやっている
論文とかあったのですが
CUDAで計算するのとの違いはあるんですか？
あと
DirectXを使いながらCUDAを利用することは
DirectXで手一杯のことをやっていたらCPUで計算したほうが早かったりするんですか？
お願いします（。・・。）
233:デフォルトの名無しさん
10/06/16 10:05:47
>>232
それが俺の読んだのと同じならcgを使ってシェーダプログラムで演算をしてた
理由はCUDAがまだ無かったから
DirectXで3D描画しながらだとそりゃ性能は落ちるでしょ。CUDA専用にカードを追加するなら話は別だけど
234:デフォルトの名無しさん
10/06/16 12:16:27
だいぶ前の話だけど、AviUtlのフィルタでシェーダよりCUDAの方が遅いとかって無かったっけ。
カリカリにチューンしたらCUDAが速いんだろうけれど、シェーダの方がメモリアクセスのお作法を意識しなくて良い分、初心者には性能向上が簡単かも知れない。
235:デフォルトの名無しさん
10/06/16 19:47:40
>>234
つまりn00bだったってことかよ!
236:デフォルトの名無しさん
10/06/16 23:49:14
>>234
???
CGをやるなら別だが、今更Cgは無いだろう。
初心者ならなおさらCUDAだ。
特に物理演算をやるなら。
237:デフォルトの名無しさん
10/06/17 00:34:00
PGIコンパイラのCUDA化って正直どうよ。
人力でやるのに疲れてきた。
それなりの性能が出るなら買ってもいいと思ってるんだけど。
238:デフォルトの名無しさん
10/06/17 06:18:56
ディレクティブでは全然性能上がらなくて頭来たからCで書き直したｗ
239:デフォルトの名無しさん
10/06/17 06:53:59
東工大、Teslaを利用したスパコン「TSUBAME2.0」を開発
URLﾘﾝｸ(pc.watch.impress.co.jp)
240:デフォルトの名無しさん
10/06/17 14:14:54
>>239
有用な結果を出さなかったら、次から仕分けな
241:デフォルトの名無しさん
10/06/17 18:10:08
国会議員を仕分けすりゃいいのに
1/3位にさｗ
242:デフォルトの名無しさん
10/06/17 20:46:42
あんたら、部外者か？
蓮舫のブレーンがそのスパコンの推進者なのに、仕分けされるわけがねーだろ。
243:デフォルトの名無しさん
10/06/17 21:00:10
つまり漣舫は神戸のスパコンをつぶしたくで仕分けしたのですね。
244:デフォルトの名無しさん
10/06/17 21:11:00
1度でいいからこんな超計算機をいじり倒したいものだ
245:デフォルトの名無しさん
10/06/17 21:15:30
部外者だけどCSX600なんて誰も使ってないんじゃね？
導入初期に不良品大量に掴まされてたみたいだけどｗ
246:デフォルトの名無しさん
10/06/17 21:28:53
まあそうでもしないと進まないところもあるんだろうけど、こんな財政状況なんだからバッサリ逝ってよし
247:デフォルトの名無しさん
10/06/17 21:34:09
>>242
むしろ部外者しかいないだろ
お前はしらんが
248:デフォルトの名無しさん
10/06/17 21:35:11
議員を減らせば良い
249:デフォルトの名無しさん
10/06/17 21:35:34
政治ネタうぜぇ
250:デフォルトの名無しさん
10/06/17 21:43:16
選挙だもの
民主党
251:デフォルトの名無しさん
10/06/17 22:18:46
れんほうのすぐ後ろに金田＠superπが座っていたから
京速機のヘッドである姫野＠ヘボベンチは何も言えんかったんだよ。
専門がスパコンはおろかコンピュータですらない姫野に
反論ができるわけもない。
252:デフォルトの名無しさん
10/06/17 22:21:51
π（笑）
253:デフォルトの名無しさん
10/06/18 02:21:28
>>251
何で姫野ヘボベンチっていわれているの？
ソースを見たけどコード自体は学生が書くような内容だね。
だからなの？
254:デフォルトの名無しさん
10/06/18 12:50:58
メモリ帯域をひたすら無駄食いするプログラムだよね＞姫野ベンチ
ここで「なぜか遅いです」とか質問してくるレベル。
255:デフォルトの名無しさん
10/06/18 15:19:59
かといってgemmでピーク性能測られても役に立たない分野も多く。
直線に強い車かカーブに強い車かってのと似てる。
256:デフォルトの名無しさん
10/06/18 17:38:19
πって何を測ってるの
257:デフォルトの名無しさん
10/06/18 17:42:40
今ではすっかり使われないx87
258:デフォルトの名無しさん
10/06/19 03:34:48
>>235-236
今更シェーダをお勧めはしないけど、畳み込みを考えた時に
コアレスにするためのコードって余分だよね。バッドノウハウだよね。
ってことが言いたかった。
259:デフォルトの名無しさん
10/06/21 12:53:32
東工大のスパコン、世界２位なんだ。
そういうことか。
260:デフォルトの名無しさん
10/06/21 13:55:44
東工大は64位じゃね？
URLﾘﾝｸ(www.top500.org)
261:デフォルトの名無しさん
10/06/21 16:58:25
ヒント：TSUBAME2はまだ実在しない
っ URLﾘﾝｸ(www.gsic.titech.ac.jp)
262:デフォルトの名無しさん
10/06/22 08:40:40
まだ計算どころか設置もしていないぞ
263:デフォルトの名無しさん
10/06/22 12:51:00
NECは京速機やめてこっちでやってんだな。
264:デフォルトの名無しさん
10/06/23 01:26:00
CUDAって未だにデバイスメモリをメインメモリにマップ出来ないの？
逆は出来るようだけど。
265:デフォルトの名無しさん
10/06/23 01:45:32
マップできなくてもいいけど、MPIでDMA使ってデバイスから通信チップに流しこんで欲しいわな
266:デフォルトの名無しさん
10/06/23 05:56:26
マップ出来ないといちいち転送指示出さなきゃいけないのが困る。
スクラッチで書くならいいけど、既存のを移植するとなると結構大変だから・・・。
267:デフォルトの名無しさん
10/06/23 19:31:52
テンプレートライブラリで、CUDAの転送まわりをほぼ完全に隠蔽してくれる奴あったよね。
268:デフォルトの名無しさん
10/06/24 12:55:15
thrust
269:デフォルトの名無しさん
10/06/26 00:36:06
使い方によるのかもしれないが、Mappedはかなり遅い気がする
Pinnedで転送したほうがまし
270:デフォルトの名無しさん
10/06/27 13:53:48
GTX480のベンチが日本語でも少しずつ出てきたね
倍精度と単精度両方載せてある
PGI CUDA Fortran を使用した行列積の計算とGPU性能
URLﾘﾝｸ(www.softek.co.jp)
CUDA Fortranだけど、CUDA Cでも結論が変わるものではないでしょう
271:デフォルトの名無しさん
10/06/27 13:57:56
MATMULとかだと倍精度も速いみたいだけど、レジスタやシェアードメモリの消費も倍になるから
実コードだとOccupancyが下がって思ったような速度でないんだよなぁ・・・ ＞Fermi
272:デフォルトの名無しさん
10/06/27 14:14:32
倍精度の実効効率が、
GTX285 約90%
GTX480 約21%
ずいぶん差があるのはGTX480の倍精度がドライバのレベルで1/4に抑えられている（らしい）ことが最大の原因？
アーキテクチャが違うから単純には言えないけど、21/25=84で、GTX285と同程度になるし…
273:デフォルトの名無しさん
10/06/27 14:51:36
ソース
274:デフォルトの名無しさん
10/06/27 14:53:17
>>273
>>272のは真上の祖父テックのベンチ結果から計算してのことだろ？
nbody.exeでも-fp64で性能出てなかったし
275:デフォルトの名無しさん
10/06/27 16:10:48
> ドライバのレベルで1/4に抑えられている（らしい）こと
ここのソース
276:デフォルトの名無しさん
10/06/27 16:41:24
googleを使えない奴にCUDA使えるとは思えないな
277:デフォルトの名無しさん
10/06/27 18:34:33
これと、中国のTUBASA？のLinpack結果を併せてだれか解説して
278:デフォルトの名無しさん
10/06/27 18:56:27
GTX480の「倍精度理論FP演算性能672 GFLOPS」という仕様そのものが間違っていて
実際は理論ピークが1/4の168。
別にドライバのレベルで抑えられているわけでもなんでもないが。
279:デフォルトの名無しさん
10/06/27 19:13:13
>>278
↓この記事を信じてしまうと
URLﾘﾝｸ(journal.mycom.co.jp)
GTX480の倍精度演算の理論値は672GFLOPSになっちゃうけど、
>実際は理論ピークが1/4の168。
が本当だとするなら、どこで計算がずれたのだろう…
280:デフォルトの名無しさん
10/06/27 22:22:30
ドライバなんだかそれ以外だかはわからないが、
GTX480の倍精度理論性能が168GFLOPSであることは信憑性高い、ということでおｋ？
281:デフォルトの名無しさん
10/06/27 22:24:19
ドライバで制限してるんじゃ無ければカタログスペックとの差異はどーすんの？
URLﾘﾝｸ(www.elsa-jp.co.jp)
282:デフォルトの名無しさん
10/06/27 23:07:18
>>281
GTX480の話をしているときにC2050持ち出してどうするの？
283:デフォルトの名無しさん
10/06/27 23:24:10
想像力を働かすと、
C2050よりもコア数が多くて、同一アーキテクチャ仕様と公表されているGTX480が168GFLOPSしか出ないなんておかしいぞ
ということかな
まあ、たとえ168GFLOPSだとしても、共有メモリやレジスタまで1/4にはなっていないだろうし（倍精度だと2倍食うけど）、
単位演算性能あたりのメモリバンド幅が単精度と比べて大きくなることを考えれば、実効効率が単精度よりも大きくなるのはわかる気がする
ECC無しでもOKで、メモリ１G程度で十分という人には結構おいしいかもしれんねGTX480
284:デフォルトの名無しさん
10/06/27 23:33:39
メモリしか差のなかったC1060に対して、C2050はもっと明示的な計算用チップとしてのアドバンテージをつけたかった
nVidiaの意向が反映した結果がGTX480だと思うのだけど。
285:デフォルトの名無しさん
10/06/28 07:10:21
C2050はGTX480はCUDAコアが少なく、クロックも低いにもかかわらず、
倍精度理論性能は515GFlopsとなりC2050ほうが高い。
これは、GTX480はドライバで倍精度演算の実行にペナルティが加わるようになってるから。
それでも、GTX285と比べても倍精度演算が高速だからいいじゃんというのがNVIDIAの考え。
もっとも実際にはそれほど性能は変わらないようだよ。
現実のカーネルコードの多くではこんなベンチと違ってメモリ帯域が足を引っ張るから、
当然C2050実効性能は理論性能遠く及ばない。
一方、GTX480は単位時間当たりの演算回数に対して、メモリ帯域がC2050よりも大きい分、
メモリ帯域による実効性能減が起きにくい。
結果として、カリカリにチューンしない限りは、GTX480とC2050の倍精度演算の
実効性能はそれほど大きくは変わらないようだ。
286:デフォルトの名無しさん
10/06/28 07:18:45
何が言いたいかというと、C2050と比べるとGTX480はそこそこお買い得だってことね。
ECC、メモリ容量、倍精度演算の潜在能力とメーカー出荷時のチェック品質が値段の違いなのかな。
287:デフォルトの名無しさん
10/06/28 07:23:59
>>283
>>281だけどそう言いたかった。OpenGLが速いQuadroみたいな感じなのかね？
Telsa C20xxのアホみたいな値段見てると・・・
288:デフォルトの名無しさん
10/06/28 07:46:15
CUDAクラスタを出荷してる会社の人から聞いた話だけど、
ごくまれにGTX系でCUDAカーネルを実行すると、間違った結果しか返さないボードがあるらしい。
他のベンチマークは普通に完走するし、ちゃんとディスプレイ出力もできて、
正常としか判断しようがないボードでも間違うのがあるそうだ。
原因は不明だが、不良と思われるそのボードを交換すると元気に動くらしいから、
やはり一見正常な不良品なんだろうといってたよ。
その点Teslaは強いといってたよ。検査も一般ボードよりずっと厳しい分、
おかしな計算結果を返すボードはまだみたことないといってた。
GTX480でも十分性能はいいけれど、Teslaは一枚は持っておくといざという時に安心らしい。
289:デフォルトの名無しさん
10/06/28 08:32:41
>>288
Tesla 1個の値段でGTXが何枚買えるかということを考えると
GTXを複数枚買って選別してしまえばよいということに・・・・・。
不良品の確率はベンダー依存だけど４枚買ったくらいではよほど引きが強くないと不良品を引けない。
まぁ、Teslaは仲間内で誰かが１個持っておけばいいやくらいの感じかな。
290:デフォルトの名無しさん
10/06/28 08:45:16
>>285
C2050とGTX 480は同じFermiアーキテクチャではあっても違う構造になっている。
倍精度周りでは480に対しては285と同じ実装方法でもうまくいくが、
C2050の場合は違う実装方法にしないと性能が出ないことがある。
C2050は全世代での不満としてで上がっていた安定性、及び倍精度性能の需要に応えるもので
その代わりに単精度計算では旧世代と比べて性能向上が少ない。
性能面では倍精度演算をよほど酷使する計算以外ではGTX480の方が速い。
メモリバンド幅がネックになる演算においてはC2050はGTX285にすら劣る。
291:デフォルトの名無しさん
10/06/28 08:46:52
＞C2050とGTX 480は同じFermiアーキテクチャではあっても違う構造になっている。
釣れますか？
292:デフォルトの名無しさん
10/06/28 09:06:02
>>289
たしかにその通りなんだけど、ボードの不良と結論付けるまで、
コードと何時間も格闘したり、ホストのメモリや電源など検証するの大変だよ。
小規模で4枚だけならなんとか頑張るとしても、大規模にやる場合は、
時間を金で買う意味でもTeslaをおすすめするってことだと思うよ。
ちゃんとコードが走るかはTeslaがあれば検証できる意味でも、
一枚はあるといいとのアドバイスね。
293:デフォルトの名無しさん
10/06/28 09:17:20
>>290
>C2050とGTX 480は同じFermiアーキテクチャではあっても違う構造になっている。
同じチップなんだから、当然同じ構造でしょ。ボードデザインは違うけど。
>メモリバンド幅がネックになる演算においてはC2050はGTX285にすら劣る。
C2050のグローバルメモリ帯域はGTX285とほとんど同じで146GB/sくらいのはず。
しかもFermi世代には、一次キャッシュと二次キャッシュがあるから多くの場合で
GTX285より高速にメモリアクセスできる。
適当言い過ぎ。
294:デフォルトの名無しさん
10/06/28 09:25:26
3.1でたね。 ・Support for printf() in device code なにげに便利かも。
295:デフォルトの名無しさん
10/06/28 10:45:13
>>294
これでもう、エミュレータを使う理由は激減するね。
>>292
だからそれをNVIDIAの営業某氏は「アキバ的発想はやめましょう」と揶揄するのだよね。
でも、10万円で2枚購入できるか1枚も購入できないかの差は法人ユーザには大きいだろ。
296:デフォルトの名無しさん
10/06/28 10:54:08
GeForce GTX 285Mが搭載されているPC買ったんですが、
URLﾘﾝｸ(developer.nvidia.com)
にあるノートブック用のドライバを落とせばいいんでしょうか？
297:デフォルトの名無しさん
10/06/28 11:27:48
>>295
むしろ法人こそTeslaボードのほうがいいんじゃないの？
Teslaはちゃんと使えれば、十分に値段分の価値があるよ。
個人ユーザーはTeslaがきついのはわかるけど、
法人はそれ使って値段以上の利益あげればいいわけだし。
もっともそれが可能かどうかは、試してみないとわからないのが困り者だけど・・・。
298:デフォルトの名無しさん
10/06/28 11:33:28
>>296
そうだよ
299:デフォルトの名無しさん
10/06/28 11:35:47
>>297
だから、予算がついちゃえばTeslaでいいんだけどね。
300:デフォルトの名無しさん
10/06/28 12:41:51
>>293
とりあえず実際に使ってみてから言ってください。
301:デフォルトの名無しさん
10/06/28 13:01:29
>>292
仕分けの時に有名になった長崎大だっけ
あそこが最初にやったのが不良品を検出するプログラムの作成でしょ。
配ってなかったっけ。
302:デフォルトの名無しさん
10/06/28 13:21:43
>>300
何を言いたいのか全く分からない。
C2050とGTX470はもってるが、GTX285はないよ。
C2050とGTX470では、倍精度演算の実効性能はそれほど違わない。
少なくとも俺のコードではね。
実際に使ったら何がわかるのさ。
303:デフォルトの名無しさん
10/06/28 13:23:26
>>302
クレクレ君みたいで申し訳ないが、サンプルプログラムのnbodyのベンチ結果を貰えない？
nbody -benchmark
と
nbody -benchmark -fp64
の２つで
304:デフォルトの名無しさん
10/06/28 13:30:51
ほれ
Run "nbody -benchmark [-n=<numBodies>]" to measure perfomance.
-fullscreen (run n-body simulation in fullscreen mode)
-fp64 (use double precision floating point values for simulation)
> Windowed mode
> Simulation data stored in video memory
> Single precision floating point simulation
> Compute 2.0 CUDA device: [Tesla C2050]
14336 bodies, total time for 10 iterations: 91.924 ms
= 22.358 billion interactions per second
= 447.154 single-precision GFLOP/s at 20 flops per interaction
Run "nbody -benchmark [-n=<numBodies>]" to measure perfomance.
-fullscreen (run n-body simulation in fullscreen mode)
-fp64 (use double precision floating point values for simulation)
> Windowed mode
> Simulation data stored in video memory
> Double precision floating point simulation
> Compute 2.0 CUDA device: [Tesla C2050]
14336 bodies, total time for 10 iterations: 328.314 ms
= 6.260 billion interactions per second
= 187.796 double-precision GFLOP/s at 30 flops per interaction
305:デフォルトの名無しさん
10/06/28 13:32:53
GTX470は家に戻らないといけない。
環境はWin7x64で、C2050でディスプレイ出力してる。
306:デフォルトの名無しさん
10/06/28 13:37:37
㌧
自分でGTX470でベンチ取ったら
516.879 GFlops@14336bodies
095.932 GFlops@14336bodies (-fp64)
だったんだが倍精度速くなってる・・ような・・・
307:デフォルトの名無しさん
10/06/28 13:49:17
FFTは使ってませんか。
倍精度FFTはGTX285からGTX480で倍速いんですがC2050はどうでしょう。
308:デフォルトの名無しさん
10/06/28 13:50:49
>>302
470のメモリバンド幅は133GB/s程度なのでC2050との差は小さいだろう。
480だと177GB/sもあるわけで。
309:デフォルトの名無しさん
10/06/28 13:52:57
>>307
FFTはデータサイズ次第だってば・・・。
まぁ、480のほうが速い場合もC2050のほうが速い場合もあります。
310:デフォルトの名無しさん
10/06/28 13:56:57
>>306
↓480での結果
[単精度]
15360 bodies, total time for 10 iterations: 73.826 ms
= 31.957 billion interactions per second
= 639.149 single-precision GFLOP/s at 20 flops per interaction
[倍精度]
15360 bodies, total time for 10 iterations: 595.640 ms
= 3.961 billion interactions per second
= 118.828 double-precision GFLOP/s at 30 flops per interaction
311:デフォルトの名無しさん
10/06/28 14:00:28
単精度では480が速いのに異論はない。
倍精度でもよっぽどチューンしない限り、480や470と
C2050のそれほど性能は変わらないよ。
そりゃ完全に同じじゃないけどさ。
312:デフォルトの名無しさん
10/06/28 15:16:14
ちょっと教えてほしいのですが、これで最後にundefined referenceになるのはなぜ？CentOS5.4 x86_64です
$ cat hello.cu
#include <stdio.h>
#include "hellonvcc.h"
int main(void){
int i1=0;
int i2=3;
printf("%d", nvccfunc(i1, i2));
return 0;
}
$ cat hellonvcc.c
#include "hellonvcc.h"
int nvccfunc(int i1, int i2){
return i1+i2;
}
$ cat hellonvcc.h
int nvccfunc(int i1, int i2);
$ gcc hellonvcc.c -o hellonvcc.o -c
$ nvcc hello.cu -o hello.o -c
$ nvcc hello.o hellonvcc.o
undefined reference to `nvccfunc(int, int)'
313:デフォルトの名無しさん
10/06/28 15:38:36
>>312
*.cuは実はc++としてnvccでコンパイルされるので、*.cから呼び出すためにはextern "C"にする必要があります。
つーか、c++拒否症でもなければhellonvcc.cをhellonvcc.cppに改名してc++としてコーディングした方が楽かもしれません。
314:デフォルトの名無しさん
10/06/28 15:46:11
ていうか、別にCUDA関係ないじゃん。
シンボル修飾についてちゃんと勉強したほうがいいよ。
315:デフォルトの名無しさん
10/06/28 15:48:35
メモリがネックにならないとき
Tesla/Geforceの単精度：Teslaの倍精度：Geforceの倍精度 = 4 : 2 : 1
メモリがネックになるとき
Tesla/Geforceの単精度：Teslaの倍精度：Geforceの倍精度 = 2 : 1 : 1
という感じだね
316:デフォルトの名無しさん
10/06/28 16:37:51
>>313
できました。どうもありがとうございます。
nvccはC++でコンパイルされるのでしたか。知りませんでした。勉強になります
ちなみに、Cで書かれたライブラリを呼び出したかったのですが、その場合は
ラッパを作るしかなさそうですね。
317:デフォルトの名無しさん
10/06/28 16:50:52
>>316
仮にcのライブラリのインクルードファイルが foo.h ならば、
extern "C" {
#include "foo.h"
}
でOK。
318:デフォルトの名無しさん
10/06/28 16:51:04
呼び出す側がCで呼び出される側Cならば問題ない。
呼び出す側がC++で呼び出される側がCならばextern "C"←今の状況はこれ
だから、呼び出されるライブラリがCで書かれてるならラッパはいらない。
というかまともなライブラリならヘッダに、適切なマクロが書いてあるんじゃないかな。
319:デフォルトの名無しさん
10/06/28 17:08:46
>>317
>>318
どうもありがとうございます。
危うく大変な作業になるところでした。
自作ライブラリなので「まともなライブラリのヘッダ」にはなっていなかったようです。
いじりたくないので317さんの方法で対応します
320:デフォルトの名無しさん
10/06/28 17:22:17
自作なら、>317の仮定でこうしてしまう方がいいかもね。
--
#if ! defined FOO_H
#define FOO_H
#if defined __cplusplus
extern "C" {
#endif
// 従来のfoo.hの中身
#ifdef __cplusplus
}
#endif // c++
#endif // FOO_H
--
一番外側にインクルードガード、その内側にC++の名前空間対策。
cuda.h辺りを参照。
321:デフォルトの名無しさん
10/06/28 19:23:46
このスレ住人は親切すねー
322:デフォルトの名無しさん
10/06/28 21:23:55
実習室でやる意味ないじゃん
323:デフォルトの名無しさん
10/06/29 01:31:18
みんな優しいなあ・・・・。
でもCUDAを始めるにはまだ早いのではないかね？
324:デフォルトの名無しさん
10/06/29 01:45:20
3.1にしたら、CSのOceanが初めて動作する様になった。
simpleparticlesのロード時間もほぼ一瞬にまで高速化されたけど、
代わりにn-bodyのロード時間がクソ長くなった。
CUDAやOpenCLと違ってCSのサンプルがいまいち安定しないのは何でだろ。
325:デフォルトの名無しさん
10/06/29 02:44:20
>>323
たしかにこの質問主はCUDA始めるにはちょいと早すぎるかもね。
質問内容がそもそもCUDA以前の内容だし。
326:デフォルトの名無しさん
10/06/29 11:57:31
>>324
環境はWindowsかい？
うちはWindowsだけど、たしかにDirectComputeのOceanは初めて見た気がする。
確かにDirectComputeのN-Bodyは起動が遅いね。
よく考えたら、DirectComputeがあるのはWindowsだけかな？
327:デフォルトの名無しさん
10/06/29 13:42:47
DirectComputeはDirectX系なのでWindowsしか無いのでは？
しかも今のところWindows7限定（´・ω・｀）
328:デフォルトの名無しさん
10/06/29 14:34:59
Vista
329:デフォルトの名無しさん
10/06/29 14:35:55
あ、Vistaさんもいけたんですね・・・すんません
330:デフォルトの名無しさん
10/06/29 15:40:39
くだもCLもwindowsだとリモート（リモートデスクトップ、ssh)では使えないわけだが、
DirectComputeなら使えるとかいうオチでもある？
331:デフォルトの名無しさん
10/06/29 22:57:17
ブブブーブォォオブオーブブォーブーブー
ブォ ／￣＼ーブーブーブブーンブォーオ
ンー| ＾o＾|∩==＜! プォープォーブブー
ブォ ＼＿／| | ブーブブーブォーーー
ォー ＿| |__| | ブォーーブブブープォー
ー| ＿| ォーブーーーブブォーブ
ブォ| | | ブォーブーブーブーンブー
332:デフォルトの名無しさん
10/06/29 22:59:57
リモートしたことないけどDirectComputeはDirect3Dの一部なのでDirect3Dがリモートできれば動く
333:デフォルトの名無しさん
10/06/30 11:56:37
cuda profilerで分岐数がわかりますが、
プロファイラで分岐数というのはどうやって数えているのですか？
1スレッドあたりの分岐数にしては、多すぎる結果が出るのですが・・・orz
334:デフォルトの名無しさん
10/06/30 13:26:58
PTXを見てbraを探せばいいんじゃない？
335:デフォルトの名無しさん
10/06/30 17:43:43
>>333
SM一個あたりか、TPC一個あたりの総計ではなかろうか。
336:sage
10/07/01 22:07:13
レスありがとうございます。
>>334
とりあえずptx見て数えてみますｗ
>>335
branchは1スレッドあたりの分岐数と表記があったのですが、
値が6桁なので、やはり総数なんでしょうかね？
ちなみに、CUDAZONEの質問掲示板のような所で
if～elseが一つ組み込まれているプログラムで
64スレッド走らせると、プロファイラではbranch=4となったという書き込みが
あるのですが、これは半ワープ４個（＝64÷16）が一つの分岐をしたから
全部で４回分岐ということなんでしょうか？
それ以外に４という数が導き出せないのですが＞＜
337:デフォルトの名無しさん
10/07/01 22:35:15
>>336
多分Warp単位だと思います。
if (C) A else B 文の場合、以下のようになりますが、
各Warpで分岐する方向がスレッド毎に異なれば（両方に分岐すれば）
２個×２Warpの合計４個とカウントされるかもしれません。
if (C) goto L0
do B
goto L1
L0:
do A
L1:
end.
338:デフォルトの名無しさん
10/07/02 01:24:16
>>336
CUDAのプロファイラは、一つのSMをターゲットにして動作している。
またカウンタの値は1スレッドごとではなく1ワープごとに増える。
詳しくは$CUDA_HOME/doc/CUDA_Profiler_3.0.txtを嫁
64スレッドの件は、branchカウンタはbranch instructionの数を数えるけど、ptx的にはconditional jumpもunconditional jumpも同じなので一緒に数えてるんじゃないかな
ようは、>337のgoto文を一つのbranchとして数える。
@C bra L0
B
bra.uni L1
L0:
A
L1:
みたいな
339:デフォルトの名無しさん
10/07/02 22:08:28
9800GTにするか、GT240にするか迷ってます。
性能は9800GTの方がいいので、そちらにしたいのですが、
プログラミングが趣味なので、CUDAのプログラミングを
する予定です
CUDAを考えたらどの程度違いますか？
全然違うんでしょうか？
340:デフォルトの名無しさん
10/07/02 23:28:57
日
本
語
で
お
k
341:デフォルトの名無しさん
10/07/02 23:41:05
GTX260以下はウンコ
もっと言えばFermiじゃなきゃ今から買う価値無し
342:デフォルトの名無しさん
10/07/03 02:01:16
>>340
コイツは日本語書けるのに日本語読めないのかｗ
プログラミングする前に国語の勉強した方が良いぞ
343:デフォルトの名無しさん
10/07/03 02:48:12
日本語がどうとうかはどうでもいいが、
GT9800とGT240を比べている時点でもう
「好きにしろ」
って感じだな。コアレッシングを気にするなら、
後者の方がいいのかな？GT240ってG200系？
344:デフォルトの名無しさん
10/07/03 05:19:24
一応GT200系。
GPGPUやるなら多分240の方がいいだろう。
345:デフォルトの名無しさん
10/07/03 14:36:26
gt240はcc1.2だからどっちにしろ倍精度は使えないね。
346:デフォルトの名無しさん
10/07/06 00:34:22
なんだこれ・・・
URLﾘﾝｸ(pc.watch.impress.co.jp)
347:デフォルトの名無しさん
10/07/06 01:15:34
ようはFermi終了のお知らせ?
Configurable CacheがサポートされてなかったらFermi用のコードなんかHPC向けしか書かないぞ
348:デフォルトの名無しさん
10/07/06 08:06:37
GTX260と比較してプログラムの実行速度はどうなりそうなのかしら
349:デフォルトの名無しさん
10/07/06 09:11:08
例えGT200がベースでも、AMD方式ではなくFermiのテッセレーター機構をぶち込むなら
演算器大増量に変わりは無いから速くはなるだろう。
ただ、純粋なFermiベースに比べればどうしても落ちるわな。
まあGeforceでもCUDA爆速では10倍の値段で売ってるTeslaの立場が無いし、
GPGPU用にトランジスタ増量したらゲーム性能大して向上してない割に爆熱大食らいになって
大多数の一般人には総スカン食らってるんだから、妥当っちゃ妥当なんでないかい。
350:デフォルトの名無しさん
10/07/06 12:10:30
値段や性能はGTX260と同じだとして消費電力がGTX480の６６％じゃGTX260より悪そうでこまる。
351:デフォルトの名無しさん
10/07/06 19:32:07
cc2.0で言語仕様が大幅に変わっているから、
460はcc1.4になるのかな。
352:デフォルトの名無しさん
10/07/07 07:21:03
倍精度切られてそうだけどね460
これで2分化が進めば、トップエンドはもっと思い切って汎用に振れるな
353:デフォルトの名無しさん
10/07/07 21:42:32
倍精度ないのかなあ。BOINCとかでも倍精度必要だから売るために載せて欲しい。
354:デフォルトの名無しさん
10/07/07 21:43:21
倍精度サポートはさすがにあるでしょ
355:デフォルトの名無しさん
10/07/07 21:52:10
おいおい、お絵かき特化のために旧コア使うってのに
お絵かきに関係のない倍精度なんて積むのかね
GT200でさえ、1:8の申し訳程度のものだったのに
これもTeslaのためだけだろ
356:デフォルトの名無しさん
10/07/08 02:11:08
別に倍精度なんか使わないから良いや
GTX465も有るんだから構わないだろ
357:デフォルトの名無しさん
10/07/08 06:02:49
これは長期的に見て良くないね。
もはやNVIDIAはHPCを諦めてきたのかな？
358:デフォルトの名無しさん
10/07/08 07:05:13
ね？
teslaって知ってる？
359:デフォルトの名無しさん
10/07/08 09:07:37
あのさ、Teslaはわかるけど、今まではGeForceとTeslaを共用してきたから開発費が押さえられていたわけで、
Teslaのみを今後開発して行くには、果たしてどれだけ開発費が回収できるかによるでしょ。
今のNVIDIAに両方を開発していく体力があるのかどうか疑問だけどね。
360:やんやん ◆yanyan72E.
10/07/08 13:45:08
TSUBAME2もできたことだし。
もうHPC事業はHPC事業独立でできるんでしょう。
361:デフォルトの名無しさん
10/07/08 14:29:44
長期的も何も今までもローエンドは倍精度サポートして来なかったじゃん
今まで通りGTX580を作る時はリッチに作ってGTX560を作る時は削れば良いだろ
362:デフォルトの名無しさん
10/07/08 18:57:55
TSUBAME1に鳴り物入りで採用されたCSX600は無視ですか、そうですか・・・
363:デフォルトの名無しさん
10/07/09 01:05:40
本気で広めたいなら下位モデルでも倍精度サポートしろよ
特定のハードでしか使えないとかまじバカじゃないの
364:デフォルトの名無しさん
10/07/09 01:22:51
チップ作るのって、簡単なことじゃないぞ。
なんか勘違いしている奴がいるが、
#ifdef GTX580
単精度サポート
倍精度サポート
#else
単精度サポート
#endif
みたいにコンパイルし直せば、出来ると思っているのか？
コンパイルのたびに数億飛んでいくってかんじなんだぞ。
デバッグの行程もたくさんあるしな。
365:デフォルトの名無しさん
10/07/09 02:11:20
うっかりF7押しただけで数億とな
F5押したらどうなっちゃうの
366:デフォルトの名無しさん
10/07/09 07:22:54
お絵かきに倍精度なんているの？ｗ
367:デフォルトの名無しさん
10/07/09 07:24:23
URLﾘﾝｸ(www.adrenaline.com.br)
コアはfermi系っぽいがキャッシュはないようだ
アンコアは前世代
368:デフォルトの名無しさん
10/07/12 15:55:32
倍精度きましたよ。１：１２とさらに削ってきた。
369:デフォルトの名無しさん
10/07/12 16:20:21
SMのSP数を変えてくるのはさすがに予想外だった
370:デフォルトの名無しさん
10/07/12 17:41:12
>>368
8%とは酷すぎ。25%がましな時代だったとはorz
371:デフォルトの名無しさん
10/07/12 18:14:02
25%だったことはないが・・・。
480は12.5%。
372:デフォルトの名無しさん
10/07/12 18:35:14
なんか廉価版tesla出そうな気もする
10万以下で
373:デフォルトの名無しさん
10/07/12 19:28:00
GTX460買ってきた。自分の倍精度プログラムはGTX260とまったく同じ性能だった。いいんじゃない。
374:デフォルトの名無しさん
10/07/12 19:56:24
URLﾘﾝｸ(pc.watch.impress.co.jp)
＞例えば、チップ全体での実行命令数を計算すると、GF100が480 Instruction/Core Clock(15 SM/2 IPC時)であるのに対して、
＞GF104は理論上のピークが448 Inst/Clk(7 SM/4 IPC時)で、実効が336 Inst/Clk(7 SM/3 IPC時)となる。
＞GF100は理論値と実効値の乖離が少ないはずだが、それでもGF104の効率はダイサイズ比を考えると相対的に高い。
それでもCUDA的には465の方が上なのか
URLﾘﾝｸ(img5.pcpop.com)
URLﾘﾝｸ(img5.pcpop.com)
URLﾘﾝｸ(img5.pcpop.com)
375:デフォルトの名無しさん
10/07/12 21:36:12
相対的にload/storeとレジスタ、キャッシュが少ない
376:デフォルトの名無しさん
10/07/12 22:41:09
CC2.1・・・
377:デフォルトの名無しさん
10/07/12 23:43:43
どの命令の組み合わせなら4issue出来るんだろうか
378:デフォルトの名無しさん
10/07/13 00:03:57
warp sizeはどうなるの? まさか48?
379:デフォルトの名無しさん
10/07/13 00:04:52
どう考えても32だが
380:デフォルトの名無しさん
10/07/13 13:03:45
>>375
同時に実行される命令が2->4になるだけなので、相対的な容量はあまり問題にならないような。
むしろメモリバンド幅が足りない。
381:デフォルトの名無しさん
10/07/13 15:59:40
460だとメモリ帯域も細いから、倍精度だともはやCPUで走らすのと大差無いんじゃね？
382:デフォルトの名無しさん
10/07/13 17:49:56
256bit 3208MHzの465より256bit 3600MHzの460の方がメモリ帯域は大きいはずだが
383:デフォルトの名無しさん
10/07/13 19:42:19
アイドルの消費電力が低いし、お試しに買うのにいいですよ。
384:デフォルトの名無しさん
10/07/16 06:15:22
URLﾘﾝｸ(pc.watch.impress.co.jp)
>データパスなどに制約があるため、倍精度演算命令を発行する際には、他の命令を発行できないためピーク性能が制約されている。
>その制約を除けば、GF100 CUDAコアはフルの倍精度演算能力を持つため、演算ユニット自体の実装コストはかなり高いと推測される。
SP自体は単体で倍精度扱えたのか
385:デフォルトの名無しさん
10/07/18 12:28:02
CUDAユーティリティのtimerを使ってみようとしたのですが、
error: cutil.h: そのようなファイルやディレクトリはありません
とでてしまいました。（この時点でパスの設定が間違ってるのでしょうか？）
そこでcutil.hを検索して.cuと同じフォルダにいれ（もしくはnvcc -Iでcutil.hがあった場所を入力し）
たのですが、今度は
/tmp/tmpxft_000052ed_00000000-11_timer.o: In function `main':
tmpxft_000052ed_00000000-10_timer.ii:(.text+0x413): undefined reference to `cutCreateTimer'
tmpxft_000052ed_00000000-10_timer.ii:(.text+0x41b): undefined reference to `cutStartTimer'
tmpxft_000052ed_00000000-10_timer.ii:(.text+0x48c): undefined reference to `cutStopTimer'
collect2: ld はステータス 1 で終了しました
となってしまいました。
ちゃんとinculdeできていないのでしょうか。
cutilを使わない、簡単なプログラムはコンパイルして実行することができています。
どんな問題が考えられるでしょうか。アホな質問ですがよろしくおねがいします。
386:デフォルトの名無しさん
10/07/18 12:32:31
>>385
cutilがビルド済みなら、-L <libcutil.aがあるディレクトリ> -lcutil
ビルドさえしてないなら、NVIDIA_CUDA_SDK/common辺りでビルド。
手元に環境がないので詳細は割愛。
387:385
10/07/18 13:13:36
>>386
無事に実行できました！
即レスありがとうございます。
なにが問題だったのでしょうか
libcutil.aにcutilが含まれてるということなのでしょうか
388:デフォルトの名無しさん
10/07/18 13:45:08
>>387
まぁ、簡単に言えばそういうこと。nm libcutil.a|gerp T してみると知見が得られるかもねw
389:デフォルトの名無しさん
10/07/19 18:58:30
このスレの住人は本当に出来た人が多いね。
と思ってたけど、ここはそういうスレなんだね。テンプレをあらためて見て気が付いたよ。
何となくだけど、385さんは結構前に自作ライブラリがリンク出来ないと質問してた人かな？
プログラムにはあまり慣れて無いけど、素直で向上心を感じます。多分学生さんなんですかね？
別に答える必要は無いけど、そうだと仮定してアドバイスすると、日経ＢＰから出てるプログラムはなぜ動くのかって本の８章あたりを読むと勉強になるかもしれない。
手元には無いので保障は出来ないけど、立ち読みした感じだと分かりやすく為になりそうなことが書いてある本だったよ。ひょっとするとシリーズの違う本の方がドンピシャの内容かもしれない。気が向いたら本屋で内容をチェックしてみてね。