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

このスレッドは、他のスレッドでは書き込めない超低レベル、
もしくは質問者自身何が何だが分からない質問を勇気を持って書き込むスレッドです。
CUDA使いが優しくコメントを返しますが、
お礼はCUDAの布教と初心者の救済をお願いします。

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

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

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

関連サイト
CUDA Zone
http://www.nvidia.co.jp/object/cuda_home_jp.html
フォーラム
http://forum.nvidia.co.jp/EokpControl?&event=TE0001

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

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

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

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

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

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

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

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

いちおつ

一乙、前スレの、魔二来vsTeslaのコストパフォーマンス比較面白いな
もっとやってくれ

>>4
もう結論でたよ

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

すげーな
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

>>7
これって、いまいち理解出来ないのだが、誰か説明して。

>>7
将来が危ぶまれるNVIDIA専用のCUDAと心中するつもりなら
Teslaでもいいかもね

>>8
たぶん

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

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

CPUは
（総費用）＝（価格）X（コア数）
（GFLOPS）＝（クロック数）X（コア数）X（CPU数）X（SIMD演算器数）X（１SIMD演算器当たりの倍精度演算数）

GPUは
（総費用）＝（価格）
（GFLOPS) ＝（クロック数）X（コア数）

>>10
間違えた

×（総費用）＝（価格）X（コア数）
○（総費用）＝（価格）X（CPU数）

OpenMPとかえらい簡単に並列化できるしな。
あと何覚えればいいですか？

普通のスレッド
MPI

>>10
OpteronってSSEとかが1コアあたり2個ある訳？

OpenMPもCUDA並に最適化が面倒。
その代わりMPIの知識があれば以外と簡単。

で、比較したのまだぁ？

そもそもC2050が出てないのに
比較なんてできるかよ

そもそもスレチ

秋葉にオープンしたfaithにCUDAのブースができたね
実体はunitcomの宣伝かな

GTX480でデモしてくれるならすぐにでも行くが

長崎大の部品を入れたとか書いてたね。CUDA話のできる店員さんいるのかしら。

チップセット・・・完全撤退
GPU・・・AMDの大攻勢の前に大苦戦
Tesla・・・CPUの高性能化・マルチコア化の前で微妙に

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

自作板に帰れ

SSEとかAVXとか覚えるよりも、
CUDAのお作法覚えるほうが簡単な気がしない？おれだけ？

>>24
でも将来性が最悪だね
遅かれ早かれOenpCLに移行するだろうし

CUDAとSIMDでは考え方が違うような気が。
個人的にはCUDAは割と大ざっぱに並べる。
SSEとかは細かなパズル。
自分でも何を言っているかわからんが、そういう思いでいる。

前スレで480特攻すると書いた者だが
倍精度1/4と聞いてキャンセルしてしまった…
ゲーマーだったら失敗覚悟で注文したんだけど

ZOTACまだリリース前になってますね。キャンセルできてよかったですね。
ケース待ちだった方は購入できたのでしょうか。

マザーボード破損の書き込みもあるし手を出さなくて正解だろうね

カーネル中でcoutとかは使えないの？
どこが悪いのか調べるときはたいていcoutで変数を書き出していたのだけど、
CUDAの場合はどうすればいいの？

cudaMemcpy()でCPUに送る。

カーネルの途中で表示したいんだけど

どんな手法にしろCPU介さなかったら表示はできんよ

http://gpu.fixstars.com/index.php/Cuda-gdbを使う

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

結果とは別にデバイス側に確保しておいて、
あとで見るのはだめなの？

>>32
・deviceemuでなんとか頑張る
・可能な限り詳細な出力をメモリに吐き出しておいて、ホストに戻ってから(転送して)解析する

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

なぜ仕事前提？

>>30
とゆうか、そうゆう発想する事態C++もCUDAも全然理解できてないよねｗ

CPUで考えると、「レジスタの計算の途中をCoutで見ることは出来ないか？」
と聞いているようなものだと思うのだが。
それを自前でやるにはレジスタの内容を一度メモリにコピーしてcoutの標準出力で表示させることになる。
それと一緒だと思うよ。

device emuで解決ジャン

待てればね。

エミュレーションモードってなくなるんだろ

すでに実機があるからなあ。

ダンプ用の領域を用意しておいて実行中に書き込んで後で見る
しかないんじゃね？

ubuntu 10.04でcudaが使えるようになるのはいつごろですか

GPUを搭載していないノートPCでエミュレーションモードを使ってたのに
なくなるってマジ？

マジ

リモートデスクトップでCUDA使えればいいのになー。

リナカフェいったひといないすか

サトリナのカフェなんてあるの？

このスレもしぼんだなあ
GTX480が倍精度外れた（らしい）のが効いているんだろうな

倍精度演算ユニットは少なくなっているだけで無くなったわけではないのだが、
そもそも480が入手できないことの方が原因ではないかと。

結局口だけの貧乏人しか居なかったって事だなｗ

おまえもな

http://www.amazon.co.jp/dp/479802578X/
これIPAの岡ちゃんが書いたの？
いつの間にかネットマイスターってところに転勤したのかな

もう入門書もでてるから、新刊はいらないな。
需要はあるのだろうか･･･。

つーか、ネットマイスター代表だってよ。
同名をtwitterで検索したらttp://twitter.com/kenmoi0715 が引っ掛かったけど……

流石にロリコンばらされたら職場には居にくいか……
アプリも割り放題だったしな。
たくましく生きているようで何よりww

この本買ったけど、ちょっと入門過ぎる気がするなあ。
「高速」という言葉にてっきり「最適化」ということを連想してしまった。
環境設定とかで半分以上割いているよ。
まあ、CPUコードと両方書いてあって、Webからダウンロード出来るから、
最初にでた本よりかはいいかもしれない。

自分で買うほどじゃないかな
図書館に購入希望だしとこう

本にするほどの蓄積と技術の咀嚼があるとは思えないけどどうなんでしょ。
あるとすればまだ論文レベルなんじゃないか？

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

プログラムの入門書なんてこんなもんだろ

GF100が失敗、GF104はキャッシュ削るみたいでGPGPUは絶望的の現在、今後広まる可能性が残っているかどうか…

は？

GPGPUに残されているのは個人では動画エンコードくらいか。
それなりの規模になると、マルチプロセッサのシステムを組んで、
CPUのコードを走らせた方が開発工数から考えると速いからなあ。

で、おまえは何しにここに来てる訳？

ばかだな
ただのAMD信者様だろｗ

GPGPUに将来性はないだろうね。
牧野もそういっているし。

コスト面がちょっと微妙かなあ
C10シリーズのようにアカデミックで10万！とかなら個人で使ってみようという流れになるだろうに
GTX480は倍精度1/4にしました！C20シリーズを売るためです！と公言されちゃったし

>>69
コスト面つか買う気失せることすんなよなー　て感じっすよ

　　　　　　　､--‐冖'⌒￣￣`ｰ-､
　　　　 ／⌒`　　　　　　　　　三ミヽｰ-ﾍ,_
　　　__,{　;;,, 　　 　 　　　　　　 ミミ　　　ｉ　´Z,
　　 ゝ　　　''〃//,,, 　　　 　,,..｀ミﾐ､_ノリ}j;　f彡
　　_)　　　 　 　 〃///,　,;彡'rffｯ､ｨ彡'ﾉ从iﾉ彡
　　>';;,, 　 　　　 ノ丿川j !川｜;　 :.`7ﾗ公 '>了
　_く彡川fﾞノ'ノﾉ ノ_ノノﾉイシﾉ｜ }.:　'〈八ﾐ､､;.)
　 ヽ.:.:.:.:.:.;=､彡／‐-ニ''_ｰ＜、{_,ﾉ -一ヾ`~;.;.;）
　 く .:.:.:.:.:!ﾊ.Yイ　　ぇ'无ﾃ,`ヽ｝｝}ィt于 ｀|ィ"~
　　 ):.:.:.:.:|.Y }: :!　 　 ｀二´/'　;　|丶ﾆ　 ﾉﾉ
　 　 ） :.: ﾄ､リ: :!ヾ:、　　　丶　;　|　ﾞ　　ｲ:}　　　　逆に考えるんだ
　　 ｛ .:.: ｌ　｛: : }　 ｀　　　　,.__(__,}　　　/ﾉ
　　　 ヽ !　　`'ﾞ! 　 　　　 ,.,,.`三'ﾞ､,_ 　/´　　　「単精度なら４倍」と
　 　 ,/´{　　ﾐ　l　　　　/ﾞ,:-…-〜､ ） |
　　,r{ 　 ＼ ﾐ 　＼　　 `'　'≡≡'　" ﾉ　　　　　　　　考えるんだ
__ノ 　ヽ 　　＼　 ヽ＼　　　　彡 　,ｲ_
　　 　 　＼ 　　＼ ヽ 丶. 　　　　ノ!|ヽ｀ヽ、
　　 　 　 　 ＼ 　　＼ヽ ｀¨¨¨¨´/ |ｌ　ト､　｀'ｰ-､__
　 　 　 　 　 　 ＼　　｀'ｰ-､　　/／　/:.:.}　 　　 　 ｀'ｰ､_
　　　　　　　 　 ｀､＼　　　／⌒ヽ　 /!:.:.|
　　　　　　　　　　｀､ ＼ /ヽLｆ___ハ/　　{
　　　　　　　　　　　　　 ′ /　! ヽ

長崎大のセンセが8800GTで激安スパコン組んだときも、
そのうち２、３割がGPGPU用としてはぶっ壊れていたじゃん。

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

牧野もそういっているし。
牧野もそういっているし。
牧野もそういっているし。
牧野もそういっているし。
牧野もそういっているし。




牧野乙

そもそも、+α要素であるアクセラレータが
終わるとか終わらないとか

おばかさん

金にならんかったら終わるだろ。

GPGPUはマルチスレッドな浮動小数点演算に特化したものだからな、使えるシーンはかなり限られてくる

forループの一部だけを実行する方法を教えてください

>>55
やっぱりこういう問い合わせが多いんだろうな
著者紹介にIPAに勤務したことはないって書いてあったよ

>>59
俺も買った。
全く同じ感想。
青木先生の本が理解できる人は全く買う必要なし。
>>55
本の経歴の部分に、「IPAに勤務したことはない。」と不自然に書いてあったよw

なんだ同名で迷惑している人なのか

んなわけない。
IPAを除名されたか、退職するときにそれに準じる誓約書を書かされたのだろう。

証拠は？

>>79
今それを見てワロタ。
「IPAに勤務したことはない。」
で株式会社ネットマイスターをググってみたけど、
すぐには見つからなかった。
でもこういった本を書く人が流出なんてするのだろうか？

GF100の板に出ていてみたいだけど、
GTX480でCUDA動作に不具合があるとか？
ドライバ改善待ち、みたい。
GTX480でのベンチマーク結果（単精度、倍精度問わず）が出ない理由はそれだったりするのかも？

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

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

>>84
＞いつかは出るであろう、Quadroにかすかに期待するしかないのか…
お値段は同程度のGFの5〜10倍でも?w

5月中にはGTX480が手に入る筈なんで、そしたら各種ベンチを取る予定なんだけどなんだけど。

処で、暫く更新していない間にnvccが新しくなっていてptx出力の段階では64bit実数処理のコードが出力されてて笑った。
-ptxをつけないと64bitコードは処理できないみたいなメッセージが出るけれど、ptxを眺めて最適化しているから
今後は実数定数の扱いにも気を遣わないとならないようだ。

それで思ったんだけど、ptx出力からメモリアクセスの回数や割り算の回数を数えるツールというか、
ptx出力を解析するツールを作ったら需要あるかな。私自身は自分で眺めたりgrepで数えたりしているんだけど。
# 要は、iccのレポート出力みたいな機能がnvccにあればいいんだけどね。

>>85
もちろん重要ありますよ。

【Fermi】GeForceGTX4xx総合Part31【GF100】
http://pc11.2ch.net/test/read.cgi/jisaku/1272318075/
の548以降でGTX480を含むGPUのベンチマークが出ている
LU分解のようです
倍精度に制限かかっているのはどうやら本当みたいだ

具体的にどう制限かかっているんだろ？
それによってCUDAのコードの書き方が変わってくると思う

>87
いや倍精度どころか単精度でも全然性能でてないわけだが

なんかそのベンチ遅くね
285の単精度で16Gflopsとか書いてあるけど

GTX285と比べる限り、GTX480の倍精度に期待できそうな気配はないなあ
1/4は本当くさいな

cublasのgemmを使ったベンチをだれかが提供すべきだな
もちろん、俺はやる気はないが

gemmもやって欲しいけど、gemmはピーク性能しか出てこないから俺はLU分解やFFTの方が多少は信用している。

というか一番多用しそうなFFTでどれだけ実効性能出るのかが気になる

>>94
例えばどんなサイズが？

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

> これはマルチスレッドで実行する関数の内部だよね？
そのとおり。

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

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

>>97
ｔｈｘ

GCC 4.5.0で使えるようにする設定を教えてください

for (int i=1;i<n-1;i++)
{
    a[i]=b[i]+b[i+1]+b[i-1]
}
をGPUで並列実行するコードにしたいのですが、どのように書けばよいのでしょうか？

三つのループにばらしてBLASを３回使えばいいけど、
もっと速い方法があるかな。

>>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];
}

こんな感じかな。

続き・・・

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

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

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

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

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

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

結局nがいくつかによるだけじゃない。
たぶん100はCUDAのくの時もまだ理解していないと思う。

>>104
あの程度だとコンテキストスイッチやホスト・デバイス間の転送等の
呼び出しコストに見合わないね。

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

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

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

突っ込みどころが多すぎるんだが…
面倒なので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のサイズしだいでは問題にならない可能性がある。

>>107
厳しい突っ込みどうも。
CUDA Coreと呼ばれるようになったのはSMじゃなくてSPだったのですね。
図を見れば一目瞭然ですね・・・うろ覚えで書いてすみませんでした。

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

だから、16人17脚でパン喰い競争やるようなもんだよw

この表現はわかりやすいね。

一人重い荷物を持たされると、全体の速度が遅くなる、的なｗ

青木本で最も評価するのがあの手書きの絵だww

青木本にこのスレから引用したと見られる節が幾つかある件w

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

コアレスのコアってなんのコアのことなの?

coalescedなので、
合体したという意味。
合体してメモリにアクセスするって意味だ。

正直その並びなら、texture使えばキャッシュ使えるじゃん

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

↓新明解(旧版)で合体の項を引くなよ

メモリにアクセスするときに、各スレッドが足並みそろえて
＋１ずつのアドレスを見に行かないといかんのです
すると32ワードなりまとめてとってくる。　とってくる間は
別のスレッドが実行されてる。

>>114
私は「連動した」メモリアクセスと説明している。要は16人17脚なんだけどねw

おまえら頭いいんだろうな

どうせ底辺大学生が研究で勉強してるだけでしょ？

と学歴昆布が申しております

でもGT200以降だと16人17脚でパン喰い競争って微妙に語弊あるよね。
16人17脚でパン喰い競争じゃない使い方したことないけど。

>>102のコードは
あるスレッドの遅れが、他のスレッドの足を引っ張る結果になるの？

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

>>102
if使ってる時点でありえないかと。。。

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

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

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

>>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って理解できている?

>>107

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

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

だから、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出力を眺めていると判ってくるところだね。

揚げ足をとるようだけど、右足を出した隣の人は左足を最初に出してる。
__syncthreadsしなければ隣のひとたちを見ないでばりばり走るし、
過度な比喩は逆に混乱させるだけだと思う。

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

>>126
いやｗ
nの数の条件分岐とか必要ないでしょｗｗ
Host側がコール時に調整すればいいだけの話

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

>>130
nが巨大なら、ループする必要があるでしょ。ループ継続条件は通常、nとの比較で行なわれるわけで。

>>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にロードされるって制約があるなら別だが。

>>131
だったら数回カーネルをコールすればいいのでは？？
この場合カーネル内でループしちゃうと、キャッシュ外れる避けられない。キャッシュミスは大きいよ。

>>100　は一次元差分法からの連想だろうけど、
差分法なら青木本に２次元があるから参考にしたら。
自分は３次元の差分法を作ったけど信じられない速度が出た。
まあ、比較したCPU版が遅いんだろうけどｗ

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

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

>>135
「bへのアクセス3つは16個連続していること」って用件は満たしてる。
多分アライメントも16個境界になってないと駄目ということか。
確かにそうすると、 b[i+1]とb[i-1]はアライメントに合ってないから
都度ロードが発生する訳だな。
理解しました。どうもありがとう。

CUDA Programing Guide 2.2を見てコアレスメモリアクセスをある程度理解したつもりですが、
Compute Capability 1.0/1.1では開始アドレスが64の倍数で、half-warp内のメモリアクセスする全てのスレッドが
開始アドレス + (スレッドID x 32bit)のアドレスにアクセスする場合に限って
トランザクションが1回にまとめられるということでしょうか？

CC 1.2以降はこのあたりの制限がかなり緩くなっているようではありますけど。

>>125
warpサイズを考慮せずに分岐のあるコードを書くと
両方の処理を行うことになって極端に効率が落ちると思いますが
分岐先で何もせずにreturnする場合でもかなり落ちるのでしょうか？

>>127
CUDAというかアーキテクチャの制限ではないでしょうかね？
そのままグローバルメモリにアクセスするのではなく、シェアードメモリを利用するという思想で作られていると思っています。

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

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

>出力メモリ先がはみ出るのはまずいかと思いますが、
メモリを考慮してグリッドとブロックの数を決めるより
はじめにグリッドとブロックの数を決めてそれに合わせてメモリを確保したほうが良い。
ってことが言いたいんじゃないか。

>>137
分岐の種類によって違う結果になる。
・ワープ中幾つかのスレッドだけが条件に該当する場合
条件分岐のコスト(のみ)が余計に掛かる。
・ワープ中全てのスレッドが条件に該当する場合
ワープによっては条件分岐のコストが嵩むが
ワープによっては条件分岐以降の処理をパスするので早く開放される。

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

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

>>102
シェアードメモリは使ったほうがいいと思いますよ

うーん、CUDAの初歩も知らない奴に取り敢えず雛形を提示したら、
ちょっと齧った連中が寄って集ってけちつけるという構図になったなw
合間に慣れた人が適切な指摘をしてくれているからいいけどねぇ。

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

やだ・・・なにこの失敗したCell並に難解なプログラミング・・・

大丈夫、Cellのspuみたいにインストラクションを一個ずつあれこれしなければならないほど複雑じゃない。

GTX4*シリーズ買う？
それとも買った？

前スレから

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


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

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

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

C20シリーズは何とか予算用意して買うか、どこかの計算機センタに導入されるのを待つ形になる
それの準備的なものとして、GTX4シリーズを買うと思う

完全にHPCユーザーしかGPGPUに興味を持たなくなったな

画像処理とか（？）する人はどう思っているんでしょう
GPGPUについては

CG系研究室の俺が。

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

>>150
> CG系研究室の俺が。
CGの人ってコンピュータビジョンにもアンテナ伸ばしてるものなんですか？

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

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

ATI stream とCUDA とどっちがよく使われているんでしょうかね

「実装してみました」系はCUDAだろうね。
ピークパフォーマンスたたき出すならATIだろJS

>>151
>CGの人ってコンピュータビジョンにもアンテナ伸ばしてるものなんですか？
ちょっとだけね。 CVのいくつかのツールは、こっちでも役に立つから。

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

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

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

ATIを利用した論文を目にしたことがないな、確かに。

米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を搭載していることを明らかにした。

>>154
GPGPUとしてはCUDAが主流になると思う。
だけどこれまではほんとに酷かったなあ。
ちょっとでもコンピュータサイエンスの知見を持っていたら恥ずかしくて発表出来ないレベルが多かった。
今までは流行でみんながCUDA、CUDAといっていたけど、多くの人はあくまでも手段の一つであって、
目的ではない。目的にするならコンピュータサイエンスの分野でやればよいわけで。
それにCUDAを使う場合って大抵最適化しているから、CPUとの比較も余り意味がないし。
これからはそういうのはなくなっていくだろうね。

480と285の倍精度計算のベンチマークの比較はどこかにありませんか?

俺が勘違いしていないかの確認なんだが

a[tx] = b[tx] + b[tx+1] + b[tx-1];

はハーフワープ当たりの読み出しが3回でいいんだよな?
16*3回じゃないよな?

それより俺は音声処理とかのIIRっぽい場合で、bがaに依存するプログラムをどう書けばいいのか知りたい。
aを更新する処理で更新前のaと比較、分岐していたりすると俺の頭脳では依存関係が解消出来ない。
世の中の音声処理はSIMDとかSMPに対応しているんじゃないかと思うんだが、どうやっているんだろう。

>>157
どういうこと？

>>159
うまく行くケース(GPUのバージョン依存だっけ？？)なら３回でいけるだろうけど、
最悪の場合

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

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

>>161
全部のSPがb[tx]を読む→コアレス
全部のSPがb[tx+1]を読む→コアレス
全部のSPがb[tx-1]を読む→コアレス

とはならないの？

>>161
そういう場合、
スレッド　i-1　はtx-1をグローバルからSharedに読む
スレッド　i　　　はtxをグローバルからSharedに読む
スレッド　i+1　はtx+1をグローバルからSharedに読む
　：
でひと固まりもってきて、
あとはShared上でCoaleshed?　ｷﾆｼﾅｲ!　で動けばいいんじゃね

>>162
だってアラインがあわないじゃん。

>>163
shared memory使うなら、
b[tx]:コアレス
b[tx+1]:コアレスx1, シリアライズx1
b[tx-1]:コアレスx1, シリアライズx1
かなー

GPUバージョンだっけ？これが2.0とかなら、たしか上記全てコアレスになると思うけど

>>162-163
「t*x」　になっているんでこのままだとtステップになっちゃうわけだよ
つまりもともとの配列の並べ方から考え直さないといかんよな。
x-1, x, x+1, ,,, になったら163みたいに、
・スレッド協調してSharedに持ってくる
・Shared上では好き勝手に読み書きしてよろし
・syncthreadする
・スレッド協調してGlobalに書き戻す
というやりかたが基本かな

>>165

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

>>166
まちがえたｗ

>>165

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

これだと全部コアレスになる？

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
少し前にテクスチャ使ったコードで例を示したけど、誰も反応しなかったぞ
なんで？？

どのレス？

>>171
>>125

OpenCLでなくCUDAを学ぶ価値ってあるの？

OpenCLってなにに使われてるの

OpenCLは今のところRadeon使う人用かと。

DirectComputeはどんな感じ？

ローカルメモリに対するテクスチャメモリのメリットがいまいち分からない
そもそもローカルメモリを使いこなせないから、テクスチャメモリまで手が回らない

>>175
え？

>>177
補間できるとか

これってcudaMallocHostってsizeof(構造体)でぴったりとその構造体の容量分だけちゃんと確保できてるんですか？
なんかsizeof使ってぴったり構造体分確保したら条件によってセグメントエラーが出るんですが（printfを適当なとこに放りこんだらエラーが出ずに通ったりします）
2倍の容量を確保したら問題なく通ります

>>177
ローカルメモリがLocal Memoryのことをいってるなら、役割がまったく違うからメリットもなにもない
Global Memoryのことをいってるなら、キャッシュに加えて補間、正規化アドレス、変換とかだな
キャッシュが有効に機能しているかは、プログラムによっては疑わしいんだけどね

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

>>180
構造体に問題があると思われます。
各要素がアラインメントの条件を満たすようにダミーの要素を追加して調整することをお勧めします。

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

printf挟むと動くと言ってるからどっかでスタック壊してんじゃないかと予想

まぁ、晒せ、と。

ω

研究室にGTX 480届いた．
smokeParticlesで170fpsくらい．
倍精度FFTとかはこれからやる．

FFTやってみた．
1024×1024の二次元FFTの順変換&逆変換を1024回ループで回した．

・単精度
285 → 1.428[s]
480 → 0.736[s]

・倍精度
285 → 7.917[s]
480 → 3.829[s]

コアの数の比を考えると速度は変わらない気がする．

nbody --benchmarkよろ

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

これでいい？

>>187
多分CUFFT3.0を使った複素２次元FFTだと思うのだが、
こちらの計測データとかなり違うような・・・。
こちらでは480の単精度で1回(片方向)あたり0.44ミリ秒、2*1024回分で0.91秒くらい。

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

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);

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");

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;
}

nvcc cufft.cu -o cufft_double -arch=sm_20 -lcudart,cutil32,cufft -Xcompiler /O2

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

>>189
おお、ありがとうございます。
GTX260でも370GFOPS位出るのであんまり速くないですね・・・（´・ω・｀）

>>187

>CUT_SAFE_CALL(cutStopTimer(timer));

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

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

同期をちゃんととってみた．

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

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

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

質問させてください

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

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

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

SLIというかマルチGPUで使う場合はOpenMPかなんかで複数スレッドを作って
それぞれののGPUを使えばいい、というかサンプルがSDKにあったはず。
ゲームみたいにSLI組んでシングルGPUに見せるのはまだ出来なかったはず

>>199
親切にありがとうございます。
まさにゲームのようなSLIの考えでいましたが、そういった使い方はできないということですので
自宅用に使用するか等、使い方を考えてみます。本当にありがとうございました。

>>197
もちろんコードを自分で書く。

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()なんかは動くのですが。

どうすれば良いのか、ヒントでも良いので教えてください。

と聞いたばかりですみません。
解決しました。
ちゃんとgputoolsのinstall.txtを呼んでいなかったので
パラメータを間違えていて、--enable-emulationというのを
つけてしまっていました。
お騒がせしました。

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にはならないって

>>204
言われてる1/4ってのは当初期待されていた性能に対して1/4ってことじゃなかったの？
いくらなんでもTesla世代と比較してないでしょ

当初は短精度に対して理論値で1/2の性能を出せるって話だったのだけど
Geforceでは削られて短精度に対して1/4しかでないことになっている

1/8じゃなかったの？

皆GeForceでは1/8 (Fermi版Teslaシリーズのさらに1/4)という意味で話していたけど
一部の人は誤解していたみたい。

どういうことだってばよ

けっきょくどういうことなんだってばよ！

こまけえことはいいんだよ！

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

Teslaじゃコードを書き換える手間を考えたら、たいしてコストパフォーマンスは良くないな。

？

>>100の様なコードだったら手間はかからんよ。

TeslaがGT世代の事じゃなくてGPGPU専用カードの事を言ってるなら
Teslaの利点はバグのあるコードを走らせてもOS巻き込んで死んだりしない点が一番じゃね？
開発機はTeslaが楽だわ

ドライバ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

2次元配列を使うと、遅くなりますが、
2次元ブロックも同様の問題をかかえていますか?

なんのこっちゃ

2次元配列なんてCUDAにはありません。
意図的に1次元に並べてください。

GTX480でおすすめのボードってありますか？
もしくはダメなボードはありますか？

こちらも480入手。取り敢えず動くのは確認。細かいテストはこれから。

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

リファレンスだからドスパラのでいいんじゃないの
シール張るだけらしいから

>>221-222
ありがとうございます。

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

>>223
ドライバのリビジョンは今手元にないから判らないけど、CUDA2.3当時のドライバでCUDA2.3が動いている。
CUDA3に関しては未だテストしていない。

最近CUDAを勉強しはじめたんですが、質問させてください。

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

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

できます。
一回目のcudaMalloc()の前にcudaThreadSynchronize()でもしておけば、一回目のcudaMalloc()は早くなります。
尤も、cudaThreadSynchronize()で時間が(恐らく100ms程度)掛かりますが。

>226
ありがとうございます。
早速試したところ、そのとおりの結果になりました。。。
これは諦めるしかないんですかね。

なんかCUDAは最初に呼ぶときに時間がかかるとかリファレンスかなんかに書いてあったような・・・
斜め読みしたけどみつけられんかった

>228
そんな仕様だったんですね、知らなかったです。

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

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

>>229
一回だけだとそうかも知れないけど、バッチやリアルタイムで複数ファイルを処理するとかなら
効いてくると思います。

研究にCUDAを使ってみたいと思っている学生です。＞＜
プログラマブルシェーダで物理演算をしてシミュレーションをやっている
論文とかあったのですが
CUDAで計算するのとの違いはあるんですか？
あと
DirectXを使いながらCUDAを利用することは
DirectXで手一杯のことをやっていたらCPUで計算したほうが早かったりするんですか？

お願いします（。・・。）

>>232
それが俺の読んだのと同じならcgを使ってシェーダプログラムで演算をしてた
理由はCUDAがまだ無かったから
DirectXで3D描画しながらだとそりゃ性能は落ちるでしょ。CUDA専用にカードを追加するなら話は別だけど

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

>>234
つまりn00bだったってことかよ!

>>234
???
CGをやるなら別だが、今更Cgは無いだろう。
初心者ならなおさらCUDAだ。
特に物理演算をやるなら。

PGIコンパイラのCUDA化って正直どうよ。
人力でやるのに疲れてきた。
それなりの性能が出るなら買ってもいいと思ってるんだけど。

ディレクティブでは全然性能上がらなくて頭来たからCで書き直したｗ

東工大、Teslaを利用したスパコン「TSUBAME2.0」を開発
http://pc.watch.impress.co.jp/docs/news/20100617_374718.html

>>239
有用な結果を出さなかったら、次から仕分けな

国会議員を仕分けすりゃいいのに
1/3位にさｗ

あんたら、部外者か？

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

つまり漣舫は神戸のスパコンをつぶしたくで仕分けしたのですね。

1度でいいからこんな超計算機をいじり倒したいものだ

部外者だけどCSX600なんて誰も使ってないんじゃね？
導入初期に不良品大量に掴まされてたみたいだけどｗ

まあそうでもしないと進まないところもあるんだろうけど、こんな財政状況なんだからバッサリ逝ってよし

>>242
むしろ部外者しかいないだろ
お前はしらんが

議員を減らせば良い

政治ネタうぜぇ

選挙だもの

　　　　　　　民主党

れんほうのすぐ後ろに金田＠superπが座っていたから
京速機のヘッドである姫野＠ヘボベンチは何も言えんかったんだよ。

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

π（笑）

>>251
何で姫野ヘボベンチっていわれているの？
ソースを見たけどコード自体は学生が書くような内容だね。
だからなの？

メモリ帯域をひたすら無駄食いするプログラムだよね＞姫野ベンチ
ここで「なぜか遅いです」とか質問してくるレベル。

かといってgemmでピーク性能測られても役に立たない分野も多く。
直線に強い車かカーブに強い車かってのと似てる。

πって何を測ってるの

今ではすっかり使われないx87

>>235-236
今更シェーダをお勧めはしないけど、畳み込みを考えた時に
コアレスにするためのコードって余分だよね。バッドノウハウだよね。
ってことが言いたかった。

東工大のスパコン、世界２位なんだ。
そういうことか。

東工大は64位じゃね？
http://www.top500.org/list/2010/06/100

ヒント：TSUBAME2はまだ実在しない
っ ttp://www.gsic.titech.ac.jp/tsubame2

まだ計算どころか設置もしていないぞ

NECは京速機やめてこっちでやってんだな。

CUDAって未だにデバイスメモリをメインメモリにマップ出来ないの？
逆は出来るようだけど。

マップできなくてもいいけど、MPIでDMA使ってデバイスから通信チップに流しこんで欲しいわな

マップ出来ないといちいち転送指示出さなきゃいけないのが困る。
スクラッチで書くならいいけど、既存のを移植するとなると結構大変だから・・・。

テンプレートライブラリで、CUDAの転送まわりをほぼ完全に隠蔽してくれる奴あったよね。

thrust

使い方によるのかもしれないが、Mappedはかなり遅い気がする
Pinnedで転送したほうがまし

GTX480のベンチが日本語でも少しずつ出てきたね
倍精度と単精度両方載せてある

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

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

MATMULとかだと倍精度も速いみたいだけど、レジスタやシェアードメモリの消費も倍になるから
実コードだとOccupancyが下がって思ったような速度でないんだよなぁ・・・ ＞Fermi

倍精度の実効効率が、

GTX285　約90%
GTX480 約21%

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

ソース

>>273
>>272のは真上の祖父テックのベンチ結果から計算してのことだろ？
nbody.exeでも-fp64で性能出てなかったし

> ドライバのレベルで1/4に抑えられている（らしい）こと
ここのソース

googleを使えない奴にCUDA使えるとは思えないな

これと、中国のTUBASA？のLinpack結果を併せてだれか解説して

GTX480の「倍精度理論FP演算性能672 GFLOPS」という仕様そのものが間違っていて
実際は理論ピークが1/4の168。

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

>>278
↓この記事を信じてしまうと
http://journal.mycom.co.jp/special/2010/gpu/004.html
GTX480の倍精度演算の理論値は672GFLOPSになっちゃうけど、

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

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

ドライバなんだかそれ以外だかはわからないが、
GTX480の倍精度理論性能が168GFLOPSであることは信憑性高い、ということでおｋ？

ドライバで制限してるんじゃ無ければカタログスペックとの差異はどーすんの？
ttp://www.elsa-jp.co.jp/products/hpc/tesla_c2050/index.html

>>281
GTX480の話をしているときにC2050持ち出してどうするの？

想像力を働かすと、
C2050よりもコア数が多くて、同一アーキテクチャ仕様と公表されているGTX480が168GFLOPSしか出ないなんておかしいぞ
ということかな

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

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

メモリしか差のなかったC1060に対して、C2050はもっと明示的な計算用チップとしてのアドバンテージをつけたかった
nVidiaの意向が反映した結果がGTX480だと思うのだけど。

C2050はGTX480はCUDAコアが少なく、クロックも低いにもかかわらず、
倍精度理論性能は515GFlopsとなりC2050ほうが高い。
これは、GTX480はドライバで倍精度演算の実行にペナルティが加わるようになってるから。
それでも、GTX285と比べても倍精度演算が高速だからいいじゃんというのがNVIDIAの考え。

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

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

何が言いたいかというと、C2050と比べるとGTX480はそこそこお買い得だってことね。
ECC、メモリ容量、倍精度演算の潜在能力とメーカー出荷時のチェック品質が値段の違いなのかな。

>>283
>>281だけどそう言いたかった。OpenGLが速いQuadroみたいな感じなのかね？
Telsa C20xxのアホみたいな値段見てると・・・

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

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

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

>>288

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

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

>>285

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

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

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

＞C2050とGTX 480は同じFermiアーキテクチャではあっても違う構造になっている。

釣れますか？

>>289
たしかにその通りなんだけど、ボードの不良と結論付けるまで、
コードと何時間も格闘したり、ホストのメモリや電源など検証するの大変だよ。

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

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

>>290
>C2050とGTX 480は同じFermiアーキテクチャではあっても違う構造になっている。
同じチップなんだから、当然同じ構造でしょ。ボードデザインは違うけど。

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

適当言い過ぎ。

3.1でたね。 ・Support for printf() in device code なにげに便利かも。

>>294
これでもう、エミュレータを使う理由は激減するね。

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

GeForce GTX 285Mが搭載されているPC買ったんですが、
ttp://developer.nvidia.com/object/cuda_3_1_downloads.html
にあるノートブック用のドライバを落とせばいいんでしょうか？

>>295
むしろ法人こそTeslaボードのほうがいいんじゃないの？
Teslaはちゃんと使えれば、十分に値段分の価値があるよ。
個人ユーザーはTeslaがきついのはわかるけど、
法人はそれ使って値段以上の利益あげればいいわけだし。

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

>>296
そうだよ

>>297
だから、予算がついちゃえばTeslaでいいんだけどね。

>>293
とりあえず実際に使ってみてから言ってください。

>>292
仕分けの時に有名になった長崎大だっけ
あそこが最初にやったのが不良品を検出するプログラムの作成でしょ。
配ってなかったっけ。

>>300
何を言いたいのか全く分からない。
C2050とGTX470はもってるが、GTX285はないよ。

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

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

>>302
クレクレ君みたいで申し訳ないが、サンプルプログラムのnbodyのベンチ結果を貰えない？
nbody -benchmark
と
nbody -benchmark -fp64
の２つで

ほれ
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

GTX470は家に戻らないといけない。
環境はWin7x64で、C2050でディスプレイ出力してる。

�d
自分でGTX470でベンチ取ったら

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

FFTは使ってませんか。
倍精度FFTはGTX285からGTX480で倍速いんですがC2050はどうでしょう。

>>302

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

>>307
FFTはデータサイズ次第だってば・・・。
まぁ、480のほうが速い場合もC2050のほうが速い場合もあります。

>>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

単精度では480が速いのに異論はない。
倍精度でもよっぽどチューンしない限り、480や470と
C2050のそれほど性能は変わらないよ。

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

ちょっと教えてほしいのですが、これで最後に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)'

>>312
*.cuは実はc++としてnvccでコンパイルされるので、*.cから呼び出すためにはextern "C"にする必要があります。
つーか、c++拒否症でもなければhellonvcc.cをhellonvcc.cppに改名してc++としてコーディングした方が楽かもしれません。

ていうか、別にCUDA関係ないじゃん。
シンボル修飾についてちゃんと勉強したほうがいいよ。

メモリがネックにならないとき
Tesla/Geforceの単精度：Teslaの倍精度：Geforceの倍精度 = 4 : 2 : 1
メモリがネックになるとき
Tesla/Geforceの単精度：Teslaの倍精度：Geforceの倍精度 = 2 : 1 : 1
という感じだね

>>313

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

>>316
仮にcのライブラリのインクルードファイルが foo.h ならば、
extern "C" {
#include "foo.h"
}
でOK。

呼び出す側がCで呼び出される側Cならば問題ない。
呼び出す側がC++で呼び出される側がCならばextern "C"←今の状況はこれ

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

>>317
>>318
どうもありがとうございます。
危うく大変な作業になるところでした。

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

自作なら、>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辺りを参照。

このスレ住人は親切すねー

実習室でやる意味ないじゃん

みんな優しいなあ・・・・。
でもCUDAを始めるにはまだ早いのではないかね？

3.1にしたら、CSのOceanが初めて動作する様になった。
simpleparticlesのロード時間もほぼ一瞬にまで高速化されたけど、
代わりにn-bodyのロード時間がクソ長くなった。
CUDAやOpenCLと違ってCSのサンプルがいまいち安定しないのは何でだろ。

>>323
たしかにこの質問主はCUDA始めるにはちょいと早すぎるかもね。
質問内容がそもそもCUDA以前の内容だし。

>>324
環境はWindowsかい？
うちはWindowsだけど、たしかにDirectComputeのOceanは初めて見た気がする。
確かにDirectComputeのN-Bodyは起動が遅いね。

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

DirectComputeはDirectX系なのでWindowsしか無いのでは？
しかも今のところWindows7限定（´・ω・｀）

Vista

あ、Vistaさんもいけたんですね・・・すんません

くだもCLもwindowsだとリモート（リモートデスクトップ、ssh)では使えないわけだが、
DirectComputeなら使えるとかいうオチでもある？

ブブブーブォォオブオーブブォーブーブー
ブォ ／￣＼ーブーブーブブーンブォーオ　
ンー|　　＾o＾|∩==＜!　プォープォーブブー
ブォ ＼＿／|　|　ブーブブーブォーーー
ォー ＿|　|__|　|　ブォーーブブブープォー
　ー|　　　　＿|　ォーブーーーブブォーブ
ブォ|　|　　 |　ブォーブーブーブーンブー

リモートしたことないけどDirectComputeはDirect3Dの一部なのでDirect3Dがリモートできれば動く

cuda profilerで分岐数がわかりますが、
プロファイラで分岐数というのはどうやって数えているのですか？

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

PTXを見てbraを探せばいいんじゃない？

>>333
SM一個あたりか、TPC一個あたりの総計ではなかろうか。

レスありがとうございます。
>>334
とりあえずptx見て数えてみますｗ
>>335
branchは1スレッドあたりの分岐数と表記があったのですが、
値が6桁なので、やはり総数なんでしょうかね？



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

>>336
多分Warp単位だと思います。
if (C) A else B 文の場合、以下のようになりますが、
各Warpで分岐する方向がスレッド毎に異なれば（両方に分岐すれば）
２個×２Warpの合計４個とカウントされるかもしれません。

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

>>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:
みたいな

9800GTにするか、GT240にするか迷ってます。
性能は9800GTの方がいいので、そちらにしたいのですが、
プログラミングが趣味なので、CUDAのプログラミングを
する予定です
CUDAを考えたらどの程度違いますか？
全然違うんでしょうか？

日
本
語
で
お
k

GTX260以下はウンコ
もっと言えばFermiじゃなきゃ今から買う価値無し

>>340
コイツは日本語書けるのに日本語読めないのかｗ
プログラミングする前に国語の勉強した方が良いぞ

日本語がどうとうかはどうでもいいが、
GT9800とGT240を比べている時点でもう
「好きにしろ」
って感じだな。コアレッシングを気にするなら、
後者の方がいいのかな？GT240ってG200系？

一応GT200系。
GPGPUやるなら多分240の方がいいだろう。

gt240はcc1.2だからどっちにしろ倍精度は使えないね。

なんだこれ・・・
ttp://pc.watch.impress.co.jp/docs/column/kaigai/20100706_378760.html

ようはFermi終了のお知らせ?
Configurable CacheがサポートされてなかったらFermi用のコードなんかHPC向けしか書かないぞ

GTX260と比較してプログラムの実行速度はどうなりそうなのかしら

例えGT200がベースでも、AMD方式ではなくFermiのテッセレーター機構をぶち込むなら
演算器大増量に変わりは無いから速くはなるだろう。
ただ、純粋なFermiベースに比べればどうしても落ちるわな。

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

値段や性能はGTX260と同じだとして消費電力がGTX480の６６％じゃGTX260より悪そうでこまる。

cc2.0で言語仕様が大幅に変わっているから、
460はcc1.4になるのかな。

倍精度切られてそうだけどね460
これで2分化が進めば、トップエンドはもっと思い切って汎用に振れるな

倍精度ないのかなあ。BOINCとかでも倍精度必要だから売るために載せて欲しい。

倍精度サポートはさすがにあるでしょ

おいおい、お絵かき特化のために旧コア使うってのに
お絵かきに関係のない倍精度なんて積むのかね
GT200でさえ、1:8の申し訳程度のものだったのに
これもTeslaのためだけだろ

別に倍精度なんか使わないから良いや
GTX465も有るんだから構わないだろ

これは長期的に見て良くないね。
もはやNVIDIAはHPCを諦めてきたのかな？

ね？
teslaって知ってる？

あのさ、Teslaはわかるけど、今まではGeForceとTeslaを共用してきたから開発費が押さえられていたわけで、
Teslaのみを今後開発して行くには、果たしてどれだけ開発費が回収できるかによるでしょ。
今のNVIDIAに両方を開発していく体力があるのかどうか疑問だけどね。

TSUBAME2もできたことだし。
もうHPC事業はHPC事業独立でできるんでしょう。

長期的も何も今までもローエンドは倍精度サポートして来なかったじゃん
今まで通りGTX580を作る時はリッチに作ってGTX560を作る時は削れば良いだろ

TSUBAME1に鳴り物入りで採用されたCSX600は無視ですか、そうですか・・・

本気で広めたいなら下位モデルでも倍精度サポートしろよ
特定のハードでしか使えないとかまじバカじゃないの

チップ作るのって、簡単なことじゃないぞ。
なんか勘違いしている奴がいるが、
#ifdef GTX580
　単精度サポート　
　倍精度サポート
#else
　単精度サポート
#endif
みたいにコンパイルし直せば、出来ると思っているのか？
コンパイルのたびに数億飛んでいくってかんじなんだぞ。
デバッグの行程もたくさんあるしな。

うっかりF7押しただけで数億とな
F5押したらどうなっちゃうの

お絵かきに倍精度なんているの？ｗ

http://www.adrenaline.com.br/tecnologia/noticias/5507/geforce-gtx-460-especificacoes-precos-e-estreia-no-brasil.html

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

倍精度きましたよ。１：１２とさらに削ってきた。

SMのSP数を変えてくるのはさすがに予想外だった

>>368
8%とは酷すぎ。25%がましな時代だったとはorz

25%だったことはないが・・・。
480は12.5%。

なんか廉価版tesla出そうな気もする
10万以下で

GTX460買ってきた。自分の倍精度プログラムはGTX260とまったく同じ性能だった。いいんじゃない。

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

相対的にload/storeとレジスタ、キャッシュが少ない

CC2.1・・・

どの命令の組み合わせなら4issue出来るんだろうか

warp sizeはどうなるの? まさか48?

どう考えても32だが

>>375
同時に実行される命令が2->4になるだけなので、相対的な容量はあまり問題にならないような。
むしろメモリバンド幅が足りない。

460だとメモリ帯域も細いから、倍精度だともはやCPUで走らすのと大差無いんじゃね？

256bit 3208MHzの465より256bit 3600MHzの460の方がメモリ帯域は大きいはずだが

アイドルの消費電力が低いし、お試しに買うのにいいですよ。

ttp://pc.watch.impress.co.jp/docs/column/kaigai/20100716_380986.html

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

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

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を使わない、簡単なプログラムはコンパイルして実行することができています。
どんな問題が考えられるでしょうか。アホな質問ですがよろしくおねがいします。

>>385
cutilがビルド済みなら、-L <libcutil.aがあるディレクトリ> -lcutil
ビルドさえしてないなら、NVIDIA_CUDA_SDK/common辺りでビルド。
手元に環境がないので詳細は割愛。

>>386
無事に実行できました！
即レスありがとうございます。

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

>>387
まぁ、簡単に言えばそういうこと。nm libcutil.a|gerp T してみると知見が得られるかもねw

このスレの住人は本当に出来た人が多いね。
と思ってたけど、ここはそういうスレなんだね。テンプレをあらためて見て気が付いたよ。

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

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

スレ違いかもしれませんが、
計算中がうまくいかず、期待してないところで0.0が出てきたので調べてみたら、
1.0e-45以下の小さい数字が出てくることが多々あり、どうやらfloatの扱える範囲外
なので0.0にされてしまっていたようです。

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

>>390
CPUではfloatで動くんですか

>>391
うごきません。cpuでも0.0になってしまいます。
dloubleにしたら1.0e-48とか表示できるんですが、
こういうときの常套手段ってあるんでしょうか

>>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を掛けて考える必要がある。

>>393
なるほど。気をつけます。

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

あ、っていうかfloat使う前提で書いちゃったけど、
ターゲットにGTX480を使えるか、もしくは
速度をそこまで気にしないならdoubleでもいいんじゃない？
あと部分的にdouble使ってもいいし。

伊理先生の「数値計算の常識」が参考になるよ
読んでなかったぜひ

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

精度にこだわるならCUDAはやめときな。
どうしてもと言うなら多倍長のライブラリをつくって・・・・てやるなら、x87を使うのが正解だろう。
そもそもCUDAを使うメリットとしてはCPUに比べメモリ帯域が広い、
並列度が高いと言うだけだから、これらが有効に使えないのなら、CPUでやった方がいいよ。
もはやCUDAを使ったからX倍速くなりました！と言うことに価値はなくなってきたからな。
CUDAを使ったからYX倍（1<Y<3）速くなりました！というなら価値はありそうだ。
CUDAを使ったからZYX倍速くなりました！というならそれは比較がおかしい。

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

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

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

>精度にこだわるならCUDAはやめときな
学者共に言ってやってくれ

学者はFermi版Tesla(Tesla版Fermi？)を使って倍精度で計算するので問題ない・・・・

念のため言っとくが単精度floatは32ビット整数すら誤差なく格納することができないからな!

>>402
何この人恥ずかしい・・・

下には下がいるってことがよくわかる

っていうか別に倍精度の値域が必要な訳じゃねえんだろ？
1.0e-7以下の精度が必要な訳じゃねえんだろ？
うぜえよ

だけど、もともとカオス性を備えた系のシミュレーションだとどんだけ精度があってもねえ

Fortranから呼び出す時って二次元配列の順番変わっちゃうんでしょ？
操作する時に列から回すとメモリが連続でなくなっちゃってアクセス遅くなったりする？

データの連続方向にスレッドを並べないと、滅茶苦茶遅くなります。
スレ違いにはなりますが、CPUの場合もデータ領域の大きさが大きい場合にはキャッシュ効率が全く変わってしまいます。

そういえばMSDNかどっかにサンプルがあったなあ
配列の順番変えるだけで10倍だか100倍だかのオーダーで実行速度が変わるやつ
確か方っぽは徹底的にキャッシュミスする並べ方らしかった

一次元配列も
1...1|2...2|...|n...n|
よりも
1...n|1...n|...|1...n|
の方がいいのかなあ。数字はスレッド番号ね。

>>410
CUDAの話なら実験してみるといいよ。驚くから。

2次元配列も全部1次元にしてるわ・・・

コンスタントメモリっていつ何時つかうのん？

gpu側でそれほど大きくなくていいけど速い必要のあるテーブルを参照したいとき。

たとえばnを外部から入力させてベクトルをn倍する関数を作りたいとき
コンスタントメモリはグローバルメモリでそんな早くないけど
一回キャッシュされたら全部のスレッドはnを読むからレジスタと同じくらい速くなる

て考えはおk？

>>415
パラメータで渡せないなら定数メモリはありだね。

>>415
最近のGPU使ってないから違うかもしれないけど、
コンスタントメモリって勝手にキャッシュとかされなくない？
まあ、ワープ内全スレッドからの同時アクセスなら１回で済むから
何回も読みに行かなければ問題はないだろうけど。

あれ、コンスタントメモリってキャッシュする命令があるんだっけ。

問答無用でキャッシュされます。

もしかしなくてもソートって並列計算に向いてない？？

bitなんちゃらソートとか、並列計算用のソートアルゴリズムもあるにはあります。
CUDA SDKのbitonicを参照あれ。

最近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);
}

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);
}

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;
}

main.hの中で色々なヘッダを纏めてインクルードしてます．
不勉強なので問題外のレベルのコードだと思いますが，方向性のヒントでも貰えたら幸いです．
よろしくお願いします．

CUDA化された姫野ベンチのソースってどっかで公開されていますか？

最近CUDAについて学び始めた初心者です。
現在三重対角行列のトーマス法をcudaで実装しようとしているのですが、計算式にある再帰処理に手も足も出ない状態です。
調べてみるとcuda（GPU）ではサポートしていないと出てきます。
GPUでの処理を考えるとこれは当然であることは分かるのですが、カーネル部分に何か工夫を行えば動かす事は可能となるのでしょうか？

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

再起処理を自前実装する

>>426
自分で書いてみたら？
CUDAの練習にもなるし。

cuda3.1から再帰は使える。
ただし、Fermiのみ。

GPUだと再帰処理ができないのが何故当然なのですか。

誰も当然とは書いていませんが・・・。
まぁ、動いたとしても遅いわけで。

>>428
再帰処理を自前実装するとは、再帰処理を非再帰処理に変えて実装するという意味ですか？
それとも何か書き方があるのですか？


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

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

各スレッドは各プロセッサにつきひとつ割り当てられて一度にプロセッサ数のスレッドが
スレッド÷プロセッサ回実行されるのですか？
それとも全スレッドはひとつのプロセッサに複数個割り当てられて
(ひとつのCUDAコアは同時に複数個のスレッドを処理できる？)
全スレッドが一回だけ同時に実行されるのですか？

>>434
難解な表現ですが、どちらかというと後者に近いです。

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

>>435
交互に処理ってちょっと違和感あって資料見たら、本当に交互に処理するのね・・・。
メモリ関係のレイテンシ待ちしてる間暇だから他のワープも処理してやんよって感じか。

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

>>433
再帰って要するに戻り先アドレスとローカル変数をスタックに乗っけてから
同じ関数実行しなおすわけだから、
そのスタック構造のとこだけ自分でメモリ上に作ってやればループに置き換えられる。

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

>>435
なるほど。でも待つということは前者のようにも感じます
480個の切符売り場に数万人が並ぶかんじでしょうか。隣がすいたら移ることもできる？
同じスレッドブロックの仲間たちは同じ列に並んでいるのでしょうか。それとも同時に切符を買えているのですか？

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

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ってこと？

＞それとも最大数ってことは1024*1*1や256*2*2はいいけど1*1*1024はダメですよってこと？
そうです。

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

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

環境くらい晒せや

具体的にどういうエラーが出てるのか書けよ馬鹿垂れ

GTX 480でキャッシュが使われているかどうか判定する方法はあるのでしょうか？
プロファイラにかけてみるとl1_global_load_hitが0で、
l1_global_load_missには値が入っています。
キャッシュが有効利用されていないと思っているのですが、どうなのでしょう。

>>441
ここの人たちは他と違ってかなり優しいけど、
はっきり言っておまえの質問は酷すぎる。
ゆとり世代か？
そもそもCUDA初心者と言っている時点でC言語の初心者じゃないのか？
環境構築はCUDA以前の問題だからな。

>>444
そもそもデータを再利用するような計算なのでしょうか？
一回読んで終わりならそうなる。

>>440
ありがとうございます。
二番目と三番目で表現は一緒なのに内容は違うんですね

>>438

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

>>444
__syncthreads()入れてる？

>>448
同じ人が再び列に並んできっぷを買い続けないといけないので例としては適切でないかと。

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

卓球ダブルス

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

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

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

FortranでCUDAができるというような話を聞いたんですが、もうできるのでしょうか？
それは純粋にFortranで書くのですか？
それともCで書いた関数をFortranで呼び出して使うという方法でしょうか？

金払いたくないならFortranからCを呼び出す
PGIからFortran用CUDAコンパイラが出てるからそれ買えば全部Fortranで書ける。
ただしFortran90以降の書き方じゃないとコンパイルできないっぽい。

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

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

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

ブロックごとのスレッド数って多い方がいいんでしたっけ？
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であることも多々あります。その場合なんかもったいない使い方なのでしょうか

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

>>459
ブロック1000スレッド1より
ブロック1スレッド1000のほうがいんですね
＞使うgpuのコア数分投入した方がいい。
というのはスレッド数をってことですか？（480基だったら400くらい？）

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

あ、一部説明が拙かった。
論理上のスレッド・ブロック・グリッドと、実際のプロセッサ・マルチプロセッサ(MP)・デバイスでは若干違っている。

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

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

ばかですみません

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が同じなら物理的な違いはないのでしょうか)

>論理上のブロックは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スレッドづつ順次割り当て、実行という感じでしょうか

長乱文失礼

>>461
書いていることがめちゃくちゃなのだが・・・・。
もしもネタでないのであればじっくり読み返していただきたい。

あれ?　そんなに間違ってたっけか。知識が古いからFermiでの数値と違っているかもしれないけれど。
後は用語がいい加減かな。
まぁいいや、誰かの修正を待とう。

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

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

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

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

GT200系までだと、1SM=8SP。Fermiはしらん
32スレッドで1Warpなのは1つのSMが4クロックの間、threadIdだけ変えて同じ処理を行うから
だからスレッド数は32の倍数が効率がよい

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

ブロック数に関しては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の実行モデルを再考することを勧める。

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

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

カーネル関数の引数はどこに保存されるのだっけ
たとえばint n=1000があって
kernel<<<>>>(n);
とやるとnはレジスタに入ってくれるの？
グローバルメモリに入っちゃうの？
配列と普通の変数で違ってくるみたいなレスを昔みたような気がするんだけど

シェアードメモリに入る

>>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が違ってきちゃった。
なぜーー？？

blockサイズ、threadサイズも入るから

>>471
double型は8の倍数のアドレスにしか置けないから。

>>471
それはさて、doubleでないといかんの?

GPU向けのコードを.cppファイルの中に書いてnvccでコンパイルすることはできないのでしょうか?

>>475
そのファイルの名前を.cuにすればコンパイルできます。

>>476
そうするとほかのコンパイラでコンパイルできなくなってしまいます。
CPU用コードとGPU用のコードを同じファイル内に書いて、ifdefでコンパイル時に切り替えたいのですが。

#include "hoge.cu"

>>475
nvccでオブジェクトファイル作って、
gccなりでリンクすればいいじゃん。

変に凝って一つのソースで両方のコード書くと条件コンパイルだらけで醜くなるだけだと思うなぁ。

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

さてと、明日からは本体の方のSharedMemoryの使い方の見直しでもするか。
# こちらは実装の都合でthread数を増やすとSharedMemoryが足りなくなるので32固定……

>>480
OccupancyはあくまでMP(=SM =Streaming Multiprocessor)内の占有率だから、
ブロック数がMPの数より少ないと、一部のMPがやることなくて遊んじゃうのかも？
遊んじゃうぐらいならOccupancy下げてMP埋めたほうが若干速いのかな。

つか若干ぐらいなら別にそうする必要なくて、Occupancy上げとけばいいのかもしんないね。

3.1を導入して、カーネル関数内ででprintf()したら、
combora.cu(275): error: identifier "cuprintf" is undefined
と言われました。

環境はVCでx64です。

何が間違っているのですか？

CC2.0以降のみ、printf()対応のようですね…
お騒がせしました。

>>473
あーなるほど。

>>474
doubleの欠点てたくさんあります？？

CUDAのドライバって
CUDAの専用ページにあるドライバと
使っているデバイスの最新ドライバは
同じ種類のものなのでしょうか？

GTX480をつかうのですが、
CUDAのページにある
devdriver_3.1_linux_64_256.40.run.sh
とGeForceGTX480用の最新ドライバである
NVIDIA-Linux-x86_64-256.44.run.sh
では480用の方が最新に見えるのですが、こちらを使用して大丈夫なのでしょうか？

半年くらい前にLinuxでNVの最新ドライバ入れたらCUDA動かなかった。
CUDA用ドライバなら動いた。

ふぅ、アク禁解けた。

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

>>484
doubleの欠点は、遅い、大きい、の2点に尽きる。
演算時間も掛かるし並列度も落ちるみたいだし、
データ量が倍になるから共有メモリに置けるデータ数が減るし、
HostからDeviceへの転送時間も長くなる。

総和計算はcpuのforループでやるより
gpuで工夫した並列リダクションのほうが何倍も早いものなのでしょうか？

1000*1000程度の配列x,y,zがあり、
それぞれの成分に重みをかけた和
の総和計算を重みを変えて1000通りやりたいのですが
どこを並列化したら良いのか悩んでいます。

>>488
xyzは固定かな。だったら転送量が少ないから望みはあるかも。
1000x1000を適当に分割した範囲で和を共有メモリに置いて、総和の計算は二分木でやればOK。
その総和をまた共有メモリに置いておいて、更に次のステップで全面分の総和を出すって感じかな。

CUDAのreductionサンプルは
最後の一手だけやれば、他は無理にやらなくてもいい。
どうせ誤差レベルの速度向上にしかならない。

あのサンプル、真面目に先頭からドキュメントとコードを読んでいって損したわw

同感。どのサンプルも動かすためにあれこれ余計なコードが多過ぎる。
デモサンプルと機能サンプルと実装サンプルは区別して欲しかった。

実はそういう意味で言ったのではない

>>489
たとえば10個分の和を各スレッドでやらせることを256スレッド/ブロックで391ブロックやって
そのままその中で各スレッドの和を共有メモリに置いて2個づつ合計することを8回
391個の和になるからそれぞれグローバルメモリに書き込む
次のカーネルで同じように和を取る
ということを1000回条件変えて繰り返すってことでしょうか。

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

>>490
最後の一手ってリダクション7の部分ですか？

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を使っています。なにか情報足りなかったら追記します。

>>494
そう。単純な総和のreductionなんて
N個のデータでN-1の加算しかしないのだから
コンピュートバウンドではなく、IOバウンド
IOを最小限にするようなアルゴリズムを組めば
多少演算が非効率でも誤差レベルの影響しか出てこない。

>>495
CentOSは知らないけれど、nVIDIAのGPUは型番を認識できないと汎用VGAと認識されるから1024x768位のXなら動く。
なので、認識はできていると言うか、できてないと言うか……
RedHatならKudzuのメッセージ見て対処できるんだけどCentOSはどうなんだろ。Linux板で聞いた方が早いかも。

__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
すみません。つけてないです
どのようなコンパイラオプションがいるのでしょうか？

-arch

>>501
-arch=sm_20
でできました！！

勉強不足でした。
ありがとうございます。

Compute Capability 2.0からはカーネルの引数はシェアードメモリを使わずに
コンスタントメモリに入れるようなのですが
カーネル内でコンスタントメモリの引数をシェアードメモリに読み込んだ方が早いのですか？

>>503
疑問に思ったら先ずは実測。
定数メモリは充分早いはずだから、共有メモリに移し変えるメリットはないんじゃない?

>>504
たしかに気にならない程度の速度でした。なんでコンスタントメモリになったのかなぁ
シェアードメモリを空けるため？？

ところで
extern __shared__ float array[];
__shared__ float array[128];
はどのような用途で使い分け、機能面でどう違いがあるのでしょうか。

>>505
4要素使いたい場合だと今まで苦労してたからなぁ〜
コンスタントメモリに入るのは個人的にうれしい

>>506
同意。引き数が多いと結構きつかった。

>>505
後半の質問の意味が判らない。要素数の定義の仕方の違いを言ってる?

>>507
前者は要素数(size)をkernel<<<>>>()の第三引数で指定するようなことが書いてありました。
要素数をカーネルの外で定義するか中で定義するかで
ただ単に宣言の仕方が違うだけで動作は全く一緒なのでしょうか？

>>508
あー、そういうことね。<<<>>>で指定する場合は可変長にできるメリットがありますね。
出力はどうだろ。ptxを出力して見比べてみては?

岡田賢治「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
>>510の本立ち読みしたら、巻末に「IPAの人とは関連ありません」って書いてあって笑ったｗ

>>512
雑談は余所でやれ、馬鹿たれが

うわぁ……

>>510
エラー内容を詳しく。まぁ、なんとなく落ちは見えたけど。

>>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 以上だとデータの読み出しは成功しますが、カーネルの計算をスルーしているような感じです。

>>516
そのエラー自体は、書かれている通りカーネルの実行時間が長すぎて出ています。
LinuxだったらX立ち上げないで(run level 3とか)起動すれば通るかもしれないけれど、Windowsはどうするんだろ。
本が想定しているGPUより能力が低いのが問題なのかもしれないので、ネカフェかなんかで速い端末で動かしてみるとか。
後は処理量を加減するしかないけれど、それには元のコードを見ないことにはなんとも。

CUDAのドライバは最新ですか。一つ前はなんかそんなエラーがでましたよ。

>>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 なら出ない?

intよりfloatの方が速い

もうちょっと具体的に言うとFMAがfloatでは使えるので、a=b×cを4クロックで実行できる。
FMA考えなくてもaddもmultiplyも4クロック。

intだとint addに4クロック、int multiplyに16クロックかかる。

>>519
自己レスです。
int でシェアードメモリを用いたら 1024 でも 2048 でも timeout にならずに
動きました。原因がわかってすっきりしました。
皆様、ありがとうございました。(_O_)

1000要素の配列データを配列として1回cudamemcpyするのと
要素ごとに1000回cudamemcpyするのではなんでこんなにも時間が違うのですか

>>523
・単純に一回リクエストする際に生じるオーバヘッドが馬鹿にならない。
・実はcudamemcpyは完了復帰じゃないので一回のcudamemcpyでは即復帰してしまい実際の転送時間は測れない。

>>524
>実はcudamemcpyは完了復帰じゃない
??
cudamemcpy直後にコピー先にアクセスしても信用ならない値なのですか？

もちろん信用できますのでご安心ください。

>>525
・HtoDの後はカーネル呼び出しの際に同期される。
・DtoHは完了復帰になる。但し、streamを使う場合は明示的に同期する必要がある。

>>526-527
ああ、DtoHのことしか考えてませんでした。

>>527
じゃあHtoDの場合はカーネル呼び出しまではストリームっぽくなるから
データ送るなら早めにやったほうがお得なのでしょうか

>>528
とってもお得。
ついでに言えば、カーネル呼び出しも(限定的に)即時復帰だから
カーネル呼び出し後にDtoHするまで他の処理をして時間を稼いだ方がお得。

>>527
>・HtoDの後はカーネル呼び出しの際に同期される。
こういう仕様だと送信バッファの内容を変更できないわけですが・・・。

カーネル呼び出しが即復帰するのと間違えているのでは？

処理→HtoD→カーネル　よりも　HtoD→処理→カーネル　の方が若干早かった気がしたんだが。

んじゃ、メモリコピーはストリームにしない限り完了復帰ってことか。

cudaMemcpyAsync

Warp sizeは32なんだけど Number of cores/Number of multiprocessorsが32なの。
この場合は1サイクルで1warpできるの？
その場合half warpとかそういう概念はかわってくるの？

>>533
1クロックで2個のハーフワープを処理する．
2クロックで2個のワープが処理されることになるが，上に書いたように，
1クロックで丸々1個のワープが処理されるわけではない．

岡田賢治のCUDA高速GPUプログラミングを読んで勉強中なのですが、
質問させてください

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

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

この本は正しいのでしょうか？
ブロックが複数のときはどうしたら良いのでしょうか？
よろしくおねがいします。

>>534
なるほど。考え方は特に変えなくていいわけですね。

11030 アップデートファイルのダウンロードがタイムアウトしました。
FINAL FANTASY XIVを再起動してください。 3:40012,20498,20049

threadIdx.xやblockDim.x等のビルトイン変数はどのメモリに置いてあるのでしょうか
カーネル内で何度も参照するよりも
カーネル内のローカル変数に格納してレジスタに置いて参照した方が高速ですか？

>>538
その辺は勝手に最適化しますよ。ptx出力をご覧あれ。
あーそうそう、その辺のビルドイン変数はshort幅ですが、
受けるローカル変数はshort幅だと変なptxになりますね。

>>535
確かに、ブロックが複数だと正しい結果になりませんね。
Nvidiaのサンプルも１ブロックで処理していました。
サイズが512までならcpuで計算しても遜色ないので、
苦労してgpuで計算する必要はないですね。

>>539
特に気にしなくていいわけですね。ありがとうございます。

ページロックメモリについてなのですが、
cudaMallocHost()にmalloc()と比べたデメリットってありますか
GPUのメモリと通信する配列をCPUに確保するときは
すべてcudaMallocHost()を使った方が良いのでしょうか

>>540
ですよね。
2^16くらいはソートしたいのに。
複数ブロックだとどうしたらいいんでしょう
1ブロック内で並び替えたバイトニック列をマージでリダクションするのでしょうか

CUDA Toolkitの3.1 64bit版をWindows7 x64にインストールしようとすると、
インストール途中のRegistering products...で止まってしまいます。

全く同じ構成のPCで昨日試したときは、この部分は一瞬で終わったんですが、
これはnVIDIA側の鯖が死んでいるのでしょうか？

>>541
CUDAテクニカルトレーニングには

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

とある。どのくらい割り当てるとどのていど低下するのかは（自分は）わからないし調べていない

>>541
ロックしたまま落ちると大変なことになりそうな悪寒。
その他の手段をやりつくしてからでも遅くないんでない?

>>544
少なくとも、スワップするような場合は明らかに違いが出る。
逆に言えば、スワップしない程度しか使っていなければ大差ないと思う。

>>543
うちのマシンでも同じ症状がでてる。
どうしたものか・・・。

safeモードで起動して、インストールするといけた。
問題が発生してたのはWindows7 Ultimate x64をクリーンインストールした環境。

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

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

困ってる人は参考にしてね。

複数GPU環境でGPU単体で動くプログラムを実行すると
どのGPUがデフォルトで動くのかだれかしってる？
デフォルトのGPUが選択される優先順位ってどうなってるのかな？

プログラム作成者が指定する。

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

手抜きなcudaSetDevice( 0 );のようなコードだと
常にデバイス0が選択される。

>>547
自分はウイルスバスター2010の環境では大丈夫でしたが、
2011だとRegistering products...で止まりました。
(ただし、ウイルスバスター2011が原因かは分かりません。
リアルタイム検索を止めても、インストールはできなかったので)

ダウンロードしたexeに引数 -r を付けて起動すると、
OSをsafeモードにしなくてもインストール出来ました。
-rはサイレントモードを意味するようですが、
何が違うかは分かりません。

処理時間が異常にかかる件、質問します。
下記のプログラムで「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;
}

block[]の配列がグローバルメモリにあるからかと。

トリップ検索か

>>552
ありがとうございます。
block[]をシェアードメモリにしても、大して改善しませんでした。
block[]にアクセスすると、Windows自体が激しく重くなります。

カーネル実行中は基本Windows重くなるよ。
軽くしたいなら、複数回のカーネル実行に分けな。

というか、Fermiなのかそうでないのか。

Fermiならキャッシュが効くから多少は軽いと思うけど。
いずれにしても、GPU向きじゃないコードだなぁ。

分岐が原因・・・なわけないか。

実行中のデバイスメモリの使用量はどのようにして知るのでしょうか。
どうもメモリがステップごとに増えてるような気がして時間推移をチェックしたいのですが。
linux,CUDA3.1,CC2.0です。

3.2RCが公開されたね。
いろいろな機能が追加された様子。

>>559
メモリ確保時の引き数と戻り値を自前で管理するしか。
勿論、解放のときの処置も忘れずに。
あーそうそう、メモリ確保と解放を頻繁に行なうと分断化が起き易いから要注意。
SDKのバージョンでその辺りの振る舞いが変わるから、動いていたコード(とGPUの組み合わせ)が
バージョンが替わると突然メモリが足りなくなったりする。

nsightのビルドルールでビルドすると、
Visual Profilerでエラーが出ます。

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

皆さんはどうですか？

環境はCUDA 3.1/Windows 7 64bit/Nsight 1.0/ 258.96/GTX 480/VS2008 SP1です。

CUDAってクラスというかオブジェクト指向を利用することは可能ですか?

質問です

GPGPUでは分岐処理は遅いと言われますが、どうして遅いのでしょうか？

>>564
全スレッドが同じ処理をするからif文がfalseでも勝手に先に進めず、
他のスレッドが終わるのを待つからでは。
ただ、たいていのプログラムはメモリアクセスのほうがネックになるから、
あまり気にしなくてもいいと思う。

>>565
っ Divergent blanch

CUDAを使った研究で、fortranのコンパイラでやることになったのですが、
分かりやすい教科書などないですかね。
なかなかfortranのものは見つからなかったんですが・・・

>>567
NVIDIAの糞営業にお伺い立てたら?

>>563
可能です。但し、パフォーマンスを追及すべきCUDAでパフォーマンスを
追及しにくいオブジェクト指向に固執することは無意味です。
肝心要の部分はオブジェクト指向を諦めるのが無難ですね。

ワーストシステムズに問い合わせたら資料くれるんじゃね？

>>569
やっぱりスピードを重視するとOODは不向きなんですね。ありがとうございます。
基本的にC++で書けるとのことなので現在の形のまま利用できればと思っていました。
vectorやstringといったものが使えるのであれば恐らく問題はないので設計しなおします。

動的にmallocするようなやつは無理じゃない?
STLならthrust使えばある程度いけるけど

>>567
PGIのリファレンス読め
それでCUDA CとCUDA Fortranの違いはわかる
CUDA Fortranの教科書を探す・出るのを待つくらいならCUDA Cの教科書読め
（C言語がわかりませんなんて言うなよ？）
CUDA Fortranにほとんどそのまま適用できる
有料でCUDA Fortranの講習会やってくれるところも探せばあるよ

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のプログラムを変形するか
どちらが適切でしょうか？

エラーを読めば分かるだろ？
「デバイス(GPU)の関数からホスト(CPU)の関数の呼び出しはエミュレーション以外では許されません」

っと思ったがもしかして__global__の存在を知らないのか

とりあえず解説本買って嫁と言いたい

あたらすぃツールキットをインストールするまえに、古いツールキットはアンインストールするべき？

cuda使ったＧＵＩアプリ開発には
どのＧＵＩライブラリを使うのが一番生産性が高くなりますか？

CUDAと何の関係もないから好きなの使え
もしくはGUI系のスレ行って聞け

CUDAだとC系ですよね
javaなんてものだと、ダメですよね
だと必然的にC系のライブラリになって、そこからCUDA呼び出せないとだめですよね
そういう前提が縛られてるのに、CUDAとは何も関係がないってよく言えますね
バカですか？というかGUI作ったことないCUIプログラマなんですか（＾＾＾＾＾＾；；；；；；；

javaからCが呼び出せないとでも思ってんの？

恥ずかしい>>580がいると聞いて

>>577
インストーラーがアンインストールしていいか聞いてきたと思う。

アンインストールするか聞かれなかったのでそのままインストールして、今二重にツールキットが入ってるんですが
v3.1は消すべき？

>>584
SDKじゃなくてツールキットが二重にはいってるの？
同じフォルダに？

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

と、ドライバーとランタイムのバージョンが異なるのは問題無？

いいわけないんじゃね？

疑問形でいいわけ？

いいんじゃね?

あんれ。toolkit3.2 64bit　のダウンロードができないなう。みんなダウンロードできる？

linux以外not foundになるぽい
べつにいいんじゃね？

ダウンロードしてから、v3.1消せば良かった…。

404自重

アドレス中のファイル名をcudatoolkit_3.2.7_win_64.msiにすればいけるよ
正式版を公開しようとして、途中でストップ掛かったんだろうか

他にもドライバーのバージョン表記がおかしかったり
260系ドライバーでOpenCLの動作がおかしくなってたりと、
GPGPU担当者がグロッキーになっているのではなかろうか

正式版が出るまで待つのが吉

ああ、あとインストール先やフォルダー名やビルドルールファイル名が3.0から変更されていて、
今迄のプロジェクトは設定の変更が必要になるから注意ね

>>596
kwsk

で、3.2ってなんかメリットあるの？

不具合修正
新機能

release notes読もうね

新機能は特に魅了を感じなかったんだけど、全体的に速くなったりしたの？

お。アドレス修正されたようだ

ひょっとして新品PCに新しくVisual Studio 2008をいれて、Developer Drivers　260.21、toolkit 3.2、SDK3.2いれるだけで、サンプルプログラムをビルド出来るようになった？

GPUなしでは動かない

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はなぜ使われているのでしょうか

三角関数つかってないか？ ＞lmem

CUDAで三角関数使うならSFU使わないともったいない。

>>605-606
使ってます。
ローカルメモリが使われるんですね。
use_fast_mathって精度が少し問題なんでしたっけ？

すいません。だれか教えてくれませんか。
cudaのサンプル実行したら以下のメッセージがでます。
cudaGetDeviceCount FAILED CUDA and Runtime version may be msimatched.
サンプルプログラムはDeviceQuery.cです。
コンパイルは通ります。
gccのバージョンでしょうか？
環境はmacBook cudadriber 3.2.7 toolkit3.2.8
をインストールしました。同様な現象のかたいますか？


Device

暗号化のメイン関数の処理にファイルIOのスピードがついてこれない。
fread３００Mで10秒以上かかるのは何とかできませんか？

別スレッドで先読みさせといたら？

マルチコアで解決できるかと・・・

>>608
とりあえずドライバのバージョンが3.2.8でない理由を聞いてもよいでしょうか？

>>608
gccのバージョンと思うなら、バージョンくらい書きなよ。
ここはエスパースレじゃねえぞ。

mac初めてさわる初心者なんだよ。
gcc4.2 toolkit何度もバージョン変えてもエラーがでる。
わからん

わからんじゃねぇ、ソース見れば原因一発で分かるだろ

>608はエラーメッセージが読めないのだろうか。

ドライバとランタイムのバージョンが合ってねえズラ。

質問クオリティの低下が止まらないな

英語分かんなくても、とりあえず翻訳サイトにその文章持ってく位の事はしろよ
つうか普通ミスマッチ位分かるだろ

ドライバも最新にすれば動くとでも？

分からんなら試せよ怠け者
黙ってcuda3.2rcのページに行ってドライバ落として突っ込んで来い

まぁ、開発者でないならば今回の3.2は落ち着くまで待った方が良いと思うが。

開発者でも普通はRCに飛びつかない。

プログラム初心者ですが飛びついてしまいました

CUDAはプログラム初心者向きとは間違っても言えない．
素直にC，C++，コンパイル，リンクの仕組みのあたりの基礎を
しっかり身に着けてからCUDAに挑戦してほしい．

資料が充実している分、ATI Streamに手を出すよりはマシ、かな？

CAL+ILならそうだが、OpenCLならそうでもないよ。

nvidiaと心中する勢いでCUDAを

CUDAなんてもう3年ぐらい放っておいて一般的になってから手を出したほうがいいと思う。

一般的になったら先行者利益がないだろ

CUDAでリアルタイムエンコーダな動画編集ソフトを作って
１５００円くらいで売って、１００００人に売れば１５００万円儲けられるな

よし、がんばるぞ☆

>>631
ネタとかツマランからポケモン板池

俺はnvidiaを信じているよ

AVCエンコーダはCUDAについてるしな。

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を分かっていませんが、よろしくお願いします。

プログラミングガイドの最新版を見たけど、G.3.3.2.とG4.3.にsharedメモリのbank conflictのことが書いてある。
結論から言えば、char型のアクセスは、G92アーキテクチャではconflictが起きて、G100アーキテクチャでは起きないようだ。
他にもバンク数が16から32になったり、細かい部分でかなり変更されている様子。
試していないがGTX480ではl1_shared_bank_conflictの値が取れるかもしれない。

>>636
なるほど、ありがとうございます。確かにそうかいてありますね。。。

とても参考になりました。
l1_shared_bank_conflictについては今度試してみたいと思います。

カーネル関数呼び出しの第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の配列の要素数はブロック内スレッド数と決まっているのでしょうか

>>638
アドレスはカーネルの方で陽に指定する
１次元アレーだから
マニュアルにサンプルがあるはず

CUDA by Example: An Introduction to General-Purpose GPU Programming
のPDF版のダウンロード購入ってどこでできるの？

>>639
長さ分のshared memoryを取ってカーネル内で分り当てするんですね。ありがとうございます

http://pc.watch.impress.co.jp/docs/column/ubiq/20101008_398601.html

カーネルの引数に配列をいくつか渡したいのですが、数が多いのでそれらの配列のポインタを配列に入れて配列を渡すことにしました。
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となってしまいます。
どこがいけないのでしょうか

>>643
tmp自体もdeviceに転送しないとダメ。
それがイヤなら構造体でもOK。

>>644
転送ってことはグローバルメモリに入るんですよね
構造体でいけました！
ありがとうございます！！！

>>645
どうでもいいけど、理由は把握できた? 把握できなきゃまたはまるよ。
# めんどくさいからレス無用

はじめてのCUDAプログラミング本のP.136にshared memory の総和計算での
ウォープ・ダイバージェントを少なくする方法が載っています。
分からないのですが、そもそも else 節がない if 文でウォープ・ダイバージェントは発生するのでしょうか？
自分の理解だと、32人33脚のパン喰い競争でアンパンをスルーするスレッドがいるだけで、次のドーナツ(else節)は
ないので、ウォープ・ダイバージェントは起きないと思うのです。

>>647
32人33脚の真ん中の奴が
「俺アンパン嫌いだからちょっとイギリスまで行ってヨークシャープディング食ってくる。俺が帰ってくるまで待ってろよ」
とか言い出したら嫌だろ

Warp divergentは分岐命令を異なる方向に分岐することをいうのでelse節があるかどうかは関係ありません。

>>648
そのWarp全員の足並みが揃っていても、
結局、他のWarpではイギリスに行く奴(if節を実行するスレッド)がいるから、
それを待たないといけないんじゃないですか。
>>649
なるほど。
if 文が終わったら合流しないんですか。
異なる方向に分岐と言っても、この場合、片方は何もしないので
合流するのであれば Warp divergentのコストは何もないと思うんです。

一人がイギリスに行ったら全員で一緒にイギリスに行くけど、付いていったやつはなにもしない。（ベクトルマスク）
誰も行かなければイギリス往復分の時間が浮くからそれがコストになる。

>>651
あ〜この場合、複数のWarpは同一のSMで切り替わりながら実行されるから
Warp内では足波を揃えていた方がいいんですね。
全部のWarpが並列的に走ってるのかと思ってました。

やっぱ　わからん
cudatoolkit 3.2.2.8
devdriver 3.2.8
gpucomputingsdk 3.2
mac os snow leopn
相変わらず　runtime version may be mismatchがでる
バージョン合ってるじゃん。
だれかかっこいい人教えて

だからソース見ろって言っただろカス

LD_LIBRARY_PATH?

そんなこといわないでくれよ
アニキ。
mac初心者なんだよ。
また朝鮮してみるよ

開発は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では正しく動くのに並列計算だとうまく動かないためカーネルの部分が原因とは思いますがまだ勉強中のため解明できません。
原因として考えられること、解決案を教えてください。

・GPUでやるほどの処理なのか疑問。
→実験目的なら納得。

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

・周辺コードがないからなんとも言えないが、ちゃんと転送されているのだろうか。
# thrust::copy()かな。
・同様に、endsは初期化されているのだろうか。

CUDAは並列動作時に同じアドレスへの書き込みを保証できないから
atomic命令使ってカウントするか、同数のtrue/falseフラグ配列を用意して
カーネル終了後に別途足し併せないと無理

そうだ、肝腎なことを忘れていた。
# 馬鹿でぇ＞自分
## つーか、流石に自分自身ではこんなミスはしないのに……

そもそも、複数のスレッドが同じメモリに書き出すのは無理。
ほぼ同時に走るから、1や2になるのも道理。
逆に言えば、巧く並列に実行しているからこそちゃんとカウントできない。

>>658
>>・可読性の悪さを〜
申し訳ないです。実際endsの初期化とかstudentへの値の代入などその他の処理が結構あったため切り貼りしたのが原因です。
>>・周辺コードが〜，同様に〜
この部分の動作は一応確認できていたため省いてしまいました。すいません。

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

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


あともう1つ質問なのですが
thrust::device_vector<hoge> hoge(N);
という宣言を行った場合ホスト側にメモリ領域を取らずデバイス側にメモリ領域を確保できるのでしょうか?

すいません肝心のお礼を書いていませんでした。
即レスありがとうございます。参考になりました。

>659が真の功労者。お礼はそっちに。

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

thrustは詳しくないけど、ちょっと見てきたところで判断する限り、ホスト側とデバイス側は別々にメモリを確保する必要がある。
従って、device_vectorはデバイス側だけメモリを確保すると思われ。

重ね重ねありがとうございます。この板IDが出ないので誰だかわからないですね。
実際に取り扱うデータは10万件くらいのデータ(蓄積しているのでもっと多い)の分析で
分析をC++で書いていたのと並列化に向いたアルゴリズムも存在していたため
CUDAなら乗せれると思って勉強しはじめました。基本はC,C++ライクなのでがんばってみます。

runtime versionてどう調べるの？
cudaの

devicequery
実行できないなら自分が入れたもんくらい覚えとけよ

CUDAでカーネルを10回実行しようと思った場合
以下のような書き方でいいのでしょうか?
for(int i=0; i<10; i++){
kernel<<<n,m>>>(,,,);
}

一応。
タイマーとか挟むつもりなら注意が必要だけど

>>667
カーネル呼び出しはコストが掛かるから、できれば一発呼びで済ませた方が。
その例なら、kernel<<<dim3(n, 10, 1), m>>>(...)できないだろうか。

667ではない，プログラミング初心者ですがコストっていうのはどういう意味ですか．

>>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
人の犠牲において, 人に損害をかけて.

配列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]
}
}

すいませんスレッドの場合、総計512が最大でした。ブロック数が(1000,1,1)です。

>スレッド数を10個(10,1,1)にした場合正しく計算されませんでした。
意味が分からん。んなわけないだろ

分岐はウォープ内で同じ方向だったら対して無駄にならないんじゃない？

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)

>>674
スレッド数が10だとsizeが100なのに10件しか処理しないという自明なことを書いているだけかと思った。

>>676
あー。。。なるほどサンクス

そりゃそーだろう...

>>670
この文脈だと、カーネルをコールする時のオーバーヘッドの事

nsight 1.5が出たけど、toolkitはRCのまま・・・

XML 解析中に次のエラーが発生しました:

ファイル: C:\Documents and Settings\All Users\Application Data\NVIDIA Corporation\NVIDIA GPU Computing SDK 3.2\C\src\vectorAdd\vectorAdd_vc90.vcproj
行: 22
列: 4
エラー メッセージ:
カスタム ビルド規則 ファイル 'C:\Program Files\Microsoft Visual Studio 9.0\VC\VCProjectDefaults\NvCudaRuntimeApi.rules' が見つからないか、または読み込めませんでした。
ファイル 'C:\Documents and Settings\All Users\Application Data\NVIDIA Corporation\NVIDIA GPU Computing SDK 3.2\C\src\vectorAdd\vectorAdd_vc90.vcproj' を読み込めませんでした。

こんなエラーがでてサンプルプログラム動かせないんだけど・・・

\Cを消したら動くんじゃね

じゃ、諦めるんだな。

cudaにおすすめのエディタとかあります？
色分け用に予約語が登録された拡張が公開されているやつとか

CUDA SDKの中にsyntax_highlightingというフォルダがあって
usertype.datというファイルに予約語が羅列されているから、それを元に色分けルールを作るのはどうかな。

>>683
Windowsですか？

SDKだけとLinuxの場合、どこにインストールするのがいいの？
ホームディレクトリ？？

rootでいいんじゃね？

ホームディレクトリか/usr/localあたりじゃね？

/opt でいいんじゃね？

>>684
意外と予約語少ないんですね。
基本はCの定義ファイル使えばいいのかな
>>685
macですがwindowsでもお勧めあるなら知りたいです

みなさんＣＵＤＡで何を作ってるの？

カーネル内で変数に整数を入れたいのですが、intほどの範囲はいらないというとき
例えばsigned char を使った方がresister/threadは少なくて済むのでしょうか

Fermiなのかそうでないのか・・・。

>>693
CUDA Capability Major revision number: 2
CUDA Capability Minor revision number: 0

ってFermiですよね

>>692
ptx出力を見るよろし。nvccのオプションでバージョン指定を忘れずに。

Windows Live MeshのリモートデスクトップだとCUDAも動くっていうのは既出？
試しにやってみたら動いちゃって、すごい感動したんだが。

レジスタの使用量について質問です。

占有率やMPあたりのアクティブスレッドブロック数を向上させるべく、
スレッドあたりのレジスタの使用量を減らしたいのですが、上手くいきません。

カーネル内で使用する変数の数を減らしてみても、レジスタの使用量は変化しませんでした。
↑プロファイラやcubinファイルで確認

スレッドあたりのレジスタの使用量はいったい何で決まっているのでしょうか？

変数の量を減らすより、INTやFLOATなんかの型の種類、条件分岐を減らす方がいい。
PTXを読んでれば分かる

>>698
ptxですか。
さっそく解析してみます。

同時に使う変数の数を減らすこと
簡単に減らせるようなとこは最適化で勝手に削られる
本当に減らそうと思ったら、アルゴリズムの根幹に手をつけないといけない

後は遅くなることを覚悟の上でインライン展開を禁止するとか。
固定長のmemset()相当のコードをコンパイルしたらレジスタを64個使って展開しやがったw

...ptxってどうやって読むのん？？

>>702
CUDAのドキュメントを読んだ上で読めば当たりがつかない?
それほど長くなければ解析してもいいけど。その為にはソースも必要だけど。

まず文字化けみたいなのばっかで意味のある文字列が最初の関数名しかわからないです('・ω・`)
バイナリで読むんじゃないよね？？

まさかELFで読もうとしてるんじゃないよな（；＾ω＾）

先ずアセンブラの基礎知識はある前提で、この辺の擬似命令が判れば読めるだろ。
ニモニック　概要
.file　　　　ファイル番号　ファイル名
.entry　　　関数宣言(パラメータリスト)　　　# 関数名は、c++同様マングリングされている
.param　　種類　パラメータ名　　　　　　　　# 種類は.u32, .u64など。ポインタは非負整数として扱われる
.reg　　　　種類　名乗り<本数>　　　　　　# 種類は.u16, .s32, .f32, .predなど
.loc　　　　ソースのファイル番号　行番号　　# 三番目の0が何かは知らね

アセンブラの基礎知識がないならいろんなソースを作ってコンパイルして辺りをつけてみるしか。

CUDA Toolkit 3.2 RC 2 (October 2010)
http://developer.nvidia.com/object/cuda_3_2_toolkit_rc.html

最近、CUDAを更新してからScilabのグラフィックができなくなりました。

NVIDIA GeForce 9300 / nForce 730i
CUDA 3.2RC
Scilab5.3.0b4

Scilabのグラフィック関係のコマンドを実行すると

NVOGLDC invisible: WScilex.exe - アプリケーションエラー
"0x69b81edf" の命令が "0x000000a0" のメモリを参照しました。
メモリが "read" になることはできませんでした。

となり、Scilabが終了してしまいます。CUDA 3.2RCが原因なのかな。

>>708
自己レス

CUDA 3.2RC2 にしてもダメだったので
CUDA 3.1 に戻したら直りました。

CUDA対応ZIPパスワード解析ソフト作っておくれ

>>710
無理です。あれは並列化する箇所が殆ど無い

ただのxor暗号でしょうが。
ブルートフォース並列化できるでしょ。

ちょっと検索したらRAR用はあったからZIP用も可能だろうな

>>708
あの遅いScilab使ってる人いるのかｗ
大学生の時、研究室が貧乏でMatlabが無かったから、
泣く泣くつかってたよｗ

つい先日、CUDAについて勉強し始めたものです
前スレにもあるように、現在でもvisual C++ express editionでは
64bitの開発はできないのでしょうか

クリーンがスキップされてしまい困っています
解決策があるようならお願いします
環境は
windows7 64bitです

可能
VC++2008SP1ならWindows SDK 7を入れた上で以下の小工夫を、
Visual C++ 2008 Express Edition And 64-Bit Targets ≪ Jens’ Blog
ttp://jenshuebel.wordpress.com/2009/02/12/visual-c-2008-express-edition-and-64-bit-targets/
VC++2010ならWindows SDK 7.1を入れて
オプション"プラットフォーム ツールセット"を"Windows7.1 SDK"にすれば開発可能になる

但し、VC++2010でCUDA開発が可能かどうかは知らん

compute capability2.0以降のものをFermiと呼ぶのでしょうか?

compute capabilityが同じならcudaとしての性能は基本的にまったく同じになるのでしょうか？

いつまでRCやってるんだ。

compute capabilityは計算機能の互換性があるということ
Fermiってチップのコードネームでしょ？

そうだね
CC 2.0ならFermiだけど
FermiはCC2.0以上の物も出てる

ところでGF GTX580として512sp版が出るな・・・Tesla版ﾏﾀﾞｧ?(・∀・ )っ/凵⌒☆ﾁﾝﾁﾝ

たとえばGTX480,470,465は同じフェルミチップを使ってて
おなじ計算機能(doubleがつかえるとかFMAが使えるとか)は使えるけど
性能はコア数とかメモリバンド幅がちがうせいで性能は違う。

それでフェルミのcompute capabilityは2.0と定義されているだけ。
GTX460では2.1になってる。CUDAのなんとかガイドのappendixに各型番のcompute capabilityが載ってるから
それみるといい。

>>719-721
ありがとうございます

CUDA C Programming Guide 3.2読んできました

計算機能の他にSM内のCUDAコアの数やレジスタ、シェアードメモリはCC依存
SM数、メモリバンド幅、グローバルメモリ量はチップ依存
なんですね

Fermiだと何が違うか＝CC2.xだと何が違うか
と思っておkですか

実際にはCCの番号はチップの型番への結びつきが強い

CC2.0はcudacore32個で１SM、

cc2.1はcudacore48個で１SMの実装がなされていて
スレッドの生成を３２単位で行うか４８単位で行うかとかの実装での最適化の部分が変わってくる

あとチップによって同時発行可能な命令数とかが変わるから実際のプログラミングでは
同じフェルミ系でも４８０と４６０では変わってくる

積算よりも除算のほうがやたらコスト高いという記事をどっかで目にしたんだけど
A/B*C/D*E/Fって計算があったら
(A*C*E)/(B*D*F)として除算を一回にした方が良いのでしょうか

むしろGPUに限らず当然やるべきことですか
それともコンパイラが勝手に最適化してくれるのでしょうか

GPUに限らず一般論として除算のコストは乗算よりはるかに高い。

そもそもCUDAに割り算命令はない。SFUで逆数を求めてそれを乗算する。

助産が遅いってのはGPUに限らずコンピュータの基本
出来るだけ減らすこと
だから(A*C*E)/(B*D*F)とすることは良いこと

かぶった

>>726
整数型の場合は結果が変わらない範囲での最適化は行なうが、>726の場合は結果が変わるから無理。
実数型の場合は演算の順序が変わるだけで結果が変わるから殆ど最適化できない。
従って、精度や値域を含めてプログラマが自分の責任で最適化するしかない。

闇雲に>728を信じて書き換えると、A*C*EやB*D*Fがオーバフローするかもしれないなw

それはいえてる

>>727-732
ありがとうございます
オーバーフローするほど大きな値はなさそうなので
順番変えても計算結果がほとんど変わらなかったらやってみます
速度的には雀の涙で、もっと根本的なアルゴリズム改善やスレッド管理のほうが必要なんだろうけど

2^n倍するときシフト演算の方が良いらしいとかパソコンでの数値計算における高速化の基礎みたいなんがぜんぜんない

そんなこと無いと思うけどな。
ループの外に出せる演算は外に出すとか、メモリ演算は避けて一時変数を使うとか、
CPU演算でも共通に使えるノウハウは色色ある。

>>734

いや、そういった知識を>>733が持っていないと言っているだけかと。

そしてそんなノウハウを解説しているサイトでもあれば誰か紹介してくれるかなと思って
失礼しました

私も教えてほいいです。

私がまいけるほいいです。

CPUとGPUじゃちと違うんでね？
CPUだとパイプラインとかキャッシュとかが重要になってくるんだけど・・・

>734は共通に使えるノウハウがあるといっているんだ。
共通に使えないノウハウだって勿論山ほどあるだろうさ。

それはさて、そういったサイト(或いはノウハウそのもの)を最適化スレである程度まとめてないのかな?
CPUのノウハウならx86スレがあったと思うのだが。
GPUというか、CUDAのノウハウはどうだろね。系統だって纏められたものはないような気がする。

じゃ〜どっちが知りたいんだい？
1．GPU高速化ノウハウ
2..汎用的高速化ノウハウ

CUDAのカーネル内での高速化が知りたいです

しかしCUDA本って全然出ないね。注目されてないんでなかろうか

本を書ける(or翻訳できる)日本人技術者が捕まらないと思われw
捕まえる気もないのかもしれないが。

で、カーネル内での高速化というと、演算云々もそうだけどメモリアクセスがメインじゃないのかなぁ。
なんか叩き台があれば、この辺をこうしてみたらとかなんとか指摘できる人はここにも何人かいるだろうけど。

>>743
本なんかでないのは programing guide で足りるからでしょ。

俺はLarrabeeのSIMD命令とかパイプラインのブロックダイアグラム見てから
CUDAのハード構造とかパフォーマンスの制約がよく理解できた

スカラ操作が複数あるように見せてるけど実際は32並列のSIMD演算を隠蔽したものなんだ

実はSSEを勉強すればCUDAも自然に覚えられるよ

1warpが1スレッドになったら起こしてくれ。

>>746を起こさないでくれ、死ぬほど疲れているんだ

＞＞742
くだのパフォーマンスはこれが一番でね？
http://developer.download.nvidia.com/compute/cuda/3_2/toolkit/docs/CUDA_C_Best_Practices_Guide.pdf
英語は×となしね

スタティックライブラリの作成方法知りたい奴いる？
-lcudartの指定が必要な

CUDAのバージョンが上がった時に死にそうなのであんまり考えたことなかったなぁ。

既存のSIMDプログラムを一部CUDAで代替しようと思ってスタティックを調べてたんだけど
それだと都度グローバルメモリにコピーが必要なわけであまり効果ないかもしれん

全部CUDAに書き換えるのが一番いいんだろうな

カーネル起動のコストを考えると、最低数千、できれば数百万のスレッド・ブロック
がないと意味がないだろ。SSEとは別次元。

580キタ！

>>745
4並列SIMD x 8 SP じゃなくて？？

8並列SIMD演算ユニット×4クロック同じ操作、あるいは16並列の2クロック

だから論理的には32並列のSIMD

>>755
その考えだとマルチプロセッサ内の8基のSPは何する人たちになるの？？
単純に積和演算系のSIMD命令に4クロックかかってるだけだと思ってたけど。

ひとつのSPあたり同じ命令を4回（あるいは2回）繰り返すんだよ。

Pentium III時代にSSEユニットが64ビットで、128ビットのSSE命令を2サイクルかけて処理してたのと同じ理屈。

それはfloatを1クロックで実行できるってことから考えにくいんだが。

doubleはそういう仕組みになってるようだが。

>>758
非fermi
1warp 32スレッド、SPが8個
4回まわす必要がある。

>>760
fermi以降しか興味なかったから、280のころのブロックダイアグラムみてみたら
確かにそうなってそうだ。

fermi
2warp 64スレッドをSP32個で実行
2回まわす必要がある。

>>762
いやいやいや、そこは同じ32スレッドで考えようよ
そうすればfermiなら1回ですむってなるじゃん。

ところで1SM 48coreな460とかでも1warp32スレッドなのかな。
1warp48スレッドであってほしいけど

>>763
それ違うから。
2サイクルかけて1ワープを処理する16並列のSIMDエンジンが2クラスタ分
ブロックダイアグラムよくみてごらん。

> ところで1SM 48coreな460とかでも1warp32スレッドなのかな。
> 1warp48スレッドであってほしいけど

ありえねえｗｗ
強いて言えば16SP×3クラスタとかじゃね？

>>764
ブロックダイアグラムだけじゃわからなかったよ。

ttp://pc.watch.impress.co.jp/img/pcw/docs/397/079/html/06.jpg.html
これみて理解できた。前に見たはずなのに、ずっと32スレッドを１つのエンジンで1クロックで実行できるものだと
思ってた。

>ありえねえｗｗ
>強いて言えば16SP×3クラスタとかじゃね？
考えてみれば下位製品なのに上位より多いはずないか。
ttp://pc.watch.impress.co.jp/img/pcw/docs/397/079/html/09.jpg.html
すげえ、あってる。

あーでもGF104はestimatedってかいてあるな

並列リダクションについて質問です

足し合わせたい要素が2＾nではない時（というか普通はそう？）、
どうしたら効率よく計算できるでしょうか

現在、要素数より大きい最小の2＾nを計算してその数だけスレッドを作成し、元の要素数を引き数にして
カーネル内ifでスレッド番号が要素数以上だと足し合わせる要素を0にする
とやっているのですが、効率悪いですか？

ハードが持っているSMの数を考えて
warpのメモリアクセスレイテンシが隠蔽できる程度の
warp数になるように調整して生成するスレッド数を決める。
適当に余りを調整しつつ要素数/スレッド数の数だけ
各スレッドで独立にリダクションした後、
ブロック内でリダクションかけて
生成したブロックの数まで要素数を削減する。
あとはCPUで処理するなり単一groupで処理するなりして
最終的に1つにする。

と言うのがリダクションのチュートリアルの結論

用語がおかしかった。
ブロック->groupね。
基本的に全てのgroupがハードに乗り切る程度の
スレッドしか立てなくていい。
で、まずSM毎に幾つのwarpがあればメモリアクセスレイテンシを
隠蔽できるかを考えてwarp数を調整する。
同じwarp数になるならSMに割り振られるgroup数を増やすより
groupのサイズを大きくした方が良い。
どうせメモリ帯域がネックになるし、group内の同期は最後だけだから
同期コストは気にならない。

MPIにはMPI_Reduceがあるから、これくらいシステムで用意して欲しかったね
ユーザーにここまでコーディングさせるのはひどい話だな

おれはリデュースとかいらんし（´・ω・｀）
しかもそれってライブラリが無いだけだろ？誰か（お前が）作れよ

Thrust : CUDA用STLライブラリ: 新 masafumi's Diary
http://masafumi.cocolog-nifty.com/masafumis_diary/2010/11/thrust-cudastl-.html

テンプレートベースのプログラミングって変態チックで惹かれるのう。
IntelのCｔは当初はC++の上位言語の予定が結局Array Building Brocksって名前になって
C++の「テンプレートライブラリ」に収まっちまった。
こっちはreduceは何も意識しなくていいぞ。最初からAPIレベルでサポートされてるから。


CUDAもカーネルベースのコーディングスタイルだけでなく32並列のSIMDを意識したベクタベースの
コーディングスタイルも用意して欲しいなぁ

関係ないけど、sandraでGPUによる暗号AES-EBCのベンチができるようになってるな
ttp://mousou01.blog.ocn.ne.jp/blog/2010/11/sisoftsandra-20.html

GPUで暗号アクセラレーションなんて・・・と思ったら案の定EBCか

他に比べて並列化しやすいんだっけ課？

勉強開始三日目
上のほうにあったがreduction分からんな
取説には実装は容易だが理解は難いとある
考えずに真似して実装しちゃえばおｋ？

>>777
ヒント：折り紙

本当にリダクションが必要かを考えてみるのもよい。
場合によっては、リダクションが必要な部分自体を1スレッドにしたほうが速くなる。

>>778-779
サンクス
ま、実のところリダクションってより共有メモリ、ワープ、ブロック等の関係把握が一苦労
論理的概念なのか物理的概念なのかとかもう訳分からん
しかもcudaのバージョンアップが急速なんで古い資料だともう使えなかったり
キャッチアップが大変な世界だな

グラボの中には描画に問題なくてもＣＵＤＡで数値計算すると以上な値をかえすコアが含まれるものもあるらしいと
聞くんですが、ボードの検査プログラムとか何か定番のものや、公開されているものはあるんでしょうか
ご存じの方いたらリンク等教えていただけないでしょうか。

初期の実装は整数演算を浮動小数点演算で肩代わりしてたから、精度が24bitしかない。というのに引っかかってるオチじゃない？
Linux専用でcuda用のmemtestみたいのがsourceforgeにあったと思うよ

それってHPC屋がTeslaC20x0売るためのFUDだよ
知り合いの先生が言ってたから、「じゃあECCの付いてないC10x0を売ったのは詐欺ですよね？」
って答えとけって返しといたｗ

長崎大の方が確か１割り位つかえないのがあり、
自作のチェックプログラムで弾いていると発表されて
ませんでしたっけ。
自分は５枚位で外れにはあたってないですが。

>>782
ありがとうございます。探してみます

>>783,782
長崎大の濱田先生の資料にもありましたし、CUDAZoneに投稿している方の資料にも
そういう記述がありましたので、現実にあると思います。

linuxの環境で、サンプルの中をのぞくとcppフィルとcuファイルでそれぞれmainを持ったソースがあるんだが
ﾊﾞｱちゃんからは、プログラム中にmainは一つだけじゃぞ、と教わってきただけにこの辺がよくわかりません
cudaのコンパイル関係の資料ってないですかね？（cpuのコンパイルの知識もあまりないんですけどね）

CPUのみとGPU使用のプログラムってだけだろ

それを一つの実行ファイルにする技を知らんのですよ
mainが二つあるオブジェをリンクできちゃうってこと？
makeファイルを読めんもんでよく分からん。。。

グラボ買ってきたんですが、これをCUDA専用のアクセラレータとして使って
画面描画には使わない場合（つまり別のグラボorオンボにつないで、買ってきたグラボにはディスプレイ繋がない）
ドライバ（通常インストールするドライバ）は入れずに、CUDAドライバ、toolkit,sdkだけをインストールすれば良いのでしょうか
それとも通常のドライバも必要としますか？

OSも書かずに

>>790
すみません、winXP 32bitです。

ドライバいらないわけないだろ

あ、何か勘違いしているようだが
CUDAドライバってのは「CUDAに対応しているバージョンのVGAドライバ」って意味だからな
VGAドライバと別物なわけじゃない

>>793
ありがとうございます。助かりました。
CUDAドライバ、toolkit、SDK

あ、途中で、、、
CUDAドライバ、toolkit、SDKのみのインストールでやってみようと思います。
CUDAドライバがcudaに対応させるためのadd-onなのか、ボード制御をすべて内包しているのかが不明
だったので、助かりました。

あるカーネルを1ステップに1回動かすということを数100ステップ行いたいのですが、
数ステップ進んだところでカーネルが途中で止まってしまい、エラーも出さずに戻ってこなくなります。
カーネル起動後にcudaGetLastError()がcudaSuccessでないときにエラーメッセージを表示させて止まるようにしてあるので
カーネル起動中に止まっているのだと思います。
最初の数ステップは問題なく動いているのでわけがわかりません。

どんなことが原因と考えられますか？
cudaMallocとcudaFreeは毎ステップ行わずに最初と最後のステップでだけ行うようにしています。
cudaMemcpyはカーネルの起動前と起動後に毎ステップ行っています

それカーネルにエラーがあって暴走してるんじゃね？
暴走させるとWindows巻き込んで死ぬで

>>796
素人意見だけど
毎ステップ後の出力データが期待通りの値になってるか確認。
もし間違っていたならカーネルかcudaMemcpyのプログラムのどちらか
もしくは両方が間違っている可能性あり。

デバッグの際、カーネルの途中で値の確認をしたい場合は
cuda-gdbデバッガなどを駆使する。

今の時代はNsightなんてものがあるんだから楽でいいよな。

>>799
Linux派としては歯がゆい

Windowsには当初デバッガ無かったわけだし・・・

>>800
俺もこてこてのLinux派だったけど、Nsightでとりあえず乗り換えた。
次元が違う。

ただ、Windowsはリモート関係が面倒だな。

Ｎｓｉｇｈｔいれようとしたら、vista sp3 or win7が要求されててwinXPの自分には無理だった
机上デバッグでやってやるぜ！

>>797-798
ちゃんと動いている時は出力結果は期待のものが得られました
(128*1*1)*(16464*43)でやりたいのですが、unspecified launch failureになりました
(128*1*1)*(2*43)にすると問題なく最後まで動きました
(128*1*1)*(50*43)数ステップ動いて途中で帰ってこなくなりました
最初はnspecified launch failureで調べてたんですけれど、ブッロク数を減らせば実行できたり
途中で止まったりとわからなくなったので質問しました
unspecified launch failureだったらどんなことが考えられますか？

原因不明のエラーメッセージなんじゃね？
さすがに何のコードかも分からないのにコメントできんわｗ

大抵は確保したグローバルメモリの範囲外にアクセスしている場合。

Nsightの1.5に付いてるToolkit 3.2ってのは3.2のRC版ですよね？

>>804
その数字がblock数*thread数ということなら、thread数が大き過ぎ。
今手元に資料がないけど、thread数の上限はx*y*zが512だったと思う。一番目は2バイト整数の範囲を越えてたまたまマイナス扱いになったからfailureになったけど、三番目はそうじゃないから帰って来ないと思われ。

いや、
(スレッド x * y * z)*(ブロック x * y)だろｊｋ・・・
GriddimにZなんてないし

>>808
今手元に資料がないけど、fermiはthread数の上限はx*y*zが1024だったと思う。

sharedMemoryを使う必要がない場合でレジスタ数が多くて困ってるときって
smemをレジスタとして使うって、あり？

>>811
試してみればいいじゃん。

レジスタでaaaって変数を定義してたならシェアードでは
__shared__ aaa[スレッド数];
ってやってスレッド分の配列作って別々の場所にアクセスすればＯＫじゃね？

まぁそれで実際にPTXの使用レジスタ数が減るかは分からんわ〜。

>>804
ブロック数がプログラムにどう影響するかよく見直してみよう。

ブロック数に比例して計算規模がでかくなるならメモリ不足とか？
数ステップで止まるなら、止まる直前までのデータを取って
異常な値を示していれば、そこからミスってる箇所を推測できるかも

まぁコード無しじゃ具体的なアドバイスはできんわ。頑張れｗｗ

いまさらだけど>>410-411ってどういう意味？

データの並びの話でしょ。
--
__global__ void kernel(int * foo, int * bar, int n)
{
for (int ic = 0; ic < n; ++ic) {
const int offset = threadIdx.x * blockIdx.x + ic;
foo[offset] = bar[offset];
}
}
--
とするか
--
__global__ void kernel(int * foo, int * bar, int n)
{
for (int ic = 0; ic < n; ++ic) {
const int offset = ic * n + threadIdx.x;
foo[offset] = bar[offset];
}
}
--
とするかの違いじゃないかと。

>>814
メモリアクセス図だろ
あの図を見て何も閃かないってことはCUDA初心者だな？

コアレスアクセスについて勉強してきましょう

Amazon EC2(1時間当たり2ドル10セント)
http://itpro.nikkeibp.co.jp/article/NEWS/20101115/354156/

IBM(月160万円)
http://itpro.nikkeibp.co.jp/article/NEWS/20101102/353730/

グラフィックボード入れたLinuxが数ヶ月起動しっぱだったので再起動したところ問題なく動くようになりました。

cudaFreeせずに毎回強制終了させてたので変なデータが残っちゃったのかな

C:\Documents and Settings\All Users\Application Data\NVIDIA Corporation\NVIDIA GPU Computing SDK\C\bin\win32\Release>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 580]
16384 bodies, total time for 10 iterations: 69.775 ms
= 38.472 billion interactions per second
= 769.433 single-precision GFLOP/s at 20 flops per interaction

>>189と比べて妥当ですか？

GTX580ベンチｷﾀ━━━━━ヽ（｀・ω・´）ﾉ━━━━━！！！
スペック的には妥当な性能うｐ具合なんじゃ？

580は本気出した480だからなぁ…　いまいちインパクトが薄い

やっと正式版
CUDA Toolkit 3.2 (November 2010)
http://developer.nvidia.com/object/cuda_3_2_downloads.html

総和計算で質問なのですが、同じ計算なのにやるたびに結果がかわってしまいます。
GPUを使わずに計算させた場合と比較すると、あってる時とあってない時があります。(プログラムは正しい？？)

要素数がスレッド数の上限を超えるため、カーネル関数を2段階に分けて計算しているのですが、
どちらのカーネルが原因なのか調べるために、1段階目を複数ブロックで計算した後、
各ブロック内の総和結果を0スレッド目が一時保存のためにグローバルメモリに書き込む直前に
printfで各ブロックの結果を出力させて、その合計を自分で取って最終結果と比べてみようと思ったのですが、
printf文を入れると必ず最終計算結果が正しくなってしまいます。
なおったと思いptintfをコメントアウトするとまた計算結果が毎回同じ値にならなくなります。

printf文を入れると何回やっても結果が正しくなり、原因が見つけられません
見ていない時だけイタズラされているみたいで非常に腹立たしいのですが
なにが原因か教えていただけませんでしょうか

桁落ちじゃね？
情報工学の初歩だからね

>>823
たぶん1回目と2回目の計算のうちいくらかを並列でやって、負荷がかかりすぎてカーネルが途中で終了してしまっているんだろう。
http://goo.gl/6RjqX
だから途中に同期が必要なprintfを入れると、GPUが一旦休憩してちゃんと動作してくれる。
カーネルをさらに細かく分割して、呼び出しを複数回に分けてさらに間に同期命令を挟めばなんとかなるんじゃないかな。

>>823
同期失敗だろ？

ソース出せ。
まぁ、>826に同意。
例えば、次のカーネルは失敗する。
--
__global__ void func(int * foo, int * sum)
{
* sum += foo[threadIdx.x];
}

>>822
待ってました。

そう言えば、pthreadで生成したスレッドで確保したデバイスメモリって
生成もとのthreadからは解放できないのね。
その逆に、デバイスメモリを確保してからpthreadで生成したスレッドで
そのメモリが使えるかどうか調べてみなくちゃだわ。

>>829
Driver APIというのを調べれば色々解決する。
と同時に、面倒くさい話も出てくる。

>>830
なるほど、そのためのDriverAPIか。確かサンプルもあったね。調べてみますわ。
情報THX!

>>824
桁落ちは実行するたびにおこったりおこらなかったりするのでしょうか
>>825
linuxなのですがやはりタイムアウトしたりするのでしょうか。
成功した時も失敗した時も1秒に満たない実行時間でした
>>826
>>827
リダクション時に他のスレッド担当のshared memoryにアクセスする前には必ず__syncthreads()
しています。
これまで同じ書き方でリダクションで失敗したことはありませんでした

threadfence、threadfence_blockの使い方がわからないのですが、

smem[threadIdx.x]=0;
smem[threadIdx.x]+=1;
と連続で書くときには同じスレッドからの書き込みでも
smem[threadIdx.x]=0;
threadfenceblock();
smem[threadIdx.x]+=1;
で値を確定？させなくてはならないのでしょうか

＞smem[threadIdx.x]=0;
＞smem[threadIdx.x]+=1;
大丈夫かどうかは兎も角、むちゃくちゃ遅そう。

そもそもどの程度の違いなんだよ？
数値例を出せよ

>>833
smemってレジスタと同じ速さだと思っていたので。
smem[threadIdx.x]=0;
for(){
smem[threadIdx.x]+=1;
}
のようなコードは
レジスタに足し合わせてから一回だけsmemに代入した方がいいのでしょうか

1.6769037288275420e-16:GPUでうまくいった場合
1.8356045526411666e-16:GPUでダメだった場合

1.6769037288276485e-16:CPUで計算したとき

桁がぜんぜん違ってくるってわけではないですが、有効数字1桁目〜2桁目がかわってきてしまいます

変なメモリにアクセスして、そこに残ってるデータが足されているのかと思いましたが、それだと値が正しくなることがあるのがおかしくなりますし、
やはり同期の問題なのでしょうか
printfを入れたら直るので、同じ場所に__syncthreads入れてprintf消してみましたが直りませんでした

ヒント：x86のCPUは64bit精度と見せかけて80bit精度のレジスタと計算回路になってる。
GPUとCPUの有効桁の差はこの所為じゃね？

x86のCPUは64bit精度と見せかけて80bit精度のレジスタと計算回路になってる。
GPUとCPUの有効桁の差はこの所為じゃね？

>>835
足している値の値域は? かけ離れた値がある場合、足す順序によって値が変わるけど。
例えば、1e-33の値が10000個と1e-16があった場合、小さい方から足すと1.0000000000001e-16になるけど
大きい方から足すと1e-16になる。バイナリに足す場合は恐らくその中間になるだろうし。
同じ理屈で、負の値があると見掛け上精度はどんどん劣化するね。

実行する度に結果が違うってリダクション演算間違ってるんだろ
とりあえずその部分のコード晒せや

>>837
80ビット使うのはx87のときだけだし倍精度の有効制度は53ビットだからそこまで変わってきようがない
たぶん単精度だろ

Project Echelon: The Future of NVIDIA's Silicon - Bright Side Of News*
http://www.brightsideofnews.com/news/2010/11/19/project-echelon-the-future-of-nvidias-silicon.aspx

ｇｔｓ250からgtx465に変更したら作成プログラムが遅くなった
どうやらカーネル起動コストの問題のようなんだがなんか情報或る？

>>840
いや、x86で浮動少数の演算といったらSSEでも使わない限りx87だから。
それに単精度じゃそこまで精度出ないし。
大きい数値と小さい数値で何度も何度も演算してるから余計に精度で問題が発生してるんだろ。

環境によるろ
intelコンパイラは64ビットの場合基本すべてsse演算になる

>>835
毎回値が変わるのは大体が同期ミス。
>>838の言うような場合だと、実行するたびに結果が変わるってことはないはず。

１段階目と２段階目のカーネルでブロック間の同期が怪しい。
一応cudaThreadSynchronize();を１個めと２個目のカーネルの間に入れてみては？

各要素の値を例えば全部1.0にして計算すれば、
桁落ちの影響かどうか分かるのでは？

それと1段階目だけ実行したときは答えは合ってるの？

ていうかリダクションよりカーネル内でループした方が早くね？

ていうか、(メモリ転送コストを隠蔽或いは無視できるなら)CPUにやらせた方が速くね?

アフォなの？

あんたが?

メモリ転送コストを隠蔽或いは無視できないならCPUにやらせた方が速い
ってことはあるけどな
>>848はなにかすごい発明でもしようとしているのだろうか？

あれ? 既にGPU側にあるデータの話じゃないの?
そもそもCPU側にあるならわざわざGPUでやろうとしないと思ったんだが。

「リダクション」の意味をわかってねーな

GPUのメモリ帯域＞＞＞＞＞＞＞＞＞＞＞＞PCIeの帯域

そもそも「同期」が必要なのが理解できないんだけどな。
物理的には同じ1ベクトル32要素のSIMDレジスタに載っかってるわけでしょ。

わざわざShared Memoryに保存して読み出してを繰り返すくらいなら
水平演算くらいサポートすればいいのに。

kernelベースのプログラミングモデル否定することになるけど。

同期しないと二重に計上する可能性があるだろ

メモリを経由するからでしょ。
レジスタ上で要素シャッフルできるならタイミングがずれる心配が無い。

カーネル内の除算ってみんな__fdividef(x,y)を使ってるのん？？

普通コンパイラに書き換えてもらう

あれ。
-use_fast_mathってx/yも勝手に__fdividef(x,y)に書き換えてくれるのか
sinf(x) -> __sinf(x)
とかだけかと思ってた。問答無用で可能なもの全部書き換えるのね。

不安だったらPTXを読もう。

俺は超越関数使ってないのに、use_fast_mathしたら速くなって、
何でだろうと思ったら除算が書き換えられてた。

一部掛け算も置き換えられなかったっけ? もしかしたら、use_fast_mathがなくても置き換わったかもしれない。

>>854
同一warp内だけのデータ交換なら
同期は必要ないだろ。

>>842
うちの環境だと同期とると遅くなるという現象が起きた。

Thrustライブラリのmax_elementを使用すれば配列の最大値は取得できたのですが
知りたいのはその位置(何番目が最大値か)で、効率的に最大値の位置だけを取得できればと思い思案中です。

CUBLASの中に配列を渡すと最大値の位置を返すcublasIdamax()という関数があるようですが
使い方がいまいちわかりません。どうかご教授いただけないでしょうか?もしくは自分で作成する場合の例を示していただきたいです。

>>864
総和を求めるプログラムを応用すれば自分で作成できると思う。

総和では値を足していくけど、
でかい方の数を代入するようにすれば(←ここはif文使うけど仕方ない？)
最後にはsmem1[0]に最大値が入る。

んで最大値求める配列とは別に、最大値の位置保持用の配列を用意して
上記のでかい方の数を代入する操作にあわせて、でかい数の方の"位置"を
代入してくようにすれば、こっちはsmem2[0]に最大値の位置が入るはず。

んまぁ、どっか間違ってるかもしれんし最速のアルゴリズムでもないかも
しれんけど参考程度に

それで位置がわかるのか？

>>866
頭鈍いな。
smem1は元配列の"実際の値"が入っている（例：[4,1,0,5,9,6,2,7]）
smem2は元配列の"配列番号"が入っている（例：[0,1,2,3,4,5,6,7]）

smem2はsmem1の代入操作とリンクさせる。
（例：smem1[3] = smem1[7]が実行される場合、同様にsmem2[3]=smem2[7]とする。）

この点を踏まえて、smem1[0]に最大値が入るよう操作したなら
smem2[0]には元配列の最大値がある位置が入るだろうってこと。

一言余計なやつだなぁ

sincosf(x,*s,*c)ってなあに
べつべつにやるより速いのん？？

>>868
intelCompilerだと実際に同時に求める関数が用意されていますが、
cudaの場合はptx出力を見れば判るように別々のニモニックに展開されます。
まぁ、無理して使わんでも。

>>868って自分が初心者じゃなくただのゆとりってことに気づいてないのかな？

>>865
解説ありがとうございます。
自分で作成も考えたのですが速度的にライブラリに勝てるだけの処理が書けないと思ったのでライブラリについて質問しました。

2つ目の配列要らなくないか？
1つ目の配列へのアクセスの添え字じゃだめなの？

>>872
いるとおもうー
1つ目の配列へのアクセスの添え字ってのはスレッド番号だから
最終的にはsmem1[0]に最大値が格納されるけど入っているのは値だけであって
その値が元々どの添字にあったかの情報はないのよ

最大値を元の配列から探した方が速いと思うが・・・
メモリのアクセス量が大して変わらなくても半分はシーケンシャルに済ませられるだろうし

>>874
元のreductionをきちんと最適化していたら
大きなNに対しては>>865の方が確実に速い。
globalメモリへのアクセス量がほぼ1/2になるから。

Fermiだと一所懸命smemとかtextureとか使っても、
グローバルメモリに比べてあんまり速くならないよね。
キャッシュ様様

逆にsmemとかtextureとか使って
きちんと速くなる用途だと、fermiにする意味が殆どない。
コストパフォーマンス的にも電力効率的にも。

>fermiにする意味
プログラミングが楽

textureのキャッシュヒットって具体的にどうゆうときになるか知ってる人いますか？？

空間的に近くを参照した時
~~~~~~~~~~

>>880
具体的にw
~~~~~~~~~

2次元で確保し、(5,5)を参照したとしたら
その8近傍をキャッシュしてくれるかもしれない。
　　　　　　　　　　　　　　　　~~~~~~~~~~~~~~~~~

1Dtextureで配列読ませてリードキャッシュに使ってるお
CUDAだとグローバルメモリがいつ書き換えられるかわからないからキャッシュしない
→textureにbindして書き換えが無いことを保証させて２回目のアクセスからキャッシュさせる

だと認識しているが・・・

>>883
なんで1Dと2Dがあるか考えたことないの？

あと
>CUDAだとグローバルメモリがいつ書き換えられるかわからないからキャッシュしない
つFermi

ごめん、>>883の言ってることに間違いはない。
Fermiの部分を除いて。

え…グローバルメモリって勝手に書き換わって信用ならないってこと?？

(Fermi以前だと、)グローバルメモリをキャッシュしても
他のSMが書き換えるかもしれないでしょ。

Fermiだと書き換えないのかい?

自動キャッシュされるが、グローバルメモリを書き換えるときにキャッシュクリアするようになってるんじゃない？
だからキャッシュにある内容は保証されるっていう

L1キャッシュは保証してないが・・・

>>890
そうなんか、Fermi持ってないからそのへん調べてないんだ（´・ω・｀）
誰かが発表で「Fermiはチューンしなくても速いんだぜ！」って言ってた記憶しかないわ・・・

ちょっとNVIDIAのサイト見ればすぐに分かることなのに

こんなとこで質問したり、間違った知識書き込んでる奴って大丈夫か？

プロファイラの評価項目で、gld/gst efficiencyは比較的簡単に１近くまで持っていけるけど
instruction throughputは毎度0.8〜0.9くらいで頭打ち
何処に無駄が隠れているのかしら。スレッド間同期かね。
それとももともとこんなものなのか

総和を求めるReductionのソースは以下のようなものでよいのでしょうか?
一応CUDAのチュートリアルに記載されていたものです。

for (int stride = numThreadsPerBlock / 2;
stride > 0; stride /= 2)
{
__syncthreads();
if (threadID < stride)
sresult[threadID] += sresult[threadID + stride];
}

>>894
だからシフトを使えとあれほど酸っぱく言ってるのに。。。
あとループの先頭でじゃなくて、後半で同期。
それと1warp。。。。。

>>894
CUDAテクニカルトレーニングが参考になるよ。
リダクションはvol2に図入りで解説してる。
特に初心者にとっては勉強になると思う。

>>895

>>894じゃないけど
for (int stride = numThreadsPerBlock>>1;
stride > 0; stride >>= 1)
{
if (threadID < stride)
sresult[threadID] += sresult[threadID + stride];
__syncthreads();
}

こういうこと？
最後の一行はどういう意味？
開幕しょっぱなスレッドの半分遊んでるとか
最後の1warpでは同期処理いらないからそれ以前と分けるとか言う話？

テクニカルトレーニングを読んだところfor文を使わずに
unsigned int tid = threadIdx.x;
unsigned int i = blockIdx.x*(blockSize*2) + tid;
unsigned int gridSize = blockSize*2*gridDim.x;
sdata[tid] = 0;

do { sdata[tid] += g_idata[i] + g_idata[i+blockSize]; i += gridSize; } while (i < n);
　__syncthreads();
　if (blockSize >= 512) { if (tid < 256) { sdata[tid] += sdata[tid + 256]; } __syncthreads(); }
　if (blockSize >= 256) { if (tid < 128) { sdata[tid] += sdata[tid + 128]; } __syncthreads(); }
　if (blockSize >= 128) { if (tid < 64) { sdata[tid] += sdata[tid + 64]; } __syncthreads(); }
　if (tid < 32) {
　　if (blockSize >= 64) sdata[tid] += sdata[tid + 32];
　　if (blockSize >= 32) sdata[tid] += sdata[tid + 16];
　　略
　}
　if (tid == 0) g_odata[blockIdx.x] = sdata[0];
｝
のように計算しているものが最適化されたもののようでした。まだ理解仕切れてないのですが
最大値とそのインデックスを求めるならば加算部分をif文で以下のように分岐させ、インデックスを保存する配列を追加すれば良いということでしょうか

if(sdata[tid] < sdata[tid + 32]){
　sdata[tid] = sdata[tid + 32];
}

>>898
一から十まで教えてもらわないと駄目なのか？
未熟ながらも理解してきてるんならこれくらい自分で
試行錯誤して完成させてみろよ