2ちゃんねる ★スマホ版★ ■掲示板に戻る■ 全部 1- 最新50  

■ このスレッドは過去ログ倉庫に格納されています

【GPGPU】くだすれCUDAスレ pert3【NVIDIA】

1 :デフォルトの名無しさん:2010/04/18(日) 19:48:30
このスレッドは、他のスレッドでは書き込めない超低レベル、
もしくは質問者自身何が何だが分からない質問を勇気を持って書き込むスレッドです。
CUDA使いが優しくコメントを返しますが、
お礼はCUDAの布教と初心者の救済をお願いします。

CUDA・HomePage
http://www.nvidia.com/cuda

関連スレ
GPUで汎用コンピューティングを行うスレ
http://pc11.2ch.net/test/read.cgi/tech/1167989627/
GPGPU#3
http://pc12.2ch.net/test/read.cgi/tech/1237630694/

前スレ
【GPGPU】くだすれCUDAスレ【NVIDIA】
http://pc12.2ch.net/test/read.cgi/tech/1206152032/
【GPGPU】くだすれCUDAスレ pert2【NVIDIA】
http://pc12.2ch.net/test/read.cgi/tech/1254997777/

2 :デフォルトの名無しさん:2010/04/18(日) 19:49:11
関連サイト
CUDA Zone
http://www.nvidia.co.jp/object/cuda_home_jp.html
フォーラム
http://forum.nvidia.co.jp/EokpControl?&event=TE0001

CUDAに触れてみる
http://chihara.naist.jp/people/STAFF/imura/computer/OpenGL/cuda1/disp_content

cudaさわってみた
http://gpgpu.jp/article/61432191.html

CUDA のインストール
http://blog.goo.ne.jp/sdpaninf/e/9533f75438b670a174af345f4a33bd51

NVIDIAの「GeForce 8800 GT(G92)」と次に控える64-bit GPUアーキテクチャ
http://pc.watch.impress.co.jp/docs/2007/1031/kaigai398.htm

CUDAを使う
http://tech.ckme.co.jp/cuda.shtml

NVIDIA CUDAを弄ってみた その2
http://dvd-r.sblo.jp/article/10422960.html

CUDAベンチ
http://wataco.air-nifty.com/syacho/2008/02/cuda_2044.html

KNOPPIX for CUDA
http://www.yasuoka.mech.keio.ac.jp/cuda/

3 :デフォルトの名無しさん:2010/04/18(日) 21:42:11
いちおつ

4 :デフォルトの名無しさん:2010/04/19(月) 00:54:58
一乙、前スレの、魔二来vsTeslaのコストパフォーマンス比較面白いな
もっとやってくれ

5 :デフォルトの名無しさん:2010/04/19(月) 05:45:00
>>4
もう結論でたよ

理論ピーク性能はTeslaの勝ち
大部分のプログラムでの実行性能はマニクールの勝ち
プログラミングのしやすさではマニクールが圧勝

6 :デフォルトの名無しさん:2010/04/19(月) 07:05:25
すげーな
48コア使いこなせるとかw

7 :デフォルトの名無しさん:2010/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 :デフォルトの名無しさん:2010/04/19(月) 07:38:14
>>7
これって、いまいち理解出来ないのだが、誰か説明して。

9 :デフォルトの名無しさん:2010/04/19(月) 09:51:51
>>7
将来が危ぶまれるNVIDIA専用のCUDAと心中するつもりなら
Teslaでもいいかもね

10 :デフォルトの名無しさん:2010/04/19(月) 12:33:45
>>8
たぶん

6168と6176seはマニクールの製品名でc2050はtesla

その下の2つは値段(ドル)と性能(GFLOPS)
値段はメーカーから直接購入した時の価格

CPUは
(総費用)=(価格)X(コア数)
(GFLOPS)=(クロック数)X(コア数)X(CPU数)X(SIMD演算器数)X(1SIMD演算器当たりの倍精度演算数)

GPUは
(総費用)=(価格)
(GFLOPS) =(クロック数)X(コア数)

11 :デフォルトの名無しさん:2010/04/19(月) 12:35:22
>>10
間違えた

×(総費用)=(価格)X(コア数)
○(総費用)=(価格)X(CPU数)



12 :デフォルトの名無しさん:2010/04/19(月) 14:12:24
OpenMPとかえらい簡単に並列化できるしな。
あと何覚えればいいですか?

13 :デフォルトの名無しさん:2010/04/19(月) 15:20:32
普通のスレッド
MPI

14 :デフォルトの名無しさん:2010/04/19(月) 16:22:11
>>10
OpteronってSSEとかが1コアあたり2個ある訳?


15 :デフォルトの名無しさん:2010/04/19(月) 16:24:58
OpenMPもCUDA並に最適化が面倒。
その代わりMPIの知識があれば以外と簡単。

16 :デフォルトの名無しさん:2010/04/19(月) 18:03:05
で、比較したのまだぁ?

17 :デフォルトの名無しさん:2010/04/19(月) 18:05:27
そもそもC2050が出てないのに
比較なんてできるかよ

18 :デフォルトの名無しさん:2010/04/19(月) 18:22:54
そもそもスレチ

19 :デフォルトの名無しさん:2010/04/19(月) 21:40:15
秋葉にオープンしたfaithにCUDAのブースができたね
実体はunitcomの宣伝かな

20 :デフォルトの名無しさん:2010/04/19(月) 22:03:30
GTX480でデモしてくれるならすぐにでも行くが

21 :デフォルトの名無しさん:2010/04/19(月) 22:04:09
長崎大の部品を入れたとか書いてたね。CUDA話のできる店員さんいるのかしら。

22 :デフォルトの名無しさん:2010/04/20(火) 07:05:14
チップセット・・・完全撤退
GPU・・・AMDの大攻勢の前に大苦戦
Tesla・・・CPUの高性能化・マルチコア化の前で微妙に

NVIDIAさんこれからどうするつもりなんでしょうね

23 :デフォルトの名無しさん:2010/04/20(火) 07:18:02
自作板に帰れ

24 :デフォルトの名無しさん:2010/04/20(火) 07:29:44
SSEとかAVXとか覚えるよりも、
CUDAのお作法覚えるほうが簡単な気がしない?おれだけ?

25 :デフォルトの名無しさん:2010/04/20(火) 08:09:46
>>24
でも将来性が最悪だね
遅かれ早かれOenpCLに移行するだろうし

26 :デフォルトの名無しさん:2010/04/20(火) 08:53:34
CUDAとSIMDでは考え方が違うような気が。
個人的にはCUDAは割と大ざっぱに並べる。
SSEとかは細かなパズル。
自分でも何を言っているかわからんが、そういう思いでいる。


27 :デフォルトの名無しさん:2010/04/20(火) 08:57:23
前スレで480特攻すると書いた者だが
倍精度1/4と聞いてキャンセルしてしまった…
ゲーマーだったら失敗覚悟で注文したんだけど

28 :デフォルトの名無しさん:2010/04/20(火) 09:09:12
ZOTACまだリリース前になってますね。キャンセルできてよかったですね。
ケース待ちだった方は購入できたのでしょうか。


29 :デフォルトの名無しさん:2010/04/20(火) 09:20:53
マザーボード破損の書き込みもあるし手を出さなくて正解だろうね

30 :デフォルトの名無しさん:2010/04/21(水) 18:23:51
カーネル中でcoutとかは使えないの?
どこが悪いのか調べるときはたいていcoutで変数を書き出していたのだけど、
CUDAの場合はどうすればいいの?

31 :デフォルトの名無しさん:2010/04/22(木) 02:17:10
cudaMemcpy()でCPUに送る。

32 :デフォルトの名無しさん:2010/04/22(木) 05:13:57
カーネルの途中で表示したいんだけど

33 :デフォルトの名無しさん:2010/04/22(木) 08:07:33
どんな手法にしろCPU介さなかったら表示はできんよ

34 :デフォルトの名無しさん:2010/04/22(木) 08:32:27
http://gpu.fixstars.com/index.php/Cuda-gdbを使う

自分はコーディングレビューしかしないけど

35 :デフォルトの名無しさん:2010/04/22(木) 09:45:25
結果とは別にデバイス側に確保しておいて、
あとで見るのはだめなの?

36 :デフォルトの名無しさん:2010/04/22(木) 10:12:38
>>32
・deviceemuでなんとか頑張る
・可能な限り詳細な出力をメモリに吐き出しておいて、ホストに戻ってから(転送して)解析する

それ以前に、そのレベルだとcudaの細かいノウハウを掴んでないだろ。
然るべき業者に委託したほうがよくないか?

37 :デフォルトの名無しさん:2010/04/22(木) 20:00:42
なぜ仕事前提?

38 :デフォルトの名無しさん:2010/04/22(木) 22:26:01
>>30
とゆうか、そうゆう発想する事態C++もCUDAも全然理解できてないよねw

39 :デフォルトの名無しさん:2010/04/23(金) 01:36:42
CPUで考えると、「レジスタの計算の途中をCoutで見ることは出来ないか?」
と聞いているようなものだと思うのだが。
それを自前でやるにはレジスタの内容を一度メモリにコピーしてcoutの標準出力で表示させることになる。
それと一緒だと思うよ。


40 :デフォルトの名無しさん:2010/04/23(金) 02:53:45
device emuで解決ジャン

41 :デフォルトの名無しさん:2010/04/23(金) 03:31:06
待てればね。

42 :デフォルトの名無しさん:2010/04/23(金) 03:47:59
エミュレーションモードってなくなるんだろ

43 :デフォルトの名無しさん:2010/04/23(金) 05:53:13
すでに実機があるからなあ。

44 :デフォルトの名無しさん:2010/04/23(金) 06:50:27
ダンプ用の領域を用意しておいて実行中に書き込んで後で見る
しかないんじゃね?

45 :デフォルトの名無しさん:2010/04/23(金) 07:05:41
ubuntu 10.04でcudaが使えるようになるのはいつごろですか

46 :デフォルトの名無しさん:2010/04/23(金) 08:28:37
GPUを搭載していないノートPCでエミュレーションモードを使ってたのに
なくなるってマジ?

47 :デフォルトの名無しさん:2010/04/23(金) 11:36:52
マジ

48 :デフォルトの名無しさん:2010/04/23(金) 11:42:20
リモートデスクトップでCUDA使えればいいのになー。

49 :デフォルトの名無しさん:2010/04/25(日) 16:17:32
リナカフェいったひといないすか

50 :デフォルトの名無しさん:2010/04/25(日) 16:22:48
サトリナのカフェなんてあるの?

51 :デフォルトの名無しさん:2010/04/26(月) 10:35:43
このスレもしぼんだなあ
GTX480が倍精度外れた(らしい)のが効いているんだろうな

52 :デフォルトの名無しさん:2010/04/26(月) 11:17:15
倍精度演算ユニットは少なくなっているだけで無くなったわけではないのだが、
そもそも480が入手できないことの方が原因ではないかと。


53 :デフォルトの名無しさん:2010/04/26(月) 15:43:31
結局口だけの貧乏人しか居なかったって事だなw

54 :デフォルトの名無しさん:2010/04/26(月) 17:33:22
おまえもな

55 :デフォルトの名無しさん:2010/04/26(月) 19:14:13
http://www.amazon.co.jp/dp/479802578X/
これIPAの岡ちゃんが書いたの?
いつの間にかネットマイスターってところに転勤したのかな

56 :デフォルトの名無しさん:2010/04/26(月) 19:24:07
もう入門書もでてるから、新刊はいらないな。
需要はあるのだろうか・・・。

57 :デフォルトの名無しさん:2010/04/26(月) 19:28:27
つーか、ネットマイスター代表だってよ。
同名をtwitterで検索したらttp://twitter.com/kenmoi0715 が引っ掛かったけど……

58 :デフォルトの名無しさん:2010/04/27(火) 01:07:50
流石にロリコンばらされたら職場には居にくいか……
アプリも割り放題だったしな。
たくましく生きているようで何よりww


59 :デフォルトの名無しさん:2010/04/27(火) 01:20:01
この本買ったけど、ちょっと入門過ぎる気がするなあ。
「高速」という言葉にてっきり「最適化」ということを連想してしまった。
環境設定とかで半分以上割いているよ。
まあ、CPUコードと両方書いてあって、Webからダウンロード出来るから、
最初にでた本よりかはいいかもしれない。

60 :デフォルトの名無しさん:2010/04/27(火) 08:02:14
自分で買うほどじゃないかな
図書館に購入希望だしとこう

61 :デフォルトの名無しさん:2010/04/27(火) 21:19:54
本にするほどの蓄積と技術の咀嚼があるとは思えないけどどうなんでしょ。
あるとすればまだ論文レベルなんじゃないか?

逆に、まだ広まってない今が新しいことするチャンスかっ!

62 :デフォルトの名無しさん:2010/04/27(火) 21:27:03
プログラムの入門書なんてこんなもんだろ

63 :デフォルトの名無しさん:2010/04/28(水) 04:21:24
GF100が失敗、GF104はキャッシュ削るみたいでGPGPUは絶望的の現在、今後広まる可能性が残っているかどうか…

64 :デフォルトの名無しさん:2010/04/28(水) 05:19:00
は?

65 :デフォルトの名無しさん:2010/04/28(水) 05:45:52
GPGPUに残されているのは個人では動画エンコードくらいか。
それなりの規模になると、マルチプロセッサのシステムを組んで、
CPUのコードを走らせた方が開発工数から考えると速いからなあ。


66 :デフォルトの名無しさん:2010/04/28(水) 06:14:57
で、おまえは何しにここに来てる訳?

67 :デフォルトの名無しさん:2010/04/28(水) 06:17:08
ばかだな
ただのAMD信者様だろw

68 :デフォルトの名無しさん:2010/04/28(水) 09:19:52
GPGPUに将来性はないだろうね。
牧野もそういっているし。

69 :デフォルトの名無しさん:2010/04/28(水) 09:29:19
コスト面がちょっと微妙かなあ
C10シリーズのようにアカデミックで10万!とかなら個人で使ってみようという流れになるだろうに
GTX480は倍精度1/4にしました!C20シリーズを売るためです!と公言されちゃったし

70 :デフォルトの名無しさん:2010/04/28(水) 09:52:12
>>69
コスト面つか買う気失せることすんなよなー て感じっすよ

71 :デフォルトの名無しさん:2010/04/28(水) 10:00:12
       、--‐冖'⌒ ̄ ̄`ー-、
     /⌒`         三ミヽー-ヘ,_
   __,{ ;;,,             ミミ   i ´Z,
   ゝ   ''〃//,,,      ,,..`ミミ、_ノリ}j; f彡
  _)        〃///, ,;彡'rffッ、ィ彡'ノ从iノ彡
  >';;,,       ノ丿川j !川|;  :.`7ラ公 '>了
 _く彡川f゙ノ'ノノ ノ_ノノノイシノ| }.: '〈八ミ、、;.)
  ヽ.:.:.:.:.:.;=、彡/‐-ニ''_ー<、{_,ノ -一ヾ`~;.;.;)
  く .:.:.:.:.:!ハ.Yイ  ぇ'无テ,`ヽ}}}ィt于 `|ィ"~
   ):.:.:.:.:|.Y }: :!    `二´/' ; |丶ニ  ノノ
    ) :.: ト、リ: :!ヾ:、   丶 ; | ゙  イ:}    逆に考えるんだ
   { .:.: l {: : }  `    ,.__(__,}   /ノ
    ヽ !  `'゙!       ,.,,.`三'゙、,_  /´   「単精度なら4倍」と
    ,/´{  ミ l    /゙,:-…-〜、 ) |
  ,r{   \ ミ  \   `' '≡≡' " ノ        考えるんだ
__ノ  ヽ   \  ヽ\    彡  ,イ_
      \   \ ヽ 丶.     ノ!|ヽ`ヽ、
         \   \ヽ `¨¨¨¨´/ |l ト、 `'ー-、__
            \  `'ー-、  // /:.:.}       `'ー、_
          `、\   /⌒ヽ  /!:.:.|
          `、 \ /ヽLf___ハ/  {
              ′ / ! ヽ

72 :デフォルトの名無しさん:2010/04/28(水) 10:01:53
長崎大のセンセが8800GTで激安スパコン組んだときも、
そのうち2、3割がGPGPU用としてはぶっ壊れていたじゃん。

今度のは、そのままだと半分以上壊れたものを出荷してるから、まずいってんで、
止めてるんでないの。

73 :デフォルトの名無しさん:2010/04/28(水) 12:24:26
牧野もそういっているし。
牧野もそういっているし。
牧野もそういっているし。
牧野もそういっているし。
牧野もそういっているし。




牧野乙

74 :デフォルトの名無しさん:2010/04/28(水) 12:29:43
そもそも、+α要素であるアクセラレータが
終わるとか終わらないとか

おばかさん

75 :デフォルトの名無しさん:2010/04/29(木) 01:52:54
金にならんかったら終わるだろ。

76 :デフォルトの名無しさん:2010/04/29(木) 09:55:55
GPGPUはマルチスレッドな浮動小数点演算に特化したものだからな、使えるシーンはかなり限られてくる

77 :デフォルトの名無しさん:2010/04/29(木) 10:26:40
forループの一部だけを実行する方法を教えてください

78 :デフォルトの名無しさん:2010/04/29(木) 12:24:55
>>55
やっぱりこういう問い合わせが多いんだろうな
著者紹介にIPAに勤務したことはないって書いてあったよ

79 :デフォルトの名無しさん:2010/04/29(木) 12:29:11
>>59
俺も買った。
全く同じ感想。
青木先生の本が理解できる人は全く買う必要なし。
>>55
本の経歴の部分に、「IPAに勤務したことはない。」と不自然に書いてあったよw

80 :デフォルトの名無しさん:2010/04/29(木) 13:44:48
なんだ同名で迷惑している人なのか

81 :デフォルトの名無しさん:2010/04/29(木) 15:21:40
んなわけない。
IPAを除名されたか、退職するときにそれに準じる誓約書を書かされたのだろう。


82 :デフォルトの名無しさん:2010/04/30(金) 02:18:54
証拠は?

83 :デフォルトの名無しさん:2010/04/30(金) 05:35:43
>>79
今それを見てワロタ。
「IPAに勤務したことはない。」
で株式会社ネットマイスターをググってみたけど、
すぐには見つからなかった。
でもこういった本を書く人が流出なんてするのだろうか?

84 :デフォルトの名無しさん:2010/04/30(金) 08:01:21
GF100の板に出ていてみたいだけど、
GTX480でCUDA動作に不具合があるとか?
ドライバ改善待ち、みたい。
GTX480でのベンチマーク結果(単精度、倍精度問わず)が出ない理由はそれだったりするのかも?

また、GTX460が6月に出るとか。
お値段は480/470よりも安いんだろうけど、倍精度しょぼくなった480よりもさらにしょぼいんだろうなあ

いつかは出るであろう、Quadroにかすかに期待するしかないのか…

85 :デフォルトの名無しさん:2010/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 :デフォルトの名無しさん:2010/05/01(土) 15:36:56
>>85
もちろん重要ありますよ。

87 :デフォルトの名無しさん:2010/05/01(土) 16:09:33
【Fermi】GeForceGTX4xx総合Part31【GF100】
http://pc11.2ch.net/test/read.cgi/jisaku/1272318075/
の548以降でGTX480を含むGPUのベンチマークが出ている
LU分解のようです
倍精度に制限かかっているのはどうやら本当みたいだ

88 :デフォルトの名無しさん:2010/05/01(土) 18:13:38
具体的にどう制限かかっているんだろ?
それによってCUDAのコードの書き方が変わってくると思う

89 :デフォルトの名無しさん:2010/05/01(土) 22:13:39
>87
いや倍精度どころか単精度でも全然性能でてないわけだが

90 :デフォルトの名無しさん:2010/05/01(土) 23:17:23
なんかそのベンチ遅くね
285の単精度で16Gflopsとか書いてあるけど

91 :デフォルトの名無しさん:2010/05/01(土) 23:25:48
GTX285と比べる限り、GTX480の倍精度に期待できそうな気配はないなあ
1/4は本当くさいな

92 :デフォルトの名無しさん:2010/05/02(日) 16:32:28
cublasのgemmを使ったベンチをだれかが提供すべきだな
もちろん、俺はやる気はないが

93 :デフォルトの名無しさん:2010/05/03(月) 10:55:57
gemmもやって欲しいけど、gemmはピーク性能しか出てこないから俺はLU分解やFFTの方が多少は信用している。

94 :デフォルトの名無しさん:2010/05/03(月) 12:02:20
というか一番多用しそうなFFTでどれだけ実効性能出るのかが気になる

95 :デフォルトの名無しさん:2010/05/03(月) 16:43:35
>>94
例えばどんなサイズが?


96 :デフォルトの名無しさん:2010/05/04(火) 18:08:02
http://journal.mycom.co.jp/special/2008/cuda/006.html
ここのコードについて質問なんだけど
これはマルチスレッドで実行する関数の内部だよね?
ということはAsやBsのデバイスメモリからシェアードメモリへの転送や、
最後のホストの変数Cへのデータ転送は無駄に全スレッドがやるの?
それとも同じデータの転送については1スレッドがやってる?

97 :デフォルトの名無しさん:2010/05/04(火) 19:47:43
> これはマルチスレッドで実行する関数の内部だよね?
そのとおり。

> AsやBsのデバイスメモリからシェアードメモリへの
全スレッドで一要素ずつ分担して転送してる。無駄にはやってない。
Cのシェアードメモリからデバイスメモリへの転送も同じ。
シェアードメモリ内の As と Bs は各スレッドあたり BLOCK_SIZE 回使用する。

> 最後のホストの変数Cへのデータ転送は
今回のコードには入ってない。ホスト側のコードで、ブロック転送を指示する。


98 :デフォルトの名無しさん:2010/05/04(火) 20:53:45
>>97
thx

99 :デフォルトの名無しさん:2010/05/07(金) 12:16:01
GCC 4.5.0で使えるようにする設定を教えてください

100 :デフォルトの名無しさん:2010/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 :デフォルトの名無しさん:2010/05/07(金) 21:33:16
三つのループにばらしてBLASを3回使えばいいけど、
もっと速い方法があるかな。

102 :デフォルトの名無しさん:2010/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 :デフォルトの名無しさん:2010/05/08(土) 01:11:54
続き・・・

これだけではSM1つしか使わないし、ARRAY_SIZE_Nが大きい場合はSMあたりのスレッド数上限を超えたりする。
なので実際には更に複数ブロックに分割して、複数のSMで実行されるようにしないと動作しなかったり非常に低い性能しか得られない。

G80やG92では256スレッド/ブロック程度が良いことが多いのかな。
でも変数の数によってはレジスタが足りなくなったりするだろうし、しっかり計算して決めるかパラメータを変えながら試行錯誤することになると思う。

お勧めサイトはイリノイ大学の
ttp://courses.ece.illinois.edu/ece498/al/

国内ではGeForceのアーキテクチャについて解説されているページも結構有用というか
ハードのことをある程度知っていないとかなり厳しいかと。

104 :デフォルトの名無しさん:2010/05/08(土) 02:45:33
ついでに言うと、>102ではカーネルの処理時間が短過ぎて呼び出しコストに見合わない。
ARRAY_SIZEをブロックとスレッドに分割したら更に残りをループにするようにしないと効率が悪そう。
それから、bのアクセスも競合するから待ちが入ることになりそう。
こっちについてはループを3回に分けた方がいいかもしれないし、強行した方がいいかもしれないし、
profileと睨めっこする必要があるかもね。

ハードの知識がなくてもptx出力を読めてprofileの意味を理解できればいいのだけれど、
それって結局ハードの知識を得るのと同じことだしね。

105 :デフォルトの名無しさん:2010/05/08(土) 03:14:44
結局nがいくつかによるだけじゃない。
たぶん100はCUDAのくの時もまだ理解していないと思う。


106 :デフォルトの名無しさん:2010/05/08(土) 03:16:30
>>104
あの程度だとコンテキストスイッチやホスト・デバイス間の転送等の
呼び出しコストに見合わないね。

学習用か極端にシンプルにしたコードだと思ってあまり気にしなかったけど。

bのアクセスの競合はブロックに分割してシェアードメモリを使うなら
バンクコンフリクトも起こらず、大丈夫かと思ったけど甘いかな?

ptx読めてprofileの意味を理解するにはアーキテクチャの理解が必要というか
もうその段階はとっくに超えているような気もする・・・

107 :デフォルトの名無しさん:2010/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ではハーフワープあたり無駄なメモリロードが2回走る。
CC2.0ではキャッシュが効くのでARRAY_SIZE_Nのサイズしだいでは問題にならない可能性がある。

108 :デフォルトの名無しさん:2010/05/08(土) 03:45:08
>>107
厳しい突っ込みどうも。
CUDA Coreと呼ばれるようになったのはSMじゃなくてSPだったのですね。
図を見れば一目瞭然ですね・・・うろ覚えで書いてすみませんでした。

コアレスアクセスはまだよく理解していないので勉強してきます・・・

109 :デフォルトの名無しさん:2010/05/08(土) 04:15:14
だから、16人17脚でパン喰い競争やるようなもんだよw

110 :デフォルトの名無しさん:2010/05/08(土) 05:18:45
この表現はわかりやすいね。

111 :デフォルトの名無しさん:2010/05/08(土) 06:18:59
一人重い荷物を持たされると、全体の速度が遅くなる、的なw

112 :デフォルトの名無しさん:2010/05/08(土) 08:19:25
青木本で最も評価するのがあの手書きの絵だww

113 :デフォルトの名無しさん:2010/05/08(土) 09:18:32
青木本にこのスレから引用したと見られる節が幾つかある件w

>>109
>107の例だと、自分のレーンのパンだけじゃなくて左右のパンも食べないといけないんだな。
そりゃ遅くなるわさ。

114 :デフォルトの名無しさん:2010/05/08(土) 11:04:49
コアレスのコアってなんのコアのことなの?

115 :デフォルトの名無しさん:2010/05/08(土) 11:07:02
coalescedなので、
合体したという意味。
合体してメモリにアクセスするって意味だ。

116 :デフォルトの名無しさん:2010/05/08(土) 12:01:16
正直その並びなら、texture使えばキャッシュ使えるじゃん

ただみんな、言うように並列度がどれくらい稼げるか(nに依存)で変わるから実際やってみないとわからないんじゃない??

117 :デフォルトの名無しさん:2010/05/08(土) 12:24:17
↓新明解(旧版)で合体の項を引くなよ

118 :デフォルトの名無しさん:2010/05/08(土) 13:18:40
メモリにアクセスするときに、各スレッドが足並みそろえて
+1ずつのアドレスを見に行かないといかんのです
すると32ワードなりまとめてとってくる。 とってくる間は
別のスレッドが実行されてる。

119 :デフォルトの名無しさん:2010/05/08(土) 14:52:01
>>114
私は「連動した」メモリアクセスと説明している。要は16人17脚なんだけどねw

120 :デフォルトの名無しさん:2010/05/08(土) 22:51:22
おまえら頭いいんだろうな

121 :デフォルトの名無しさん:2010/05/08(土) 22:57:18
どうせ底辺大学生が研究で勉強してるだけでしょ?

122 :デフォルトの名無しさん:2010/05/08(土) 23:11:19
と学歴昆布が申しております

123 :デフォルトの名無しさん:2010/05/09(日) 15:32:57
でもGT200以降だと16人17脚でパン喰い競争って微妙に語弊あるよね。
16人17脚でパン喰い競争じゃない使い方したことないけど。

124 :デフォルトの名無しさん:2010/05/09(日) 16:44:58
>>102のコードは
あるスレッドの遅れが、他のスレッドの足を引っ張る結果になるの?

コード内で参照する領域に依存関係がないのに
他のスレッドに影響するの?
なんだか納得出来ないなあ

125 :デフォルトの名無しさん:2010/05/09(日) 17:42:10
>>102
if使ってる時点でありえないかと。。。

適当に(関数名とか間違ってるかも)書くけど、

int tx = threadIdx.x;
a[tx] = text1D(tx-1) + text1D(tx) + text1D(tx+1);

で、thread数とblock数を最適化しておけば、おkでは??
出力もコアレスだしさ

126 :デフォルトの名無しさん:2010/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 :デフォルトの名無しさん:2010/05/09(日) 20:59:34
>>107

>a[tx] = b[tx] + b[tx+1] + b[tx-1];
前提としているGPUがわからんが、根本的にコアレスメモリアクセスになっていない。
本当にわからんなあ。例えばコアレスアクセスになるにはどう書けばいいの?

つーかCUDAって根本的に間違ってるよなぁ。
C言語で書けることが利点?なのに、ベタにCで書くと
コアレスじゃないとか何とかで遅いというw
じゃあ、遅くないfloatN_tでも用意してそれ使わせればいいのに、と考えると
HLSLとかの方がいいんじゃねえの?って結論になるw


128 :デフォルトの名無しさん:2010/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 :デフォルトの名無しさん:2010/05/09(日) 21:18:40
揚げ足をとるようだけど、右足を出した隣の人は左足を最初に出してる。
__syncthreadsしなければ隣のひとたちを見ないでばりばり走るし、
過度な比喩は逆に混乱させるだけだと思う。

>>100みたいなのは計算式を丸ごと変えないといけないな。
プログラムより、数学的な知識が必要だと思う。

130 :デフォルトの名無しさん:2010/05/09(日) 21:19:54
>>126
いやw
nの数の条件分岐とか必要ないでしょww
Host側がコール時に調整すればいいだけの話

CUDAの場合、無駄な命令入れるよりかは、少し出力メモリ先がはみ出ててもコアレスならおk

131 :デフォルトの名無しさん:2010/05/09(日) 21:23:42
>>130
nが巨大なら、ループする必要があるでしょ。ループ継続条件は通常、nとの比較で行なわれるわけで。

132 :デフォルトの名無しさん:2010/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 :デフォルトの名無しさん:2010/05/09(日) 21:32:13
>>131
だったら数回カーネルをコールすればいいのでは??
この場合カーネル内でループしちゃうと、キャッシュ外れる避けられない。キャッシュミスは大きいよ。

134 :デフォルトの名無しさん:2010/05/09(日) 21:39:07
>>100 は一次元差分法からの連想だろうけど、
差分法なら青木本に2次元があるから参考にしたら。
自分は3次元の差分法を作ったけど信じられない速度が出た。
まあ、比較したCPU版が遅いんだろうけどw

135 :デフォルトの名無しさん:2010/05/09(日) 22:06:07
>>132
CC1.1までを前提にするけど、コアレスであるためには
スレッドがアクセスするメモリがハーフワープ(16個)連続して配置されている必要がある。
SMがロードする際16個まとめてグローバルメモリをロードして各スレッドに割り当てるから。
だからハーフワープあたりのグローバルメモリのアクセスは1回になる。

じゃあコアレスじゃないとどうなるかというと、
キャッシュなんてないないから各SPは個別に毎回グローバルメモリに読みにいく。
だから遅い。

136 :デフォルトの名無しさん:2010/05/09(日) 22:20:52
>>135
「bへのアクセス3つは16個連続していること」って用件は満たしてる。
多分アライメントも16個境界になってないと駄目ということか。
確かにそうすると、 b[i+1]とb[i-1]はアライメントに合ってないから
都度ロードが発生する訳だな。
理解しました。どうもありがとう。

137 :デフォルトの名無しさん:2010/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 :デフォルトの名無しさん:2010/05/10(月) 00:51:19
>>127
CUDAというかアーキテクチャの制限ではないでしょうかね?
そのままグローバルメモリにアクセスするのではなく、シェアードメモリを利用するという思想で作られていると思っています。

>>130
出力メモリ先がはみ出るのはまずいかと思いますが、分岐が複雑になる場合は
分岐を行わずに配列に0を入れておくとかの方が良さそうですね。

>>134
あの本はそういったことも載っているのですか。
やはり一度は読むべきなのでしょうかね・・・

139 :デフォルトの名無しさん:2010/05/10(月) 02:36:11
>出力メモリ先がはみ出るのはまずいかと思いますが、
メモリを考慮してグリッドとブロックの数を決めるより
はじめにグリッドとブロックの数を決めてそれに合わせてメモリを確保したほうが良い。
ってことが言いたいんじゃないか。

140 :デフォルトの名無しさん:2010/05/10(月) 07:11:13
>>137
分岐の種類によって違う結果になる。
・ワープ中幾つかのスレッドだけが条件に該当する場合
条件分岐のコスト(のみ)が余計に掛かる。
・ワープ中全てのスレッドが条件に該当する場合
ワープによっては条件分岐のコストが嵩むが
ワープによっては条件分岐以降の処理をパスするので早く開放される。

尚、if文の中身が複雑でループ内にあるがループには依存しない場合、
条件判断処理をループ外で行ない結果のBoolean値だけをレジスタに保存しておく最適化は行なわれる。

条件分岐は怖くありませーん。下手に複雑な演算するよりお徳でーす。

141 :デフォルトの名無しさん:2010/05/10(月) 11:26:17
>>102
シェアードメモリは使ったほうがいいと思いますよ

142 :デフォルトの名無しさん:2010/05/10(月) 11:47:48
うーん、CUDAの初歩も知らない奴に取り敢えず雛形を提示したら、
ちょっと齧った連中が寄って集ってけちつけるという構図になったなw
合間に慣れた人が適切な指摘をしてくれているからいいけどねぇ。

つーか、CUDAがそれだけ「注目」されていると言うことなんだろうけど。

143 :デフォルトの名無しさん:2010/05/10(月) 13:31:41
やだ・・・なにこの失敗したCell並に難解なプログラミング・・・

144 :デフォルトの名無しさん:2010/05/10(月) 14:51:44
大丈夫、Cellのspuみたいにインストラクションを一個ずつあれこれしなければならないほど複雑じゃない。

145 :デフォルトの名無しさん:2010/05/10(月) 17:57:02
GTX4*シリーズ買う?
それとも買った?

146 :デフォルトの名無しさん:2010/05/10(月) 23:14:36
前スレから

303 名前:デフォルトの名無しさん[sage] 投稿日:2009/12/05(土) 19:32:04
倍精度性能はGT200世代はおまけで1個演算機がついてるだけだから期待するだけ無駄だよ
Fermi世代から各コアに演算機付けるって言ってるけど、
一般向けには倍精度削ってGF100の名前で出すとか言ってるからどうなるか分からん


>一般向けには倍精度削ってGF100の名前で出すとか言ってるからどうなるか分からん

GTX480で倍精度1/4とかなった(ようだけど)
去年の12月の時点でこんな話が出ていたんだね

で、GTX4*買うかはやっぱり倍精度が気になるわけで、様子見の状態

147 :デフォルトの名無しさん:2010/05/12(水) 09:25:19
C20シリーズは何とか予算用意して買うか、どこかの計算機センタに導入されるのを待つ形になる
それの準備的なものとして、GTX4シリーズを買うと思う

148 :デフォルトの名無しさん:2010/05/14(金) 10:29:43
完全にHPCユーザーしかGPGPUに興味を持たなくなったな

149 :デフォルトの名無しさん:2010/05/14(金) 11:10:38
画像処理とか(?)する人はどう思っているんでしょう
GPGPUについては

150 :デフォルトの名無しさん:2010/05/14(金) 12:19:47
CG系研究室の俺が。

GPGPU(=CUDA)は今や当たり前の技術です。
なんか論文で新しい手法を提案するとき、実行時間の項にCPUでの計算時間だけ書いていたら、
「GPUでやったらどうなのよ?」と査読者から普通にコメントつけられるレベルです。
GPUに載せづらい既存の手法を、ちょっと改良してGPU適合なアルゴリズムにしました、的な論文が
3年ぐらい前から大量に出ています。

151 :デフォルトの名無しさん:2010/05/14(金) 13:02:33
>>150
> CG系研究室の俺が。
CGの人ってコンピュータビジョンにもアンテナ伸ばしてるものなんですか?

> 「GPUでやったらどうなのよ?」と査読者から普通にコメントつけられるレベルです。
*ただしリアルタイム応用に限る。

> 3年ぐらい前から大量に出ています。
出過ぎたせいで最近は実装してみましたレベルの論文はむしろ通りにくいですよ。

152 :デフォルトの名無しさん:2010/05/14(金) 18:40:04
ATI stream とCUDA とどっちがよく使われているんでしょうかね

153 :デフォルトの名無しさん:2010/05/14(金) 19:19:01
「実装してみました」系はCUDAだろうね。
ピークパフォーマンスたたき出すならATIだろJS

154 :150:2010/05/14(金) 20:07:19
>>151
>CGの人ってコンピュータビジョンにもアンテナ伸ばしてるものなんですか?
ちょっとだけね。 CVのいくつかのツールは、こっちでも役に立つから。

>出過ぎたせいで最近は実装してみましたレベルの論文はむしろ通りにくいですよ。
あるあるwwww

>>152-153
俺の狭い視界だけで言うと、ATI streamは使っている人を見たことがない。
俺は昔試しに使ってみたけど、ちょっと複雑なことをやろうとするとすぐに
Direct3D8でやったようなシェーダアセンブリを書かなくちゃならなくなって辟易した。

今のところGPGPUではCUDA一択だと思います。
研究者の皆さんもツールの使い方にそれほど時間を掛けられるわけでもないので、
ほぼC/C++と同様に書けるCUDAが主流になるのは仕方ないと思います。

155 :デフォルトの名無しさん:2010/05/14(金) 20:19:13
ATIを利用した論文を目にしたことがないな、確かに。

156 :デフォルトの名無しさん:2010/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 :デフォルトの名無しさん:2010/05/15(土) 01:12:46
>>154
GPGPUとしてはCUDAが主流になると思う。
だけどこれまではほんとに酷かったなあ。
ちょっとでもコンピュータサイエンスの知見を持っていたら恥ずかしくて発表出来ないレベルが多かった。
今までは流行でみんながCUDA、CUDAといっていたけど、多くの人はあくまでも手段の一つであって、
目的ではない。目的にするならコンピュータサイエンスの分野でやればよいわけで。
それにCUDAを使う場合って大抵最適化しているから、CPUとの比較も余り意味がないし。
これからはそういうのはなくなっていくだろうね。




158 :& ◆6kLsuKGM9vPI :2010/05/15(土) 02:51:39
480と285の倍精度計算のベンチマークの比較はどこかにありませんか?

159 :デフォルトの名無しさん:2010/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 :デフォルトの名無しさん:2010/05/15(土) 15:40:25
>>157
どういうこと?

161 :デフォルトの名無しさん:2010/05/15(土) 16:14:37
>>159
うまく行くケース(GPUのバージョン依存だっけ??)なら3回でいけるだろうけど、
最悪の場合

b[tx]:コアレス
b[tx+1]:シリアライズ
b[tx-1]:シリアライズ

ってなるから、1+16+16 = 33 じゃね??

162 :デフォルトの名無しさん:2010/05/15(土) 16:39:09
>>161
全部のSPがb[tx]を読む→コアレス
全部のSPがb[tx+1]を読む→コアレス
全部のSPがb[tx-1]を読む→コアレス

とはならないの?

163 :デフォルトの名無しさん:2010/05/15(土) 17:00:15
>>161
そういう場合、
スレッド i-1 はtx-1をグローバルからSharedに読む
スレッド i   はtxをグローバルからSharedに読む
スレッド i+1 はtx+1をグローバルからSharedに読む
 :
でひと固まりもってきて、
あとはShared上でCoaleshed? キニシナイ! で動けばいいんじゃね

164 :デフォルトの名無しさん:2010/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 :デフォルトの名無しさん:2010/05/16(日) 02:31:46
>>162-163
「t*x」 になっているんでこのままだとtステップになっちゃうわけだよ
つまりもともとの配列の並べ方から考え直さないといかんよな。
x-1, x, x+1, ,,, になったら163みたいに、
・スレッド協調してSharedに持ってくる
・Shared上では好き勝手に読み書きしてよろし
・syncthreadする
・スレッド協調してGlobalに書き戻す
というやりかたが基本かな

166 :デフォルトの名無しさん:2010/05/16(日) 09:56:22
>>165

>・Shared上では好き勝手に読み書きしてよろし
今扱ってる問題だと場合によっては、bank conconflictがあったりする。。。

167 :デフォルトの名無しさん:2010/05/16(日) 09:57:09
>>166
まちがえたw

>>165

>・Shared上では好き勝手に読み書きしてよろし
今扱ってる問題だと大丈夫だけど、場合によってはbank conconflictがあったりする。。。

168 :デフォルトの名無しさん:2010/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 :デフォルトの名無しさん:2010/05/16(日) 11:06:33
テクスチャキャッシュを使うのが面倒なくていいと思うのは俺だけか?

170 :デフォルトの名無しさん:2010/05/16(日) 11:33:22
>>169
少し前にテクスチャ使ったコードで例を示したけど、誰も反応しなかったぞ
なんで??

171 :デフォルトの名無しさん:2010/05/16(日) 14:23:32
どのレス?

172 :デフォルトの名無しさん:2010/05/16(日) 18:59:40
>>171
>>125

173 :デフォルトの名無しさん:2010/05/16(日) 19:31:54
OpenCLでなくCUDAを学ぶ価値ってあるの?

174 :デフォルトの名無しさん:2010/05/16(日) 19:34:06
OpenCLってなにに使われてるの

175 :デフォルトの名無しさん:2010/05/16(日) 19:38:18
OpenCLは今のところRadeon使う人用かと。


176 :デフォルトの名無しさん:2010/05/16(日) 20:33:59
DirectComputeはどんな感じ?

177 :デフォルトの名無しさん:2010/05/16(日) 21:33:22
ローカルメモリに対するテクスチャメモリのメリットがいまいち分からない
そもそもローカルメモリを使いこなせないから、テクスチャメモリまで手が回らない

178 :デフォルトの名無しさん:2010/05/16(日) 23:49:15
>>175
え?

179 :デフォルトの名無しさん:2010/05/17(月) 00:41:56
>>177
補間できるとか

180 :デフォルトの名無しさん:2010/05/17(月) 20:52:47
これってcudaMallocHostってsizeof(構造体)でぴったりとその構造体の容量分だけちゃんと確保できてるんですか?
なんかsizeof使ってぴったり構造体分確保したら条件によってセグメントエラーが出るんですが(printfを適当なとこに放りこんだらエラーが出ずに通ったりします)
2倍の容量を確保したら問題なく通ります


181 :デフォルトの名無しさん:2010/05/17(月) 23:08:13
>>177
ローカルメモリがLocal Memoryのことをいってるなら、役割がまったく違うからメリットもなにもない
Global Memoryのことをいってるなら、キャッシュに加えて補間、正規化アドレス、変換とかだな
キャッシュが有効に機能しているかは、プログラムによっては疑わしいんだけどね

あとコーディングがちょっとめんどくさい

182 :デフォルトの名無しさん:2010/05/18(火) 09:35:45
>>180
構造体に問題があると思われます。
各要素がアラインメントの条件を満たすようにダミーの要素を追加して調整することをお勧めします。

くらいしかアドバイスしようがないな。


183 :デフォルトの名無しさん:2010/05/18(火) 10:19:15
printf挟むと動くと言ってるからどっかでスタック壊してんじゃないかと予想

184 :デフォルトの名無しさん:2010/05/19(水) 14:37:47
まぁ、晒せ、と。

185 :デフォルトの名無しさん:2010/05/20(木) 11:59:21
ω

186 :デフォルトの名無しさん:2010/05/21(金) 09:13:45
研究室にGTX 480届いた.
smokeParticlesで170fpsくらい.
倍精度FFTとかはこれからやる.

187 :デフォルトの名無しさん:2010/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 :デフォルトの名無しさん:2010/05/21(金) 11:33:50
nbody --benchmarkよろ

189 :デフォルトの名無しさん:2010/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 :デフォルトの名無しさん:2010/05/21(金) 13:05:20
>>187
多分CUFFT3.0を使った複素2次元FFTだと思うのだが、
こちらの計測データとかなり違うような・・・。
こちらでは480の単精度で1回(片方向)あたり0.44ミリ秒、2*1024回分で0.91秒くらい。

ちなみに頑張ると倍精度は2.6秒くらいでいける。

191 :187:2010/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:2010/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:2010/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:2010/05/21(金) 13:21:21
nvcc cufft.cu -o cufft_double -arch=sm_20 -lcudart,cutil32,cufft -Xcompiler /O2

↑のソースをこれでコンパイルした.

195 :デフォルトの名無しさん:2010/05/21(金) 13:47:54
>>189
おお、ありがとうございます。
GTX260でも370GFOPS位出るのであんまり速くないですね・・・(´・ω・`)

196 :デフォルトの名無しさん:2010/05/21(金) 13:54:24
>>187

>CUT_SAFE_CALL(cutStopTimer(timer));

これの前にcudaThreadSynchronize();を入れないといけない予感。

k の反復回数を512にした時のほぼ倍の時間になるかどうかをチェックしてみて。



197 :187:2010/05/21(金) 15:00:01
同期をちゃんととってみた.

・単精度
285 → 1.710[s]
480 → 0.875[s]

・倍精度
285 → 8.438[s]
480 → 4.076[s]

>>190
倍精度を頑張るってどう頑張るんですか?

198 :デフォルトの名無しさん:2010/05/21(金) 15:28:32
質問させてください

CPU...:Intel Core i7-860(2.80GHz)
メモリ........:4G
M/B .......:ASUS P7P55D Delux
VGA  .....:FX580
電源..........:850W (80Plus)
の機材で現在ps4とpremireで地方の小さなブライダルの写真や映像をやっているのですが

友人からFX580を譲っていただけるとのことで2枚差しで使ってみようかと思うのですが
SLIは全く知識がなく、2枚差しの効果があるのかがわかりません。

どこで聞いたらよいのかわからずスレチかもしれませんが先輩方、アドバイスお願いします。

199 :デフォルトの名無しさん:2010/05/21(金) 20:45:35
SLIというかマルチGPUで使う場合はOpenMPかなんかで複数スレッドを作って
それぞれののGPUを使えばいい、というかサンプルがSDKにあったはず。
ゲームみたいにSLI組んでシングルGPUに見せるのはまだ出来なかったはず

200 :デフォルトの名無しさん:2010/05/21(金) 21:36:18
>>199
親切にありがとうございます。
まさにゲームのようなSLIの考えでいましたが、そういった使い方はできないということですので
自宅用に使用するか等、使い方を考えてみます。本当にありがとうございました。


201 :デフォルトの名無しさん:2010/05/21(金) 22:56:10
>>197
もちろんコードを自分で書く。


202 :デフォルトの名無しさん:2010/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:2010/05/26(水) 17:15:17
と聞いたばかりですみません。
解決しました。
ちゃんとgputoolsのinstall.txtを呼んでいなかったので
パラメータを間違えていて、--enable-emulationというのを
つけてしまっていました。
お騒がせしました。

204 :デフォルトの名無しさん:2010/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 :デフォルトの名無しさん:2010/05/26(水) 23:38:38
>>204
言われてる1/4ってのは当初期待されていた性能に対して1/4ってことじゃなかったの?
いくらなんでもTesla世代と比較してないでしょ

206 :デフォルトの名無しさん:2010/05/27(木) 03:41:07
当初は短精度に対して理論値で1/2の性能を出せるって話だったのだけど
Geforceでは削られて短精度に対して1/4しかでないことになっている

207 :デフォルトの名無しさん:2010/05/27(木) 05:53:57
1/8じゃなかったの?

208 :デフォルトの名無しさん:2010/05/27(木) 06:06:21
皆GeForceでは1/8 (Fermi版Teslaシリーズのさらに1/4)という意味で話していたけど
一部の人は誤解していたみたい。

209 :デフォルトの名無しさん:2010/05/27(木) 20:51:44
どういうことだってばよ

210 :デフォルトの名無しさん:2010/05/28(金) 00:28:01
けっきょくどういうことなんだってばよ!

211 :デフォルトの名無しさん:2010/05/28(金) 08:37:10
こまけえことはいいんだよ!

>>205
俺の理解はこれなんだが
GTX480も在庫が増えて値段も下がってきたね
俺、5万円切ったら本気出すんだ…

212 :デフォルトの名無しさん:2010/05/28(金) 17:09:48
Teslaじゃコードを書き換える手間を考えたら、たいしてコストパフォーマンスは良くないな。


213 :デフォルトの名無しさん:2010/05/28(金) 21:44:15


214 :デフォルトの名無しさん:2010/05/29(土) 05:36:25
>>100の様なコードだったら手間はかからんよ。

215 :デフォルトの名無しさん:2010/05/29(土) 12:56:20
TeslaがGT世代の事じゃなくてGPGPU専用カードの事を言ってるなら
Teslaの利点はバグのあるコードを走らせてもOS巻き込んで死んだりしない点が一番じゃね?
開発機はTeslaが楽だわ

216 :デフォルトの名無しさん:2010/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 :デフォルトの名無しさん:2010/06/03(木) 06:46:24
2次元配列を使うと、遅くなりますが、
2次元ブロックも同様の問題をかかえていますか?

218 :デフォルトの名無しさん:2010/06/04(金) 17:03:58
なんのこっちゃ

219 :デフォルトの名無しさん:2010/06/05(土) 01:05:17
2次元配列なんてCUDAにはありません。
意図的に1次元に並べてください。

220 :デフォルトの名無しさん:2010/06/05(土) 20:17:51
GTX480でおすすめのボードってありますか?
もしくはダメなボードはありますか?

221 :デフォルトの名無しさん:2010/06/05(土) 21:43:43
こちらも480入手。取り敢えず動くのは確認。細かいテストはこれから。

>>220
今のところどこも経験値低いだろうからリファレンスボードのままだろうし、
ELSAが無難じゃね? 高いかもしれないけど。
ヒートパイプの所為で更に場所を喰うので要注意。

222 :デフォルトの名無しさん:2010/06/05(土) 21:52:51
リファレンスだからドスパラのでいいんじゃないの
シール張るだけらしいから

223 :220:2010/06/07(月) 11:06:57
>>221-222
ありがとうございます。

elsaのGTX480のページに対応OSがwindowsしか書いてないんですが、
Linuxでは動かないのでしょうか?
GeforceのCUDA用ドライバをいれたら普通に動くんでしょうか?

224 :221:2010/06/07(月) 11:31:39
>>223
ドライバのリビジョンは今手元にないから判らないけど、CUDA2.3当時のドライバでCUDA2.3が動いている。
CUDA3に関しては未だテストしていない。

225 :デフォルトの名無しさん:2010/06/14(月) 13:16:15
最近CUDAを勉強しはじめたんですが、質問させてください。

入力画像をグレースケール化して出力する、というプログラムを作ったところOpenCVだけで作ったプログラムに比べて、
プログラム全体の処理時間は速くなりませんでした。どうやらプログラムで最初に実行されるcudaMallocで100ms近く掛かっているようです。
次の行にも同じくcudaMallocがあるんですが、これは1msほどで完了します。
これは改善できませんか?

OSはUbuntu 9.04、GPUは9500GT、cudaのバージョンは2.3です。

226 :デフォルトの名無しさん:2010/06/14(月) 13:20:32
できます。
一回目のcudaMalloc()の前にcudaThreadSynchronize()でもしておけば、一回目のcudaMalloc()は早くなります。
尤も、cudaThreadSynchronize()で時間が(恐らく100ms程度)掛かりますが。

227 :デフォルトの名無しさん:2010/06/14(月) 13:29:32
>226
ありがとうございます。
早速試したところ、そのとおりの結果になりました。。。
これは諦めるしかないんですかね。

228 :デフォルトの名無しさん:2010/06/14(月) 13:29:33
なんかCUDAは最初に呼ぶときに時間がかかるとかリファレンスかなんかに書いてあったような・・・
斜め読みしたけどみつけられんかった

229 :デフォルトの名無しさん:2010/06/14(月) 13:41:43
>228
そんな仕様だったんですね、知らなかったです。

書き換えたもとの処理がたいして時間が掛からない場合、CUDAの利点を活かせないんですね。

230 :226:2010/06/14(月) 13:45:06
>>227
cudaは元々単発プロセス向きじゃないので、諦めて常駐型のプロセスか処理時間の掛かるプロセスに使いましょう。
具体的に何をやっているかは判らないけど、cudaのAPIを呼ぶ初回は必ず時間が掛かるのよね。
例えばmain()の先頭ででも一回syncするだけでいいから、私の用途では問題にならないんだけど。

231 :デフォルトの名無しさん:2010/06/14(月) 13:48:59
>>229
一回だけだとそうかも知れないけど、バッチやリアルタイムで複数ファイルを処理するとかなら
効いてくると思います。

232 :デフォルトの名無しさん:2010/06/16(水) 10:00:49
研究にCUDAを使ってみたいと思っている学生です。><
プログラマブルシェーダで物理演算をしてシミュレーションをやっている
論文とかあったのですが
CUDAで計算するのとの違いはあるんですか?
あと
DirectXを使いながらCUDAを利用することは
DirectXで手一杯のことをやっていたらCPUで計算したほうが早かったりするんですか?

お願いします(。・・。)

233 :デフォルトの名無しさん:2010/06/16(水) 10:05:47
>>232
それが俺の読んだのと同じならcgを使ってシェーダプログラムで演算をしてた
理由はCUDAがまだ無かったから
DirectXで3D描画しながらだとそりゃ性能は落ちるでしょ。CUDA専用にカードを追加するなら話は別だけど

234 :デフォルトの名無しさん:2010/06/16(水) 12:16:27
だいぶ前の話だけど、AviUtlのフィルタでシェーダよりCUDAの方が遅いとかって無かったっけ。
カリカリにチューンしたらCUDAが速いんだろうけれど、シェーダの方がメモリアクセスのお作法を意識しなくて良い分、初心者には性能向上が簡単かも知れない。

235 :デフォルトの名無しさん:2010/06/16(水) 19:47:40
>>234
つまりn00bだったってことかよ!

236 :デフォルトの名無しさん:2010/06/16(水) 23:49:14
>>234
???
CGをやるなら別だが、今更Cgは無いだろう。
初心者ならなおさらCUDAだ。
特に物理演算をやるなら。



237 :デフォルトの名無しさん:2010/06/17(木) 00:34:00
PGIコンパイラのCUDA化って正直どうよ。
人力でやるのに疲れてきた。
それなりの性能が出るなら買ってもいいと思ってるんだけど。

238 :デフォルトの名無しさん:2010/06/17(木) 06:18:56
ディレクティブでは全然性能上がらなくて頭来たからCで書き直したw

239 :デフォルトの名無しさん:2010/06/17(木) 06:53:59
東工大、Teslaを利用したスパコン「TSUBAME2.0」を開発
http://pc.watch.impress.co.jp/docs/news/20100617_374718.html

240 :デフォルトの名無しさん:2010/06/17(木) 14:14:54
>>239
有用な結果を出さなかったら、次から仕分けな

241 :デフォルトの名無しさん:2010/06/17(木) 18:10:08
国会議員を仕分けすりゃいいのに
1/3位にさw

242 :デフォルトの名無しさん:2010/06/17(木) 20:46:42
あんたら、部外者か?

蓮舫のブレーンがそのスパコンの推進者なのに、仕分けされるわけがねーだろ。

243 :デフォルトの名無しさん:2010/06/17(木) 21:00:10
つまり漣舫は神戸のスパコンをつぶしたくで仕分けしたのですね。

244 :デフォルトの名無しさん:2010/06/17(木) 21:11:00
1度でいいからこんな超計算機をいじり倒したいものだ

245 :デフォルトの名無しさん:2010/06/17(木) 21:15:30
部外者だけどCSX600なんて誰も使ってないんじゃね?
導入初期に不良品大量に掴まされてたみたいだけどw

246 :デフォルトの名無しさん:2010/06/17(木) 21:28:53
まあそうでもしないと進まないところもあるんだろうけど、こんな財政状況なんだからバッサリ逝ってよし

247 :デフォルトの名無しさん:2010/06/17(木) 21:34:09
>>242
むしろ部外者しかいないだろ
お前はしらんが

248 :デフォルトの名無しさん:2010/06/17(木) 21:35:11
議員を減らせば良い

249 :デフォルトの名無しさん:2010/06/17(木) 21:35:34
政治ネタうぜぇ

250 :デフォルトの名無しさん:2010/06/17(木) 21:43:16
選挙だもの

       民主党

251 :デフォルトの名無しさん:2010/06/17(木) 22:18:46
れんほうのすぐ後ろに金田@superπが座っていたから
京速機のヘッドである姫野@ヘボベンチは何も言えんかったんだよ。

専門がスパコンはおろかコンピュータですらない姫野に
反論ができるわけもない。

252 :デフォルトの名無しさん:2010/06/17(木) 22:21:51
π(笑)

253 :デフォルトの名無しさん:2010/06/18(金) 02:21:28
>>251
何で姫野ヘボベンチっていわれているの?
ソースを見たけどコード自体は学生が書くような内容だね。
だからなの?

254 :デフォルトの名無しさん:2010/06/18(金) 12:50:58
メモリ帯域をひたすら無駄食いするプログラムだよね>姫野ベンチ
ここで「なぜか遅いです」とか質問してくるレベル。

255 :デフォルトの名無しさん:2010/06/18(金) 15:19:59
かといってgemmでピーク性能測られても役に立たない分野も多く。
直線に強い車かカーブに強い車かってのと似てる。

256 :デフォルトの名無しさん:2010/06/18(金) 17:38:19
πって何を測ってるの

257 :デフォルトの名無しさん:2010/06/18(金) 17:42:40
今ではすっかり使われないx87

258 :デフォルトの名無しさん:2010/06/19(土) 03:34:48
>>235-236
今更シェーダをお勧めはしないけど、畳み込みを考えた時に
コアレスにするためのコードって余分だよね。バッドノウハウだよね。
ってことが言いたかった。

259 :デフォルトの名無しさん:2010/06/21(月) 12:53:32
東工大のスパコン、世界2位なんだ。
そういうことか。

260 :デフォルトの名無しさん:2010/06/21(月) 13:55:44
東工大は64位じゃね?
http://www.top500.org/list/2010/06/100

261 :デフォルトの名無しさん:2010/06/21(月) 16:58:25
ヒント:TSUBAME2はまだ実在しない
っ ttp://www.gsic.titech.ac.jp/tsubame2

262 :デフォルトの名無しさん:2010/06/22(火) 08:40:40
まだ計算どころか設置もしていないぞ

263 :デフォルトの名無しさん:2010/06/22(火) 12:51:00
NECは京速機やめてこっちでやってんだな。

264 :デフォルトの名無しさん:2010/06/23(水) 01:26:00
CUDAって未だにデバイスメモリをメインメモリにマップ出来ないの?
逆は出来るようだけど。

265 :デフォルトの名無しさん:2010/06/23(水) 01:45:32
マップできなくてもいいけど、MPIでDMA使ってデバイスから通信チップに流しこんで欲しいわな

266 :デフォルトの名無しさん:2010/06/23(水) 05:56:26
マップ出来ないといちいち転送指示出さなきゃいけないのが困る。
スクラッチで書くならいいけど、既存のを移植するとなると結構大変だから・・・。

267 :デフォルトの名無しさん:2010/06/23(水) 19:31:52
テンプレートライブラリで、CUDAの転送まわりをほぼ完全に隠蔽してくれる奴あったよね。

268 :デフォルトの名無しさん:2010/06/24(木) 12:55:15
thrust

269 :デフォルトの名無しさん:2010/06/26(土) 00:36:06
使い方によるのかもしれないが、Mappedはかなり遅い気がする
Pinnedで転送したほうがまし

270 :デフォルトの名無しさん:2010/06/27(日) 13:53:48
GTX480のベンチが日本語でも少しずつ出てきたね
倍精度と単精度両方載せてある

PGI CUDA Fortran を使用した行列積の計算とGPU性能
ttp://www.softek.co.jp/SPG/Pgi/TIPS/public/accel/cuf-matmul.html

CUDA Fortranだけど、CUDA Cでも結論が変わるものではないでしょう

271 :デフォルトの名無しさん:2010/06/27(日) 13:57:56
MATMULとかだと倍精度も速いみたいだけど、レジスタやシェアードメモリの消費も倍になるから
実コードだとOccupancyが下がって思ったような速度でないんだよなぁ・・・ >Fermi

272 :デフォルトの名無しさん:2010/06/27(日) 14:14:32
倍精度の実効効率が、

GTX285 約90%
GTX480 約21%

ずいぶん差があるのはGTX480の倍精度がドライバのレベルで1/4に抑えられている(らしい)ことが最大の原因?
アーキテクチャが違うから単純には言えないけど、21/25=84で、GTX285と同程度になるし…

273 :デフォルトの名無しさん:2010/06/27(日) 14:51:36
ソース

274 :デフォルトの名無しさん:2010/06/27(日) 14:53:17
>>273
>>272のは真上の祖父テックのベンチ結果から計算してのことだろ?
nbody.exeでも-fp64で性能出てなかったし

275 :デフォルトの名無しさん:2010/06/27(日) 16:10:48
> ドライバのレベルで1/4に抑えられている(らしい)こと
ここのソース

276 :デフォルトの名無しさん:2010/06/27(日) 16:41:24
googleを使えない奴にCUDA使えるとは思えないな

277 :デフォルトの名無しさん:2010/06/27(日) 18:34:33
これと、中国のTUBASA?のLinpack結果を併せてだれか解説して


278 :デフォルトの名無しさん:2010/06/27(日) 18:56:27
GTX480の「倍精度理論FP演算性能672 GFLOPS」という仕様そのものが間違っていて
実際は理論ピークが1/4の168。

別にドライバのレベルで抑えられているわけでもなんでもないが。

279 :デフォルトの名無しさん:2010/06/27(日) 19:13:13
>>278
↓この記事を信じてしまうと
http://journal.mycom.co.jp/special/2010/gpu/004.html
GTX480の倍精度演算の理論値は672GFLOPSになっちゃうけど、

>実際は理論ピークが1/4の168。

が本当だとするなら、どこで計算がずれたのだろう…

280 :デフォルトの名無しさん:2010/06/27(日) 22:22:30
ドライバなんだかそれ以外だかはわからないが、
GTX480の倍精度理論性能が168GFLOPSであることは信憑性高い、ということでおk?

281 :デフォルトの名無しさん:2010/06/27(日) 22:24:19
ドライバで制限してるんじゃ無ければカタログスペックとの差異はどーすんの?
ttp://www.elsa-jp.co.jp/products/hpc/tesla_c2050/index.html

282 :デフォルトの名無しさん:2010/06/27(日) 23:07:18
>>281
GTX480の話をしているときにC2050持ち出してどうするの?

283 :デフォルトの名無しさん:2010/06/27(日) 23:24:10
想像力を働かすと、
C2050よりもコア数が多くて、同一アーキテクチャ仕様と公表されているGTX480が168GFLOPSしか出ないなんておかしいぞ
ということかな

まあ、たとえ168GFLOPSだとしても、共有メモリやレジスタまで1/4にはなっていないだろうし(倍精度だと2倍食うけど)、
単位演算性能あたりのメモリバンド幅が単精度と比べて大きくなることを考えれば、実効効率が単精度よりも大きくなるのはわかる気がする

ECC無しでもOKで、メモリ1G程度で十分という人には結構おいしいかもしれんねGTX480

284 :デフォルトの名無しさん:2010/06/27(日) 23:33:39
メモリしか差のなかったC1060に対して、C2050はもっと明示的な計算用チップとしてのアドバンテージをつけたかった
nVidiaの意向が反映した結果がGTX480だと思うのだけど。

285 :デフォルトの名無しさん:2010/06/28(月) 07:10:21
C2050はGTX480はCUDAコアが少なく、クロックも低いにもかかわらず、
倍精度理論性能は515GFlopsとなりC2050ほうが高い。
これは、GTX480はドライバで倍精度演算の実行にペナルティが加わるようになってるから。
それでも、GTX285と比べても倍精度演算が高速だからいいじゃんというのがNVIDIAの考え。

もっとも実際にはそれほど性能は変わらないようだよ。

現実のカーネルコードの多くではこんなベンチと違ってメモリ帯域が足を引っ張るから、
当然C2050実効性能は理論性能遠く及ばない。
一方、GTX480は単位時間当たりの演算回数に対して、メモリ帯域がC2050よりも大きい分、
メモリ帯域による実効性能減が起きにくい。
結果として、カリカリにチューンしない限りは、GTX480とC2050の倍精度演算の
実効性能はそれほど大きくは変わらないようだ。

286 :デフォルトの名無しさん:2010/06/28(月) 07:18:45
何が言いたいかというと、C2050と比べるとGTX480はそこそこお買い得だってことね。
ECC、メモリ容量、倍精度演算の潜在能力とメーカー出荷時のチェック品質が値段の違いなのかな。

287 :デフォルトの名無しさん:2010/06/28(月) 07:23:59
>>283
>>281だけどそう言いたかった。OpenGLが速いQuadroみたいな感じなのかね?
Telsa C20xxのアホみたいな値段見てると・・・

288 :デフォルトの名無しさん:2010/06/28(月) 07:46:15
CUDAクラスタを出荷してる会社の人から聞いた話だけど、
ごくまれにGTX系でCUDAカーネルを実行すると、間違った結果しか返さないボードがあるらしい。
他のベンチマークは普通に完走するし、ちゃんとディスプレイ出力もできて、
正常としか判断しようがないボードでも間違うのがあるそうだ。
原因は不明だが、不良と思われるそのボードを交換すると元気に動くらしいから、
やはり一見正常な不良品なんだろうといってたよ。

その点Teslaは強いといってたよ。検査も一般ボードよりずっと厳しい分、
おかしな計算結果を返すボードはまだみたことないといってた。

GTX480でも十分性能はいいけれど、Teslaは一枚は持っておくといざという時に安心らしい。

289 :デフォルトの名無しさん:2010/06/28(月) 08:32:41
>>288

Tesla 1個の値段でGTXが何枚買えるかということを考えると
GTXを複数枚買って選別してしまえばよいということに・・・・・。
不良品の確率はベンダー依存だけど4枚買ったくらいではよほど引きが強くないと不良品を引けない。

まぁ、Teslaは仲間内で誰かが1個持っておけばいいやくらいの感じかな。

290 :デフォルトの名無しさん:2010/06/28(月) 08:45:16
>>285

C2050とGTX 480は同じFermiアーキテクチャではあっても違う構造になっている。
倍精度周りでは480に対しては285と同じ実装方法でもうまくいくが、
C2050の場合は違う実装方法にしないと性能が出ないことがある。

C2050は全世代での不満としてで上がっていた安定性、及び倍精度性能の需要に応えるもので
その代わりに単精度計算では旧世代と比べて性能向上が少ない。

性能面では倍精度演算をよほど酷使する計算以外ではGTX480の方が速い。
メモリバンド幅がネックになる演算においてはC2050はGTX285にすら劣る。


291 :デフォルトの名無しさん:2010/06/28(月) 08:46:52
>C2050とGTX 480は同じFermiアーキテクチャではあっても違う構造になっている。

釣れますか?


292 :デフォルトの名無しさん:2010/06/28(月) 09:06:02
>>289
たしかにその通りなんだけど、ボードの不良と結論付けるまで、
コードと何時間も格闘したり、ホストのメモリや電源など検証するの大変だよ。

小規模で4枚だけならなんとか頑張るとしても、大規模にやる場合は、
時間を金で買う意味でもTeslaをおすすめするってことだと思うよ。

ちゃんとコードが走るかはTeslaがあれば検証できる意味でも、
一枚はあるといいとのアドバイスね。

293 :デフォルトの名無しさん:2010/06/28(月) 09:17:20
>>290
>C2050とGTX 480は同じFermiアーキテクチャではあっても違う構造になっている。
同じチップなんだから、当然同じ構造でしょ。ボードデザインは違うけど。

>メモリバンド幅がネックになる演算においてはC2050はGTX285にすら劣る。
C2050のグローバルメモリ帯域はGTX285とほとんど同じで146GB/sくらいのはず。
しかもFermi世代には、一次キャッシュと二次キャッシュがあるから多くの場合で
GTX285より高速にメモリアクセスできる。

適当言い過ぎ。

294 :デフォルトの名無しさん:2010/06/28(月) 09:25:26
3.1でたね。 ・Support for printf() in device code なにげに便利かも。

295 :デフォルトの名無しさん:2010/06/28(月) 10:45:13
>>294
これでもう、エミュレータを使う理由は激減するね。

>>292
だからそれをNVIDIAの営業某氏は「アキバ的発想はやめましょう」と揶揄するのだよね。
でも、10万円で2枚購入できるか1枚も購入できないかの差は法人ユーザには大きいだろ。

296 :デフォルトの名無しさん:2010/06/28(月) 10:54:08
GeForce GTX 285Mが搭載されているPC買ったんですが、
ttp://developer.nvidia.com/object/cuda_3_1_downloads.html
にあるノートブック用のドライバを落とせばいいんでしょうか?

297 :デフォルトの名無しさん:2010/06/28(月) 11:27:48
>>295
むしろ法人こそTeslaボードのほうがいいんじゃないの?
Teslaはちゃんと使えれば、十分に値段分の価値があるよ。
個人ユーザーはTeslaがきついのはわかるけど、
法人はそれ使って値段以上の利益あげればいいわけだし。

もっともそれが可能かどうかは、試してみないとわからないのが困り者だけど・・・。

298 :デフォルトの名無しさん:2010/06/28(月) 11:33:28
>>296
そうだよ

299 :デフォルトの名無しさん:2010/06/28(月) 11:35:47
>>297
だから、予算がついちゃえばTeslaでいいんだけどね。

300 :デフォルトの名無しさん:2010/06/28(月) 12:41:51
>>293
とりあえず実際に使ってみてから言ってください。

301 :デフォルトの名無しさん:2010/06/28(月) 13:01:29
>>292
仕分けの時に有名になった長崎大だっけ
あそこが最初にやったのが不良品を検出するプログラムの作成でしょ。
配ってなかったっけ。

302 :デフォルトの名無しさん:2010/06/28(月) 13:21:43
>>300
何を言いたいのか全く分からない。
C2050とGTX470はもってるが、GTX285はないよ。

C2050とGTX470では、倍精度演算の実効性能はそれほど違わない。
少なくとも俺のコードではね。

実際に使ったら何がわかるのさ。

303 :デフォルトの名無しさん:2010/06/28(月) 13:23:26
>>302
クレクレ君みたいで申し訳ないが、サンプルプログラムのnbodyのベンチ結果を貰えない?
nbody -benchmark

nbody -benchmark -fp64
の2つで

304 :デフォルトの名無しさん:2010/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 :デフォルトの名無しさん:2010/06/28(月) 13:32:53
GTX470は家に戻らないといけない。
環境はWin7x64で、C2050でディスプレイ出力してる。

306 :デフォルトの名無しさん:2010/06/28(月) 13:37:37
d
自分でGTX470でベンチ取ったら

516.879 GFlops@14336bodies
095.932 GFlops@14336bodies (-fp64)
だったんだが倍精度速くなってる・・ような・・・

307 :デフォルトの名無しさん:2010/06/28(月) 13:49:17
FFTは使ってませんか。
倍精度FFTはGTX285からGTX480で倍速いんですがC2050はどうでしょう。

308 :デフォルトの名無しさん:2010/06/28(月) 13:50:49
>>302

470のメモリバンド幅は133GB/s程度なのでC2050との差は小さいだろう。
480だと177GB/sもあるわけで。


309 :デフォルトの名無しさん:2010/06/28(月) 13:52:57
>>307
FFTはデータサイズ次第だってば・・・。
まぁ、480のほうが速い場合もC2050のほうが速い場合もあります。

310 :デフォルトの名無しさん:2010/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 :デフォルトの名無しさん:2010/06/28(月) 14:00:28
単精度では480が速いのに異論はない。
倍精度でもよっぽどチューンしない限り、480や470と
C2050のそれほど性能は変わらないよ。

そりゃ完全に同じじゃないけどさ。

312 :デフォルトの名無しさん:2010/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 :デフォルトの名無しさん:2010/06/28(月) 15:38:36
>>312
*.cuは実はc++としてnvccでコンパイルされるので、*.cから呼び出すためにはextern "C"にする必要があります。
つーか、c++拒否症でもなければhellonvcc.cをhellonvcc.cppに改名してc++としてコーディングした方が楽かもしれません。

314 :デフォルトの名無しさん:2010/06/28(月) 15:46:11
ていうか、別にCUDA関係ないじゃん。
シンボル修飾についてちゃんと勉強したほうがいいよ。


315 :デフォルトの名無しさん:2010/06/28(月) 15:48:35
メモリがネックにならないとき
Tesla/Geforceの単精度:Teslaの倍精度:Geforceの倍精度 = 4 : 2 : 1
メモリがネックになるとき
Tesla/Geforceの単精度:Teslaの倍精度:Geforceの倍精度 = 2 : 1 : 1
という感じだね

316 :デフォルトの名無しさん:2010/06/28(月) 16:37:51
>>313

できました。どうもありがとうございます。
nvccはC++でコンパイルされるのでしたか。知りませんでした。勉強になります
ちなみに、Cで書かれたライブラリを呼び出したかったのですが、その場合は
ラッパを作るしかなさそうですね。

317 :デフォルトの名無しさん:2010/06/28(月) 16:50:52
>>316
仮にcのライブラリのインクルードファイルが foo.h ならば、
extern "C" {
#include "foo.h"
}
でOK。

318 :デフォルトの名無しさん:2010/06/28(月) 16:51:04
呼び出す側がCで呼び出される側Cならば問題ない。
呼び出す側がC++で呼び出される側がCならばextern "C"←今の状況はこれ

だから、呼び出されるライブラリがCで書かれてるならラッパはいらない。
というかまともなライブラリならヘッダに、適切なマクロが書いてあるんじゃないかな。

319 :デフォルトの名無しさん:2010/06/28(月) 17:08:46
>>317
>>318
どうもありがとうございます。
危うく大変な作業になるところでした。

自作ライブラリなので「まともなライブラリのヘッダ」にはなっていなかったようです。
いじりたくないので317さんの方法で対応します

320 :デフォルトの名無しさん:2010/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 :デフォルトの名無しさん:2010/06/28(月) 19:23:46
このスレ住人は親切すねー

322 :デフォルトの名無しさん:2010/06/28(月) 21:23:55
実習室でやる意味ないじゃん

323 :デフォルトの名無しさん:2010/06/29(火) 01:31:18
みんな優しいなあ・・・・。
でもCUDAを始めるにはまだ早いのではないかね?

324 :デフォルトの名無しさん:2010/06/29(火) 01:45:20
3.1にしたら、CSのOceanが初めて動作する様になった。
simpleparticlesのロード時間もほぼ一瞬にまで高速化されたけど、
代わりにn-bodyのロード時間がクソ長くなった。
CUDAやOpenCLと違ってCSのサンプルがいまいち安定しないのは何でだろ。

325 :デフォルトの名無しさん:2010/06/29(火) 02:44:20
>>323
たしかにこの質問主はCUDA始めるにはちょいと早すぎるかもね。
質問内容がそもそもCUDA以前の内容だし。

326 :デフォルトの名無しさん:2010/06/29(火) 11:57:31
>>324
環境はWindowsかい?
うちはWindowsだけど、たしかにDirectComputeのOceanは初めて見た気がする。
確かにDirectComputeのN-Bodyは起動が遅いね。

よく考えたら、DirectComputeがあるのはWindowsだけかな?

327 :デフォルトの名無しさん:2010/06/29(火) 13:42:47
DirectComputeはDirectX系なのでWindowsしか無いのでは?
しかも今のところWindows7限定(´・ω・`)

328 :デフォルトの名無しさん:2010/06/29(火) 14:34:59
Vista

329 :デフォルトの名無しさん:2010/06/29(火) 14:35:55
あ、Vistaさんもいけたんですね・・・すんません

330 :デフォルトの名無しさん:2010/06/29(火) 15:40:39
くだもCLもwindowsだとリモート(リモートデスクトップ、ssh)では使えないわけだが、
DirectComputeなら使えるとかいうオチでもある?

331 :デフォルトの名無しさん:2010/06/29(火) 22:57:17
ブブブーブォォオブオーブブォーブーブー
ブォ / ̄\ーブーブーブブーンブォーオ 
ンー|  ^o^|∩==<! プォープォーブブー
ブォ \_/| | ブーブブーブォーーー
ォー _| |__| | ブォーーブブブープォー
 ー|    _| ォーブーーーブブォーブ
ブォ| |   | ブォーブーブーブーンブー


332 :デフォルトの名無しさん:2010/06/29(火) 22:59:57
リモートしたことないけどDirectComputeはDirect3Dの一部なのでDirect3Dがリモートできれば動く

333 :デフォルトの名無しさん:2010/06/30(水) 11:56:37
cuda profilerで分岐数がわかりますが、
プロファイラで分岐数というのはどうやって数えているのですか?

1スレッドあたりの分岐数にしては、多すぎる結果が出るのですが・・・orz

334 :デフォルトの名無しさん:2010/06/30(水) 13:26:58
PTXを見てbraを探せばいいんじゃない?

335 :デフォルトの名無しさん:2010/06/30(水) 17:43:43
>>333
SM一個あたりか、TPC一個あたりの総計ではなかろうか。

336 :sage:2010/07/01(木) 22:07:13
レスありがとうございます。
>>334
とりあえずptx見て数えてみますw
>>335
branchは1スレッドあたりの分岐数と表記があったのですが、
値が6桁なので、やはり総数なんでしょうかね?



ちなみに、CUDAZONEの質問掲示板のような所で
if〜elseが一つ組み込まれているプログラムで
64スレッド走らせると、プロファイラではbranch=4となったという書き込みが
あるのですが、これは半ワープ4個(=64÷16)が一つの分岐をしたから
全部で4回分岐ということなんでしょうか?
それ以外に4という数が導き出せないのですが><

337 :デフォルトの名無しさん:2010/07/01(木) 22:35:15
>>336
多分Warp単位だと思います。
if (C) A else B 文の場合、以下のようになりますが、
各Warpで分岐する方向がスレッド毎に異なれば(両方に分岐すれば)
2個×2Warpの合計4個とカウントされるかもしれません。

if (C) goto L0
do B
goto L1
L0:
do A
L1:
end.

338 :デフォルトの名無しさん:2010/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 :デフォルトの名無しさん:2010/07/02(金) 22:08:28
9800GTにするか、GT240にするか迷ってます。
性能は9800GTの方がいいので、そちらにしたいのですが、
プログラミングが趣味なので、CUDAのプログラミングを
する予定です
CUDAを考えたらどの程度違いますか?
全然違うんでしょうか?

340 :デフォルトの名無しさん:2010/07/02(金) 23:28:57





k


341 :デフォルトの名無しさん:2010/07/02(金) 23:41:05
GTX260以下はウンコ
もっと言えばFermiじゃなきゃ今から買う価値無し

342 :デフォルトの名無しさん:2010/07/03(土) 02:01:16
>>340
コイツは日本語書けるのに日本語読めないのかw
プログラミングする前に国語の勉強した方が良いぞ

343 :デフォルトの名無しさん:2010/07/03(土) 02:48:12
日本語がどうとうかはどうでもいいが、
GT9800とGT240を比べている時点でもう
「好きにしろ」
って感じだな。コアレッシングを気にするなら、
後者の方がいいのかな?GT240ってG200系?

344 :デフォルトの名無しさん:2010/07/03(土) 05:19:24
一応GT200系。
GPGPUやるなら多分240の方がいいだろう。

345 :デフォルトの名無しさん:2010/07/03(土) 14:36:26
gt240はcc1.2だからどっちにしろ倍精度は使えないね。


346 :デフォルトの名無しさん:2010/07/06(火) 00:34:22
なんだこれ・・・
ttp://pc.watch.impress.co.jp/docs/column/kaigai/20100706_378760.html

347 :デフォルトの名無しさん:2010/07/06(火) 01:15:34
ようはFermi終了のお知らせ?
Configurable CacheがサポートされてなかったらFermi用のコードなんかHPC向けしか書かないぞ

348 :デフォルトの名無しさん:2010/07/06(火) 08:06:37
GTX260と比較してプログラムの実行速度はどうなりそうなのかしら


349 :デフォルトの名無しさん:2010/07/06(火) 09:11:08
例えGT200がベースでも、AMD方式ではなくFermiのテッセレーター機構をぶち込むなら
演算器大増量に変わりは無いから速くはなるだろう。
ただ、純粋なFermiベースに比べればどうしても落ちるわな。

まあGeforceでもCUDA爆速では10倍の値段で売ってるTeslaの立場が無いし、
GPGPU用にトランジスタ増量したらゲーム性能大して向上してない割に爆熱大食らいになって
大多数の一般人には総スカン食らってるんだから、妥当っちゃ妥当なんでないかい。


350 :デフォルトの名無しさん:2010/07/06(火) 12:10:30
値段や性能はGTX260と同じだとして消費電力がGTX480の66%じゃGTX260より悪そうでこまる。

351 :デフォルトの名無しさん:2010/07/06(火) 19:32:07
cc2.0で言語仕様が大幅に変わっているから、
460はcc1.4になるのかな。

352 :デフォルトの名無しさん:2010/07/07(水) 07:21:03
倍精度切られてそうだけどね460
これで2分化が進めば、トップエンドはもっと思い切って汎用に振れるな

353 :デフォルトの名無しさん:2010/07/07(水) 21:42:32
倍精度ないのかなあ。BOINCとかでも倍精度必要だから売るために載せて欲しい。

354 :デフォルトの名無しさん:2010/07/07(水) 21:43:21
倍精度サポートはさすがにあるでしょ

355 :デフォルトの名無しさん:2010/07/07(水) 21:52:10
おいおい、お絵かき特化のために旧コア使うってのに
お絵かきに関係のない倍精度なんて積むのかね
GT200でさえ、1:8の申し訳程度のものだったのに
これもTeslaのためだけだろ

356 :デフォルトの名無しさん:2010/07/08(木) 02:11:08
別に倍精度なんか使わないから良いや
GTX465も有るんだから構わないだろ

357 :デフォルトの名無しさん:2010/07/08(木) 06:02:49
これは長期的に見て良くないね。
もはやNVIDIAはHPCを諦めてきたのかな?

358 :デフォルトの名無しさん:2010/07/08(木) 07:05:13
ね?
teslaって知ってる?

359 :デフォルトの名無しさん:2010/07/08(木) 09:07:37
あのさ、Teslaはわかるけど、今まではGeForceとTeslaを共用してきたから開発費が押さえられていたわけで、
Teslaのみを今後開発して行くには、果たしてどれだけ開発費が回収できるかによるでしょ。
今のNVIDIAに両方を開発していく体力があるのかどうか疑問だけどね。

360 :やんやん ◆yanyan72E. :2010/07/08(木) 13:45:08
TSUBAME2もできたことだし。
もうHPC事業はHPC事業独立でできるんでしょう。

361 :デフォルトの名無しさん:2010/07/08(木) 14:29:44
長期的も何も今までもローエンドは倍精度サポートして来なかったじゃん
今まで通りGTX580を作る時はリッチに作ってGTX560を作る時は削れば良いだろ

362 :デフォルトの名無しさん:2010/07/08(木) 18:57:55
TSUBAME1に鳴り物入りで採用されたCSX600は無視ですか、そうですか・・・

363 :デフォルトの名無しさん:2010/07/09(金) 01:05:40
本気で広めたいなら下位モデルでも倍精度サポートしろよ
特定のハードでしか使えないとかまじバカじゃないの

364 :デフォルトの名無しさん:2010/07/09(金) 01:22:51
チップ作るのって、簡単なことじゃないぞ。
なんか勘違いしている奴がいるが、
#ifdef GTX580
 単精度サポート 
 倍精度サポート
#else
 単精度サポート
#endif
みたいにコンパイルし直せば、出来ると思っているのか?
コンパイルのたびに数億飛んでいくってかんじなんだぞ。
デバッグの行程もたくさんあるしな。


365 :デフォルトの名無しさん:2010/07/09(金) 02:11:20
うっかりF7押しただけで数億とな
F5押したらどうなっちゃうの

366 :デフォルトの名無しさん:2010/07/09(金) 07:22:54
お絵かきに倍精度なんているの?w

367 :デフォルトの名無しさん:2010/07/09(金) 07:24:23
http://www.adrenaline.com.br/tecnologia/noticias/5507/geforce-gtx-460-especificacoes-precos-e-estreia-no-brasil.html

コアはfermi系っぽいがキャッシュはないようだ
アンコアは前世代

368 :デフォルトの名無しさん:2010/07/12(月) 15:55:32
倍精度きましたよ。1:12とさらに削ってきた。


369 :デフォルトの名無しさん:2010/07/12(月) 16:20:21
SMのSP数を変えてくるのはさすがに予想外だった

370 :デフォルトの名無しさん:2010/07/12(月) 17:41:12
>>368
8%とは酷すぎ。25%がましな時代だったとはorz


371 :デフォルトの名無しさん:2010/07/12(月) 18:14:02
25%だったことはないが・・・。
480は12.5%。

372 :デフォルトの名無しさん:2010/07/12(月) 18:35:14
なんか廉価版tesla出そうな気もする
10万以下で

373 :デフォルトの名無しさん:2010/07/12(月) 19:28:00
GTX460買ってきた。自分の倍精度プログラムはGTX260とまったく同じ性能だった。いいんじゃない。

374 :デフォルトの名無しさん:2010/07/12(月) 19:56:24
ttp://pc.watch.impress.co.jp/docs/column/kaigai/20100712_380148.html
>例えば、チップ全体での実行命令数を計算すると、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の方が上なのか
ttp://img5.pcpop.com/ArticleImages/0x0/1/1569/001569032.jpg
ttp://img5.pcpop.com/ArticleImages/0x0/1/1569/001569039.jpg
ttp://img5.pcpop.com/ArticleImages/0x0/1/1569/001569040.jpg

375 :デフォルトの名無しさん:2010/07/12(月) 21:36:12
相対的にload/storeとレジスタ、キャッシュが少ない

376 :デフォルトの名無しさん:2010/07/12(月) 22:41:09
CC2.1・・・

377 :デフォルトの名無しさん:2010/07/12(月) 23:43:43
どの命令の組み合わせなら4issue出来るんだろうか

378 :デフォルトの名無しさん:2010/07/13(火) 00:03:57
warp sizeはどうなるの? まさか48?

379 :デフォルトの名無しさん:2010/07/13(火) 00:04:52
どう考えても32だが

380 :デフォルトの名無しさん:2010/07/13(火) 13:03:45
>>375
同時に実行される命令が2->4になるだけなので、相対的な容量はあまり問題にならないような。
むしろメモリバンド幅が足りない。


381 :デフォルトの名無しさん:2010/07/13(火) 15:59:40
460だとメモリ帯域も細いから、倍精度だともはやCPUで走らすのと大差無いんじゃね?

382 :デフォルトの名無しさん:2010/07/13(火) 17:49:56
256bit 3208MHzの465より256bit 3600MHzの460の方がメモリ帯域は大きいはずだが

383 :デフォルトの名無しさん:2010/07/13(火) 19:42:19
アイドルの消費電力が低いし、お試しに買うのにいいですよ。


384 :デフォルトの名無しさん:2010/07/16(金) 06:15:22
ttp://pc.watch.impress.co.jp/docs/column/kaigai/20100716_380986.html

>データパスなどに制約があるため、倍精度演算命令を発行する際には、他の命令を発行できないためピーク性能が制約されている。
>その制約を除けば、GF100 CUDAコアはフルの倍精度演算能力を持つため、演算ユニット自体の実装コストはかなり高いと推測される。

SP自体は単体で倍精度扱えたのか

385 :デフォルトの名無しさん:2010/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 :デフォルトの名無しさん:2010/07/18(日) 12:32:31
>>385
cutilがビルド済みなら、-L <libcutil.aがあるディレクトリ> -lcutil
ビルドさえしてないなら、NVIDIA_CUDA_SDK/common辺りでビルド。
手元に環境がないので詳細は割愛。

387 :385:2010/07/18(日) 13:13:36
>>386
無事に実行できました!
即レスありがとうございます。

なにが問題だったのでしょうか
libcutil.aにcutilが含まれてるということなのでしょうか

388 :デフォルトの名無しさん:2010/07/18(日) 13:45:08
>>387
まぁ、簡単に言えばそういうこと。nm libcutil.a|gerp T してみると知見が得られるかもねw

389 :デフォルトの名無しさん:2010/07/19(月) 18:58:30
このスレの住人は本当に出来た人が多いね。
と思ってたけど、ここはそういうスレなんだね。テンプレをあらためて見て気が付いたよ。

何となくだけど、385さんは結構前に自作ライブラリがリンク出来ないと質問してた人かな?
プログラムにはあまり慣れて無いけど、素直で向上心を感じます。多分学生さんなんですかね?
別に答える必要は無いけど、そうだと仮定してアドバイスすると、日経BPから出てるプログラムはなぜ動くのかって本の8章あたりを読むと勉強になるかもしれない。

手元には無いので保障は出来ないけど、立ち読みした感じだと分かりやすく為になりそうなことが書いてある本だったよ。ひょっとするとシリーズの違う本の方がドンピシャの内容かもしれない。気が向いたら本屋で内容をチェックしてみてね。

390 :デフォルトの名無しさん:2010/07/21(水) 11:49:31
スレ違いかもしれませんが、
計算中がうまくいかず、期待してないところで0.0が出てきたので調べてみたら、
1.0e-45以下の小さい数字が出てくることが多々あり、どうやらfloatの扱える範囲外
なので0.0にされてしまっていたようです。

double(お勧めされてない?)を使う、
全部に1.0e+20くらいかけといて結果を見るときに脳内で割る
とかそんな解決方法でしょうか?


391 :デフォルトの名無しさん:2010/07/21(水) 13:06:12
>>390
CPUではfloatで動くんですか

392 :390:2010/07/21(水) 14:04:49
>>391
うごきません。cpuでも0.0になってしまいます。
dloubleにしたら1.0e-48とか表示できるんですが、
こういうときの常套手段ってあるんでしょうか

393 :デフォルトの名無しさん:2010/07/21(水) 15:08:40
>>392
そりゃあ、一部の計算するときだけ単位を変えればいいでしょ。
ただ掛け算とかすると単位も掛け算しちゃうのでそこらへん気をつける必要が
あるけど。

X * Y = Z
という計算でZが小さくなりすぎるなら

x * u = X
y * u = Y
が成り立つようにしておいて、uは単位。例えば0.001とか。

そんで計算機上では
x * y = z を求める。

でも実際は単位込で以下の計算なので
Z = (x * u) * (y * u) = z * (u*u)
最後のzの結果を見るときはu*uを掛けて考える必要がある。

394 :390:2010/07/21(水) 15:43:36
>>393
なるほど。気をつけます。

しかも情報落ちまでしてました…
doubleより有効数字小さいのか
doubleで作ったプログラムそのまま移植できないんですね。

395 :デフォルトの名無しさん:2010/07/21(水) 16:13:21
あ、っていうかfloat使う前提で書いちゃったけど、
ターゲットにGTX480を使えるか、もしくは
速度をそこまで気にしないならdoubleでもいいんじゃない?
あと部分的にdouble使ってもいいし。

396 :デフォルトの名無しさん:2010/07/21(水) 17:58:51
伊理先生の「数値計算の常識」が参考になるよ
読んでなかったぜひ

397 :デフォルトの名無しさん:2010/07/22(木) 01:27:32
>>394
そりゃぁ、doubleは「倍精度」なんだから。で、余程cuda向きの演算でない限りdoubleでXeonに勝つのは至難の技だから、
ドラスティックにアルゴリズムやデータ構造を見直さないとダメかもね。
# そして、見直した結果CPUのままでも速くなってしまうのはよくある話

398 :デフォルトの名無しさん:2010/07/22(木) 02:40:04
精度にこだわるならCUDAはやめときな。
どうしてもと言うなら多倍長のライブラリをつくって・・・・てやるなら、x87を使うのが正解だろう。
そもそもCUDAを使うメリットとしてはCPUに比べメモリ帯域が広い、
並列度が高いと言うだけだから、これらが有効に使えないのなら、CPUでやった方がいいよ。
もはやCUDAを使ったからX倍速くなりました!と言うことに価値はなくなってきたからな。
CUDAを使ったからYX倍(1<Y<3)速くなりました!というなら価値はありそうだ。
CUDAを使ったからZYX倍速くなりました!というならそれは比較がおかしい。

399 :デフォルトの名無しさん:2010/07/22(木) 04:01:35

ttp://journal.mycom.co.jp/articles/2010/07/21/fermi_cache/index.html
NVIDIAのFermiで新設されたキャッシュは効いているのか

>また、成瀬氏の実測ではGTX 480の1次キャッシュのアクセスレーテンシは約70ns、2次キャッシュのアクセスレーテンシは約250ns程度とのことで、
>CPUのキャッシュと比べるとこれでもキャッシュ? という程度の速度であるが、
>多数ワープを切り替えて実行する超マルチスレッド実行であるのでレイテンシを隠ぺいでき効果が出ているのであろう。

予想通り、パイプラインが深すぎてレイテンシがでかすぎるもよう

400 :デフォルトの名無しさん:2010/07/22(木) 10:06:37
>精度にこだわるならCUDAはやめときな
学者共に言ってやってくれ

401 :デフォルトの名無しさん:2010/07/22(木) 11:55:37
学者はFermi版Tesla(Tesla版Fermi?)を使って倍精度で計算するので問題ない・・・・

402 :デフォルトの名無しさん:2010/07/22(木) 12:45:29
念のため言っとくが単精度floatは32ビット整数すら誤差なく格納することができないからな!

403 :デフォルトの名無しさん:2010/07/22(木) 19:48:10
>>402
何この人恥ずかしい・・・

404 :デフォルトの名無しさん:2010/07/22(木) 21:49:08
下には下がいるってことがよくわかる

405 :デフォルトの名無しさん:2010/07/23(金) 00:02:17
っていうか別に倍精度の値域が必要な訳じゃねえんだろ?
1.0e-7以下の精度が必要な訳じゃねえんだろ?
うぜえよ

406 :デフォルトの名無しさん:2010/07/23(金) 00:22:22
だけど、もともとカオス性を備えた系のシミュレーションだとどんだけ精度があってもねえ

407 :デフォルトの名無しさん:2010/07/24(土) 16:46:26
Fortranから呼び出す時って二次元配列の順番変わっちゃうんでしょ?
操作する時に列から回すとメモリが連続でなくなっちゃってアクセス遅くなったりする?

408 :デフォルトの名無しさん:2010/07/24(土) 21:04:41
データの連続方向にスレッドを並べないと、滅茶苦茶遅くなります。
スレ違いにはなりますが、CPUの場合もデータ領域の大きさが大きい場合にはキャッシュ効率が全く変わってしまいます。

409 :デフォルトの名無しさん:2010/07/24(土) 21:30:40
そういえばMSDNかどっかにサンプルがあったなあ
配列の順番変えるだけで10倍だか100倍だかのオーダーで実行速度が変わるやつ
確か方っぽは徹底的にキャッシュミスする並べ方らしかった

410 :デフォルトの名無しさん:2010/07/24(土) 22:16:57
一次元配列も
1...1|2...2|...|n...n|
よりも
1...n|1...n|...|1...n|
の方がいいのかなあ。数字はスレッド番号ね。

411 :デフォルトの名無しさん:2010/07/24(土) 23:51:35
>>410
CUDAの話なら実験してみるといいよ。驚くから。

412 :デフォルトの名無しさん:2010/07/25(日) 00:01:44
2次元配列も全部1次元にしてるわ・・・

413 :デフォルトの名無しさん:2010/07/25(日) 01:24:39
コンスタントメモリっていつ何時つかうのん?


414 :デフォルトの名無しさん:2010/07/25(日) 02:21:28
gpu側でそれほど大きくなくていいけど速い必要のあるテーブルを参照したいとき。

415 :デフォルトの名無しさん:2010/07/25(日) 10:07:32
たとえばnを外部から入力させてベクトルをn倍する関数を作りたいとき
コンスタントメモリはグローバルメモリでそんな早くないけど
一回キャッシュされたら全部のスレッドはnを読むからレジスタと同じくらい速くなる

て考えはおk?

416 :デフォルトの名無しさん:2010/07/25(日) 13:31:18
>>415
パラメータで渡せないなら定数メモリはありだね。

417 :デフォルトの名無しさん:2010/07/25(日) 14:30:13
>>415
最近のGPU使ってないから違うかもしれないけど、
コンスタントメモリって勝手にキャッシュとかされなくない?
まあ、ワープ内全スレッドからの同時アクセスなら1回で済むから
何回も読みに行かなければ問題はないだろうけど。

418 :デフォルトの名無しさん:2010/07/25(日) 15:53:11
あれ、コンスタントメモリってキャッシュする命令があるんだっけ。

419 :デフォルトの名無しさん:2010/07/25(日) 16:37:58
問答無用でキャッシュされます。

420 :デフォルトの名無しさん:2010/07/25(日) 18:43:06
もしかしなくてもソートって並列計算に向いてない??

421 :デフォルトの名無しさん:2010/07/25(日) 20:00:43
bitなんちゃらソートとか、並列計算用のソートアルゴリズムもあるにはあります。
CUDA SDKのbitonicを参照あれ。

422 :デフォルトの名無しさん:2010/07/25(日) 20:30:28
最近OpenGLの勉強を始めたんですが,
CUDAプログラミングガイドの3.2.8.1にあるサンプルを実行するには
どれくらいコードを追加すればいいのでしょうか?
初期化などを追加しても実行中に停止してしまいます.
ソースを載せますのでご教授お願いします.

#include "main.h"

#define Wwidth 512
#define Wheight 512

GLuint positionsVBO;
struct cudaGraphicsResource* positionsVBO_CUDA;

float timer = 0.;

__global__ void createVertices(float4* positions, float timer, unsigned int width, unsigned int height)
{
unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;

// uv座標を計算する
float u = x / (float)width;
float v = y / (float)height;
u = u * 2.0f - 1.0f;

// 単純なsin波形を計算する
float freq = 4.0f;
float w = sinf(u * freq + timer) * cosf(v * freq + timer) * 0.5f;
// 位置を書き込む
positions[y * width + x] = make_float4(u, w, v, 1.0f);
}

423 :デフォルトの名無しさん:2010/07/25(日) 20:31:22
void display()
{
// CUDAにより書き込まれたバッファオブジェクトをマップする
float4* positions;
cudaGraphicsMapResources(1, &positionsVBO_CUDA, 0);
size_t num_bytes;
cudaGraphicsResourceGetMappedPointer((void**)&positions, &num_bytes, positionsVBO_CUDA);

// カーネル関数を起動する
dim3 dimBlock(16, 16, 1);
dim3 dimGrid(Wwidth / dimBlock.x, Wheight / dimBlock.y, 1);
createVertices <<< dimGrid, dimBlock >>> (positions, timer, Wwidth, Wheight);

// バッファオブジェクトをアンマップする
cudaGraphicsUnmapResources(1, &positionsVBO_CUDA, 0);
// バッファオブジェクトからレンダリングする
glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT);
glBindBuffer(GL_ARRAY_BUFFER, positionsVBO);
glVertexPointer(4, GL_FLOAT, 0, 0);
glEnableClientState(GL_VERTEX_ARRAY);
glDrawArrays(GL_POINTS, 0, Wwidth * Wheight);
glDisableClientState(GL_VERTEX_ARRAY);
// バッファを交換する
glutSwapBuffers();
timer += 0.01;
glutPostRedisplay();
}
void deleteVBO()
{
cudaGraphicsUnregisterResource(positionsVBO_CUDA);
glDeleteBuffers(1, &positionsVBO);
}

424 :デフォルトの名無しさん:2010/07/25(日) 20:32:03
int main(int argc, char **argv)
{
// 明示的にデバイスを設定する
cudaGLSetGLDevice(0);
// OpenGLとGLUT初期化する
glutInit(&argc, argv);
glutInitDisplayMode(GLUT_RGBA | GLUT_DOUBLE);
glutInitWindowPosition(702, 128);
glutInitWindowSize(Wwidth, Wheight);
glutCreateWindow("Sample");

glutSetOption(GLUT_ACTION_ON_WINDOW_CLOSE, GLUT_ACTION_GLUTMAINLOOP_RETURNS);

glutDisplayFunc(display);

// バッファオブジェクトを生成し,CUDAに登録する
glGenBuffers(1, &positionsVBO);
glBindBuffer(GL_ARRAY_BUFFER, positionsVBO);
unsigned int size = Wwidth * Wheight * 4 * sizeof(float);
glBufferData(GL_ARRAY_BUFFER, size, 0, GL_DYNAMIC_DRAW);
glBindBuffer(GL_ARRAY_BUFFER, 0);
cudaGraphicsGLRegisterBuffer(&positionsVBO_CUDA, positionsVBO, cudaGraphicsMapFlagsWriteDiscard);

// レンダリングループを起動する
glutMainLoop();

deleteVBO();

printf("Exit...\n");

return 0;
}

425 :422:2010/07/25(日) 20:35:43
main.hの中で色々なヘッダを纏めてインクルードしてます.
不勉強なので問題外のレベルのコードだと思いますが,方向性のヒントでも貰えたら幸いです.
よろしくお願いします.

426 :デフォルトの名無しさん:2010/07/25(日) 20:58:06
CUDA化された姫野ベンチのソースってどっかで公開されていますか?

427 :デフォルトの名無しさん:2010/07/27(火) 02:26:16
最近CUDAについて学び始めた初心者です。
現在三重対角行列のトーマス法をcudaで実装しようとしているのですが、計算式にある再帰処理に手も足も出ない状態です。
調べてみるとcuda(GPU)ではサポートしていないと出てきます。
GPUでの処理を考えるとこれは当然であることは分かるのですが、カーネル部分に何か工夫を行えば動かす事は可能となるのでしょうか?

もし行えるのであれば、効率等は考えなくても良いので良ければご教示お願いします。


428 :デフォルトの名無しさん:2010/07/27(火) 02:28:13
再起処理を自前実装する

429 :デフォルトの名無しさん:2010/07/27(火) 04:51:48
>>426
自分で書いてみたら?
CUDAの練習にもなるし。

430 :デフォルトの名無しさん:2010/07/27(火) 08:25:33
cuda3.1から再帰は使える。
ただし、Fermiのみ。

431 :デフォルトの名無しさん:2010/07/27(火) 11:14:48
GPUだと再帰処理ができないのが何故当然なのですか。

432 :デフォルトの名無しさん:2010/07/27(火) 11:18:10
誰も当然とは書いていませんが・・・。
まぁ、動いたとしても遅いわけで。

433 :427:2010/07/27(火) 20:43:18
>>428
再帰処理を自前実装するとは、再帰処理を非再帰処理に変えて実装するという意味ですか?
それとも何か書き方があるのですか?


>>430
3.1は再帰使えるのですか?!
普通に書けるのなら、1度試してみたいです。

>>431
すみません。初心者なもので、GPUでは再帰はできないって認識で固定していてしまいました。
並列で動かす各スレッドのどれが終わっているか分からないので,再帰はできないという認識でした。




434 :デフォルトの名無しさん:2010/07/27(火) 21:20:51
各スレッドは各プロセッサにつきひとつ割り当てられて一度にプロセッサ数のスレッドが
スレッド÷プロセッサ回実行されるのですか?
それとも全スレッドはひとつのプロセッサに複数個割り当てられて
(ひとつのCUDAコアは同時に複数個のスレッドを処理できる?)
全スレッドが一回だけ同時に実行されるのですか?

435 :デフォルトの名無しさん:2010/07/27(火) 22:08:50
>>434
難解な表現ですが、どちらかというと後者に近いです。

CUDAコアは複数個のスレッドを交互に処理します。
またCUDAコアが管理できるスレッド数には限界があり、
スレッドが多い場合には一部のスレッドのみが交互に実行され、
それ以外のスレッドは実行中のどれかのスレッド(スレッドブロック単位)が完了するのを待ちます。


436 :デフォルトの名無しさん:2010/07/28(水) 05:09:13
>>435
交互に処理ってちょっと違和感あって資料見たら、本当に交互に処理するのね・・・。
メモリ関係のレイテンシ待ちしてる間暇だから他のワープも処理してやんよって感じか。

それにしてもFermiはL1/L2キャッシュもついてるし、便利ね。

437 :デフォルトの名無しさん:2010/07/28(水) 05:19:46
>>433
再帰って要するに戻り先アドレスとローカル変数をスタックに乗っけてから
同じ関数実行しなおすわけだから、
そのスタック構造のとこだけ自分でメモリ上に作ってやればループに置き換えられる。

結果的には大抵、ローカル変数を全部配列化して、再帰レベルカウントを配列インデックスに指定するような方法で
対応できる。CUDAの場合それをどこのメモリに置くかを適切に決めないといけないけど。

438 :デフォルトの名無しさん:2010/07/28(水) 08:27:25
>>435
なるほど。でも待つということは前者のようにも感じます
480個の切符売り場に数万人が並ぶかんじでしょうか。隣がすいたら移ることもできる?
同じスレッドブロックの仲間たちは同じ列に並んでいるのでしょうか。それとも同時に切符を買えているのですか?

また、同じwarpの仲間たちではどうでしょうか。halfwarp単位で同時にメモリアクセスってことは同時に切符を買えている??

439 :デフォルトの名無しさん:2010/07/28(水) 09:28:55
Maximum number of threads per block: 1024
Maximum sizes of each dimension of a block: 1024 x 1024 x 64
Maximum sizes of each dimension of a grid: 65535 x 65535 x 1

1ブロックに1024スレッドってのはまあわかるんだけど、
下の二つはどういうこと??
xyzの各次元の最大数ってことならブロックあたり1024*1024*64スレッドってあきらかに1024をオーバーするよね。
それとも最大数ってことは1024*1*1や256*2*2はいいけど1*1*1024はダメですよってこと?
最終的にスレッドの上限は65536*1024ってこと?

440 :デフォルトの名無しさん:2010/07/28(水) 10:18:24
>それとも最大数ってことは1024*1*1や256*2*2はいいけど1*1*1024はダメですよってこと?
そうです。

>最終的にスレッドの上限は65536*1024ってこと?
こちらは65535*65535*1024ですね。


441 :デフォルトの名無しさん:2010/07/28(水) 12:42:03
CUDAを最近始めた初心者です。
既存のプログラムをまずは利用しようと思い
ttp://kgxpx834.blog58.fc2.com/blog-entry-6.html
を参考にCUDAの開発環境を作成した後
ttp://forums.nvidia.com/index.php?showtopic=107474
のCUDAMCML.zipをダウンロードし.cuと.hのソースファイルを取り込んで
コンパイルしようとしたところ大量のエラーを吐かれました。
ヘッダーファイルだけはコンパイルできるのですがほかのファイルがコンパイルできません。
当然ビルドもできません。
環境設定で他にやることがあるのでしょうか?
よろしくお願いします。

442 :デフォルトの名無しさん:2010/07/28(水) 12:46:48
環境くらい晒せや

443 :デフォルトの名無しさん:2010/07/28(水) 15:04:09
具体的にどういうエラーが出てるのか書けよ馬鹿垂れ

444 :デフォルトの名無しさん:2010/07/28(水) 20:41:02
GTX 480でキャッシュが使われているかどうか判定する方法はあるのでしょうか?
プロファイラにかけてみるとl1_global_load_hitが0で、
l1_global_load_missには値が入っています。
キャッシュが有効利用されていないと思っているのですが、どうなのでしょう。

445 :デフォルトの名無しさん:2010/07/28(水) 23:15:00
>>441
ここの人たちは他と違ってかなり優しいけど、
はっきり言っておまえの質問は酷すぎる。
ゆとり世代か?
そもそもCUDA初心者と言っている時点でC言語の初心者じゃないのか?
環境構築はCUDA以前の問題だからな。

446 :デフォルトの名無しさん:2010/07/29(木) 00:36:28
>>444
そもそもデータを再利用するような計算なのでしょうか?
一回読んで終わりならそうなる。

447 :439:2010/07/29(木) 10:27:45
>>440
ありがとうございます。
二番目と三番目で表現は一緒なのに内容は違うんですね

448 :デフォルトの名無しさん:2010/07/29(木) 20:19:44
>>438

>>436をみたかんじ
切符をたのんで財布からお金をだしてるあいだに窓口は次の人の買う切符を聞いてるかんじかね

449 :デフォルトの名無しさん:2010/07/29(木) 20:39:06
>>444
__syncthreads()入れてる?

450 :デフォルトの名無しさん:2010/07/29(木) 21:49:51
>>448
同じ人が再び列に並んできっぷを買い続けないといけないので例としては適切でないかと。

あれだ・・・・火縄銃を3人交代で打つ戦法・・・何の例だか分からなくなってきたが。

451 :デフォルトの名無しさん:2010/07/29(木) 23:25:15
卓球ダブルス

452 :444:2010/07/29(木) 23:30:21
>>446
そうなのですか。
最も今回のコードが1回の読み込み(ブロック内ですよね?)にあたるか
よく分からないのですが…
コアレスメモリアクセスにしてもキャッシュにヒットしないのは
ブロック単位で計測されているからでしょうか?
それともコードが間違ってて本当にキャッシュが効いていないのでしょうか?
キャッシュが何か良く分からなくなってきました。

>>449
色々なとこに入れてみましたけど効果はありませんでした。

453 :デフォルトの名無しさん:2010/07/30(金) 04:17:09
>>452
キャッシュとはデータが再利用されなれば意味がないよ。
プリフェチと勘違いしてない?
あとコアレッシングにしても速さが変わらないとなると、L2が効いているのかも。
いずれにせよ、まずはShared Memoryで書いて、効果が出るかどうかを調べることだね。
つか、そんな質問する位ならコードさらせや。

454 :デフォルトの名無しさん:2010/07/30(金) 21:05:15
FortranでCUDAができるというような話を聞いたんですが、もうできるのでしょうか?
それは純粋にFortranで書くのですか?
それともCで書いた関数をFortranで呼び出して使うという方法でしょうか?

455 :デフォルトの名無しさん:2010/07/30(金) 22:08:11
金払いたくないならFortranからCを呼び出す
PGIからFortran用CUDAコンパイラが出てるからそれ買えば全部Fortranで書ける。
ただしFortran90以降の書き方じゃないとコンパイルできないっぽい。

興味あったら15日体験版あるから試してみたら?

456 :454:2010/07/31(土) 12:46:48
>>455
ありがとうございます。
FortranからCを呼び出しても実行速度やらはかわらないのでしょうか。
プログラミングの情報はCでのCUDAの方が多いようなのでどちらも同じなら
C呼び出しの方を使いたいと思うのですが

457 :デフォルトの名無しさん:2010/08/02(月) 10:20:03
>>456
FortranからCを呼び出す部分は通常演算の核になる部分から離れているので余り問題にならないはず。
と言うより、それが問題になるようなら演算の核になる部分が小さすぎるからCUDAに向かない実装と言うことになる。
一度CUDAのカーネルを呼び出したら、最低でも100msは戻ってこないくらい演算を集中させないと
カーネル呼び出しのコスト自体がネックになってしまう。

458 :デフォルトの名無しさん:2010/08/02(月) 12:01:59
ブロックごとのスレッド数って多い方がいいんでしたっけ?
array[3000][N][A][B][C]の配列で、N,A,B,Cは結構変動するのですが、
3000*Nブロック(x=3000,y=N)のA*B*Cスレッド(x=A,y=B,z=C)にしようと思っているのですが、
A,B,Cがそれぞれ1であることも多々あります。その場合なんかもったいない使い方なのでしょうか

459 :デフォルトの名無しさん:2010/08/02(月) 12:14:31
>>458
実際にやってみれば判るけど、1スレッドしかなくても一つのサブプロセッサが動く。
つまりブロック数が同じとき、スレッド数が1でも2でも同じ時間が掛かる。
また、並列して投入できるカーネル関数がないなら空いているサブプロセッサがあっても使われない。
従って、使うgpuのコア数分投入した方がいい。
# 尤も、GUI表示と兼用ならその分は少なくても無駄は出ないけれど。

460 :458:2010/08/02(月) 15:08:29
>>459
ブロック1000スレッド1より
ブロック1スレッド1000のほうがいんですね
>使うgpuのコア数分投入した方がいい。
というのはスレッド数をってことですか?(480基だったら400くらい?)

ひとつのコアにつき同じブロック内のスレッドがひとつづつバーーっと割り当てられて
ブロック内のスレッドは各々のコアが同時に処理するのでしたっけ

461 :デフォルトの名無しさん:2010/08/02(月) 15:37:48
あ、一部説明が拙かった。
論理上のスレッド・ブロック・グリッドと、実際のプロセッサ・マルチプロセッサ(MP)・デバイスでは若干違っている。

論理上のスレッドは実際のMP内のプロセッサに1vs1で割り当てられる。
スレッド数が32を超える場合は複数のMPに順に割り当てられる。
但し、同時に使用できるMP数を越える場合は実行が保留され、MP単位で随時処理される。
従って、スレッド数は32の倍数でデバイス辺りのプロセッサ数を上限とする程度でいい。
尚、共有メモリはブロック内で共有という仕様なので、MPに収まらない範囲で利用すると
グローバルメモリに退避することになるのでパフォーマンスが低下する。このため、共有メモリを使うならスレッド数は多くない方がいい。
私の場合、32、64、96、128、256辺りにしていることが多い。

論理上のブロックはMPに1vs1で割り当てられる。
ブロック間では共有できる資源がないのでMPの空きに応じて適宜割り当てられる。
当然、一度に処理できなければ随時処理される。
従って、デバイス辺りプロセッサ数をスレッド数で割った数より大きければ特に上限はない。
VIP-Wikiでは1000という例が挙がっている。
ブロック数*スレッド数が必要な処理数に満たないときはループを組んで処理することになるが、
カーネル内でループを作るのを避けてブロック数を増やすのも選択肢の一つ。


462 :458:2010/08/02(月) 17:52:24
ばかですみません

MP内のコアは8個でしたっけ
32スレッドが一つのwarpという単位でランダムな順番、割り振りで(この場合の32スレッド自体は連続なID?)
MP1内の8個のコアで同時期に順次に処理
33スレッド目から64スレッド目まではMP2内に割り当てられ、処理
MPが15個の場合、32*15+1スレッド目からは同時に割り当てられるMPがないため保留して
どっかのMPが空き次第、(連続の?)32スレッド単位で空いたMPに割り当てられ、処理
つまりこの場合1ブロックが32*15スレッドまでなら一回で同時に割り当てられるのでそれ以下が望ましい

共有メモリはブロック内で共有するが、物理的にはMP内に一つあって8コアで共有する形なので
1ブロックが32スレッドを超えるとMP1からMP2の共有メモリにアクセスする場合
グローバルメモリを介さなくてはならない(このアクセスは1ブロック内の共有メモリアクセスなので論理的には速いはず)
あれ?つまり1ブロックは32スレッド以下が望ましいってことに??

ここまではこんな解釈でよろしいのでしょうか??

(あと、blockDim.x,y,zは利便性のためだけで、大事なのはスレッド数であって
x,y,zをどうとろうがx*y*zが同じなら物理的な違いはないのでしょうか)


463 :458:2010/08/02(月) 18:19:08
>論理上のブロックはMPに1vs1で割り当てられる。
ここもよくわからなかったのですが、
どんなにスレッド数が少ないブロックでも一つのMP内には1ブロックしか存在できないという意味でしょうか
2スレッドのブロックが4つあったとき、8スレッドしかないのでスレッド数的には1MPでおさまるはずだけど
実際は4MP使うよと。

64スレッドのブロックが4つあったときはMP1とMP2に1ブロック目が割り当てられ、
計8MP使う

32*14スレッドのブロックが2つある場合、1ブロック目がMP14まで割り当てられ、
2ブロック目の最初32スレッドだけがMP15に割り当てられ、あとは保留
MPが空いたら2ブロック目の続きを32スレッドづつ順次割り当て、実行という感じでしょうか

長乱文失礼


464 :デフォルトの名無しさん:2010/08/02(月) 18:20:33
>>461
書いていることがめちゃくちゃなのだが・・・・。
もしもネタでないのであればじっくり読み返していただきたい。

465 :461:2010/08/02(月) 19:22:16
あれ? そんなに間違ってたっけか。知識が古いからFermiでの数値と違っているかもしれないけれど。
後は用語がいい加減かな。
まぁいいや、誰かの修正を待とう。

あー、共有メモリの下りは明らかにおかしいな。MP間で同期を取ると遅くなるけどグローバルメモリに退避しているわけじゃないか。

それはさて、ブロックは論理的に独立だからMP内で共存できないのはあっていると思う。
少なくとも、ブロック数が充分あるときにスレッド数1とスレッド数4とで所要時間が変わらなかった記憶がある。

それと、スレッドブロックのxyz次元の差はなかったと思うけど、これは実験結果を残してないから記憶も曖昧だな。

ってことで、滅茶苦茶だそうだから以降は自粛。

466 :デフォルトの名無しさん:2010/08/02(月) 19:49:55
GT200系までだと、1SM=8SP。Fermiはしらん
32スレッドで1Warpなのは1つのSMが4クロックの間、threadIdだけ変えて同じ処理を行うから
だからスレッド数は32の倍数が効率がよい

32を超えると32単位で実行順は保証されないが、1ブロック内のスレッドは同一のSMで実行される
だからSM数≦ブロック数になるのがよい

467 :デフォルトの名無しさん:2010/08/02(月) 22:03:19
ブロック数に関してはSM数の倍数になるのがいい。
これは1つのブロックは1つのSMに割り当てられて実行されるため。
ただしSM数<<ブロック数だったら無理に合わせる必要はあまりない

スレッド数は32の倍数がよいが、具体的に何倍がいいかはコードによってかわる。
それを計るのがプロファイラやnvcc --ptxas-options=-vでえられるOccupancy。
これは基本的にはレジスタやShared Memoryの使用率を元に、SM内でブロックを同時起動した際に合計でどれぐらいレジスタ、Shared Memoryの無駄が少ないか?を指標化したもの。

1ブロックあたりの使用率が高い場合(=スレッド数が多すぎる)はOccupancyが低くなり、1つのSM内で同時に実行できるブロックが少なくなる。
同時に起動できるブロックが少ないと、あるブロックが同期待ち、終了待ちをしている時にSMの実行時間に無駄が生じる。

一応、ブロックあたり192スレッド(6 warps)以上、同時に起動できるブロック数は2以上がいいとされている。
なので>>458への回答としては、なるべく多い方がいいが、32の倍数でかつOccupancyが高くなる数にするのがよい。

>>459-465の話(特に共有メモリ云々)はめちゃくちゃなので、CUDAの実行モデルを再考することを勧める。

468 :458:2010/08/04(水) 10:22:43
Occupancy調べてみました。
tをスレッド数とするとwarp数w=[t/32]+1、
レジスタ数r、シェアードメモリsを出力して
copute capability2.0の場合は
Min([48/w],[32768/(r*t)],[49152/s],8)
ってかんじですかね。

>>459-467
お騒がせしました。ありがとうございます。
もう一度ハードウェア周りについて勉強してみます。

469 :デフォルトの名無しさん:2010/08/04(水) 14:44:24
カーネル関数の引数はどこに保存されるのだっけ
たとえばint n=1000があって
kernel<<<>>>(n);
とやるとnはレジスタに入ってくれるの?
グローバルメモリに入っちゃうの?
配列と普通の変数で違ってくるみたいなレスを昔みたような気がするんだけど

470 :デフォルトの名無しさん:2010/08/04(水) 15:57:37
シェアードメモリに入る

471 :469:2010/08/04(水) 17:12:34
>>470
ありがとう。
よく考えたら-cubinで自分で調べられたね。すみません。
で、やってみたんだけど
最初smem=40のカーネルに引数intをひとつ加えるとたしかに4増えて44になった。
そこにさらに引数doubleをひとつ加えてみたら8増えて52になるはずが12増えて56になっちゃった
あれっと思ってsmem=40の状態に戻してdoubleを加えるとちゃんと8増えて48になった。
さらにintを加えると52に。
つまり40の状態から引数をint→doubleって加えるのとdouble→intって加えるのでは
使用smemが違ってきちゃった。
なぜーー??


472 :デフォルトの名無しさん:2010/08/04(水) 17:54:00
blockサイズ、threadサイズも入るから

130 KB
■ このスレッドは過去ログ倉庫に格納されています

★スマホ版★ 掲示板に戻る 全部 前100 次100 最新50

read.cgi ver 05.02.02 2014/06/23 Mango Mangüé ★
FOX ★ DSO(Dynamic Shared Object)