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

[表示 : 全て最新50 1-99 101- 201- 301- 401- 501- 601- 2chのread.cgiへ]
Update time : 05/09 23:27 / Filesize : 186 KB / Number-of Response : 679
[このスレッドの書き込みを削除する]
[＋板最近立ったスレ＆熱いスレ一覧 : ＋板最近立ったスレ／記者別一覧] [類似スレッド一覧]


↑キャッシュ検索、類似スレ動作を修正しました、ご迷惑をお掛けしました

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

CUDA・HomePage
www.nvidia.com/cuda

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

前スレ
【GPGPU】くだすれCUDAスレ【NVIDIA】
pc12.2ch.net/test/read.cgi/tech/1206152032/
【GPGPU】くだすれCUDAスレ pert2【NVIDIA】
pc12.2ch.net/test/read.cgi/tech/1254997777/
2 名前：デフォルトの名無しさん mailto:sage [2010/04/18(日) 19:49:11 ]: 関連サイト
CUDA Zone
www.nvidia.co.jp/object/cuda_home_jp.html
フォーラム
forum.nvidia.co.jp/EokpControl?&event=TE0001

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

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

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

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

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

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

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

KNOPPIX for CUDA
www.yasuoka.mech.keio.ac.jp/cuda/
3 名前：デフォルトの名無しさん mailto:sage [2010/04/18(日) 21:42:11 ]: いちおつ
4 名前：デフォルトの名無しさん mailto:sage [2010/04/19(月) 00:54:58 ]: 一乙、前スレの、魔二来vsTeslaのコストパフォーマンス比較面白いな
もっとやってくれ
5 名前：デフォルトの名無しさん mailto:sage [2010/04/19(月) 05:45:00 ]: >>4
もう結論でたよ

理論ピーク性能はTeslaの勝ち
大部分のプログラムでの実行性能はマニクールの勝ち
プログラミングのしやすさではマニクールが圧勝
6 名前：デフォルトの名無しさん mailto:sage [2010/04/19(月) 07:05:25 ]: すげーな
48コア使いこなせるとかｗ
7 名前：デフォルトの名無しさん mailto:sage [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 名前：デフォルトの名無しさん mailto:sage [2010/04/19(月) 07:38:14 ]: >>7
これって、いまいち理解出来ないのだが、誰か説明して。
9 名前：デフォルトの名無しさん mailto:sage [2010/04/19(月) 09:51:51 ]: >>7
将来が危ぶまれるNVIDIA専用のCUDAと心中するつもりなら
Teslaでもいいかもね
10 名前：デフォルトの名無しさん mailto:sage [2010/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 名前：デフォルトの名無しさん mailto:sage [2010/04/19(月) 12:35:22 ]: >>10
間違えた

×（総費用）＝（価格）X（コア数）
○（総費用）＝（価格）X（CPU数）
12 名前：デフォルトの名無しさん mailto:sage [2010/04/19(月) 14:12:24 ]: OpenMPとかえらい簡単に並列化できるしな。
あと何覚えればいいですか？
13 名前：デフォルトの名無しさん mailto:sage [2010/04/19(月) 15:20:32 ]: 普通のスレッド
MPI
14 名前：デフォルトの名無しさん mailto:sage [2010/04/19(月) 16:22:11 ]: >>10
OpteronってSSEとかが1コアあたり2個ある訳？
15 名前：デフォルトの名無しさん mailto:sage [2010/04/19(月) 16:24:58 ]: OpenMPもCUDA並に最適化が面倒。
その代わりMPIの知識があれば以外と簡単。
16 名前：デフォルトの名無しさん mailto:sage [2010/04/19(月) 18:03:05 ]: で、比較したのまだぁ？
17 名前：デフォルトの名無しさん mailto:sage [2010/04/19(月) 18:05:27 ]: そもそもC2050が出てないのに
比較なんてできるかよ
18 名前：デフォルトの名無しさん mailto:sage [2010/04/19(月) 18:22:54 ]: そもそもスレチ
19 名前：デフォルトの名無しさん [2010/04/19(月) 21:40:15 ]: 秋葉にオープンしたfaithにCUDAのブースができたね
実体はunitcomの宣伝かな
20 名前：デフォルトの名無しさん mailto:sage [2010/04/19(月) 22:03:30 ]: GTX480でデモしてくれるならすぐにでも行くが
21 名前：デフォルトの名無しさん mailto:sage [2010/04/19(月) 22:04:09 ]: 長崎大の部品を入れたとか書いてたね。CUDA話のできる店員さんいるのかしら。
22 名前：デフォルトの名無しさん mailto:sage [2010/04/20(火) 07:05:14 ]: チップセット・・・完全撤退
GPU・・・AMDの大攻勢の前に大苦戦
Tesla・・・CPUの高性能化・マルチコア化の前で微妙に

NVIDIAさんこれからどうするつもりなんでしょうね
23 名前：デフォルトの名無しさん mailto:sage [2010/04/20(火) 07:18:02 ]: 自作板に帰れ
24 名前：デフォルトの名無しさん mailto:sage [2010/04/20(火) 07:29:44 ]: SSEとかAVXとか覚えるよりも、
CUDAのお作法覚えるほうが簡単な気がしない？おれだけ？
25 名前：デフォルトの名無しさん mailto:sage [2010/04/20(火) 08:09:46 ]: >>24
でも将来性が最悪だね
遅かれ早かれOenpCLに移行するだろうし
26 名前：デフォルトの名無しさん mailto:sage [2010/04/20(火) 08:53:34 ]: CUDAとSIMDでは考え方が違うような気が。
個人的にはCUDAは割と大ざっぱに並べる。
SSEとかは細かなパズル。
自分でも何を言っているかわからんが、そういう思いでいる。
27 名前：デフォルトの名無しさん mailto:sage [2010/04/20(火) 08:57:23 ]: 前スレで480特攻すると書いた者だが
倍精度1/4と聞いてキャンセルしてしまった…
ゲーマーだったら失敗覚悟で注文したんだけど
28 名前：デフォルトの名無しさん mailto:sage [2010/04/20(火) 09:09:12 ]: ZOTACまだリリース前になってますね。キャンセルできてよかったですね。
ケース待ちだった方は購入できたのでしょうか。
29 名前：デフォルトの名無しさん mailto:sage [2010/04/20(火) 09:20:53 ]: マザーボード破損の書き込みもあるし手を出さなくて正解だろうね
30 名前：デフォルトの名無しさん mailto:sage [2010/04/21(水) 18:23:51 ]: カーネル中でcoutとかは使えないの？
どこが悪いのか調べるときはたいていcoutで変数を書き出していたのだけど、
CUDAの場合はどうすればいいの？
31 名前：デフォルトの名無しさん mailto:sage [2010/04/22(木) 02:17:10 ]: cudaMemcpy()でCPUに送る。
32 名前：デフォルトの名無しさん [2010/04/22(木) 05:13:57 ]: カーネルの途中で表示したいんだけど
33 名前：デフォルトの名無しさん mailto:sage [2010/04/22(木) 08:07:33 ]: どんな手法にしろCPU介さなかったら表示はできんよ
34 名前：デフォルトの名無しさん mailto:sage [2010/04/22(木) 08:32:27 ]: gpu.fixstars.com/index.php/Cuda-gdbを使う

自分はコーディングレビューしかしないけど
35 名前：デフォルトの名無しさん mailto:sage [2010/04/22(木) 09:45:25 ]: 結果とは別にデバイス側に確保しておいて、
あとで見るのはだめなの？
36 名前：デフォルトの名無しさん mailto:sage [2010/04/22(木) 10:12:38 ]: >>32
・deviceemuでなんとか頑張る
・可能な限り詳細な出力をメモリに吐き出しておいて、ホストに戻ってから(転送して)解析する

それ以前に、そのレベルだとcudaの細かいノウハウを掴んでないだろ。
然るべき業者に委託したほうがよくないか?
37 名前：デフォルトの名無しさん mailto:sage [2010/04/22(木) 20:00:42 ]: なぜ仕事前提？
38 名前：デフォルトの名無しさん [2010/04/22(木) 22:26:01 ]: >>30
とゆうか、そうゆう発想する事態C++もCUDAも全然理解できてないよねｗ
39 名前：デフォルトの名無しさん mailto:sage [2010/04/23(金) 01:36:42 ]: CPUで考えると、「レジスタの計算の途中をCoutで見ることは出来ないか？」
と聞いているようなものだと思うのだが。
それを自前でやるにはレジスタの内容を一度メモリにコピーしてcoutの標準出力で表示させることになる。
それと一緒だと思うよ。
40 名前：デフォルトの名無しさん mailto:sage [2010/04/23(金) 02:53:45 ]: device emuで解決ジャン
41 名前：デフォルトの名無しさん mailto:sage [2010/04/23(金) 03:31:06 ]: 待てればね。
42 名前：デフォルトの名無しさん mailto:sage [2010/04/23(金) 03:47:59 ]: エミュレーションモードってなくなるんだろ
43 名前：デフォルトの名無しさん mailto:sage [2010/04/23(金) 05:53:13 ]: すでに実機があるからなあ。
44 名前：デフォルトの名無しさん mailto:sage [2010/04/23(金) 06:50:27 ]: ダンプ用の領域を用意しておいて実行中に書き込んで後で見る
しかないんじゃね？
45 名前：デフォルトの名無しさん [2010/04/23(金) 07:05:41 ]: ubuntu 10.04でcudaが使えるようになるのはいつごろですか
46 名前：デフォルトの名無しさん mailto:sage [2010/04/23(金) 08:28:37 ]: GPUを搭載していないノートPCでエミュレーションモードを使ってたのに
なくなるってマジ？
47 名前：デフォルトの名無しさん mailto:sage [2010/04/23(金) 11:36:52 ]: マジ
48 名前：デフォルトの名無しさん mailto:sage [2010/04/23(金) 11:42:20 ]: リモートデスクトップでCUDA使えればいいのになー。
49 名前：デフォルトの名無しさん mailto:sage [2010/04/25(日) 16:17:32 ]: リナカフェいったひといないすか
50 名前：デフォルトの名無しさん mailto:sage [2010/04/25(日) 16:22:48 ]: サトリナのカフェなんてあるの？
51 名前：デフォルトの名無しさん mailto:sage [2010/04/26(月) 10:35:43 ]: このスレもしぼんだなあ
GTX480が倍精度外れた（らしい）のが効いているんだろうな
52 名前：デフォルトの名無しさん mailto:sage [2010/04/26(月) 11:17:15 ]: 倍精度演算ユニットは少なくなっているだけで無くなったわけではないのだが、
そもそも480が入手できないことの方が原因ではないかと。
53 名前：デフォルトの名無しさん mailto:sage [2010/04/26(月) 15:43:31 ]: 結局口だけの貧乏人しか居なかったって事だなｗ
54 名前：デフォルトの名無しさん mailto:sage [2010/04/26(月) 17:33:22 ]: おまえもな
55 名前：デフォルトの名無しさん mailto:sage [2010/04/26(月) 19:14:13 ]: www.amazon.co.jp/dp/479802578X/
これIPAの岡ちゃんが書いたの？
いつの間にかネットマイスターってところに転勤したのかな
56 名前：デフォルトの名無しさん mailto:sage [2010/04/26(月) 19:24:07 ]: もう入門書もでてるから、新刊はいらないな。
需要はあるのだろうか･･･。
57 名前：デフォルトの名無しさん mailto:sage [2010/04/26(月) 19:28:27 ]: つーか、ネットマイスター代表だってよ。
同名をtwitterで検索したらttp://twitter.com/kenmoi0715 が引っ掛かったけど……
58 名前：デフォルトの名無しさん mailto:sage [2010/04/27(火) 01:07:50 ]: 流石にロリコンばらされたら職場には居にくいか……
アプリも割り放題だったしな。
たくましく生きているようで何よりww
59 名前：デフォルトの名無しさん mailto:sage [2010/04/27(火) 01:20:01 ]: この本買ったけど、ちょっと入門過ぎる気がするなあ。
「高速」という言葉にてっきり「最適化」ということを連想してしまった。
環境設定とかで半分以上割いているよ。
まあ、CPUコードと両方書いてあって、Webからダウンロード出来るから、
最初にでた本よりかはいいかもしれない。
60 名前：デフォルトの名無しさん mailto:sage [2010/04/27(火) 08:02:14 ]: 自分で買うほどじゃないかな
図書館に購入希望だしとこう
61 名前：デフォルトの名無しさん mailto:sage [2010/04/27(火) 21:19:54 ]: 本にするほどの蓄積と技術の咀嚼があるとは思えないけどどうなんでしょ。
あるとすればまだ論文レベルなんじゃないか？

逆に、まだ広まってない今が新しいことするチャンスかっ！
62 名前：デフォルトの名無しさん mailto:sage [2010/04/27(火) 21:27:03 ]: プログラムの入門書なんてこんなもんだろ
63 名前：デフォルトの名無しさん mailto:sage [2010/04/28(水) 04:21:24 ]: GF100が失敗、GF104はキャッシュ削るみたいでGPGPUは絶望的の現在、今後広まる可能性が残っているかどうか…
64 名前：デフォルトの名無しさん mailto:sage [2010/04/28(水) 05:19:00 ]: は？
65 名前：デフォルトの名無しさん mailto:sage [2010/04/28(水) 05:45:52 ]: GPGPUに残されているのは個人では動画エンコードくらいか。
それなりの規模になると、マルチプロセッサのシステムを組んで、
CPUのコードを走らせた方が開発工数から考えると速いからなあ。
66 名前：デフォルトの名無しさん mailto:sage [2010/04/28(水) 06:14:57 ]: で、おまえは何しにここに来てる訳？
67 名前：デフォルトの名無しさん mailto:sage [2010/04/28(水) 06:17:08 ]: ばかだな
ただのAMD信者様だろｗ
68 名前：デフォルトの名無しさん mailto:sage [2010/04/28(水) 09:19:52 ]: GPGPUに将来性はないだろうね。
牧野もそういっているし。
69 名前：デフォルトの名無しさん mailto:sage [2010/04/28(水) 09:29:19 ]: コスト面がちょっと微妙かなあ
C10シリーズのようにアカデミックで10万！とかなら個人で使ってみようという流れになるだろうに
GTX480は倍精度1/4にしました！C20シリーズを売るためです！と公言されちゃったし
70 名前：デフォルトの名無しさん mailto:sage [2010/04/28(水) 09:52:12 ]: >>69
コスト面つか買う気失せることすんなよなー　て感じっすよ
71 名前：デフォルトの名無しさん mailto:sage [2010/04/28(水) 10:00:12 ]: 　　　　　　　､--‐冖'⌒￣￣`ｰ-､
　　　　／⌒`　　　　　　　　　三ミヽｰ-ﾍ,_
　　　__,{　;;,, 　　　　　　　　　ミミ　　　ｉ　´Z,
　　ゝ　　　''〃//,,, 　　　　,,..｀ミﾐ､_ノリ}j;　f彡
　　_)　　　　　〃///,　,;彡'rffｯ､ｨ彡'ﾉ从iﾉ彡
　　>';;,, 　　　　ノ丿川j !川｜;　 :.`7ﾗ公 '>了
　_く彡川fﾞノ'ノﾉノ_ノノﾉイシﾉ｜ }.:　'〈八ﾐ､､;.)
　ヽ.:.:.:.:.:.;=､彡／‐-ニ''_ｰ＜、{_,ﾉ -一ヾ`~;.;.;）
　く .:.:.:.:.:!ﾊ.Yイ　　ぇ'无ﾃ,`ヽ｝｝}ィt于｀|ィ"~
　　 ):.:.:.:.:|.Y }: :!　　｀二´/'　;　|丶ﾆ　ﾉﾉ
　　） :.: ﾄ､リ: :!ヾ:、　　　丶　;　|　ﾞ　　ｲ:}　　　　逆に考えるんだ
　　｛ .:.: ｌ　｛: : }　｀　　　　,.__(__,}　　　/ﾉ
　　　ヽ !　　`'ﾞ! 　　　　 ,.,,.`三'ﾞ､,_ 　/´　　　「単精度なら４倍」と
　　 ,/´{　　ﾐ　l　　　　/ﾞ,:-…-～､） |
　　,r{ 　＼ﾐ　＼　　 `'　'≡≡'　" ﾉ　　　　　　　　考えるんだ
__ノ　ヽ　　＼　ヽ＼　　　　彡　,ｲ_
　　　　＼　　＼ヽ丶. 　　　　ノ!|ヽ｀ヽ、
　　　　　＼　　＼ヽ｀¨¨¨¨´/ |ｌ　ト､　｀'ｰ-､__
　　　　　　＼　　｀'ｰ-､　　/／　/:.:.}　　　　｀'ｰ､_
　　　　　　　　｀､＼　　　／⌒ヽ　 /!:.:.|
　　　　　　　　　　｀､＼ /ヽLｆ___ハ/　　{
　　　　　　　　　　　　　 ′ /　! ヽ
72 名前：デフォルトの名無しさん [2010/04/28(水) 10:01:53 ]: 長崎大のセンセが8800GTで激安スパコン組んだときも、
そのうち２、３割がGPGPU用としてはぶっ壊れていたじゃん。

今度のは、そのままだと半分以上壊れたものを出荷してるから、まずいってんで、
止めてるんでないの。
73 名前：デフォルトの名無しさん [2010/04/28(水) 12:24:26 ]: 牧野もそういっているし。
牧野もそういっているし。
牧野もそういっているし。
牧野もそういっているし。
牧野もそういっているし。

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

おばかさん
75 名前：デフォルトの名無しさん mailto:sage [2010/04/29(木) 01:52:54 ]: 金にならんかったら終わるだろ。
76 名前：デフォルトの名無しさん mailto:sage [2010/04/29(木) 09:55:55 ]: GPGPUはマルチスレッドな浮動小数点演算に特化したものだからな、使えるシーンはかなり限られてくる
77 名前：デフォルトの名無しさん [2010/04/29(木) 10:26:40 ]: forループの一部だけを実行する方法を教えてください
78 名前：デフォルトの名無しさん mailto:sage [2010/04/29(木) 12:24:55 ]: >>55
やっぱりこういう問い合わせが多いんだろうな
著者紹介にIPAに勤務したことはないって書いてあったよ
79 名前：デフォルトの名無しさん mailto:sage [2010/04/29(木) 12:29:11 ]: >>59
俺も買った。
全く同じ感想。
青木先生の本が理解できる人は全く買う必要なし。
>>55
本の経歴の部分に、「IPAに勤務したことはない。」と不自然に書いてあったよw
80 名前：デフォルトの名無しさん mailto:sage [2010/04/29(木) 13:44:48 ]: なんだ同名で迷惑している人なのか
81 名前：デフォルトの名無しさん mailto:sage [2010/04/29(木) 15:21:40 ]: んなわけない。
IPAを除名されたか、退職するときにそれに準じる誓約書を書かされたのだろう。
82 名前：デフォルトの名無しさん mailto:sage [2010/04/30(金) 02:18:54 ]: 証拠は？
83 名前：デフォルトの名無しさん mailto:sage [2010/04/30(金) 05:35:43 ]: >>79
今それを見てワロタ。
「IPAに勤務したことはない。」
で株式会社ネットマイスターをググってみたけど、
すぐには見つからなかった。
でもこういった本を書く人が流出なんてするのだろうか？
84 名前：デフォルトの名無しさん mailto:sage [2010/04/30(金) 08:01:21 ]: GF100の板に出ていてみたいだけど、
GTX480でCUDA動作に不具合があるとか？
ドライバ改善待ち、みたい。
GTX480でのベンチマーク結果（単精度、倍精度問わず）が出ない理由はそれだったりするのかも？

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

いつかは出るであろう、Quadroにかすかに期待するしかないのか…
85 名前：デフォルトの名無しさん mailto:sage [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 名前：デフォルトの名無しさん mailto:sage [2010/05/01(土) 15:36:56 ]: >>85
もちろん重要ありますよ。
87 名前：デフォルトの名無しさん mailto:sage [2010/05/01(土) 16:09:33 ]: 【Fermi】GeForceGTX4xx総合Part31【GF100】
pc11.2ch.net/test/read.cgi/jisaku/1272318075/
の548以降でGTX480を含むGPUのベンチマークが出ている
LU分解のようです
倍精度に制限かかっているのはどうやら本当みたいだ
88 名前：デフォルトの名無しさん mailto:sage [2010/05/01(土) 18:13:38 ]: 具体的にどう制限かかっているんだろ？
それによってCUDAのコードの書き方が変わってくると思う
89 名前：デフォルトの名無しさん mailto:sage [2010/05/01(土) 22:13:39 ]: >87
いや倍精度どころか単精度でも全然性能でてないわけだが
90 名前：デフォルトの名無しさん mailto:sage [2010/05/01(土) 23:17:23 ]: なんかそのベンチ遅くね
285の単精度で16Gflopsとか書いてあるけど
91 名前：デフォルトの名無しさん mailto:sage [2010/05/01(土) 23:25:48 ]: GTX285と比べる限り、GTX480の倍精度に期待できそうな気配はないなあ
1/4は本当くさいな
92 名前：デフォルトの名無しさん mailto:sage [2010/05/02(日) 16:32:28 ]: cublasのgemmを使ったベンチをだれかが提供すべきだな
もちろん、俺はやる気はないが
93 名前：デフォルトの名無しさん mailto:sage [2010/05/03(月) 10:55:57 ]: gemmもやって欲しいけど、gemmはピーク性能しか出てこないから俺はLU分解やFFTの方が多少は信用している。
94 名前：デフォルトの名無しさん mailto:sage [2010/05/03(月) 12:02:20 ]: というか一番多用しそうなFFTでどれだけ実効性能出るのかが気になる
95 名前：デフォルトの名無しさん mailto:sage [2010/05/03(月) 16:43:35 ]: >>94
例えばどんなサイズが？
96 名前：デフォルトの名無しさん [2010/05/04(火) 18:08:02 ]: journal.mycom.co.jp/special/2008/cuda/006.html
ここのコードについて質問なんだけど
これはマルチスレッドで実行する関数の内部だよね？
ということはAsやBsのデバイスメモリからシェアードメモリへの転送や、
最後のホストの変数Cへのデータ転送は無駄に全スレッドがやるの？
それとも同じデータの転送については１スレッドがやってる？
97 名前：デフォルトの名無しさん mailto:sage [2010/05/04(火) 19:47:43 ]: > これはマルチスレッドで実行する関数の内部だよね？
そのとおり。

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

> 最後のホストの変数Cへのデータ転送は
今回のコードには入ってない。ホスト側のコードで、ブロック転送を指示する。
98 名前：デフォルトの名無しさん mailto:sage [2010/05/04(火) 20:53:45 ]: >>97
ｔｈｘ
99 名前：デフォルトの名無しさん [2010/05/07(金) 12:16:01 ]: GCC 4.5.0で使えるようにする設定を教えてください
100 名前：デフォルトの名無しさん mailto:sage [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を３回使えばいいけど、
もっと速い方法があるかな。
102 名前：デフォルトの名無しさん mailto:sage [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 名前：デフォルトの名無しさん mailto:sage [2010/05/08(土) 01:11:54 ]: 続き・・・

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

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

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

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

ハードの知識がなくてもptx出力を読めてprofileの意味を理解できればいいのだけれど、
それって結局ハードの知識を得るのと同じことだしね。
105 名前：デフォルトの名無しさん mailto:sage [2010/05/08(土) 03:14:44 ]: 結局nがいくつかによるだけじゃない。
たぶん100はCUDAのくの時もまだ理解していないと思う。
106 名前：デフォルトの名無しさん mailto:sage [2010/05/08(土) 03:16:30 ]: >>104
あの程度だとコンテキストスイッチやホスト・デバイス間の転送等の
呼び出しコストに見合わないね。

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

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

ptx読めてprofileの意味を理解するにはアーキテクチャの理解が必要というか
もうその段階はとっくに超えているような気もする・・・
107 名前：デフォルトの名無しさん mailto:sage [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ではハーフワープあたり無駄なメモリロードが２回走る。
CC2.0ではキャッシュが効くのでARRAY_SIZE_Nのサイズしだいでは問題にならない可能性がある。
108 名前：デフォルトの名無しさん mailto:sage [2010/05/08(土) 03:45:08 ]: >>107
厳しい突っ込みどうも。
CUDA Coreと呼ばれるようになったのはSMじゃなくてSPだったのですね。
図を見れば一目瞭然ですね・・・うろ覚えで書いてすみませんでした。

コアレスアクセスはまだよく理解していないので勉強してきます・・・
109 名前：デフォルトの名無しさん mailto:sage [2010/05/08(土) 04:15:14 ]: だから、16人17脚でパン喰い競争やるようなもんだよw
110 名前：デフォルトの名無しさん mailto:sage [2010/05/08(土) 05:18:45 ]: この表現はわかりやすいね。
111 名前：デフォルトの名無しさん mailto:sage [2010/05/08(土) 06:18:59 ]: 一人重い荷物を持たされると、全体の速度が遅くなる、的なｗ
112 名前：デフォルトの名無しさん mailto:sage [2010/05/08(土) 08:19:25 ]: 青木本で最も評価するのがあの手書きの絵だww
113 名前：デフォルトの名無しさん mailto:sage [2010/05/08(土) 09:18:32 ]: 青木本にこのスレから引用したと見られる節が幾つかある件w

>>109
>107の例だと、自分のレーンのパンだけじゃなくて左右のパンも食べないといけないんだな。
そりゃ遅くなるわさ。
114 名前：デフォルトの名無しさん [2010/05/08(土) 11:04:49 ]: コアレスのコアってなんのコアのことなの?
115 名前：デフォルトの名無しさん mailto:sage [2010/05/08(土) 11:07:02 ]: coalescedなので、
合体したという意味。
合体してメモリにアクセスするって意味だ。
116 名前：デフォルトの名無しさん mailto:sage [2010/05/08(土) 12:01:16 ]: 正直その並びなら、texture使えばキャッシュ使えるじゃん

ただみんな、言うように並列度がどれくらい稼げるか(nに依存)で変わるから実際やってみないとわからないんじゃない？？
117 名前：デフォルトの名無しさん [2010/05/08(土) 12:24:17 ]: ↓新明解(旧版)で合体の項を引くなよ
118 名前：デフォルトの名無しさん mailto:sage [2010/05/08(土) 13:18:40 ]: メモリにアクセスするときに、各スレッドが足並みそろえて
＋１ずつのアドレスを見に行かないといかんのです
すると32ワードなりまとめてとってくる。　とってくる間は
別のスレッドが実行されてる。
119 名前：デフォルトの名無しさん mailto:sage [2010/05/08(土) 14:52:01 ]: >>114
私は「連動した」メモリアクセスと説明している。要は16人17脚なんだけどねw
120 名前：デフォルトの名無しさん mailto:sage [2010/05/08(土) 22:51:22 ]: おまえら頭いいんだろうな
121 名前：デフォルトの名無しさん mailto:sage [2010/05/08(土) 22:57:18 ]: どうせ底辺大学生が研究で勉強してるだけでしょ？
122 名前：デフォルトの名無しさん mailto:sage [2010/05/08(土) 23:11:19 ]: と学歴昆布が申しております
123 名前：デフォルトの名無しさん mailto:sage [2010/05/09(日) 15:32:57 ]: でもGT200以降だと16人17脚でパン喰い競争って微妙に語弊あるよね。
16人17脚でパン喰い競争じゃない使い方したことないけど。
124 名前：デフォルトの名無しさん mailto:sage [2010/05/09(日) 16:44:58 ]: >>102のコードは
あるスレッドの遅れが、他のスレッドの足を引っ張る結果になるの？

コード内で参照する領域に依存関係がないのに
他のスレッドに影響するの？
なんだか納得出来ないなあ
125 名前：デフォルトの名無しさん mailto:sage [2010/05/09(日) 17:42:10 ]: >>102
if使ってる時点でありえないかと。。。

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

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

で、thread数とblock数を最適化しておけば、おｋでは？？
出力もコアレスだしさ
126 名前：デフォルトの名無しさん mailto:sage [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 名前：デフォルトの名無しさん mailto:sage [2010/05/09(日) 20:59:34 ]: >>107

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

つーかCUDAって根本的に間違ってるよなぁ。
C言語で書けることが利点？なのに、ベタにCで書くと
コアレスじゃないとか何とかで遅いというｗ
じゃあ、遅くないfloatN_tでも用意してそれ使わせればいいのに、と考えると
HLSLとかの方がいいんじゃねえの？って結論になるｗ
128 名前：デフォルトの名無しさん mailto:sage [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 名前：デフォルトの名無しさん mailto:sage [2010/05/09(日) 21:18:40 ]: 揚げ足をとるようだけど、右足を出した隣の人は左足を最初に出してる。
__syncthreadsしなければ隣のひとたちを見ないでばりばり走るし、
過度な比喩は逆に混乱させるだけだと思う。

>>100みたいなのは計算式を丸ごと変えないといけないな。
プログラムより、数学的な知識が必要だと思う。
130 名前：デフォルトの名無しさん mailto:sage [2010/05/09(日) 21:19:54 ]: >>126
いやｗ
nの数の条件分岐とか必要ないでしょｗｗ
Host側がコール時に調整すればいいだけの話

CUDAの場合、無駄な命令入れるよりかは、少し出力メモリ先がはみ出ててもコアレスならおｋ
131 名前：デフォルトの名無しさん mailto:sage [2010/05/09(日) 21:23:42 ]: >>130
nが巨大なら、ループする必要があるでしょ。ループ継続条件は通常、nとの比較で行なわれるわけで。
132 名前：デフォルトの名無しさん mailto:sage [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 名前：デフォルトの名無しさん mailto:sage [2010/05/09(日) 21:32:13 ]: >>131
だったら数回カーネルをコールすればいいのでは？？
この場合カーネル内でループしちゃうと、キャッシュ外れる避けられない。キャッシュミスは大きいよ。
134 名前：デフォルトの名無しさん [2010/05/09(日) 21:39:07 ]: >>100　は一次元差分法からの連想だろうけど、
差分法なら青木本に２次元があるから参考にしたら。
自分は３次元の差分法を作ったけど信じられない速度が出た。
まあ、比較したCPU版が遅いんだろうけどｗ
135 名前：デフォルトの名無しさん mailto:sage [2010/05/09(日) 22:06:07 ]: >>132
CC1.1までを前提にするけど、コアレスであるためには
スレッドがアクセスするメモリがハーフワープ(16個)連続して配置されている必要がある。
SMがロードする際16個まとめてグローバルメモリをロードして各スレッドに割り当てるから。
だからハーフワープあたりのグローバルメモリのアクセスは1回になる。

じゃあコアレスじゃないとどうなるかというと、
キャッシュなんてないないから各SPは個別に毎回グローバルメモリに読みにいく。
だから遅い。
136 名前：デフォルトの名無しさん mailto:sage [2010/05/09(日) 22:20:52 ]: >>135
「bへのアクセス3つは16個連続していること」って用件は満たしてる。
多分アライメントも16個境界になってないと駄目ということか。
確かにそうすると、 b[i+1]とb[i-1]はアライメントに合ってないから
都度ロードが発生する訳だな。
理解しました。どうもありがとう。
137 名前：デフォルトの名無しさん mailto:sage [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 名前：デフォルトの名無しさん mailto:sage [2010/05/10(月) 00:51:19 ]: >>127
CUDAというかアーキテクチャの制限ではないでしょうかね？
そのままグローバルメモリにアクセスするのではなく、シェアードメモリを利用するという思想で作られていると思っています。

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

>>134
あの本はそういったことも載っているのですか。
やはり一度は読むべきなのでしょうかね・・・
139 名前：デフォルトの名無しさん mailto:sage [2010/05/10(月) 02:36:11 ]: >出力メモリ先がはみ出るのはまずいかと思いますが、
メモリを考慮してグリッドとブロックの数を決めるより
はじめにグリッドとブロックの数を決めてそれに合わせてメモリを確保したほうが良い。
ってことが言いたいんじゃないか。
140 名前：デフォルトの名無しさん mailto:sage [2010/05/10(月) 07:11:13 ]: >>137
分岐の種類によって違う結果になる。
・ワープ中幾つかのスレッドだけが条件に該当する場合
条件分岐のコスト(のみ)が余計に掛かる。
・ワープ中全てのスレッドが条件に該当する場合
ワープによっては条件分岐のコストが嵩むが
ワープによっては条件分岐以降の処理をパスするので早く開放される。

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

条件分岐は怖くありませーん。下手に複雑な演算するよりお徳でーす。
141 名前：デフォルトの名無しさん mailto:sage [2010/05/10(月) 11:26:17 ]: >>102
シェアードメモリは使ったほうがいいと思いますよ
142 名前：デフォルトの名無しさん mailto:sage [2010/05/10(月) 11:47:48 ]: うーん、CUDAの初歩も知らない奴に取り敢えず雛形を提示したら、
ちょっと齧った連中が寄って集ってけちつけるという構図になったなw
合間に慣れた人が適切な指摘をしてくれているからいいけどねぇ。

つーか、CUDAがそれだけ「注目」されていると言うことなんだろうけど。
143 名前：デフォルトの名無しさん mailto:sage [2010/05/10(月) 13:31:41 ]: やだ・・・なにこの失敗したCell並に難解なプログラミング・・・
144 名前：デフォルトの名無しさん mailto:sage [2010/05/10(月) 14:51:44 ]: 大丈夫、Cellのspuみたいにインストラクションを一個ずつあれこれしなければならないほど複雑じゃない。
145 名前：デフォルトの名無しさん mailto:sage [2010/05/10(月) 17:57:02 ]: GTX4*シリーズ買う？
それとも買った？
146 名前：デフォルトの名無しさん mailto:sage [2010/05/10(月) 23:14:36 ]: 前スレから

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

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

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

で、GTX4*買うかはやっぱり倍精度が気になるわけで、様子見の状態
147 名前：デフォルトの名無しさん mailto:sage [2010/05/12(水) 09:25:19 ]: C20シリーズは何とか予算用意して買うか、どこかの計算機センタに導入されるのを待つ形になる
それの準備的なものとして、GTX4シリーズを買うと思う
148 名前：デフォルトの名無しさん mailto:sage [2010/05/14(金) 10:29:43 ]: 完全にHPCユーザーしかGPGPUに興味を持たなくなったな
149 名前：デフォルトの名無しさん mailto:sage [2010/05/14(金) 11:10:38 ]: 画像処理とか（？）する人はどう思っているんでしょう
GPGPUについては
150 名前：デフォルトの名無しさん mailto:sage [2010/05/14(金) 12:19:47 ]: CG系研究室の俺が。

GPGPU(=CUDA)は今や当たり前の技術です。
なんか論文で新しい手法を提案するとき、実行時間の項にCPUでの計算時間だけ書いていたら、
「GPUでやったらどうなのよ？」と査読者から普通にコメントつけられるレベルです。
GPUに載せづらい既存の手法を、ちょっと改良してGPU適合なアルゴリズムにしました、的な論文が
3年ぐらい前から大量に出ています。
151 名前：デフォルトの名無しさん mailto:sage [2010/05/14(金) 13:02:33 ]: >>150
> CG系研究室の俺が。
CGの人ってコンピュータビジョンにもアンテナ伸ばしてるものなんですか？

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

> 3年ぐらい前から大量に出ています。
出過ぎたせいで最近は実装してみましたレベルの論文はむしろ通りにくいですよ。
152 名前：デフォルトの名無しさん mailto:sage [2010/05/14(金) 18:40:04 ]: ATI stream とCUDA とどっちがよく使われているんでしょうかね
153 名前：デフォルトの名無しさん [2010/05/14(金) 19:19:01 ]: 「実装してみました」系はCUDAだろうね。
ピークパフォーマンスたたき出すならATIだろJS
154 名前：150 mailto:sage [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 名前：デフォルトの名無しさん mailto:sage [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 名前：デフォルトの名無しさん mailto:sage [2010/05/15(土) 01:12:46 ]: >>154
GPGPUとしてはCUDAが主流になると思う。
だけどこれまではほんとに酷かったなあ。
ちょっとでもコンピュータサイエンスの知見を持っていたら恥ずかしくて発表出来ないレベルが多かった。
今までは流行でみんながCUDA、CUDAといっていたけど、多くの人はあくまでも手段の一つであって、
目的ではない。目的にするならコンピュータサイエンスの分野でやればよいわけで。
それにCUDAを使う場合って大抵最適化しているから、CPUとの比較も余り意味がないし。
これからはそういうのはなくなっていくだろうね。
158 名前：& ◆6kLsuKGM9vPI [2010/05/15(土) 02:51:39 ]: 480と285の倍精度計算のベンチマークの比較はどこかにありませんか?
159 名前：デフォルトの名無しさん mailto:sage [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 名前：デフォルトの名無しさん mailto:sage [2010/05/15(土) 15:40:25 ]: >>157
どういうこと？
161 名前：デフォルトの名無しさん mailto:sage [2010/05/15(土) 16:14:37 ]: >>159
うまく行くケース(GPUのバージョン依存だっけ？？)なら３回でいけるだろうけど、
最悪の場合

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

ってなるから、1+16+16 = 33　じゃね？？
162 名前：デフォルトの名無しさん mailto:sage [2010/05/15(土) 16:39:09 ]: >>161
全部のSPがb[tx]を読む→コアレス
全部のSPがb[tx+1]を読む→コアレス
全部のSPがb[tx-1]を読む→コアレス

とはならないの？
163 名前：デフォルトの名無しさん mailto:sage [2010/05/15(土) 17:00:15 ]: >>161
そういう場合、
スレッド　i-1　はtx-1をグローバルからSharedに読む
スレッド　i　　　はtxをグローバルからSharedに読む
スレッド　i+1　はtx+1をグローバルからSharedに読む
　：
でひと固まりもってきて、
あとはShared上でCoaleshed?　ｷﾆｼﾅｲ!　で動けばいいんじゃね
164 名前：デフォルトの名無しさん mailto:sage [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 名前：デフォルトの名無しさん mailto:sage [2010/05/16(日) 02:31:46 ]: >>162-163
「t*x」　になっているんでこのままだとtステップになっちゃうわけだよ
つまりもともとの配列の並べ方から考え直さないといかんよな。
x-1, x, x+1, ,,, になったら163みたいに、
・スレッド協調してSharedに持ってくる
・Shared上では好き勝手に読み書きしてよろし
・syncthreadする
・スレッド協調してGlobalに書き戻す
というやりかたが基本かな
166 名前：デフォルトの名無しさん mailto:sage [2010/05/16(日) 09:56:22 ]: >>165

>・Shared上では好き勝手に読み書きしてよろし
今扱ってる問題だと場合によっては、bank conconflictがあったりする。。。
167 名前：デフォルトの名無しさん mailto:sage [2010/05/16(日) 09:57:09 ]: >>166
まちがえたｗ

>>165

>・Shared上では好き勝手に読み書きしてよろし
今扱ってる問題だと大丈夫だけど、場合によってはbank conconflictがあったりする。。。
168 名前：デフォルトの名無しさん mailto:sage [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 名前：デフォルトの名無しさん mailto:sage [2010/05/16(日) 11:06:33 ]: テクスチャキャッシュを使うのが面倒なくていいと思うのは俺だけか?
170 名前：デフォルトの名無しさん mailto:sage [2010/05/16(日) 11:33:22 ]: >>169
少し前にテクスチャ使ったコードで例を示したけど、誰も反応しなかったぞ
なんで？？
171 名前：デフォルトの名無しさん mailto:sage [2010/05/16(日) 14:23:32 ]: どのレス？
172 名前：デフォルトの名無しさん mailto:sage [2010/05/16(日) 18:59:40 ]: >>171
>>125
173 名前：デフォルトの名無しさん mailto:sage [2010/05/16(日) 19:31:54 ]: OpenCLでなくCUDAを学ぶ価値ってあるの？
174 名前：デフォルトの名無しさん mailto:sage [2010/05/16(日) 19:34:06 ]: OpenCLってなにに使われてるの
175 名前：デフォルトの名無しさん mailto:sage [2010/05/16(日) 19:38:18 ]: OpenCLは今のところRadeon使う人用かと。
176 名前：デフォルトの名無しさん mailto:sage [2010/05/16(日) 20:33:59 ]: DirectComputeはどんな感じ？
177 名前：デフォルトの名無しさん mailto:sage [2010/05/16(日) 21:33:22 ]: ローカルメモリに対するテクスチャメモリのメリットがいまいち分からない
そもそもローカルメモリを使いこなせないから、テクスチャメモリまで手が回らない
178 名前：デフォルトの名無しさん mailto:sage [2010/05/16(日) 23:49:15 ]: >>175
え？
179 名前：デフォルトの名無しさん mailto:sage [2010/05/17(月) 00:41:56 ]: >>177
補間できるとか
180 名前：デフォルトの名無しさん mailto:sage [2010/05/17(月) 20:52:47 ]: これってcudaMallocHostってsizeof(構造体)でぴったりとその構造体の容量分だけちゃんと確保できてるんですか？
なんかsizeof使ってぴったり構造体分確保したら条件によってセグメントエラーが出るんですが（printfを適当なとこに放りこんだらエラーが出ずに通ったりします）
2倍の容量を確保したら問題なく通ります
181 名前：デフォルトの名無しさん mailto:sage [2010/05/17(月) 23:08:13 ]: >>177
ローカルメモリがLocal Memoryのことをいってるなら、役割がまったく違うからメリットもなにもない
Global Memoryのことをいってるなら、キャッシュに加えて補間、正規化アドレス、変換とかだな
キャッシュが有効に機能しているかは、プログラムによっては疑わしいんだけどね

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

くらいしかアドバイスしようがないな。
183 名前：デフォルトの名無しさん mailto:sage [2010/05/18(火) 10:19:15 ]: printf挟むと動くと言ってるからどっかでスタック壊してんじゃないかと予想
184 名前：デフォルトの名無しさん mailto:sage [2010/05/19(水) 14:37:47 ]: まぁ、晒せ、と。
185 名前：デフォルトの名無しさん mailto:sage [2010/05/20(木) 11:59:21 ]: ω
186 名前：デフォルトの名無しさん mailto:sage [2010/05/21(金) 09:13:45 ]: 研究室にGTX 480届いた．
smokeParticlesで170fpsくらい．
倍精度FFTとかはこれからやる．
187 名前：デフォルトの名無しさん mailto:sage [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 名前：デフォルトの名無しさん mailto:sage [2010/05/21(金) 11:33:50 ]: nbody --benchmarkよろ
189 名前：デフォルトの名無しさん mailto:sage [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 名前：デフォルトの名無しさん mailto:sage [2010/05/21(金) 13:05:20 ]: >>187
多分CUFFT3.0を使った複素２次元FFTだと思うのだが、
こちらの計測データとかなり違うような・・・。
こちらでは480の単精度で1回(片方向)あたり0.44ミリ秒、2*1024回分で0.91秒くらい。

ちなみに頑張ると倍精度は2.6秒くらいでいける。
191 名前：187 mailto:sage [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 mailto:sage [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 mailto:sage [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 mailto:sage [2010/05/21(金) 13:21:21 ]: nvcc cufft.cu -o cufft_double -arch=sm_20 -lcudart,cutil32,cufft -Xcompiler /O2

↑のソースをこれでコンパイルした．
195 名前：デフォルトの名無しさん mailto:sage [2010/05/21(金) 13:47:54 ]: >>189
おお、ありがとうございます。
GTX260でも370GFOPS位出るのであんまり速くないですね・・・（´・ω・｀）
196 名前：デフォルトの名無しさん mailto:sage [2010/05/21(金) 13:54:24 ]: >>187

>CUT_SAFE_CALL(cutStopTimer(timer));

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

k の反復回数を512にした時のほぼ倍の時間になるかどうかをチェックしてみて。
197 名前：187 mailto:sage [2010/05/21(金) 15:00:01 ]: 同期をちゃんととってみた．

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

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

>>190
倍精度を頑張るってどう頑張るんですか？
198 名前：デフォルトの名無しさん mailto:sage [2010/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 名前：デフォルトの名無しさん mailto:sage [2010/05/21(金) 20:45:35 ]: SLIというかマルチGPUで使う場合はOpenMPかなんかで複数スレッドを作って
それぞれののGPUを使えばいい、というかサンプルがSDKにあったはず。
ゲームみたいにSLI組んでシングルGPUに見せるのはまだ出来なかったはず
200 名前：デフォルトの名無しさん mailto:sage [2010/05/21(金) 21:36:18 ]: >>199
親切にありがとうございます。
まさにゲームのようなSLIの考えでいましたが、そういった使い方はできないということですので
自宅用に使用するか等、使い方を考えてみます。本当にありがとうございました。
201 名前：デフォルトの名無しさん mailto:sage [2010/05/21(金) 22:56:10 ]: >>197
もちろんコードを自分で書く。
202 名前：デフォルトの名無しさん mailto:sage [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 mailto:sage [2010/05/26(水) 17:15:17 ]: と聞いたばかりですみません。
解決しました。
ちゃんとgputoolsのinstall.txtを呼んでいなかったので
パラメータを間違えていて、--enable-emulationというのを
つけてしまっていました。
お騒がせしました。
204 名前：デフォルトの名無しさん mailto:sage [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 名前：デフォルトの名無しさん mailto:sage [2010/05/26(水) 23:38:38 ]: >>204
言われてる1/4ってのは当初期待されていた性能に対して1/4ってことじゃなかったの？
いくらなんでもTesla世代と比較してないでしょ
206 名前：デフォルトの名無しさん mailto:sage [2010/05/27(木) 03:41:07 ]: 当初は短精度に対して理論値で1/2の性能を出せるって話だったのだけど
Geforceでは削られて短精度に対して1/4しかでないことになっている
207 名前：デフォルトの名無しさん mailto:sage [2010/05/27(木) 05:53:57 ]: 1/8じゃなかったの？
208 名前：デフォルトの名無しさん mailto:sage [2010/05/27(木) 06:06:21 ]: 皆GeForceでは1/8 (Fermi版Teslaシリーズのさらに1/4)という意味で話していたけど
一部の人は誤解していたみたい。
209 名前：デフォルトの名無しさん mailto:sage [2010/05/27(木) 20:51:44 ]: どういうことだってばよ
210 名前：デフォルトの名無しさん mailto:sage [2010/05/28(金) 00:28:01 ]: けっきょくどういうことなんだってばよ！
211 名前：デフォルトの名無しさん mailto:sage [2010/05/28(金) 08:37:10 ]: こまけえことはいいんだよ！

>>205
俺の理解はこれなんだが
GTX480も在庫が増えて値段も下がってきたね
俺、5万円切ったら本気出すんだ…
212 名前：デフォルトの名無しさん mailto:sage [2010/05/28(金) 17:09:48 ]: Teslaじゃコードを書き換える手間を考えたら、たいしてコストパフォーマンスは良くないな。
213 名前：デフォルトの名無しさん mailto:sage [2010/05/28(金) 21:44:15 ]: ？
214 名前：デフォルトの名無しさん mailto:sage [2010/05/29(土) 05:36:25 ]: >>100の様なコードだったら手間はかからんよ。
215 名前：デフォルトの名無しさん mailto:sage [2010/05/29(土) 12:56:20 ]: TeslaがGT世代の事じゃなくてGPGPU専用カードの事を言ってるなら
Teslaの利点はバグのあるコードを走らせてもOS巻き込んで死んだりしない点が一番じゃね？
開発機はTeslaが楽だわ
216 名前：デフォルトの名無しさん mailto:sage [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 名前：デフォルトの名無しさん mailto:sage [2010/06/05(土) 01:05:17 ]: 2次元配列なんてCUDAにはありません。
意図的に1次元に並べてください。
220 名前：デフォルトの名無しさん mailto:sage [2010/06/05(土) 20:17:51 ]: GTX480でおすすめのボードってありますか？
もしくはダメなボードはありますか？
221 名前：デフォルトの名無しさん mailto:sage [2010/06/05(土) 21:43:43 ]: こちらも480入手。取り敢えず動くのは確認。細かいテストはこれから。

>>220
今のところどこも経験値低いだろうからリファレンスボードのままだろうし、
ELSAが無難じゃね? 高いかもしれないけど。
ヒートパイプの所為で更に場所を喰うので要注意。
222 名前：デフォルトの名無しさん mailto:sage [2010/06/05(土) 21:52:51 ]: リファレンスだからドスパラのでいいんじゃないの
シール張るだけらしいから
223 名前：220 mailto:sage [2010/06/07(月) 11:06:57 ]: >>221-222
ありがとうございます。

elsaのGTX480のページに対応OSがwindowsしか書いてないんですが、
Linuxでは動かないのでしょうか？
GeforceのCUDA用ドライバをいれたら普通に動くんでしょうか？
224 名前：221 mailto:sage [2010/06/07(月) 11:31:39 ]: >>223
ドライバのリビジョンは今手元にないから判らないけど、CUDA2.3当時のドライバでCUDA2.3が動いている。
CUDA3に関しては未だテストしていない。
225 名前：デフォルトの名無しさん mailto:sage [2010/06/14(月) 13:16:15 ]: 最近CUDAを勉強しはじめたんですが、質問させてください。

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

OSはUbuntu 9.04、GPUは9500GT、cudaのバージョンは2.3です。
226 名前：デフォルトの名無しさん mailto:sage [2010/06/14(月) 13:20:32 ]: できます。
一回目のcudaMalloc()の前にcudaThreadSynchronize()でもしておけば、一回目のcudaMalloc()は早くなります。
尤も、cudaThreadSynchronize()で時間が(恐らく100ms程度)掛かりますが。
227 名前：デフォルトの名無しさん mailto:sage [2010/06/14(月) 13:29:32 ]: >226
ありがとうございます。
早速試したところ、そのとおりの結果になりました。。。
これは諦めるしかないんですかね。
228 名前：デフォルトの名無しさん mailto:sage [2010/06/14(月) 13:29:33 ]: なんかCUDAは最初に呼ぶときに時間がかかるとかリファレンスかなんかに書いてあったような・・・
斜め読みしたけどみつけられんかった
229 名前：デフォルトの名無しさん mailto:sage [2010/06/14(月) 13:41:43 ]: >228
そんな仕様だったんですね、知らなかったです。

書き換えたもとの処理がたいして時間が掛からない場合、CUDAの利点を活かせないんですね。
230 名前：226 mailto:sage [2010/06/14(月) 13:45:06 ]: >>227
cudaは元々単発プロセス向きじゃないので、諦めて常駐型のプロセスか処理時間の掛かるプロセスに使いましょう。
具体的に何をやっているかは判らないけど、cudaのAPIを呼ぶ初回は必ず時間が掛かるのよね。
例えばmain()の先頭ででも一回syncするだけでいいから、私の用途では問題にならないんだけど。
231 名前：デフォルトの名無しさん mailto:sage [2010/06/14(月) 13:48:59 ]: >>229
一回だけだとそうかも知れないけど、バッチやリアルタイムで複数ファイルを処理するとかなら
効いてくると思います。
232 名前：デフォルトの名無しさん mailto:sage [2010/06/16(水) 10:00:49 ]: 研究にCUDAを使ってみたいと思っている学生です。＞＜
プログラマブルシェーダで物理演算をしてシミュレーションをやっている
論文とかあったのですが
CUDAで計算するのとの違いはあるんですか？
あと
DirectXを使いながらCUDAを利用することは
DirectXで手一杯のことをやっていたらCPUで計算したほうが早かったりするんですか？

お願いします（。・・。）
233 名前：デフォルトの名無しさん mailto:sage [2010/06/16(水) 10:05:47 ]: >>232
それが俺の読んだのと同じならcgを使ってシェーダプログラムで演算をしてた
理由はCUDAがまだ無かったから
DirectXで3D描画しながらだとそりゃ性能は落ちるでしょ。CUDA専用にカードを追加するなら話は別だけど
234 名前：デフォルトの名無しさん mailto:sage [2010/06/16(水) 12:16:27 ]: だいぶ前の話だけど、AviUtlのフィルタでシェーダよりCUDAの方が遅いとかって無かったっけ。
カリカリにチューンしたらCUDAが速いんだろうけれど、シェーダの方がメモリアクセスのお作法を意識しなくて良い分、初心者には性能向上が簡単かも知れない。
235 名前：デフォルトの名無しさん [2010/06/16(水) 19:47:40 ]: >>234
つまりn00bだったってことかよ!
236 名前：デフォルトの名無しさん mailto:sage [2010/06/16(水) 23:49:14 ]: >>234
???
CGをやるなら別だが、今更Cgは無いだろう。
初心者ならなおさらCUDAだ。
特に物理演算をやるなら。
237 名前：デフォルトの名無しさん mailto:sage [2010/06/17(木) 00:34:00 ]: PGIコンパイラのCUDA化って正直どうよ。
人力でやるのに疲れてきた。
それなりの性能が出るなら買ってもいいと思ってるんだけど。
238 名前：デフォルトの名無しさん mailto:sage [2010/06/17(木) 06:18:56 ]: ディレクティブでは全然性能上がらなくて頭来たからCで書き直したｗ
239 名前：デフォルトの名無しさん mailto:sage [2010/06/17(木) 06:53:59 ]: 東工大、Teslaを利用したスパコン「TSUBAME2.0」を開発
pc.watch.impress.co.jp/docs/news/20100617_374718.html
240 名前：デフォルトの名無しさん mailto:sage [2010/06/17(木) 14:14:54 ]: >>239
有用な結果を出さなかったら、次から仕分けな
241 名前：デフォルトの名無しさん mailto:sage [2010/06/17(木) 18:10:08 ]: 国会議員を仕分けすりゃいいのに
1/3位にさｗ
242 名前：デフォルトの名無しさん mailto:sage [2010/06/17(木) 20:46:42 ]: あんたら、部外者か？

蓮舫のブレーンがそのスパコンの推進者なのに、仕分けされるわけがねーだろ。
243 名前：デフォルトの名無しさん mailto:sage [2010/06/17(木) 21:00:10 ]: つまり漣舫は神戸のスパコンをつぶしたくで仕分けしたのですね。
244 名前：デフォルトの名無しさん mailto:sage [2010/06/17(木) 21:11:00 ]: 1度でいいからこんな超計算機をいじり倒したいものだ
245 名前：デフォルトの名無しさん mailto:sage [2010/06/17(木) 21:15:30 ]: 部外者だけどCSX600なんて誰も使ってないんじゃね？
導入初期に不良品大量に掴まされてたみたいだけどｗ
246 名前：デフォルトの名無しさん mailto:sage [2010/06/17(木) 21:28:53 ]: まあそうでもしないと進まないところもあるんだろうけど、こんな財政状況なんだからバッサリ逝ってよし
247 名前：デフォルトの名無しさん mailto:sage [2010/06/17(木) 21:34:09 ]: >>242
むしろ部外者しかいないだろ
お前はしらんが
248 名前：デフォルトの名無しさん mailto:sage [2010/06/17(木) 21:35:11 ]: 議員を減らせば良い
249 名前：デフォルトの名無しさん mailto:sage [2010/06/17(木) 21:35:34 ]: 政治ネタうぜぇ
250 名前：デフォルトの名無しさん mailto:sage [2010/06/17(木) 21:43:16 ]: 選挙だもの

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

専門がスパコンはおろかコンピュータですらない姫野に
反論ができるわけもない。
252 名前：デフォルトの名無しさん mailto:sage [2010/06/17(木) 22:21:51 ]: π（笑）
253 名前：デフォルトの名無しさん mailto:sage [2010/06/18(金) 02:21:28 ]: >>251
何で姫野ヘボベンチっていわれているの？
ソースを見たけどコード自体は学生が書くような内容だね。
だからなの？
254 名前：デフォルトの名無しさん mailto:sage [2010/06/18(金) 12:50:58 ]: メモリ帯域をひたすら無駄食いするプログラムだよね＞姫野ベンチ
ここで「なぜか遅いです」とか質問してくるレベル。
255 名前：デフォルトの名無しさん mailto:sage [2010/06/18(金) 15:19:59 ]: かといってgemmでピーク性能測られても役に立たない分野も多く。
直線に強い車かカーブに強い車かってのと似てる。
256 名前：デフォルトの名無しさん mailto:sage [2010/06/18(金) 17:38:19 ]: πって何を測ってるの
257 名前：デフォルトの名無しさん mailto:sage [2010/06/18(金) 17:42:40 ]: 今ではすっかり使われないx87
258 名前：デフォルトの名無しさん mailto:sage [2010/06/19(土) 03:34:48 ]: >>235-236
今更シェーダをお勧めはしないけど、畳み込みを考えた時に
コアレスにするためのコードって余分だよね。バッドノウハウだよね。
ってことが言いたかった。
259 名前：デフォルトの名無しさん mailto:sage [2010/06/21(月) 12:53:32 ]: 東工大のスパコン、世界２位なんだ。
そういうことか。
260 名前：デフォルトの名無しさん mailto:sage [2010/06/21(月) 13:55:44 ]: 東工大は64位じゃね？
www.top500.org/list/2010/06/100
261 名前：デフォルトの名無しさん mailto:sage [2010/06/21(月) 16:58:25 ]: ヒント：TSUBAME2はまだ実在しない
っ ttp://www.gsic.titech.ac.jp/tsubame2
262 名前：デフォルトの名無しさん mailto:sage [2010/06/22(火) 08:40:40 ]: まだ計算どころか設置もしていないぞ
263 名前：デフォルトの名無しさん mailto:sage [2010/06/22(火) 12:51:00 ]: NECは京速機やめてこっちでやってんだな。
264 名前：デフォルトの名無しさん mailto:sage [2010/06/23(水) 01:26:00 ]: CUDAって未だにデバイスメモリをメインメモリにマップ出来ないの？
逆は出来るようだけど。
265 名前：デフォルトの名無しさん mailto:sage [2010/06/23(水) 01:45:32 ]: マップできなくてもいいけど、MPIでDMA使ってデバイスから通信チップに流しこんで欲しいわな
266 名前：デフォルトの名無しさん mailto:sage [2010/06/23(水) 05:56:26 ]: マップ出来ないといちいち転送指示出さなきゃいけないのが困る。
スクラッチで書くならいいけど、既存のを移植するとなると結構大変だから・・・。
267 名前：デフォルトの名無しさん mailto:sage [2010/06/23(水) 19:31:52 ]: テンプレートライブラリで、CUDAの転送まわりをほぼ完全に隠蔽してくれる奴あったよね。
268 名前：デフォルトの名無しさん mailto:sage [2010/06/24(木) 12:55:15 ]: thrust
269 名前：デフォルトの名無しさん mailto:sage [2010/06/26(土) 00:36:06 ]: 使い方によるのかもしれないが、Mappedはかなり遅い気がする
Pinnedで転送したほうがまし
270 名前：デフォルトの名無しさん mailto:sage [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 名前：デフォルトの名無しさん mailto:sage [2010/06/27(日) 13:57:56 ]: MATMULとかだと倍精度も速いみたいだけど、レジスタやシェアードメモリの消費も倍になるから
実コードだとOccupancyが下がって思ったような速度でないんだよなぁ・・・＞Fermi
272 名前：デフォルトの名無しさん mailto:sage [2010/06/27(日) 14:14:32 ]: 倍精度の実効効率が、

GTX285　約90%
GTX480 約21%

ずいぶん差があるのはGTX480の倍精度がドライバのレベルで1/4に抑えられている（らしい）ことが最大の原因？
アーキテクチャが違うから単純には言えないけど、21/25=84で、GTX285と同程度になるし…
273 名前：デフォルトの名無しさん mailto:sage [2010/06/27(日) 14:51:36 ]: ソース
274 名前：デフォルトの名無しさん mailto:sage [2010/06/27(日) 14:53:17 ]: >>273
>>272のは真上の祖父テックのベンチ結果から計算してのことだろ？
nbody.exeでも-fp64で性能出てなかったし
275 名前：デフォルトの名無しさん mailto:sage [2010/06/27(日) 16:10:48 ]: > ドライバのレベルで1/4に抑えられている（らしい）こと
ここのソース
276 名前：デフォルトの名無しさん mailto:sage [2010/06/27(日) 16:41:24 ]: googleを使えない奴にCUDA使えるとは思えないな
277 名前：デフォルトの名無しさん mailto:sage [2010/06/27(日) 18:34:33 ]: これと、中国のTUBASA？のLinpack結果を併せてだれか解説して
278 名前：デフォルトの名無しさん mailto:sage [2010/06/27(日) 18:56:27 ]: GTX480の「倍精度理論FP演算性能672 GFLOPS」という仕様そのものが間違っていて
実際は理論ピークが1/4の168。

別にドライバのレベルで抑えられているわけでもなんでもないが。
279 名前：デフォルトの名無しさん mailto:sage [2010/06/27(日) 19:13:13 ]: >>278
↓この記事を信じてしまうと
journal.mycom.co.jp/special/2010/gpu/004.html
GTX480の倍精度演算の理論値は672GFLOPSになっちゃうけど、

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

が本当だとするなら、どこで計算がずれたのだろう…
280 名前：デフォルトの名無しさん mailto:sage [2010/06/27(日) 22:22:30 ]: ドライバなんだかそれ以外だかはわからないが、
GTX480の倍精度理論性能が168GFLOPSであることは信憑性高い、ということでおｋ？
281 名前：デフォルトの名無しさん mailto:sage [2010/06/27(日) 22:24:19 ]: ドライバで制限してるんじゃ無ければカタログスペックとの差異はどーすんの？
ttp://www.elsa-jp.co.jp/products/hpc/tesla_c2050/index.html
282 名前：デフォルトの名無しさん mailto:sage [2010/06/27(日) 23:07:18 ]: >>281
GTX480の話をしているときにC2050持ち出してどうするの？
283 名前：デフォルトの名無しさん mailto:sage [2010/06/27(日) 23:24:10 ]: 想像力を働かすと、
C2050よりもコア数が多くて、同一アーキテクチャ仕様と公表されているGTX480が168GFLOPSしか出ないなんておかしいぞ
ということかな

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

ECC無しでもOKで、メモリ１G程度で十分という人には結構おいしいかもしれんねGTX480
284 名前：デフォルトの名無しさん mailto:sage [2010/06/27(日) 23:33:39 ]: メモリしか差のなかったC1060に対して、C2050はもっと明示的な計算用チップとしてのアドバンテージをつけたかった
nVidiaの意向が反映した結果がGTX480だと思うのだけど。
285 名前：デフォルトの名無しさん mailto:sage [2010/06/28(月) 07:10:21 ]: C2050はGTX480はCUDAコアが少なく、クロックも低いにもかかわらず、
倍精度理論性能は515GFlopsとなりC2050ほうが高い。
これは、GTX480はドライバで倍精度演算の実行にペナルティが加わるようになってるから。
それでも、GTX285と比べても倍精度演算が高速だからいいじゃんというのがNVIDIAの考え。

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

現実のカーネルコードの多くではこんなベンチと違ってメモリ帯域が足を引っ張るから、
当然C2050実効性能は理論性能遠く及ばない。
一方、GTX480は単位時間当たりの演算回数に対して、メモリ帯域がC2050よりも大きい分、
メモリ帯域による実効性能減が起きにくい。
結果として、カリカリにチューンしない限りは、GTX480とC2050の倍精度演算の
実効性能はそれほど大きくは変わらないようだ。
286 名前：デフォルトの名無しさん mailto:sage [2010/06/28(月) 07:18:45 ]: 何が言いたいかというと、C2050と比べるとGTX480はそこそこお買い得だってことね。
ECC、メモリ容量、倍精度演算の潜在能力とメーカー出荷時のチェック品質が値段の違いなのかな。
287 名前：デフォルトの名無しさん mailto:sage [2010/06/28(月) 07:23:59 ]: >>283
>>281だけどそう言いたかった。OpenGLが速いQuadroみたいな感じなのかね？
Telsa C20xxのアホみたいな値段見てると・・・
288 名前：デフォルトの名無しさん mailto:sage [2010/06/28(月) 07:46:15 ]: CUDAクラスタを出荷してる会社の人から聞いた話だけど、
ごくまれにGTX系でCUDAカーネルを実行すると、間違った結果しか返さないボードがあるらしい。
他のベンチマークは普通に完走するし、ちゃんとディスプレイ出力もできて、
正常としか判断しようがないボードでも間違うのがあるそうだ。
原因は不明だが、不良と思われるそのボードを交換すると元気に動くらしいから、
やはり一見正常な不良品なんだろうといってたよ。

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

GTX480でも十分性能はいいけれど、Teslaは一枚は持っておくといざという時に安心らしい。
289 名前：デフォルトの名無しさん mailto:sage [2010/06/28(月) 08:32:41 ]: >>288

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

まぁ、Teslaは仲間内で誰かが１個持っておけばいいやくらいの感じかな。
290 名前：デフォルトの名無しさん mailto:sage [2010/06/28(月) 08:45:16 ]: >>285

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

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

性能面では倍精度演算をよほど酷使する計算以外ではGTX480の方が速い。
メモリバンド幅がネックになる演算においてはC2050はGTX285にすら劣る。
291 名前：デフォルトの名無しさん mailto:sage [2010/06/28(月) 08:46:52 ]: ＞C2050とGTX 480は同じFermiアーキテクチャではあっても違う構造になっている。

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

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

ちゃんとコードが走るかはTeslaがあれば検証できる意味でも、
一枚はあるといいとのアドバイスね。
293 名前：デフォルトの名無しさん mailto:sage [2010/06/28(月) 09:17:20 ]: >>290
>C2050とGTX 480は同じFermiアーキテクチャではあっても違う構造になっている。
同じチップなんだから、当然同じ構造でしょ。ボードデザインは違うけど。

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

適当言い過ぎ。
294 名前：デフォルトの名無しさん mailto:sage [2010/06/28(月) 09:25:26 ]: 3.1でたね。・Support for printf() in device code なにげに便利かも。
295 名前：デフォルトの名無しさん mailto:sage [2010/06/28(月) 10:45:13 ]: >>294
これでもう、エミュレータを使う理由は激減するね。

>>292
だからそれをNVIDIAの営業某氏は「アキバ的発想はやめましょう」と揶揄するのだよね。
でも、10万円で2枚購入できるか1枚も購入できないかの差は法人ユーザには大きいだろ。
296 名前：デフォルトの名無しさん mailto:sage [2010/06/28(月) 10:54:08 ]: GeForce GTX 285Mが搭載されているPC買ったんですが、
ttp://developer.nvidia.com/object/cuda_3_1_downloads.html
にあるノートブック用のドライバを落とせばいいんでしょうか？
297 名前：デフォルトの名無しさん mailto:sage [2010/06/28(月) 11:27:48 ]: >>295
むしろ法人こそTeslaボードのほうがいいんじゃないの？
Teslaはちゃんと使えれば、十分に値段分の価値があるよ。
個人ユーザーはTeslaがきついのはわかるけど、
法人はそれ使って値段以上の利益あげればいいわけだし。

もっともそれが可能かどうかは、試してみないとわからないのが困り者だけど・・・。
298 名前：デフォルトの名無しさん mailto:sage [2010/06/28(月) 11:33:28 ]: >>296
そうだよ
299 名前：デフォルトの名無しさん mailto:sage [2010/06/28(月) 11:35:47 ]: >>297
だから、予算がついちゃえばTeslaでいいんだけどね。
300 名前：デフォルトの名無しさん mailto:sage [2010/06/28(月) 12:41:51 ]: >>293
とりあえず実際に使ってみてから言ってください。
301 名前：デフォルトの名無しさん mailto:sage [2010/06/28(月) 13:01:29 ]: >>292
仕分けの時に有名になった長崎大だっけ
あそこが最初にやったのが不良品を検出するプログラムの作成でしょ。
配ってなかったっけ。
302 名前：デフォルトの名無しさん mailto:sage [2010/06/28(月) 13:21:43 ]: >>300
何を言いたいのか全く分からない。
C2050とGTX470はもってるが、GTX285はないよ。

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

実際に使ったら何がわかるのさ。
303 名前：デフォルトの名無しさん mailto:sage [2010/06/28(月) 13:23:26 ]: >>302
クレクレ君みたいで申し訳ないが、サンプルプログラムのnbodyのベンチ結果を貰えない？
nbody -benchmark
と
nbody -benchmark -fp64
の２つで
304 名前：デフォルトの名無しさん mailto:sage [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 名前：デフォルトの名無しさん mailto:sage [2010/06/28(月) 13:32:53 ]: GTX470は家に戻らないといけない。
環境はWin7x64で、C2050でディスプレイ出力してる。
306 名前：デフォルトの名無しさん mailto:sage [2010/06/28(月) 13:37:37 ]: ㌧
自分でGTX470でベンチ取ったら

516.879 GFlops@14336bodies
095.932 GFlops@14336bodies (-fp64)
だったんだが倍精度速くなってる・・ような・・・
307 名前：デフォルトの名無しさん mailto:sage [2010/06/28(月) 13:49:17 ]: FFTは使ってませんか。
倍精度FFTはGTX285からGTX480で倍速いんですがC2050はどうでしょう。
308 名前：デフォルトの名無しさん mailto:sage [2010/06/28(月) 13:50:49 ]: >>302

470のメモリバンド幅は133GB/s程度なのでC2050との差は小さいだろう。
480だと177GB/sもあるわけで。
309 名前：デフォルトの名無しさん mailto:sage [2010/06/28(月) 13:52:57 ]: >>307
FFTはデータサイズ次第だってば・・・。
まぁ、480のほうが速い場合もC2050のほうが速い場合もあります。
310 名前：デフォルトの名無しさん mailto:sage [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 名前：デフォルトの名無しさん mailto:sage [2010/06/28(月) 14:00:28 ]: 単精度では480が速いのに異論はない。
倍精度でもよっぽどチューンしない限り、480や470と
C2050のそれほど性能は変わらないよ。

そりゃ完全に同じじゃないけどさ。
312 名前：デフォルトの名無しさん mailto:sage [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 名前：デフォルトの名無しさん mailto:sage [2010/06/28(月) 15:38:36 ]: >>312
*.cuは実はc++としてnvccでコンパイルされるので、*.cから呼び出すためにはextern "C"にする必要があります。
つーか、c++拒否症でもなければhellonvcc.cをhellonvcc.cppに改名してc++としてコーディングした方が楽かもしれません。
314 名前：デフォルトの名無しさん mailto:sage [2010/06/28(月) 15:46:11 ]: ていうか、別にCUDA関係ないじゃん。
シンボル修飾についてちゃんと勉強したほうがいいよ。
315 名前：デフォルトの名無しさん mailto:sage [2010/06/28(月) 15:48:35 ]: メモリがネックにならないとき
Tesla/Geforceの単精度：Teslaの倍精度：Geforceの倍精度 = 4 : 2 : 1
メモリがネックになるとき
Tesla/Geforceの単精度：Teslaの倍精度：Geforceの倍精度 = 2 : 1 : 1
という感じだね
316 名前：デフォルトの名無しさん mailto:sage [2010/06/28(月) 16:37:51 ]: >>313

できました。どうもありがとうございます。
nvccはC++でコンパイルされるのでしたか。知りませんでした。勉強になります
ちなみに、Cで書かれたライブラリを呼び出したかったのですが、その場合は
ラッパを作るしかなさそうですね。
317 名前：デフォルトの名無しさん mailto:sage [2010/06/28(月) 16:50:52 ]: >>316
仮にcのライブラリのインクルードファイルが foo.h ならば、
extern "C" {
#include "foo.h"
}
でOK。
318 名前：デフォルトの名無しさん mailto:sage [2010/06/28(月) 16:51:04 ]: 呼び出す側がCで呼び出される側Cならば問題ない。
呼び出す側がC++で呼び出される側がCならばextern "C"←今の状況はこれ

だから、呼び出されるライブラリがCで書かれてるならラッパはいらない。
というかまともなライブラリならヘッダに、適切なマクロが書いてあるんじゃないかな。
319 名前：デフォルトの名無しさん mailto:sage [2010/06/28(月) 17:08:46 ]: >>317
>>318
どうもありがとうございます。
危うく大変な作業になるところでした。

自作ライブラリなので「まともなライブラリのヘッダ」にはなっていなかったようです。
いじりたくないので317さんの方法で対応します
320 名前：デフォルトの名無しさん mailto:sage [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 名前：デフォルトの名無しさん mailto:sage [2010/06/28(月) 19:23:46 ]: このスレ住人は親切すねー
322 名前：デフォルトの名無しさん mailto:sage [2010/06/28(月) 21:23:55 ]: 実習室でやる意味ないじゃん
323 名前：デフォルトの名無しさん mailto:sage [2010/06/29(火) 01:31:18 ]: みんな優しいなあ・・・・。
でもCUDAを始めるにはまだ早いのではないかね？
324 名前：デフォルトの名無しさん mailto:sage [2010/06/29(火) 01:45:20 ]: 3.1にしたら、CSのOceanが初めて動作する様になった。
simpleparticlesのロード時間もほぼ一瞬にまで高速化されたけど、
代わりにn-bodyのロード時間がクソ長くなった。
CUDAやOpenCLと違ってCSのサンプルがいまいち安定しないのは何でだろ。
325 名前：デフォルトの名無しさん mailto:sage [2010/06/29(火) 02:44:20 ]: >>323
たしかにこの質問主はCUDA始めるにはちょいと早すぎるかもね。
質問内容がそもそもCUDA以前の内容だし。
326 名前：デフォルトの名無しさん mailto:sage [2010/06/29(火) 11:57:31 ]: >>324
環境はWindowsかい？
うちはWindowsだけど、たしかにDirectComputeのOceanは初めて見た気がする。
確かにDirectComputeのN-Bodyは起動が遅いね。

よく考えたら、DirectComputeがあるのはWindowsだけかな？
327 名前：デフォルトの名無しさん mailto:sage [2010/06/29(火) 13:42:47 ]: DirectComputeはDirectX系なのでWindowsしか無いのでは？
しかも今のところWindows7限定（´・ω・｀）
328 名前：デフォルトの名無しさん mailto:sage [2010/06/29(火) 14:34:59 ]: Vista
329 名前：デフォルトの名無しさん mailto:sage [2010/06/29(火) 14:35:55 ]: あ、Vistaさんもいけたんですね・・・すんません
330 名前：デフォルトの名無しさん mailto:sage [2010/06/29(火) 15:40:39 ]: くだもCLもwindowsだとリモート（リモートデスクトップ、ssh)では使えないわけだが、
DirectComputeなら使えるとかいうオチでもある？
331 名前：デフォルトの名無しさん [2010/06/29(火) 22:57:17 ]: ブブブーブォォオブオーブブォーブーブー
ブォ／￣＼ーブーブーブブーンブォーオ　
ンー|　　＾o＾|∩==＜!　プォープォーブブー
ブォ＼＿／|　|　ブーブブーブォーーー
ォー＿|　|__|　|　ブォーーブブブープォー
　ー|　　　　＿|　ォーブーーーブブォーブ
ブォ|　|　　 |　ブォーブーブーブーンブー
332 名前：デフォルトの名無しさん mailto:sage [2010/06/29(火) 22:59:57 ]: リモートしたことないけどDirectComputeはDirect3Dの一部なのでDirect3Dがリモートできれば動く
333 名前：デフォルトの名無しさん [2010/06/30(水) 11:56:37 ]: cuda profilerで分岐数がわかりますが、
プロファイラで分岐数というのはどうやって数えているのですか？

1スレッドあたりの分岐数にしては、多すぎる結果が出るのですが・・・orz
334 名前：デフォルトの名無しさん mailto:sage [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見て数えてみますｗ
>>335
branchは1スレッドあたりの分岐数と表記があったのですが、
値が6桁なので、やはり総数なんでしょうかね？

ちなみに、CUDAZONEの質問掲示板のような所で
if～elseが一つ組み込まれているプログラムで
64スレッド走らせると、プロファイラではbranch=4となったという書き込みが
あるのですが、これは半ワープ４個（＝64÷16）が一つの分岐をしたから
全部で４回分岐ということなんでしょうか？
それ以外に４という数が導き出せないのですが＞＜
337 名前：デフォルトの名無しさん mailto:sage [2010/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 名前：デフォルトの名無しさん mailto:sage [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 名前：デフォルトの名無しさん mailto:sage [2010/07/02(金) 22:08:28 ]: 9800GTにするか、GT240にするか迷ってます。
性能は9800GTの方がいいので、そちらにしたいのですが、
プログラミングが趣味なので、CUDAのプログラミングを
する予定です
CUDAを考えたらどの程度違いますか？
全然違うんでしょうか？
340 名前：デフォルトの名無しさん [2010/07/02(金) 23:28:57 ]: 日
本
語
で
お
k
341 名前：デフォルトの名無しさん mailto:sage [2010/07/02(金) 23:41:05 ]: GTX260以下はウンコ
もっと言えばFermiじゃなきゃ今から買う価値無し
342 名前：デフォルトの名無しさん mailto:sage [2010/07/03(土) 02:01:16 ]: >>340
コイツは日本語書けるのに日本語読めないのかｗ
プログラミングする前に国語の勉強した方が良いぞ
343 名前：デフォルトの名無しさん mailto:sage [2010/07/03(土) 02:48:12 ]: 日本語がどうとうかはどうでもいいが、
GT9800とGT240を比べている時点でもう
「好きにしろ」
って感じだな。コアレッシングを気にするなら、
後者の方がいいのかな？GT240ってG200系？
344 名前：デフォルトの名無しさん mailto:sage [2010/07/03(土) 05:19:24 ]: 一応GT200系。
GPGPUやるなら多分240の方がいいだろう。
345 名前：デフォルトの名無しさん mailto:sage [2010/07/03(土) 14:36:26 ]: gt240はcc1.2だからどっちにしろ倍精度は使えないね。
346 名前：デフォルトの名無しさん mailto:sage [2010/07/06(火) 00:34:22 ]: なんだこれ・・・
ttp://pc.watch.impress.co.jp/docs/column/kaigai/20100706_378760.html
347 名前：デフォルトの名無しさん mailto:sage [2010/07/06(火) 01:15:34 ]: ようはFermi終了のお知らせ?
Configurable CacheがサポートされてなかったらFermi用のコードなんかHPC向けしか書かないぞ
348 名前：デフォルトの名無しさん mailto:sage [2010/07/06(火) 08:06:37 ]: GTX260と比較してプログラムの実行速度はどうなりそうなのかしら
349 名前：デフォルトの名無しさん mailto:sage [2010/07/06(火) 09:11:08 ]: 例えGT200がベースでも、AMD方式ではなくFermiのテッセレーター機構をぶち込むなら
演算器大増量に変わりは無いから速くはなるだろう。
ただ、純粋なFermiベースに比べればどうしても落ちるわな。

まあGeforceでもCUDA爆速では10倍の値段で売ってるTeslaの立場が無いし、
GPGPU用にトランジスタ増量したらゲーム性能大して向上してない割に爆熱大食らいになって
大多数の一般人には総スカン食らってるんだから、妥当っちゃ妥当なんでないかい。
350 名前：デフォルトの名無しさん mailto:sage [2010/07/06(火) 12:10:30 ]: 値段や性能はGTX260と同じだとして消費電力がGTX480の６６％じゃGTX260より悪そうでこまる。
351 名前：デフォルトの名無しさん [2010/07/06(火) 19:32:07 ]: cc2.0で言語仕様が大幅に変わっているから、
460はcc1.4になるのかな。
352 名前：デフォルトの名無しさん mailto:sage [2010/07/07(水) 07:21:03 ]: 倍精度切られてそうだけどね460
これで2分化が進めば、トップエンドはもっと思い切って汎用に振れるな
353 名前：デフォルトの名無しさん mailto:sage [2010/07/07(水) 21:42:32 ]: 倍精度ないのかなあ。BOINCとかでも倍精度必要だから売るために載せて欲しい。
354 名前：デフォルトの名無しさん mailto:sage [2010/07/07(水) 21:43:21 ]: 倍精度サポートはさすがにあるでしょ
355 名前：デフォルトの名無しさん mailto:sage [2010/07/07(水) 21:52:10 ]: おいおい、お絵かき特化のために旧コア使うってのに
お絵かきに関係のない倍精度なんて積むのかね
GT200でさえ、1:8の申し訳程度のものだったのに
これもTeslaのためだけだろ
356 名前：デフォルトの名無しさん mailto:sage [2010/07/08(木) 02:11:08 ]: 別に倍精度なんか使わないから良いや
GTX465も有るんだから構わないだろ
357 名前：デフォルトの名無しさん mailto:sage [2010/07/08(木) 06:02:49 ]: これは長期的に見て良くないね。
もはやNVIDIAはHPCを諦めてきたのかな？
358 名前：デフォルトの名無しさん mailto:sage [2010/07/08(木) 07:05:13 ]: ね？
teslaって知ってる？
359 名前：デフォルトの名無しさん mailto:sage [2010/07/08(木) 09:07:37 ]: あのさ、Teslaはわかるけど、今まではGeForceとTeslaを共用してきたから開発費が押さえられていたわけで、
Teslaのみを今後開発して行くには、果たしてどれだけ開発費が回収できるかによるでしょ。
今のNVIDIAに両方を開発していく体力があるのかどうか疑問だけどね。
360 名前：やんやん ◆yanyan72E. mailto:sage [2010/07/08(木) 13:45:08 ]: TSUBAME2もできたことだし。
もうHPC事業はHPC事業独立でできるんでしょう。
361 名前：デフォルトの名無しさん mailto:sage [2010/07/08(木) 14:29:44 ]: 長期的も何も今までもローエンドは倍精度サポートして来なかったじゃん
今まで通りGTX580を作る時はリッチに作ってGTX560を作る時は削れば良いだろ
362 名前：デフォルトの名無しさん mailto:sage [2010/07/08(木) 18:57:55 ]: TSUBAME1に鳴り物入りで採用されたCSX600は無視ですか、そうですか・・・
363 名前：デフォルトの名無しさん mailto:sage [2010/07/09(金) 01:05:40 ]: 本気で広めたいなら下位モデルでも倍精度サポートしろよ
特定のハードでしか使えないとかまじバカじゃないの
364 名前：デフォルトの名無しさん mailto:sage [2010/07/09(金) 01:22:51 ]: チップ作るのって、簡単なことじゃないぞ。
なんか勘違いしている奴がいるが、
#ifdef GTX580
　単精度サポート　
　倍精度サポート
#else
　単精度サポート
#endif
みたいにコンパイルし直せば、出来ると思っているのか？
コンパイルのたびに数億飛んでいくってかんじなんだぞ。
デバッグの行程もたくさんあるしな。
365 名前：デフォルトの名無しさん mailto:sage [2010/07/09(金) 02:11:20 ]: うっかりF7押しただけで数億とな
F5押したらどうなっちゃうの
366 名前：デフォルトの名無しさん mailto:sage [2010/07/09(金) 07:22:54 ]: お絵かきに倍精度なんているの？ｗ
367 名前：デフォルトの名無しさん mailto:sage [2010/07/09(金) 07:24:23 ]: www.adrenaline.com.br/tecnologia/noticias/5507/geforce-gtx-460-especificacoes-precos-e-estreia-no-brasil.html

コアはfermi系っぽいがキャッシュはないようだ
アンコアは前世代
368 名前：デフォルトの名無しさん mailto:sage [2010/07/12(月) 15:55:32 ]: 倍精度きましたよ。１：１２とさらに削ってきた。
369 名前：デフォルトの名無しさん mailto:sage [2010/07/12(月) 16:20:21 ]: SMのSP数を変えてくるのはさすがに予想外だった
370 名前：デフォルトの名無しさん mailto:sage [2010/07/12(月) 17:41:12 ]: >>368
8%とは酷すぎ。25%がましな時代だったとはorz
371 名前：デフォルトの名無しさん mailto:sage [2010/07/12(月) 18:14:02 ]: 25%だったことはないが・・・。
480は12.5%。
372 名前：デフォルトの名無しさん mailto:sage [2010/07/12(月) 18:35:14 ]: なんか廉価版tesla出そうな気もする
10万以下で
373 名前：デフォルトの名無しさん mailto:sage [2010/07/12(月) 19:28:00 ]: GTX460買ってきた。自分の倍精度プログラムはGTX260とまったく同じ性能だった。いいんじゃない。
374 名前：デフォルトの名無しさん mailto:sage [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 名前：デフォルトの名無しさん mailto:sage [2010/07/12(月) 21:36:12 ]: 相対的にload/storeとレジスタ、キャッシュが少ない
376 名前：デフォルトの名無しさん mailto:sage [2010/07/12(月) 22:41:09 ]: CC2.1・・・
377 名前：デフォルトの名無しさん mailto:sage [2010/07/12(月) 23:43:43 ]: どの命令の組み合わせなら4issue出来るんだろうか
378 名前：デフォルトの名無しさん mailto:sage [2010/07/13(火) 00:03:57 ]: warp sizeはどうなるの? まさか48?
379 名前：デフォルトの名無しさん mailto:sage [2010/07/13(火) 00:04:52 ]: どう考えても32だが
380 名前：デフォルトの名無しさん mailto:sage [2010/07/13(火) 13:03:45 ]: >>375
同時に実行される命令が2->4になるだけなので、相対的な容量はあまり問題にならないような。
むしろメモリバンド幅が足りない。
381 名前：デフォルトの名無しさん mailto:sage [2010/07/13(火) 15:59:40 ]: 460だとメモリ帯域も細いから、倍精度だともはやCPUで走らすのと大差無いんじゃね？
382 名前：デフォルトの名無しさん mailto:sage [2010/07/13(火) 17:49:56 ]: 256bit 3208MHzの465より256bit 3600MHzの460の方がメモリ帯域は大きいはずだが
383 名前：デフォルトの名無しさん mailto:sage [2010/07/13(火) 19:42:19 ]: アイドルの消費電力が低いし、お試しに買うのにいいですよ。
384 名前：デフォルトの名無しさん mailto:sage [2010/07/16(金) 06:15:22 ]: ttp://pc.watch.impress.co.jp/docs/column/kaigai/20100716_380986.html

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

SP自体は単体で倍精度扱えたのか
385 名前：デフォルトの名無しさん mailto:sage [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 名前：デフォルトの名無しさん mailto:sage [2010/07/18(日) 12:32:31 ]: >>385
cutilがビルド済みなら、-L <libcutil.aがあるディレクトリ> -lcutil
ビルドさえしてないなら、NVIDIA_CUDA_SDK/common辺りでビルド。
手元に環境がないので詳細は割愛。
387 名前：385 mailto:sage [2010/07/18(日) 13:13:36 ]: >>386
無事に実行できました！
即レスありがとうございます。

なにが問題だったのでしょうか
libcutil.aにcutilが含まれてるということなのでしょうか
388 名前：デフォルトの名無しさん mailto:sage [2010/07/18(日) 13:45:08 ]: >>387
まぁ、簡単に言えばそういうこと。nm libcutil.a|gerp T してみると知見が得られるかもねw
389 名前：デフォルトの名無しさん mailto:sage [2010/07/19(月) 18:58:30 ]: このスレの住人は本当に出来た人が多いね。
と思ってたけど、ここはそういうスレなんだね。テンプレをあらためて見て気が付いたよ。

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

手元には無いので保障は出来ないけど、立ち読みした感じだと分かりやすく為になりそうなことが書いてある本だったよ。ひょっとするとシリーズの違う本の方がドンピシャの内容かもしれない。気が向いたら本屋で内容をチェックしてみてね。
390 名前：デフォルトの名無しさん mailto:sage [2010/07/21(水) 11:49:31 ]: スレ違いかもしれませんが、
計算中がうまくいかず、期待してないところで0.0が出てきたので調べてみたら、
1.0e-45以下の小さい数字が出てくることが多々あり、どうやらfloatの扱える範囲外
なので0.0にされてしまっていたようです。

double(お勧めされてない？)を使う、
全部に1.0e+20くらいかけといて結果を見るときに脳内で割る
とかそんな解決方法でしょうか？
391 名前：デフォルトの名無しさん mailto:sage [2010/07/21(水) 13:06:12 ]: >>390
CPUではfloatで動くんですか
392 名前：390 mailto:sage [2010/07/21(水) 14:04:49 ]: >>391
うごきません。cpuでも0.0になってしまいます。
dloubleにしたら1.0e-48とか表示できるんですが、
こういうときの常套手段ってあるんでしょうか
393 名前：デフォルトの名無しさん mailto:sage [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 mailto:sage [2010/07/21(水) 15:43:36 ]: >>393
なるほど。気をつけます。

しかも情報落ちまでしてました…
doubleより有効数字小さいのか
doubleで作ったプログラムそのまま移植できないんですね。
395 名前：デフォルトの名無しさん mailto:sage [2010/07/21(水) 16:13:21 ]: あ、っていうかfloat使う前提で書いちゃったけど、
ターゲットにGTX480を使えるか、もしくは
速度をそこまで気にしないならdoubleでもいいんじゃない？
あと部分的にdouble使ってもいいし。
396 名前：デフォルトの名無しさん mailto:sage [2010/07/21(水) 17:58:51 ]: 伊理先生の「数値計算の常識」が参考になるよ
読んでなかったぜひ
397 名前：デフォルトの名無しさん mailto:sage [2010/07/22(木) 01:27:32 ]: >>394
そりゃぁ、doubleは「倍精度」なんだから。で、余程cuda向きの演算でない限りdoubleでXeonに勝つのは至難の技だから、
ドラスティックにアルゴリズムやデータ構造を見直さないとダメかもね。
# そして、見直した結果CPUのままでも速くなってしまうのはよくある話
398 名前：デフォルトの名無しさん mailto:sage [2010/07/22(木) 02:40:04 ]: 精度にこだわるならCUDAはやめときな。
どうしてもと言うなら多倍長のライブラリをつくって・・・・てやるなら、x87を使うのが正解だろう。
そもそもCUDAを使うメリットとしてはCPUに比べメモリ帯域が広い、
並列度が高いと言うだけだから、これらが有効に使えないのなら、CPUでやった方がいいよ。
もはやCUDAを使ったからX倍速くなりました！と言うことに価値はなくなってきたからな。
CUDAを使ったからYX倍（1<Y<3）速くなりました！というなら価値はありそうだ。
CUDAを使ったからZYX倍速くなりました！というならそれは比較がおかしい。
399 名前：デフォルトの名無しさん mailto:sage [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 名前：デフォルトの名無しさん mailto:sage [2010/07/22(木) 10:06:37 ]: >精度にこだわるならCUDAはやめときな
学者共に言ってやってくれ
401 名前：デフォルトの名無しさん mailto:sage [2010/07/22(木) 11:55:37 ]: 学者はFermi版Tesla(Tesla版Fermi？)を使って倍精度で計算するので問題ない・・・・
402 名前：デフォルトの名無しさん [2010/07/22(木) 12:45:29 ]: 念のため言っとくが単精度floatは32ビット整数すら誤差なく格納することができないからな!
403 名前：デフォルトの名無しさん mailto:sage [2010/07/22(木) 19:48:10 ]: >>402
何この人恥ずかしい・・・
404 名前：デフォルトの名無しさん [2010/07/22(木) 21:49:08 ]: 下には下がいるってことがよくわかる
405 名前：デフォルトの名無しさん mailto:sage [2010/07/23(金) 00:02:17 ]: っていうか別に倍精度の値域が必要な訳じゃねえんだろ？
1.0e-7以下の精度が必要な訳じゃねえんだろ？
うぜえよ
406 名前：デフォルトの名無しさん mailto:sage [2010/07/23(金) 00:22:22 ]: だけど、もともとカオス性を備えた系のシミュレーションだとどんだけ精度があってもねえ
407 名前：デフォルトの名無しさん mailto:sage [2010/07/24(土) 16:46:26 ]: Fortranから呼び出す時って二次元配列の順番変わっちゃうんでしょ？
操作する時に列から回すとメモリが連続でなくなっちゃってアクセス遅くなったりする？
408 名前：デフォルトの名無しさん mailto:sage [2010/07/24(土) 21:04:41 ]: データの連続方向にスレッドを並べないと、滅茶苦茶遅くなります。
スレ違いにはなりますが、CPUの場合もデータ領域の大きさが大きい場合にはキャッシュ効率が全く変わってしまいます。
409 名前：デフォルトの名無しさん mailto:sage [2010/07/24(土) 21:30:40 ]: そういえばMSDNかどっかにサンプルがあったなあ
配列の順番変えるだけで10倍だか100倍だかのオーダーで実行速度が変わるやつ
確か方っぽは徹底的にキャッシュミスする並べ方らしかった
410 名前：デフォルトの名無しさん mailto:sage [2010/07/24(土) 22:16:57 ]: 一次元配列も
1...1|2...2|...|n...n|
よりも
1...n|1...n|...|1...n|
の方がいいのかなあ。数字はスレッド番号ね。
411 名前：デフォルトの名無しさん mailto:sage [2010/07/24(土) 23:51:35 ]: >>410
CUDAの話なら実験してみるといいよ。驚くから。
412 名前：デフォルトの名無しさん mailto:sage [2010/07/25(日) 00:01:44 ]: 2次元配列も全部1次元にしてるわ・・・
413 名前：デフォルトの名無しさん mailto:sage [2010/07/25(日) 01:24:39 ]: コンスタントメモリっていつ何時つかうのん？
414 名前：デフォルトの名無しさん mailto:sage [2010/07/25(日) 02:21:28 ]: gpu側でそれほど大きくなくていいけど速い必要のあるテーブルを参照したいとき。
415 名前：デフォルトの名無しさん mailto:sage [2010/07/25(日) 10:07:32 ]: たとえばnを外部から入力させてベクトルをn倍する関数を作りたいとき
コンスタントメモリはグローバルメモリでそんな早くないけど
一回キャッシュされたら全部のスレッドはnを読むからレジスタと同じくらい速くなる

て考えはおk？
416 名前：デフォルトの名無しさん mailto:sage [2010/07/25(日) 13:31:18 ]: >>415
パラメータで渡せないなら定数メモリはありだね。
417 名前：デフォルトの名無しさん mailto:sage [2010/07/25(日) 14:30:13 ]: >>415
最近のGPU使ってないから違うかもしれないけど、
コンスタントメモリって勝手にキャッシュとかされなくない？
まあ、ワープ内全スレッドからの同時アクセスなら１回で済むから
何回も読みに行かなければ問題はないだろうけど。
418 名前：デフォルトの名無しさん mailto:sage [2010/07/25(日) 15:53:11 ]: あれ、コンスタントメモリってキャッシュする命令があるんだっけ。
419 名前：デフォルトの名無しさん mailto:sage [2010/07/25(日) 16:37:58 ]: 問答無用でキャッシュされます。
420 名前：デフォルトの名無しさん mailto:sage [2010/07/25(日) 18:43:06 ]: もしかしなくてもソートって並列計算に向いてない？？
421 名前：デフォルトの名無しさん mailto:sage [2010/07/25(日) 20:00:43 ]: bitなんちゃらソートとか、並列計算用のソートアルゴリズムもあるにはあります。
CUDA SDKのbitonicを参照あれ。
422 名前：デフォルトの名無しさん mailto:sage [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 名前：デフォルトの名無しさん mailto:sage [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 名前：デフォルトの名無しさん mailto:sage [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 mailto:sage [2010/07/25(日) 20:35:43 ]: main.hの中で色々なヘッダを纏めてインクルードしてます．
不勉強なので問題外のレベルのコードだと思いますが，方向性のヒントでも貰えたら幸いです．
よろしくお願いします．
426 名前：デフォルトの名無しさん mailto:sage [2010/07/25(日) 20:58:06 ]: CUDA化された姫野ベンチのソースってどっかで公開されていますか？
427 名前：デフォルトの名無しさん mailto:sage [2010/07/27(火) 02:26:16 ]: 最近CUDAについて学び始めた初心者です。
現在三重対角行列のトーマス法をcudaで実装しようとしているのですが、計算式にある再帰処理に手も足も出ない状態です。
調べてみるとcuda（GPU）ではサポートしていないと出てきます。
GPUでの処理を考えるとこれは当然であることは分かるのですが、カーネル部分に何か工夫を行えば動かす事は可能となるのでしょうか？

もし行えるのであれば、効率等は考えなくても良いので良ければご教示お願いします。
428 名前：デフォルトの名無しさん mailto:sage [2010/07/27(火) 02:28:13 ]: 再起処理を自前実装する
429 名前：デフォルトの名無しさん mailto:sage [2010/07/27(火) 04:51:48 ]: >>426
自分で書いてみたら？
CUDAの練習にもなるし。
430 名前：デフォルトの名無しさん mailto:sage [2010/07/27(火) 08:25:33 ]: cuda3.1から再帰は使える。
ただし、Fermiのみ。
431 名前：デフォルトの名無しさん mailto:sage [2010/07/27(火) 11:14:48 ]: GPUだと再帰処理ができないのが何故当然なのですか。
432 名前：デフォルトの名無しさん mailto:sage [2010/07/27(火) 11:18:10 ]: 誰も当然とは書いていませんが・・・。
まぁ、動いたとしても遅いわけで。
433 名前：427 mailto:sage [2010/07/27(火) 20:43:18 ]: >>428
再帰処理を自前実装するとは、再帰処理を非再帰処理に変えて実装するという意味ですか？
それとも何か書き方があるのですか？

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

>>431
すみません。初心者なもので、GPUでは再帰はできないって認識で固定していてしまいました。
並列で動かす各スレッドのどれが終わっているか分からないので，再帰はできないという認識でした。
434 名前：デフォルトの名無しさん mailto:sage [2010/07/27(火) 21:20:51 ]: 各スレッドは各プロセッサにつきひとつ割り当てられて一度にプロセッサ数のスレッドが
スレッド÷プロセッサ回実行されるのですか？
それとも全スレッドはひとつのプロセッサに複数個割り当てられて
(ひとつのCUDAコアは同時に複数個のスレッドを処理できる？)
全スレッドが一回だけ同時に実行されるのですか？
435 名前：デフォルトの名無しさん mailto:sage [2010/07/27(火) 22:08:50 ]: >>434
難解な表現ですが、どちらかというと後者に近いです。

CUDAコアは複数個のスレッドを交互に処理します。
またCUDAコアが管理できるスレッド数には限界があり、
スレッドが多い場合には一部のスレッドのみが交互に実行され、
それ以外のスレッドは実行中のどれかのスレッド（スレッドブロック単位）が完了するのを待ちます。
436 名前：デフォルトの名無しさん mailto:sage [2010/07/28(水) 05:09:13 ]: >>435
交互に処理ってちょっと違和感あって資料見たら、本当に交互に処理するのね・・・。
メモリ関係のレイテンシ待ちしてる間暇だから他のワープも処理してやんよって感じか。

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

結果的には大抵、ローカル変数を全部配列化して、再帰レベルカウントを配列インデックスに指定するような方法で
対応できる。CUDAの場合それをどこのメモリに置くかを適切に決めないといけないけど。
438 名前：デフォルトの名無しさん mailto:sage [2010/07/28(水) 08:27:25 ]: >>435
なるほど。でも待つということは前者のようにも感じます
480個の切符売り場に数万人が並ぶかんじでしょうか。隣がすいたら移ることもできる？
同じスレッドブロックの仲間たちは同じ列に並んでいるのでしょうか。それとも同時に切符を買えているのですか？

また、同じwarpの仲間たちではどうでしょうか。halfwarp単位で同時にメモリアクセスってことは同時に切符を買えている？？
439 名前：デフォルトの名無しさん mailto:sage [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 名前：デフォルトの名無しさん mailto:sage [2010/07/28(水) 10:18:24 ]: ＞それとも最大数ってことは1024*1*1や256*2*2はいいけど1*1*1024はダメですよってこと？
そうです。

＞最終的にスレッドの上限は65536*1024ってこと？
こちらは65535*65535*1024ですね。
441 名前：デフォルトの名無しさん mailto:sage [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 名前：デフォルトの名無しさん mailto:sage [2010/07/28(水) 12:46:48 ]: 環境くらい晒せや
443 名前：デフォルトの名無しさん mailto:sage [2010/07/28(水) 15:04:09 ]: 具体的にどういうエラーが出てるのか書けよ馬鹿垂れ
444 名前：デフォルトの名無しさん mailto:sage [2010/07/28(水) 20:41:02 ]: GTX 480でキャッシュが使われているかどうか判定する方法はあるのでしょうか？
プロファイラにかけてみるとl1_global_load_hitが0で、
l1_global_load_missには値が入っています。
キャッシュが有効利用されていないと思っているのですが、どうなのでしょう。
445 名前：デフォルトの名無しさん mailto:sage [2010/07/28(水) 23:15:00 ]: >>441
ここの人たちは他と違ってかなり優しいけど、
はっきり言っておまえの質問は酷すぎる。
ゆとり世代か？
そもそもCUDA初心者と言っている時点でC言語の初心者じゃないのか？
環境構築はCUDA以前の問題だからな。
446 名前：デフォルトの名無しさん mailto:sage [2010/07/29(木) 00:36:28 ]: >>444
そもそもデータを再利用するような計算なのでしょうか？
一回読んで終わりならそうなる。
447 名前：439 mailto:sage [2010/07/29(木) 10:27:45 ]: >>440
ありがとうございます。
二番目と三番目で表現は一緒なのに内容は違うんですね
448 名前：デフォルトの名無しさん mailto:sage [2010/07/29(木) 20:19:44 ]: >>438

>>436をみたかんじ
切符をたのんで財布からお金をだしてるあいだに窓口は次の人の買う切符を聞いてるかんじかね
449 名前：デフォルトの名無しさん mailto:sage [2010/07/29(木) 20:39:06 ]: >>444
__syncthreads()入れてる？
450 名前：デフォルトの名無しさん mailto:sage [2010/07/29(木) 21:49:51 ]: >>448
同じ人が再び列に並んできっぷを買い続けないといけないので例としては適切でないかと。

あれだ・・・・火縄銃を３人交代で打つ戦法・・・何の例だか分からなくなってきたが。
451 名前：デフォルトの名無しさん mailto:sage [2010/07/29(木) 23:25:15 ]: 卓球ダブルス
452 名前：444 mailto:sage [2010/07/29(木) 23:30:21 ]: >>446
そうなのですか。
最も今回のコードが1回の読み込み(ブロック内ですよね?)にあたるか
よく分からないのですが…
コアレスメモリアクセスにしてもキャッシュにヒットしないのは
ブロック単位で計測されているからでしょうか?
それともコードが間違ってて本当にキャッシュが効いていないのでしょうか?
キャッシュが何か良く分からなくなってきました。

>>449
色々なとこに入れてみましたけど効果はありませんでした。
453 名前：デフォルトの名無しさん mailto:sage [2010/07/30(金) 04:17:09 ]: >>452
キャッシュとはデータが再利用されなれば意味がないよ。
プリフェチと勘違いしてない？
あとコアレッシングにしても速さが変わらないとなると、L2が効いているのかも。
いずれにせよ、まずはShared Memoryで書いて、効果が出るかどうかを調べることだね。
つか、そんな質問する位ならコードさらせや。
454 名前：デフォルトの名無しさん mailto:sage [2010/07/30(金) 21:05:15 ]: FortranでCUDAができるというような話を聞いたんですが、もうできるのでしょうか？
それは純粋にFortranで書くのですか？
それともCで書いた関数をFortranで呼び出して使うという方法でしょうか？
455 名前：デフォルトの名無しさん mailto:sage [2010/07/30(金) 22:08:11 ]: 金払いたくないならFortranからCを呼び出す
PGIからFortran用CUDAコンパイラが出てるからそれ買えば全部Fortranで書ける。
ただしFortran90以降の書き方じゃないとコンパイルできないっぽい。

興味あったら15日体験版あるから試してみたら？
456 名前：454 mailto:sage [2010/07/31(土) 12:46:48 ]: >>455
ありがとうございます。
FortranからCを呼び出しても実行速度やらはかわらないのでしょうか。
プログラミングの情報はCでのCUDAの方が多いようなのでどちらも同じなら
C呼び出しの方を使いたいと思うのですが
457 名前：デフォルトの名無しさん mailto:sage [2010/08/02(月) 10:20:03 ]: >>456
FortranからCを呼び出す部分は通常演算の核になる部分から離れているので余り問題にならないはず。
と言うより、それが問題になるようなら演算の核になる部分が小さすぎるからCUDAに向かない実装と言うことになる。
一度CUDAのカーネルを呼び出したら、最低でも100msは戻ってこないくらい演算を集中させないと
カーネル呼び出しのコスト自体がネックになってしまう。
458 名前：デフォルトの名無しさん mailto:sage [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 名前：デフォルトの名無しさん mailto:sage [2010/08/02(月) 12:14:31 ]: >>458
実際にやってみれば判るけど、1スレッドしかなくても一つのサブプロセッサが動く。
つまりブロック数が同じとき、スレッド数が1でも2でも同じ時間が掛かる。
また、並列して投入できるカーネル関数がないなら空いているサブプロセッサがあっても使われない。
従って、使うgpuのコア数分投入した方がいい。
# 尤も、GUI表示と兼用ならその分は少なくても無駄は出ないけれど。
460 名前：458 mailto:sage [2010/08/02(月) 15:08:29 ]: >>459
ブロック1000スレッド1より
ブロック1スレッド1000のほうがいんですね
＞使うgpuのコア数分投入した方がいい。
というのはスレッド数をってことですか？（480基だったら400くらい？）

ひとつのコアにつき同じブロック内のスレッドがひとつづつバーーっと割り当てられて
ブロック内のスレッドは各々のコアが同時に処理するのでしたっけ
461 名前：デフォルトの名無しさん mailto:sage [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 mailto:sage [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 mailto:sage [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 名前：デフォルトの名無しさん mailto:sage [2010/08/02(月) 18:20:33 ]: >>461
書いていることがめちゃくちゃなのだが・・・・。
もしもネタでないのであればじっくり読み返していただきたい。
465 名前：461 mailto:sage [2010/08/02(月) 19:22:16 ]: あれ?　そんなに間違ってたっけか。知識が古いからFermiでの数値と違っているかもしれないけれど。
後は用語がいい加減かな。
まぁいいや、誰かの修正を待とう。

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

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

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

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

32を超えると32単位で実行順は保証されないが、1ブロック内のスレッドは同一のSMで実行される
だからSM数≦ブロック数になるのがよい
467 名前：デフォルトの名無しさん mailto:sage [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 mailto:sage [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 名前：デフォルトの名無しさん mailto:sage [2010/08/04(水) 14:44:24 ]: カーネル関数の引数はどこに保存されるのだっけ
たとえばint n=1000があって
kernel<<<>>>(n);
とやるとnはレジスタに入ってくれるの？
グローバルメモリに入っちゃうの？
配列と普通の変数で違ってくるみたいなレスを昔みたような気がするんだけど
470 名前：デフォルトの名無しさん mailto:sage [2010/08/04(水) 15:57:37 ]: シェアードメモリに入る
471 名前：469 mailto:sage [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 名前：デフォルトの名無しさん mailto:sage [2010/08/04(水) 17:54:00 ]: blockサイズ、threadサイズも入るから
473 名前：デフォルトの名無しさん mailto:sage [2010/08/09(月) 10:00:28 ]: >>471
double型は8の倍数のアドレスにしか置けないから。
474 名前：デフォルトの名無しさん mailto:sage [2010/08/09(月) 10:49:19 ]: >>471
それはさて、doubleでないといかんの?
475 名前：デフォルトの名無しさん [2010/08/10(火) 10:50:57 ]: GPU向けのコードを.cppファイルの中に書いてnvccでコンパイルすることはできないのでしょうか?
476 名前：デフォルトの名無しさん mailto:sage [2010/08/10(火) 11:18:14 ]: >>475
そのファイルの名前を.cuにすればコンパイルできます。
477 名前：デフォルトの名無しさん [2010/08/11(水) 11:11:33 ]: >>476
そうするとほかのコンパイラでコンパイルできなくなってしまいます。
CPU用コードとGPU用のコードを同じファイル内に書いて、ifdefでコンパイル時に切り替えたいのですが。
478 名前：デフォルトの名無しさん mailto:sage [2010/08/11(水) 11:13:22 ]: #include "hoge.cu"
479 名前：デフォルトの名無しさん mailto:sage [2010/08/11(水) 14:22:27 ]: >>475
nvccでオブジェクトファイル作って、
gccなりでリンクすればいいじゃん。
480 名前：デフォルトの名無しさん mailto:sage [2010/08/11(水) 18:44:41 ]: 変に凝って一つのソースで両方のコード書くと条件コンパイルだらけで醜くなるだけだと思うなぁ。

それはさて、thread数が32固定でOccupancyが0.167の関数があったから>467を踏まえて実験してみた。
なるほど、thread数を192にしてやるとOccupancyは1.00になった。
問題は、ある程度計算量が多いときは192の方が数割速くなったけど計算量が少ないときに若干遅くなったこと。
ループ回数のチェックはしていないからなんとも言えないけれど、これはちょっと気になるところ。
# 尤も、問題の関数はトータルの計算時間の極一部でしかないから調査は保留だなぁ。

さてと、明日からは本体の方のSharedMemoryの使い方の見直しでもするか。
# こちらは実装の都合でthread数を増やすとSharedMemoryが足りなくなるので32固定……
481 名前：デフォルトの名無しさん mailto:sage [2010/08/12(木) 03:35:42 ]: >>480
OccupancyはあくまでMP(=SM =Streaming Multiprocessor)内の占有率だから、
ブロック数がMPの数より少ないと、一部のMPがやることなくて遊んじゃうのかも？
遊んじゃうぐらいならOccupancy下げてMP埋めたほうが若干速いのかな。

つか若干ぐらいなら別にそうする必要なくて、Occupancy上げとけばいいのかもしんないね。
482 名前：デフォルトの名無しさん [2010/08/15(日) 21:59:10 ]: 3.1を導入して、カーネル関数内ででprintf()したら、
combora.cu(275): error: identifier "cuprintf" is undefined
と言われました。

環境はVCでx64です。

何が間違っているのですか？
483 名前：デフォルトの名無しさん mailto:sage [2010/08/15(日) 22:21:56 ]: CC2.0以降のみ、printf()対応のようですね…
お騒がせしました。
484 名前：471 mailto:sage [2010/08/17(火) 11:49:24 ]: >>473
あーなるほど。

>>474
doubleの欠点てたくさんあります？？
485 名前：デフォルトの名無しさん mailto:sage [2010/08/20(金) 14:10:18 ]: CUDAのドライバって
CUDAの専用ページにあるドライバと
使っているデバイスの最新ドライバは
同じ種類のものなのでしょうか？

GTX480をつかうのですが、
CUDAのページにある
devdriver_3.1_linux_64_256.40.run.sh
とGeForceGTX480用の最新ドライバである
NVIDIA-Linux-x86_64-256.44.run.sh
では480用の方が最新に見えるのですが、こちらを使用して大丈夫なのでしょうか？
486 名前：デフォルトの名無しさん mailto:sage [2010/08/20(金) 18:20:26 ]: 半年くらい前にLinuxでNVの最新ドライバ入れたらCUDA動かなかった。
CUDA用ドライバなら動いた。
487 名前：デフォルトの名無しさん mailto:sage [2010/08/21(土) 20:55:17 ]: ふぅ、アク禁解けた。

>>485
開発チームが違うのか、CUDA対応が常に遅れるので専用を拾わないとダメ。

>>484
doubleの欠点は、遅い、大きい、の2点に尽きる。
演算時間も掛かるし並列度も落ちるみたいだし、
データ量が倍になるから共有メモリに置けるデータ数が減るし、
HostからDeviceへの転送時間も長くなる。
488 名前：デフォルトの名無しさん mailto:sage [2010/08/22(日) 20:25:57 ]: 総和計算はcpuのforループでやるより
gpuで工夫した並列リダクションのほうが何倍も早いものなのでしょうか？

1000*1000程度の配列x,y,zがあり、
それぞれの成分に重みをかけた和
の総和計算を重みを変えて1000通りやりたいのですが
どこを並列化したら良いのか悩んでいます。
489 名前：デフォルトの名無しさん mailto:sage [2010/08/22(日) 22:32:26 ]: >>488
xyzは固定かな。だったら転送量が少ないから望みはあるかも。
1000x1000を適当に分割した範囲で和を共有メモリに置いて、総和の計算は二分木でやればOK。
その総和をまた共有メモリに置いておいて、更に次のステップで全面分の総和を出すって感じかな。
490 名前：デフォルトの名無しさん mailto:sage [2010/08/22(日) 22:36:18 ]: CUDAのreductionサンプルは
最後の一手だけやれば、他は無理にやらなくてもいい。
どうせ誤差レベルの速度向上にしかならない。
491 名前：デフォルトの名無しさん mailto:sage [2010/08/23(月) 03:04:29 ]: あのサンプル、真面目に先頭からドキュメントとコードを読んでいって損したわw
492 名前：デフォルトの名無しさん mailto:sage [2010/08/23(月) 08:49:02 ]: 同感。どのサンプルも動かすためにあれこれ余計なコードが多過ぎる。
デモサンプルと機能サンプルと実装サンプルは区別して欲しかった。
493 名前：デフォルトの名無しさん mailto:sage [2010/08/23(月) 08:54:17 ]: 実はそういう意味で言ったのではない
494 名前：488 mailto:sage [2010/08/23(月) 11:09:01 ]: >>489
たとえば10個分の和を各スレッドでやらせることを256スレッド/ブロックで391ブロックやって
そのままその中で各スレッドの和を共有メモリに置いて2個づつ合計することを8回
391個の和になるからそれぞれグローバルメモリに書き込む
次のカーネルで同じように和を取る
ということを1000回条件変えて繰り返すってことでしょうか。

変える条件は何にも依存しない場合、10条件繰り返すよりも
まとめて10条件分計算して100回繰り返した方が速いかも？

>>490
最後の一手ってリダクション7の部分ですか？
495 名前：デフォルトの名無しさん mailto:sage [2010/08/23(月) 16:46:39 ]: cuda2.3を使っていて3.1にしようと思い、
centOS5.3→5.5にしたところ、デバイスの型番が認識できなくなりました。
/sbin/lspciを実行すると
03:00.0 VGA compatible controller: nVidia Corporation Unknown device 06c0 (rev a3)
となります。ドライバをいくつか変えてもダメでした
そこでOSを5.3に戻してもと入れていたドライバを入れたところ、同じ現象になってしまいました。
当然devicequeryも実行できません

Xは起動しているので認識自体はしてるはずなのですが。

どんな原因が考えられますか？
ドライバ入れた後なにか特殊な操作が必要でしたっけ？
型番認識自体はドライバとかなしでもできるものなのではないのですか

ASUS の ENGTX480/2DI/1536MD5を使っています。なにか情報足りなかったら追記します。
496 名前：490 mailto:sage [2010/08/23(月) 19:34:59 ]: >>494
そう。単純な総和のreductionなんて
N個のデータでN-1の加算しかしないのだから
コンピュートバウンドではなく、IOバウンド
IOを最小限にするようなアルゴリズムを組めば
多少演算が非効率でも誤差レベルの影響しか出てこない。
497 名前：デフォルトの名無しさん mailto:sage [2010/08/23(月) 22:05:43 ]: >>495
CentOSは知らないけれど、nVIDIAのGPUは型番を認識できないと汎用VGAと認識されるから1024x768位のXなら動く。
なので、認識はできていると言うか、できてないと言うか……
RedHatならKudzuのメッセージ見て対処できるんだけどCentOSはどうなんだろ。Linux板で聞いた方が早いかも。
498 名前：デフォルトの名無しさん mailto:sage [2010/08/24(火) 15:57:59 ]: __global__でprintf()ってどうやるのん
そのまま書いたら
calling a host function from a __device__/__global__ function is not allowed

って出た。

CUDA3.1ってデバイスでprintf()使えるんだよね？？

Device 0: "GeForce GTX 480"
CUDA Driver Version: 3.10
CUDA Runtime Version: 3.10
CUDA Capability Major revision number: 2
CUDA Capability Minor revision number: 0

なんか間違ってますか？
499 名前：デフォルトの名無しさん mailto:sage [2010/08/24(火) 16:13:31 ]: コンパイラオプションをつけているのでしょうか？
500 名前：498 mailto:sage [2010/08/24(火) 16:35:31 ]: >>499
すみません。つけてないです
どのようなコンパイラオプションがいるのでしょうか？
501 名前：デフォルトの名無しさん mailto:sage [2010/08/24(火) 18:36:42 ]: -arch
502 名前：498 mailto:sage [2010/08/24(火) 18:44:29 ]: >>501
-arch=sm_20
でできました！！

勉強不足でした。
ありがとうございます。
503 名前：デフォルトの名無しさん mailto:sage [2010/08/26(木) 02:45:20 ]: Compute Capability 2.0からはカーネルの引数はシェアードメモリを使わずに
コンスタントメモリに入れるようなのですが
カーネル内でコンスタントメモリの引数をシェアードメモリに読み込んだ方が早いのですか？
504 名前：デフォルトの名無しさん mailto:sage [2010/08/26(木) 08:01:25 ]: >>503
疑問に思ったら先ずは実測。
定数メモリは充分早いはずだから、共有メモリに移し変えるメリットはないんじゃない?
505 名前：503 mailto:sage [2010/08/26(木) 21:31:43 ]: >>504
たしかに気にならない程度の速度でした。なんでコンスタントメモリになったのかなぁ
シェアードメモリを空けるため？？

ところで
extern __shared__ float array[];
__shared__ float array[128];
はどのような用途で使い分け、機能面でどう違いがあるのでしょうか。
506 名前：デフォルトの名無しさん mailto:sage [2010/08/26(木) 22:18:37 ]: >>505
4要素使いたい場合だと今まで苦労してたからなぁ～
コンスタントメモリに入るのは個人的にうれしい
507 名前：デフォルトの名無しさん mailto:sage [2010/08/27(金) 01:22:18 ]: >>506
同意。引き数が多いと結構きつかった。

>>505
後半の質問の意味が判らない。要素数の定義の仕方の違いを言ってる?
508 名前：503 mailto:sage [2010/08/27(金) 01:33:44 ]: >>507
前者は要素数(size)をkernel<<<>>>()の第三引数で指定するようなことが書いてありました。
要素数をカーネルの外で定義するか中で定義するかで
ただ単に宣言の仕方が違うだけで動作は全く一緒なのでしょうか？
509 名前：デフォルトの名無しさん mailto:sage [2010/08/27(金) 02:06:51 ]: >>508
あー、そういうことね。<<<>>>で指定する場合は可変長にできるメリットがありますね。
出力はどうだろ。ptxを出力して見比べてみては?
510 名前：デフォルトの名無しさん mailto:sage [2010/08/27(金) 19:24:42 ]: 岡田賢治「CUDA高速GPUプログラミング」秀和システム
の p.140 のプログラム 07-02 をそのまま実行すると
画面がいったん消え、57行目(カーネル後、最初のcutilSafeCall)
でエラーだと表示されます。

BLOCK_SIZE を 16 から 32 (または 64, 128, 256, 512) に
変更すると、なぜか無事に実行します。

同じ状況になった人はいませんでしょうか。

OSはWindowsXP 32bit環境でCUDA 3.0です。
ＧＰＵはオンボードのGeForce9300mgpu, グラフィック兼用です。

block(16, 16)で1ブロックあたり 256 スレッドとなり、
512 スレッド以下なので、16 で正解だと思ったのですが、
2次元以上でも一辺が1 Warp=32 の倍数で 512 以下のほうがOK
ということでしょうか。
511 名前：デフォルトの名無しさん mailto:sage [2010/08/27(金) 19:42:27 ]: まぴょーん
512 名前：デフォルトの名無しさん mailto:sage [2010/08/27(金) 20:59:31 ]: >>511
>>510の本立ち読みしたら、巻末に「IPAの人とは関連ありません」って書いてあって笑ったｗ
513 名前：デフォルトの名無しさん mailto:sage [2010/08/27(金) 21:26:38 ]: >>512
雑談は余所でやれ、馬鹿たれが
514 名前：デフォルトの名無しさん mailto:sage [2010/08/27(金) 21:59:08 ]: うわぁ……
515 名前：デフォルトの名無しさん mailto:sage [2010/08/27(金) 22:29:50 ]: >>510
エラー内容を詳しく。まぁ、なんとなく落ちは見えたけど。
516 名前：デフォルトの名無しさん mailto:sage [2010/08/29(日) 00:50:01 ]: >>515
しばらく沈黙して
cudamatrix.cu(57) : cudaSafeCall() Runtime API error : the launch timed out and was terminated.
と出ます。

確認のため 07-01 と 07-02 の結果の一部を出力してみたらおかしかったので、
概算してみたところ初期設定が大きすぎて int の制限を超えていることがわかりました。
30-31行目の 1024*1024 を 1024 に代えて確認してみました。
BLOCK_SIZE=16 だと読み出しが不安定ですが計算結果はデバイスに残っているようです。
32 以上だとデータの読み出しは成功しますが、カーネルの計算をスルーしているような感じです。
517 名前：デフォルトの名無しさん mailto:sage [2010/08/29(日) 04:46:02 ]: >>516
そのエラー自体は、書かれている通りカーネルの実行時間が長すぎて出ています。
LinuxだったらX立ち上げないで(run level 3とか)起動すれば通るかもしれないけれど、Windowsはどうするんだろ。
本が想定しているGPUより能力が低いのが問題なのかもしれないので、ネカフェかなんかで速い端末で動かしてみるとか。
後は処理量を加減するしかないけれど、それには元のコードを見ないことにはなんとも。
518 名前：デフォルトの名無しさん mailto:sage [2010/08/29(日) 07:45:18 ]: CUDAのドライバは最新ですか。一つ前はなんかそんなエラーがでましたよ。
519 名前：デフォルトの名無しさん mailto:sage [2010/08/29(日) 19:50:22 ]: >>517
なるほど、実行時間が長すぎると出るエラーですか。

p.206のベンチマークでも同じプログラムで GeForce9400M/256MB
では timeout になっているようなので、MATRIX_SIZE を減らすしか
ないようですね。確かに、１辺 512 ならOKでした。

ただ、同様の行列積算の float版が ASCII. technologies 2009/12, p.46
に載っていたので試してみたら、シェアードメモリを使っても使わなくても
SIZE=1024 でも 2048 でもちゃんと動きました。
# GPUは float 演算のほうが速い?

>>518
3.0 です。3.1 なら出ない?
520 名前：デフォルトの名無しさん mailto:sage [2010/08/29(日) 19:57:54 ]: intよりfloatの方が速い
521 名前：デフォルトの名無しさん mailto:sage [2010/08/29(日) 20:04:26 ]: もうちょっと具体的に言うとFMAがfloatでは使えるので、a=b×cを4クロックで実行できる。
FMA考えなくてもaddもmultiplyも4クロック。

intだとint addに4クロック、int multiplyに16クロックかかる。
522 名前：デフォルトの名無しさん mailto:sage [2010/08/29(日) 21:02:16 ]: >>519
自己レスです。
int でシェアードメモリを用いたら 1024 でも 2048 でも timeout にならずに
動きました。原因がわかってすっきりしました。
皆様、ありがとうございました。(_O_)
523 名前：デフォルトの名無しさん mailto:sage [2010/08/30(月) 13:19:17 ]: 1000要素の配列データを配列として1回cudamemcpyするのと
要素ごとに1000回cudamemcpyするのではなんでこんなにも時間が違うのですか
524 名前：デフォルトの名無しさん mailto:sage [2010/08/30(月) 13:55:15 ]: >>523
・単純に一回リクエストする際に生じるオーバヘッドが馬鹿にならない。
・実はcudamemcpyは完了復帰じゃないので一回のcudamemcpyでは即復帰してしまい実際の転送時間は測れない。
525 名前：523 mailto:sage [2010/08/30(月) 15:20:19 ]: >>524
>実はcudamemcpyは完了復帰じゃない
??
cudamemcpy直後にコピー先にアクセスしても信用ならない値なのですか？
526 名前：デフォルトの名無しさん mailto:sage [2010/08/30(月) 15:39:19 ]: もちろん信用できますのでご安心ください。
527 名前：デフォルトの名無しさん mailto:sage [2010/08/30(月) 15:45:36 ]: >>525
・HtoDの後はカーネル呼び出しの際に同期される。
・DtoHは完了復帰になる。但し、streamを使う場合は明示的に同期する必要がある。
528 名前：523 mailto:sage [2010/08/30(月) 16:00:55 ]: >>526-527
ああ、DtoHのことしか考えてませんでした。

>>527
じゃあHtoDの場合はカーネル呼び出しまではストリームっぽくなるから
データ送るなら早めにやったほうがお得なのでしょうか
529 名前：デフォルトの名無しさん mailto:sage [2010/08/30(月) 16:17:11 ]: >>528
とってもお得。
ついでに言えば、カーネル呼び出しも(限定的に)即時復帰だから
カーネル呼び出し後にDtoHするまで他の処理をして時間を稼いだ方がお得。
530 名前：デフォルトの名無しさん mailto:sage [2010/08/30(月) 16:36:39 ]: >>527
>・HtoDの後はカーネル呼び出しの際に同期される。
こういう仕様だと送信バッファの内容を変更できないわけですが・・・。

カーネル呼び出しが即復帰するのと間違えているのでは？
531 名前：デフォルトの名無しさん mailto:sage [2010/08/30(月) 17:16:46 ]: 処理→HtoD→カーネル　よりも　HtoD→処理→カーネル　の方が若干早かった気がしたんだが。

んじゃ、メモリコピーはストリームにしない限り完了復帰ってことか。
532 名前：デフォルトの名無しさん [2010/08/31(火) 08:28:18 ]: cudaMemcpyAsync
533 名前：デフォルトの名無しさん mailto:sage [2010/08/31(火) 13:16:50 ]: Warp sizeは32なんだけど Number of cores/Number of multiprocessorsが32なの。
この場合は1サイクルで1warpできるの？
その場合half warpとかそういう概念はかわってくるの？
534 名前：デフォルトの名無しさん mailto:sage [2010/08/31(火) 13:35:12 ]: >>533
1クロックで2個のハーフワープを処理する．
2クロックで2個のワープが処理されることになるが，上に書いたように，
1クロックで丸々1個のワープが処理されるわけではない．
535 名前：デフォルトの名無しさん mailto:sage [2010/08/31(火) 19:35:42 ]: 岡田賢治のCUDA高速GPUプログラミングを読んで勉強中なのですが、
質問させてください

バイトニックソートのGPUプログラム(ttp://www.shuwasystem.co.jp/support/7980html/2578.html)
でカーネル関数を10万回繰り返しているのは1回だけだと短すぎるからであって、実際にソートしてるのは
最初の1回だけ、と本に書いてあったのでforをはずして1回だけカーネルを実行し、結果コピー後に
printfで値を見たところ、正しくソートされていませんでした。

値をいくつか変えて試したところ、カーネル1回の実行ですべて正しくソートされるのは
ブロックが1つの時だけのような気がします。
が、本のサンプルでは4ブロックになっています。

この本は正しいのでしょうか？
ブロックが複数のときはどうしたら良いのでしょうか？
よろしくおねがいします。
536 名前：533 mailto:sage [2010/09/01(水) 10:37:12 ]: >>534
なるほど。考え方は特に変えなくていいわけですね。
537 名前：デフォルトの名無しさん [2010/09/01(水) 10:38:57 ]: 11030 アップデートファイルのダウンロードがタイムアウトしました。
FINAL FANTASY XIVを再起動してください。 3:40012,20498,20049
538 名前：デフォルトの名無しさん mailto:sage [2010/09/02(木) 10:33:26 ]: threadIdx.xやblockDim.x等のビルトイン変数はどのメモリに置いてあるのでしょうか
カーネル内で何度も参照するよりも
カーネル内のローカル変数に格納してレジスタに置いて参照した方が高速ですか？
539 名前：デフォルトの名無しさん mailto:sage [2010/09/02(木) 10:56:57 ]: >>538
その辺は勝手に最適化しますよ。ptx出力をご覧あれ。
あーそうそう、その辺のビルドイン変数はshort幅ですが、
受けるローカル変数はshort幅だと変なptxになりますね。
540 名前：デフォルトの名無しさん mailto:sage [2010/09/02(木) 12:58:42 ]: >>535
確かに、ブロックが複数だと正しい結果になりませんね。
Nvidiaのサンプルも１ブロックで処理していました。
サイズが512までならcpuで計算しても遜色ないので、
苦労してgpuで計算する必要はないですね。
541 名前：538 mailto:sage [2010/09/02(木) 13:16:50 ]: >>539
特に気にしなくていいわけですね。ありがとうございます。

ページロックメモリについてなのですが、
cudaMallocHost()にmalloc()と比べたデメリットってありますか
GPUのメモリと通信する配列をCPUに確保するときは
すべてcudaMallocHost()を使った方が良いのでしょうか
542 名前：デフォルトの名無しさん mailto:sage [2010/09/04(土) 21:31:02 ]: >>540
ですよね。
2^16くらいはソートしたいのに。
複数ブロックだとどうしたらいいんでしょう
1ブロック内で並び替えたバイトニック列をマージでリダクションするのでしょうか
543 名前：デフォルトの名無しさん mailto:sage [2010/09/06(月) 23:37:20 ]: CUDA Toolkitの3.1 64bit版をWindows7 x64にインストールしようとすると、
インストール途中のRegistering products...で止まってしまいます。

全く同じ構成のPCで昨日試したときは、この部分は一瞬で終わったんですが、
これはnVIDIA側の鯖が死んでいるのでしょうか？
544 名前：デフォルトの名無しさん mailto:sage [2010/09/07(火) 10:44:43 ]: >>541
CUDAテクニカルトレーニングには

　ページロックのメモリを多く割り当てすぎると、システム全体のパフォーマンスが低下する場合がある

とある。どのくらい割り当てるとどのていど低下するのかは（自分は）わからないし調べていない
545 名前：デフォルトの名無しさん mailto:sage [2010/09/07(火) 21:39:04 ]: >>541
ロックしたまま落ちると大変なことになりそうな悪寒。
その他の手段をやりつくしてからでも遅くないんでない?

>>544
少なくとも、スワップするような場合は明らかに違いが出る。
逆に言えば、スワップしない程度しか使っていなければ大差ないと思う。
546 名前：デフォルトの名無しさん mailto:sage [2010/09/09(木) 09:34:02 ]: >>543
うちのマシンでも同じ症状がでてる。
どうしたものか・・・。
547 名前：546 mailto:sage [2010/09/09(木) 10:56:43 ]: safeモードで起動して、インストールするといけた。
問題が発生してたのはWindows7 Ultimate x64をクリーンインストールした環境。

検証としてほかのWindows環境で試してみたよ(クリーン環境じゃない)
Windows Vista Ultimate x64、　Windows 7 Enterprise x64,
Windows Server 2008 R2 x64では問題は生じなかった。

ウイルス対策ソフトはServer以外ウイルスバスター2011がインストール済、
プロセッサはインテル、AMD両方あるので、おいらの結論としては、
Windows 7 Ultimateのサービスのどれかと相性が悪かった気がする。

困ってる人は参考にしてね。
548 名前：デフォルトの名無しさん mailto:sage [2010/09/09(木) 21:31:06 ]: 複数GPU環境でGPU単体で動くプログラムを実行すると
どのGPUがデフォルトで動くのかだれかしってる？
デフォルトのGPUが選択される優先順位ってどうなってるのかな？
549 名前：デフォルトの名無しさん mailto:sage [2010/09/09(木) 21:48:40 ]: プログラム作成者が指定する。

SDKのnbodyサンプルをみると、
devID = cutGetMaxGflopsDeviceId();
cudaSetDevice( devID );
でデフォルトのデバイスを設定してる。
つまり、一番ピーク性能がよいものを選択するよう書かれてる。

手抜きなcudaSetDevice( 0 );のようなコードだと
常にデバイス0が選択される。
550 名前：デフォルトの名無しさん mailto:sage [2010/09/10(金) 00:36:48 ]: >>547
自分はウイルスバスター2010の環境では大丈夫でしたが、
2011だとRegistering products...で止まりました。
(ただし、ウイルスバスター2011が原因かは分かりません。
リアルタイム検索を止めても、インストールはできなかったので)

ダウンロードしたexeに引数 -r を付けて起動すると、
OSをsafeモードにしなくてもインストール出来ました。
-rはサイレントモードを意味するようですが、
何が違うかは分かりません。
551 名前：デフォルトの名無しさん mailto:sage [2010/09/10(金) 23:52:37 ]: 処理時間が異常にかかる件、質問します。
下記のプログラムで「c |= block[6 * i + j];」の行があると、
すごく速度が遅くなります。
環境はWindow7+CUDA3.1+VisualStudio2008です。
原因を教えてください。
for(int i = 0; i < 11; i++){
char c = 0;
for(int j = 0; j < 6; j++){
c <<= 1;
c |= block[6 * i + j];
}
c += '.';
if(c > '9') c += 7;
if(c > 'Z') c += 6;
cudaTrip[(65535 * blockIdx.y + blockIdx.x) * 10 + i] = c;
}
552 名前：デフォルトの名無しさん mailto:sage [2010/09/11(土) 00:07:06 ]: block[]の配列がグローバルメモリにあるからかと。
553 名前：デフォルトの名無しさん mailto:sage [2010/09/11(土) 01:00:53 ]: トリップ検索か
554 名前：デフォルトの名無しさん mailto:sage [2010/09/11(土) 01:10:45 ]: >>552
ありがとうございます。
block[]をシェアードメモリにしても、大して改善しませんでした。
block[]にアクセスすると、Windows自体が激しく重くなります。
555 名前：デフォルトの名無しさん mailto:sage [2010/09/11(土) 01:41:19 ]: カーネル実行中は基本Windows重くなるよ。
軽くしたいなら、複数回のカーネル実行に分けな。
556 名前：デフォルトの名無しさん mailto:sage [2010/09/11(土) 02:00:57 ]: というか、Fermiなのかそうでないのか。
557 名前：デフォルトの名無しさん mailto:sage [2010/09/11(土) 08:31:14 ]: Fermiならキャッシュが効くから多少は軽いと思うけど。
いずれにしても、GPU向きじゃないコードだなぁ。
558 名前：デフォルトの名無しさん mailto:sage [2010/09/12(日) 02:40:52 ]: 分岐が原因・・・なわけないか。
559 名前：デフォルトの名無しさん mailto:sage [2010/09/15(水) 13:37:11 ]: 実行中のデバイスメモリの使用量はどのようにして知るのでしょうか。
どうもメモリがステップごとに増えてるような気がして時間推移をチェックしたいのですが。
linux,CUDA3.1,CC2.0です。
560 名前：デフォルトの名無しさん mailto:sage [2010/09/15(水) 17:46:14 ]: 3.2RCが公開されたね。
いろいろな機能が追加された様子。
561 名前：デフォルトの名無しさん mailto:sage [2010/09/15(水) 23:19:57 ]: >>559
メモリ確保時の引き数と戻り値を自前で管理するしか。
勿論、解放のときの処置も忘れずに。
あーそうそう、メモリ確保と解放を頻繁に行なうと分断化が起き易いから要注意。
SDKのバージョンでその辺りの振る舞いが変わるから、動いていたコード(とGPUの組み合わせ)が
バージョンが替わると突然メモリが足りなくなったりする。
562 名前：デフォルトの名無しさん mailto:sage [2010/09/16(木) 00:30:05 ]: nsightのビルドルールでビルドすると、
Visual Profilerでエラーが出ます。

SDK(サンプル)付属のビルドルールでビルドすると、
エラーは出ません。

皆さんはどうですか？

環境はCUDA 3.1/Windows 7 64bit/Nsight 1.0/ 258.96/GTX 480/VS2008 SP1です。
563 名前：デフォルトの名無しさん mailto:sage [2010/09/16(木) 13:24:38 ]: CUDAってクラスというかオブジェクト指向を利用することは可能ですか?
564 名前：デフォルトの名無しさん mailto:sage [2010/09/16(木) 13:29:48 ]: 質問です

GPGPUでは分岐処理は遅いと言われますが、どうして遅いのでしょうか？
565 名前：デフォルトの名無しさん mailto:sage [2010/09/16(木) 14:06:34 ]: >>564
全スレッドが同じ処理をするからif文がfalseでも勝手に先に進めず、
他のスレッドが終わるのを待つからでは。
ただ、たいていのプログラムはメモリアクセスのほうがネックになるから、
あまり気にしなくてもいいと思う。
566 名前：デフォルトの名無しさん mailto:sage [2010/09/16(木) 14:43:44 ]: >>565
っ Divergent blanch
567 名前：デフォルトの名無しさん mailto:sage [2010/09/16(木) 14:58:01 ]: CUDAを使った研究で、fortranのコンパイラでやることになったのですが、
分かりやすい教科書などないですかね。
なかなかfortranのものは見つからなかったんですが・・・
568 名前：デフォルトの名無しさん mailto:sage [2010/09/17(金) 02:47:42 ]: >>567
NVIDIAの糞営業にお伺い立てたら?
569 名前：デフォルトの名無しさん mailto:sage [2010/09/17(金) 02:49:38 ]: >>563
可能です。但し、パフォーマンスを追及すべきCUDAでパフォーマンスを
追及しにくいオブジェクト指向に固執することは無意味です。
肝心要の部分はオブジェクト指向を諦めるのが無難ですね。
570 名前：デフォルトの名無しさん mailto:sage [2010/09/17(金) 02:51:15 ]: ワーストシステムズに問い合わせたら資料くれるんじゃね？
571 名前：デフォルトの名無しさん mailto:sage [2010/09/17(金) 12:44:21 ]: >>569
やっぱりスピードを重視するとOODは不向きなんですね。ありがとうございます。
基本的にC++で書けるとのことなので現在の形のまま利用できればと思っていました。
vectorやstringといったものが使えるのであれば恐らく問題はないので設計しなおします。
572 名前：デフォルトの名無しさん mailto:sage [2010/09/17(金) 19:44:47 ]: 動的にmallocするようなやつは無理じゃない?
STLならthrust使えばある程度いけるけど
573 名前：デフォルトの名無しさん mailto:sage [2010/09/18(土) 17:06:00 ]: >>567
PGIのリファレンス読め
それでCUDA CとCUDA Fortranの違いはわかる
CUDA Fortranの教科書を探す・出るのを待つくらいならCUDA Cの教科書読め
（C言語がわかりませんなんて言うなよ？）
CUDA Fortranにほとんどそのまま適用できる
有料でCUDA Fortranの講習会やってくれるところも探せばあるよ
574 名前：sage [2010/09/21(火) 18:46:38 ]: CUDAを始めたばかりの初心者です．
現在既存のCのプログラムを並列化・速度を気にせずCUDAで
動作させようとしているのですが，
関数内の関数の変換がうまくいきません．
void A(～，～){B(～,～)}
void B(～，～){C(～，～)}
intC(～，～){}
・・・
main{A(～，～)}　　　　　　　　みたいなものです．

エラーでは
calling a host function from a __device__/__global__ function is only allowed in device emulation mode
と出ます．

このままプログラムを修正するか，一度Cのプログラムを変形するか
どちらが適切でしょうか？
575 名前：デフォルトの名無しさん mailto:sage [2010/09/21(火) 20:42:10 ]: エラーを読めば分かるだろ？
「デバイス(GPU)の関数からホスト(CPU)の関数の呼び出しはエミュレーション以外では許されません」

っと思ったがもしかして__global__の存在を知らないのか
576 名前：デフォルトの名無しさん mailto:sage [2010/09/21(火) 21:23:33 ]: とりあえず解説本買って嫁と言いたい
577 名前：デフォルトの名無しさん mailto:sage [2010/09/22(水) 08:46:15 ]: あたらすぃツールキットをインストールするまえに、古いツールキットはアンインストールするべき？
578 名前：デフォルトの名無しさん mailto:sage [2010/09/22(水) 09:39:00 ]: cuda使ったＧＵＩアプリ開発には
どのＧＵＩライブラリを使うのが一番生産性が高くなりますか？
579 名前：デフォルトの名無しさん mailto:sage [2010/09/22(水) 10:15:58 ]: CUDAと何の関係もないから好きなの使え
もしくはGUI系のスレ行って聞け
580 名前：デフォルトの名無しさん mailto:sage [2010/09/22(水) 10:22:05 ]: CUDAだとC系ですよね
javaなんてものだと、ダメですよね
だと必然的にC系のライブラリになって、そこからCUDA呼び出せないとだめですよね
そういう前提が縛られてるのに、CUDAとは何も関係がないってよく言えますね
バカですか？というかGUI作ったことないCUIプログラマなんですか（＾＾＾＾＾＾；；；；；；；
581 名前：デフォルトの名無しさん mailto:sage [2010/09/22(水) 10:42:04 ]: javaからCが呼び出せないとでも思ってんの？
582 名前：デフォルトの名無しさん mailto:sage [2010/09/22(水) 10:44:40 ]: 恥ずかしい>>580がいると聞いて
583 名前：デフォルトの名無しさん mailto:sage [2010/09/22(水) 11:06:59 ]: >>577
インストーラーがアンインストールしていいか聞いてきたと思う。
584 名前：デフォルトの名無しさん mailto:sage [2010/09/22(水) 16:47:33 ]: アンインストールするか聞かれなかったのでそのままインストールして、今二重にツールキットが入ってるんですが
v3.1は消すべき？
585 名前：デフォルトの名無しさん mailto:sage [2010/09/22(水) 17:02:28 ]: >>584
SDKじゃなくてツールキットが二重にはいってるの？
同じフォルダに？
586 名前：デフォルトの名無しさん mailto:sage [2010/09/22(水) 19:54:37 ]: C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA
にはv3.2しかないみたい。ならいいのかしら…。

C:\CUDA
に同様のファイル・フォルダが残ってたり、コントロールパネルには
NVIDIA CUDA Toolkit
NVIDIA CUDA Toolkit v3.2 (64 bit)
となってて

deviceQueryを実行すると
CUDA Driver Version: 3.20
CUDA Runtime Version: 3.10

と、ドライバーとランタイムのバージョンが異なるのは問題無？
587 名前：デフォルトの名無しさん mailto:sage [2010/09/22(水) 19:57:12 ]: いいわけないんじゃね？
588 名前：デフォルトの名無しさん mailto:sage [2010/09/22(水) 21:09:39 ]: 疑問形でいいわけ？
589 名前：デフォルトの名無しさん mailto:sage [2010/09/23(木) 02:19:24 ]: いいんじゃね?
590 名前：デフォルトの名無しさん mailto:sage [2010/09/23(木) 05:49:31 ]: あんれ。toolkit3.2 64bit　のダウンロードができないなう。みんなダウンロードできる？
591 名前：デフォルトの名無しさん mailto:sage [2010/09/23(木) 13:23:29 ]: linux以外not foundになるぽい
べつにいいんじゃね？
592 名前：デフォルトの名無しさん mailto:sage [2010/09/23(木) 16:41:48 ]: ダウンロードしてから、v3.1消せば良かった…。
593 名前：デフォルトの名無しさん mailto:sage [2010/09/23(木) 18:19:26 ]: 404自重
594 名前：デフォルトの名無しさん mailto:sage [2010/09/24(金) 01:51:15 ]: アドレス中のファイル名をcudatoolkit_3.2.7_win_64.msiにすればいけるよ
正式版を公開しようとして、途中でストップ掛かったんだろうか

他にもドライバーのバージョン表記がおかしかったり
260系ドライバーでOpenCLの動作がおかしくなってたりと、
GPGPU担当者がグロッキーになっているのではなかろうか
595 名前：デフォルトの名無しさん mailto:sage [2010/09/24(金) 02:24:10 ]: 正式版が出るまで待つのが吉
596 名前：デフォルトの名無しさん mailto:sage [2010/09/24(金) 02:48:22 ]: ああ、あとインストール先やフォルダー名やビルドルールファイル名が3.0から変更されていて、
今迄のプロジェクトは設定の変更が必要になるから注意ね
597 名前：デフォルトの名無しさん mailto:sage [2010/09/24(金) 08:46:54 ]: >>596
kwsk
598 名前：デフォルトの名無しさん mailto:sage [2010/09/24(金) 10:09:36 ]: で、3.2ってなんかメリットあるの？
599 名前：デフォルトの名無しさん mailto:sage [2010/09/24(金) 10:10:55 ]: 不具合修正
新機能

release notes読もうね
600 名前：デフォルトの名無しさん mailto:sage [2010/09/24(金) 10:34:00 ]: 新機能は特に魅了を感じなかったんだけど、全体的に速くなったりしたの？
601 名前：デフォルトの名無しさん mailto:sage [2010/09/24(金) 11:10:37 ]: お。アドレス修正されたようだ
602 名前：デフォルトの名無しさん mailto:sage [2010/09/24(金) 12:31:43 ]: ひょっとして新品PCに新しくVisual Studio 2008をいれて、Developer Drivers　260.21、toolkit 3.2、SDK3.2いれるだけで、サンプルプログラムをビルド出来るようになった？
603 名前：デフォルトの名無しさん mailto:sage [2010/09/24(金) 23:31:12 ]: GPUなしでは動かない
604 名前：デフォルトの名無しさん mailto:sage [2010/09/25(土) 20:55:21 ]: ptxas infoの見方がよくわかりません。

ptxas info : Used 14 registers, 48 bytes cmem[0], 368 bytes cmem[2], 4 bytes cmem[16]
ptxas info : Used 19 registers, 3072+0 bytes smem, 40 bytes cmem[0], 368 bytes cmem[2]
ptxas info : Used 27 registers, 18432+0 bytes smem, 120 bytes cmem[0], 368 bytes cmem[2]
ptxas info : Used 40 registers, 18432+0 bytes smem, 128 bytes cmem[0], 368 bytes cmem[2], 12 bytes cmem[16]
ptxas info : Used 23 registers, 1024+0 bytes smem, 104 bytes cmem[0], 368 bytes cmem[2]
ptxas info : Used 55 registers, 8+0 bytes lmem, 4096+0 bytes smem, 128 bytes cmem[0], 368 bytes cmem[2], 28 bytes cmem[16]
ptxas info : Used 54 registers, 8+0 bytes lmem, 88 bytes cmem[0], 368 bytes cmem[2], 12 bytes cmem[16]
ptxas info : Used 12 registers, 56 bytes cmem[0], 368 bytes cmem[2]

1024+0 bytes smemのような+はなにを表しているのでしょうか
368 bytes cmem[2]のようなcmem[]の中身は何を表しているのでしょうか
ローカルメモリ使うほどレジスタ使ったつもりはないのですがlmemはなぜ使われているのでしょうか
605 名前：デフォルトの名無しさん mailto:sage [2010/09/25(土) 22:05:04 ]: 三角関数つかってないか？＞lmem
606 名前：デフォルトの名無しさん mailto:sage [2010/09/26(日) 00:06:40 ]: CUDAで三角関数使うならSFU使わないともったいない。
607 名前：デフォルトの名無しさん mailto:sage [2010/09/27(月) 10:54:22 ]: >>605-606
使ってます。
ローカルメモリが使われるんですね。
use_fast_mathって精度が少し問題なんでしたっけ？
608 名前：デフォルトの名無しさん [2010/09/27(月) 15:43:35 ]: すいません。だれか教えてくれませんか。
cudaのサンプル実行したら以下のメッセージがでます。
cudaGetDeviceCount FAILED CUDA and Runtime version may be msimatched.
サンプルプログラムはDeviceQuery.cです。
コンパイルは通ります。
gccのバージョンでしょうか？
環境はmacBook cudadriber 3.2.7 toolkit3.2.8
をインストールしました。同様な現象のかたいますか？

Device
609 名前：デフォルトの名無しさん [2010/09/27(月) 17:21:34 ]: 暗号化のメイン関数の処理にファイルIOのスピードがついてこれない。
fread３００Mで10秒以上かかるのは何とかできませんか？
610 名前：デフォルトの名無しさん mailto:sage [2010/09/27(月) 17:23:04 ]: 別スレッドで先読みさせといたら？
611 名前：デフォルトの名無しさん [2010/09/27(月) 17:36:05 ]: マルチコアで解決できるかと・・・
612 名前：デフォルトの名無しさん mailto:sage [2010/09/27(月) 19:28:24 ]: >>608
とりあえずドライバのバージョンが3.2.8でない理由を聞いてもよいでしょうか？
613 名前：デフォルトの名無しさん mailto:sage [2010/09/27(月) 20:34:04 ]: >>608
gccのバージョンと思うなら、バージョンくらい書きなよ。
ここはエスパースレじゃねえぞ。
614 名前：デフォルトの名無しさん [2010/09/28(火) 00:15:33 ]: mac初めてさわる初心者なんだよ。
gcc4.2 toolkit何度もバージョン変えてもエラーがでる。
わからん
615 名前：デフォルトの名無しさん mailto:sage [2010/09/28(火) 00:21:22 ]: わからんじゃねぇ、ソース見れば原因一発で分かるだろ
616 名前：デフォルトの名無しさん mailto:sage [2010/09/28(火) 00:31:17 ]: >608はエラーメッセージが読めないのだろうか。
617 名前：デフォルトの名無しさん mailto:sage [2010/09/28(火) 00:41:02 ]: ドライバとランタイムのバージョンが合ってねえズラ。
618 名前：デフォルトの名無しさん mailto:sage [2010/09/28(火) 01:09:27 ]: 質問クオリティの低下が止まらないな
619 名前：デフォルトの名無しさん mailto:sage [2010/09/28(火) 03:25:56 ]: 英語分かんなくても、とりあえず翻訳サイトにその文章持ってく位の事はしろよ
つうか普通ミスマッチ位分かるだろ
620 名前：デフォルトの名無しさん mailto:sage [2010/09/28(火) 13:16:06 ]: ドライバも最新にすれば動くとでも？
621 名前：デフォルトの名無しさん mailto:sage [2010/09/28(火) 14:47:07 ]: 分からんなら試せよ怠け者
黙ってcuda3.2rcのページに行ってドライバ落として突っ込んで来い
622 名前：デフォルトの名無しさん mailto:sage [2010/09/29(水) 03:26:30 ]: まぁ、開発者でないならば今回の3.2は落ち着くまで待った方が良いと思うが。
623 名前：デフォルトの名無しさん mailto:sage [2010/09/29(水) 04:04:14 ]: 開発者でも普通はRCに飛びつかない。
624 名前：デフォルトの名無しさん mailto:sage [2010/09/29(水) 09:14:17 ]: プログラム初心者ですが飛びついてしまいました
625 名前：デフォルトの名無しさん mailto:sage [2010/09/29(水) 19:49:11 ]: CUDAはプログラム初心者向きとは間違っても言えない．
素直にC，C++，コンパイル，リンクの仕組みのあたりの基礎を
しっかり身に着けてからCUDAに挑戦してほしい．
626 名前：デフォルトの名無しさん mailto:sage [2010/09/29(水) 23:11:44 ]: 資料が充実している分、ATI Streamに手を出すよりはマシ、かな？
627 名前：デフォルトの名無しさん mailto:sage [2010/09/30(木) 00:18:55 ]: CAL+ILならそうだが、OpenCLならそうでもないよ。
628 名前：デフォルトの名無しさん mailto:sage [2010/10/01(金) 23:56:14 ]: nvidiaと心中する勢いでCUDAを
629 名前：デフォルトの名無しさん mailto:sage [2010/10/02(土) 12:00:43 ]: CUDAなんてもう3年ぐらい放っておいて一般的になってから手を出したほうがいいと思う。
630 名前：デフォルトの名無しさん mailto:sage [2010/10/02(土) 12:11:19 ]: 一般的になったら先行者利益がないだろ
631 名前：デフォルトの名無しさん mailto:sage [2010/10/02(土) 12:44:00 ]: CUDAでリアルタイムエンコーダな動画編集ソフトを作って
１５００円くらいで売って、１００００人に売れば１５００万円儲けられるな

よし、がんばるぞ☆
632 名前：デフォルトの名無しさん mailto:sage [2010/10/02(土) 12:54:51 ]: >>631
ネタとかツマランからポケモン板池
633 名前：デフォルトの名無しさん mailto:sage [2010/10/02(土) 16:59:01 ]: 俺はnvidiaを信じているよ
634 名前：デフォルトの名無しさん mailto:sage [2010/10/02(土) 20:10:50 ]: AVCエンコーダはCUDAについてるしな。
635 名前：デフォルトの名無しさん mailto:sage [2010/10/03(日) 15:30:40 ]: CPU AMD Phenom II X4 955
GPU GTX260
OS windows7 32bitにおいて
char型のsharedメモリを用意し、連続アクセスするとバンクコンフリクトが起こっているようなのですが
これは1バンクが32bitだから8bitのchar型だと4-wayバンクコンフリクトが起こっているという認識で良いでしょうか。
バンクコンフリクトが起こっているかどうかはint型のsharedメモリを用意し速度検証し、int型の方が早いことを確認しました。
またvisual profiler でwarp sirialize項目を確認したところ、char型のみバンクコンフリクトが確認されましたので4-wayかなと思っています。

また
CPU Intel Corei 7 930
GPU GTX480
OS windows7 64bit
において同じことをGTX480でためすとどうもバンクコンフリクトが起こっているように思えません。
調べたところint型の場合と比べてchar型の方が実行時間が短いからそう思いました。
ちなみにこちらもvisual profilerで確認したかったのですが、warp sirialize項目がなぜかないため断念しました。

質問としては、
GTX480ではGTX260とバンクコンフリクトの条件が違うのでしょうか。
またGTX480でのvisual profilerにてwarp sirializeの項目がないのでどうにかして測定できないでしょうか。
あまりcudaやvisual profilerを分かっていませんが、よろしくお願いします。
636 名前：デフォルトの名無しさん mailto:sage [2010/10/03(日) 16:10:39 ]: プログラミングガイドの最新版を見たけど、G.3.3.2.とG4.3.にsharedメモリのbank conflictのことが書いてある。
結論から言えば、char型のアクセスは、G92アーキテクチャではconflictが起きて、G100アーキテクチャでは起きないようだ。
他にもバンク数が16から32になったり、細かい部分でかなり変更されている様子。
試していないがGTX480ではl1_shared_bank_conflictの値が取れるかもしれない。
637 名前：デフォルトの名無しさん mailto:sage [2010/10/03(日) 19:47:32 ]: >>636
なるほど、ありがとうございます。確かにそうかいてありますね。。。

とても参考になりました。
l1_shared_bank_conflictについては今度試してみたいと思います。
638 名前：デフォルトの名無しさん mailto:sage [2010/10/06(水) 17:44:22 ]: カーネル関数呼び出しの第3引数にextern __shared__のシェアードメモリの
サイズ指定がありますが、
float s[256]
だったら4*256=1024byteですよね。
float s1[256]
float s2[256]
と2種類欲しければ2048といれればいのでしょうか？
その場合
float s1[128]
float s2[256]
と欲しければ4*128+4*256=1536を引き数にした時に
float s1[256]
float s2[128]
と解釈してしまう可能性もあるのでしょうか？(extern __shared__ s1[],s2[]としか書かないため)
それともshared memoryの配列の要素数はブロック内スレッド数と決まっているのでしょうか
639 名前：デフォルトの名無しさん [2010/10/06(水) 18:17:15 ]: >>638
アドレスはカーネルの方で陽に指定する
１次元アレーだから
マニュアルにサンプルがあるはず
640 名前：デフォルトの名無しさん mailto:sage [2010/10/07(木) 05:59:58 ]: CUDA by Example: An Introduction to General-Purpose GPU Programming
のPDF版のダウンロード購入ってどこでできるの？
641 名前：638 mailto:sage [2010/10/07(木) 10:11:52 ]: >>639
長さ分のshared memoryを取ってカーネル内で分り当てするんですね。ありがとうございます
642 名前：デフォルトの名無しさん mailto:sage [2010/10/08(金) 00:03:40 ]: pc.watch.impress.co.jp/docs/column/ubiq/20101008_398601.html
643 名前：デフォルトの名無しさん mailto:sage [2010/10/08(金) 15:07:14 ]: カーネルの引数に配列をいくつか渡したいのですが、数が多いのでそれらの配列のポインタを配列に入れて配列を渡すことにしました。
float *tmp[3],*tmp0,*tmp1,*tmp2;
tmp0,1,2をcudamalloc
tmp[0]=tmp0;tmp[1]=tmp1;tmp[2]=tmp2;
kernel<<<>>>(tmp)
kernel内でtmp[0][0]=...
のように使ったのですが、unspecified launch failureとなってしまいます。
どこがいけないのでしょうか
644 名前：デフォルトの名無しさん mailto:sage [2010/10/08(金) 16:46:17 ]: >>643
tmp自体もdeviceに転送しないとダメ。
それがイヤなら構造体でもOK。
645 名前：デフォルトの名無しさん mailto:sage [2010/10/08(金) 18:09:39 ]: >>644
転送ってことはグローバルメモリに入るんですよね
構造体でいけました！
ありがとうございます！！！
646 名前：デフォルトの名無しさん mailto:sage [2010/10/08(金) 19:51:11 ]: >>645
どうでもいいけど、理由は把握できた? 把握できなきゃまたはまるよ。
# めんどくさいからレス無用
647 名前：デフォルトの名無しさん [2010/10/10(日) 02:07:27 ]: はじめてのCUDAプログラミング本のP.136にshared memory の総和計算での
ウォープ・ダイバージェントを少なくする方法が載っています。
分からないのですが、そもそも else 節がない if 文でウォープ・ダイバージェントは発生するのでしょうか？
自分の理解だと、32人33脚のパン喰い競争でアンパンをスルーするスレッドがいるだけで、次のドーナツ(else節)は
ないので、ウォープ・ダイバージェントは起きないと思うのです。
648 名前：デフォルトの名無しさん mailto:sage [2010/10/10(日) 08:15:18 ]: >>647
32人33脚の真ん中の奴が
「俺アンパン嫌いだからちょっとイギリスまで行ってヨークシャープディング食ってくる。俺が帰ってくるまで待ってろよ」
とか言い出したら嫌だろ
649 名前：デフォルトの名無しさん mailto:sage [2010/10/10(日) 11:49:00 ]: Warp divergentは分岐命令を異なる方向に分岐することをいうのでelse節があるかどうかは関係ありません。
650 名前：デフォルトの名無しさん mailto:sage [2010/10/10(日) 12:22:05 ]: >>648
そのWarp全員の足並みが揃っていても、
結局、他のWarpではイギリスに行く奴(if節を実行するスレッド)がいるから、
それを待たないといけないんじゃないですか。
>>649
なるほど。
if 文が終わったら合流しないんですか。
異なる方向に分岐と言っても、この場合、片方は何もしないので
合流するのであれば Warp divergentのコストは何もないと思うんです。
651 名前：デフォルトの名無しさん mailto:sage [2010/10/10(日) 12:42:42 ]: 一人がイギリスに行ったら全員で一緒にイギリスに行くけど、付いていったやつはなにもしない。（ベクトルマスク）
誰も行かなければイギリス往復分の時間が浮くからそれがコストになる。
652 名前：デフォルトの名無しさん mailto:sage [2010/10/10(日) 12:59:44 ]: >>651
あ～この場合、複数のWarpは同一のSMで切り替わりながら実行されるから
Warp内では足波を揃えていた方がいいんですね。
全部のWarpが並列的に走ってるのかと思ってました。
653 名前：デフォルトの名無しさん [2010/10/12(火) 03:17:49 ]: やっぱ　わからん
cudatoolkit 3.2.2.8
devdriver 3.2.8
gpucomputingsdk 3.2
mac os snow leopn
相変わらず　runtime version may be mismatchがでる
バージョン合ってるじゃん。
だれかかっこいい人教えて
654 名前：デフォルトの名無しさん mailto:sage [2010/10/12(火) 03:19:32 ]: だからソース見ろって言っただろカス
655 名前：デフォルトの名無しさん mailto:sage [2010/10/12(火) 06:52:14 ]: LD_LIBRARY_PATH?
656 名前：デフォルトの名無しさん [2010/10/12(火) 23:48:44 ]: そんなこといわないでくれよ
アニキ。
mac初心者なんだよ。
また朝鮮してみるよ
657 名前：デフォルトの名無しさん [2010/10/13(水) 15:26:01 ]: 開発はGT220を使いWindows7で行っています。

現在構造体とthrustのテストを行っているのですがうまくいきません。
構造体はStudentという名前で学籍番号,クラス,点数を持っています。学生は全部で100人いるため以下のように宣言をしました。

thrust::host_vector<Student> host_student(100);
thrust::device_vector<Student> device_student(100);

この部分とvectorにデータを入れる部分はデータの出力を見る限り問題はなかったようです。
問題は点数が60点以上と未満に分けるという処理を行おうとした際にEmuDebugでは正しい結果(80:20)が出たのですが普通に実行した場合不正な値(2:1)が出ます。
結果は以下に宣言したends[0]に60点以上の人数、ends[1]に60点未満の人数を入れるつもりです。

thrust::host_vector<int> host_ends(2);
thrust::device_vector<int> ends(2);

以下のようにカーネルを実行し

cudaKernel<<< 1, 100 >>>(thrust::raw_pointer_cast(&*device_student.begin()) ,thrust::raw_pointer_cast(&*ends.begin()));

カーネルでは以下のように記述しています。
__global__ void checkS(Student *stu, int *ends){
// スレッドID
int id = threadIdx.x;
if(stu[id].point > 59){//60以上ならば
ends[0] += 1;
}else ends[1] += 1;
}
Windowsなのでデバッグが正しく出来ていないため挙動が不明です。
EmuDebugでは正しく動くのに並列計算だとうまく動かないためカーネルの部分が原因とは思いますがまだ勉強中のため解明できません。
原因として考えられること、解決案を教えてください。
658 名前：デフォルトの名無しさん mailto:sage [2010/10/13(水) 15:43:50 ]: ・GPUでやるほどの処理なのか疑問。
→実験目的なら納得。

・可読性の悪さをコメントで補うという厄介なコードになっている。読み易いように以下のように書けば充分。
・同様に、elseの前後で複文と単文を混ぜるのは宜しくない。後でコードを追加する場合の備えも兼ねて複文の方が宜しかろう。
・ついでに、c/c++では「一つ増やす」ことを明示的に書けるから積極的に使った方が読み易い。
# この場合、「数え上げ」のために1増やすのであって、1という数値自体に意味があるわけではない。
--
if (stu[id].point >= 60) {
++ends[0];
} else {
++ends[1];
}
--

・周辺コードがないからなんとも言えないが、ちゃんと転送されているのだろうか。
# thrust::copy()かな。
・同様に、endsは初期化されているのだろうか。
659 名前：デフォルトの名無しさん mailto:sage [2010/10/13(水) 15:57:16 ]: CUDAは並列動作時に同じアドレスへの書き込みを保証できないから
atomic命令使ってカウントするか、同数のtrue/falseフラグ配列を用意して
カーネル終了後に別途足し併せないと無理
660 名前：デフォルトの名無しさん mailto:sage [2010/10/13(水) 16:06:49 ]: そうだ、肝腎なことを忘れていた。
# 馬鹿でぇ＞自分
## つーか、流石に自分自身ではこんなミスはしないのに……

そもそも、複数のスレッドが同じメモリに書き出すのは無理。
ほぼ同時に走るから、1や2になるのも道理。
逆に言えば、巧く並列に実行しているからこそちゃんとカウントできない。
661 名前：デフォルトの名無しさん mailto:sage [2010/10/13(水) 16:47:31 ]: >>658
>>・可読性の悪さを～
申し訳ないです。実際endsの初期化とかstudentへの値の代入などその他の処理が結構あったため切り貼りしたのが原因です。
>>・周辺コードが～，同様に～
この部分の動作は一応確認できていたため省いてしまいました。すいません。

>>660
自分も並列につまり同時に計算しているのに同じ場所に加算を行えるのか疑問でしたがやはりムリだったんですね。
確かに逆に考えれば動いているともいえそうです。

実は今は多少違う方法で解決したため問題はなさそうですがEmuで成功した理由がわからなかった+後学のために質問しました。

あともう1つ質問なのですが
thrust::device_vector<hoge> hoge(N);
という宣言を行った場合ホスト側にメモリ領域を取らずデバイス側にメモリ領域を確保できるのでしょうか?
662 名前：デフォルトの名無しさん mailto:sage [2010/10/13(水) 16:52:00 ]: すいません肝心のお礼を書いていませんでした。
即レスありがとうございます。参考になりました。
663 名前：デフォルトの名無しさん mailto:sage [2010/10/13(水) 17:00:45 ]: >659が真の功労者。お礼はそっちに。

emuで成功するのは、マルチスレッドでエミュレーションしないでシングルスレッドでエミュレーションしているからだと思われ。
マルチスレッドで並列化しても同じように起きる典型例だからね。

thrustは詳しくないけど、ちょっと見てきたところで判断する限り、ホスト側とデバイス側は別々にメモリを確保する必要がある。
従って、device_vectorはデバイス側だけメモリを確保すると思われ。
664 名前：デフォルトの名無しさん mailto:sage [2010/10/13(水) 17:09:35 ]: 重ね重ねありがとうございます。この板IDが出ないので誰だかわからないですね。
実際に取り扱うデータは10万件くらいのデータ(蓄積しているのでもっと多い)の分析で
分析をC++で書いていたのと並列化に向いたアルゴリズムも存在していたため
CUDAなら乗せれると思って勉強しはじめました。基本はC,C++ライクなのでがんばってみます。
665 名前：デフォルトの名無しさん [2010/10/14(木) 00:37:01 ]: runtime versionてどう調べるの？
cudaの
666 名前：デフォルトの名無しさん mailto:sage [2010/10/14(木) 00:42:13 ]: devicequery
実行できないなら自分が入れたもんくらい覚えとけよ
667 名前：デフォルトの名無しさん mailto:sage [2010/10/14(木) 18:01:10 ]: CUDAでカーネルを10回実行しようと思った場合
以下のような書き方でいいのでしょうか?
for(int i=0; i<10; i++){
kernel<<<n,m>>>(,,,);
}
668 名前：デフォルトの名無しさん mailto:sage [2010/10/14(木) 18:22:25 ]: 一応。
タイマーとか挟むつもりなら注意が必要だけど
669 名前：デフォルトの名無しさん mailto:sage [2010/10/14(木) 21:32:45 ]: >>667
カーネル呼び出しはコストが掛かるから、できれば一発呼びで済ませた方が。
その例なら、kernel<<<dim3(n, 10, 1), m>>>(...)できないだろうか。
670 名前：デフォルトの名無しさん mailto:sage [2010/10/15(金) 02:11:53 ]: 667ではない，プログラミング初心者ですがコストっていうのはどういう意味ですか．
671 名前：2の方だね mailto:sage [2010/10/15(金) 11:36:33 ]: >>670
cost
[名][U][C]
1 代価；原価；支出, 費用, 経費, コスト(running costs)；((?s))（家庭・企業などの）諸経費；((?s))《法律》訴訟費用
the high ［the low］ cost of electricity
高い［低い］電力コスト
the total cost of the scheme
計画の総経費
the cost of production ［＝production costs］
生産費
sell at ［below］ cost
原価［原価以下］で売る
without cost
無料で
cover the cost
元をとる
cut ［reduce］ costs
経費を切り詰める
My car was repaired at a cost of ￥50,000.
車の修理代は5万円だった. ⇒PRICE[類語]
2 ((しばしば the ?))（時間・労働などの）犠牲, 損失
at a person's cost
人の犠牲において, 人に損害をかけて.
672 名前：デフォルトの名無しさん mailto:sage [2010/10/15(金) 13:12:33 ]: 配列aには100以下の数値がランダムで100個格納されています。
配列bの中身をaと同じにしたい場合以下のように書けばできました。結果も正しく出力できました。
__global__ void changeA(int *a, int *b){
int id = threadIdx.x;
a[id]=b[id]
}
上記は100スレッドで並列計算を行いましたがスレッド数を10個(10,1,1)にした場合正しく計算されませんでした。
ブロック数を10個(10,1,1)スレッド数を10個(10,1,1)にしてidを
int id = blockIdx.x * blockDim.x + threadIdx.x;
とした場合正しく計算できました。こうなる原因は大体わかっているのですが
スレッド数が少ない場合、またはスレッド数が膨大な場合たとえば(1000,1,1)のようにした場合で正しい結果を得ることができるのでしょうか?
スレッドが膨大な場合size(この場合100)を渡しidがsizeを超えれば処理をしないというものです。
この場合無駄ができそうなのでよい方法はないでしょうか?

__global__ void changeA(int *a, int *b, int size){
int id = blockIdx.x * blockDim.x + threadIdx.x;
if(id =< size){
a[id]=b[id]
}
}
673 名前：デフォルトの名無しさん mailto:sage [2010/10/15(金) 13:15:30 ]: すいませんスレッドの場合、総計512が最大でした。ブロック数が(1000,1,1)です。
674 名前：デフォルトの名無しさん mailto:sage [2010/10/15(金) 14:23:27 ]: >スレッド数を10個(10,1,1)にした場合正しく計算されませんでした。
意味が分からん。んなわけないだろ

分岐はウォープ内で同じ方向だったら対して無駄にならないんじゃない？
675 名前：デフォルトの名無しさん mailto:sage [2010/10/15(金) 14:23:37 ]: a[id] = b[id]じゃ説明が逆だろうけどそれはそれとして。
単にコピーしているだけとして、これでいいと思う。
--
__global__ void changeA(int * a, int * b, int size)
{
for (int idx = blockIdx.x * blockDim.x + threadIdx.x; idx < size; idx += gridDim.x * blockDim.x) {
a[id] = b[id];
}
}
--
勿論gridDim.x * blockDim.xがsizeよりも大きければ>672と同じになるし、sizeが大きくても破綻しない。
ブロック数が過剰な場合に無駄は出ないかという点だけど、それについてはホスト側で判定すればいいだけじゃないかと。
--
const int NumThreads = 512; // 実際には192以上の32の倍数で調整すべきかも
const int NumBlocks = std::min((size + NumThreads - 1) / NumThreads * NumThreads, 1000); // block数の最大を1000としてみた
changeA<<<dim3(NumBlocks, 1, 1), dim3(NumThreads, 1, 1)>>>(a, b, size)
676 名前：デフォルトの名無しさん mailto:sage [2010/10/15(金) 14:25:43 ]: >>674
スレッド数が10だとsizeが100なのに10件しか処理しないという自明なことを書いているだけかと思った。
677 名前：デフォルトの名無しさん mailto:sage [2010/10/15(金) 17:56:27 ]: >>676
あー。。。なるほどサンクス

そりゃそーだろう...
678 名前：デフォルトの名無しさん mailto:sage [2010/10/15(金) 20:22:23 ]: >>670
この文脈だと、カーネルをコールする時のオーバーヘッドの事

[ 新着レスの取得/表示 (agate) ] / [ 携帯版 ]

read.cgi ver5.27 [feat.BBS2 +1.6] / e.0.2 (02/09/03) / eucaly.net products.
担当:undef