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

1 名前：デフォルトの名無しさん mailto:sage [2009/10/08(木) 19:29:37 ] このスレッドは、他のスレッドでは書き込めない超低レベル、 もしくは質問者自身何が何だが分からない質問を勇気を持って書き込むスレッドです。 CUDA使いが優しくコメントを返しますが、 お礼はCUDAの布教と初心者の救済をお願いします。 CUDA・HomePage www.nvidia.com/cuda 関連スレ GPUで汎用コンピューティングを行うスレ pc11.2ch.net/test/read.cgi/tech/1167989627/ GPGPU#3 pc12.2ch.net/test/read.cgi/tech/1237630694/ 352 名前：デフォルトの名無しさん mailto:sage [2009/12/10(木) 19:27:33 ] >質問1. ブロックあたりのスレッド数を512にすると計算結果が正しくなりません。256なら正しいです。なぜでしょう？ カーネルコールの後でエラーキャッチのコード入れてる？エラー出てない？ 353 名前：デフォルトの名無しさん mailto:sage [2009/12/10(木) 21:05:48 ] 質問の傾向をみると、 CUDAって、普通のCのように自由度高く書けるけど、 実際は、サンプルソースからは読み取れない あれは駄目、こうしたら駄目、的な制約が多すぎるのかな？ DXAPI+HLSLのような、自由度の少ない環境の方が むしろ、良質なソースが書けるのかも。 354 名前：デフォルトの名無しさん mailto:sage [2009/12/10(木) 21:54:43 ] 参考資料 Device 0: "GeForce GTX 295" CUDA Driver Version: 3.0 CUDA Runtime Version: 2.30 CUDA Capability Major revision number: 1 CUDA Capability Minor revision number: 3 Total amount of global memory: 939524096 bytes Number of multiprocessors: 30 Number of cores: 240 Total amount of constant memory: 65536 bytes Total amount of shared memory per block: 16384 bytes Total number of registers available per block: 16384 Warp size: 32 Maximum number of threads per block: 512 Maximum sizes of each dimension of a block: 512 x 512 x 64 Maximum sizes of each dimension of a grid: 65535 x 65535 x 1 Maximum memory pitch: 262144 bytes Texture alignment: 256 bytes Clock rate: 1.24 GHz Concurrent copy and execution: Yes Run time limit on kernels: No Integrated: No Support host page-locked memory mapping: Yes Compute mode: Default (multiple host threads can use this device simultaneously) 355 名前：デフォルトの名無しさん mailto:sage [2009/12/10(木) 22:17:49 ] >>353 CUDAをただの言語だと思っていてGPUのアーキテクチャを理解していないからかと。 後半については同意。限られた記述で実装可能なアプリしか書かないからというのがより正確かも。 356 名前：343 mailto:sage [2009/12/11(金) 00:35:49 ] >>351 スレッド数512、デバイスメモリ使用でやってみましたが、やはり正しくありませんでした。 Emuではちゃんと動くので、原因はよく分かりません。 コンスタントも使って多のでそっちが原因かも知れません。 分からん。もういい。 >>343 の質問2については、multithreading.cppをプロジェクトに追加してないだけでした。 俺死ね。 357 名前：デフォルトの名無しさん mailto:sage [2009/12/11(金) 01:23:14 ] >353 テクスチャはまじで鬼門。 凝った使い方をしようとすると、コンパイル通っても動かない・動いても結果間違ってるがざらにある。 358 名前：デフォルトの名無しさん mailto:sage [2009/12/11(金) 01:42:00 ] >>357 あるある。 deviceemuともまた動作違うこと多いしね。 NEXUS出ればデバッグもうちょっと楽になりそう。 359 名前：デフォルトの名無しさん mailto:sage [2009/12/11(金) 09:51:22 ] >>343 歩留まり向上のため、重い処理をさせると計算が狂うことがあるらしい どこかのPDFで見た 360 名前：デフォルトの名無しさん mailto:sage [2009/12/11(金) 10:57:49 ] >>359 それちょっと表現が微妙に誤解を招くような・・・。 出荷されている状態で、計算がまともじゃないコアがたまにまぎれこんでいる。 グラフィック用途と違って、メーカーの想定より負荷が偏った計算をさせることがあるから、 やや耐性の低いコアの耐性がちょっと足りない部分がエラーを起こす。 高負荷プログラムなんかを使ってそういうコアというか製品をはじいてやれば、あとはほぼ安定稼働する、という話だと思う。 361 名前：デフォルトの名無しさん mailto:sage [2009/12/11(金) 11:17:51 ] はじめてのCUDAがいつまでたっても届かない 一体どこに発注したんだよ上の人 362 名前：デフォルトの名無しさん mailto:sage [2009/12/11(金) 12:42:26 ] >>360 CUDAのスパコンで1/6が不良品だったと言ってたが 363 名前：デフォルトの名無しさん mailto:sage [2009/12/11(金) 12:44:30 ] えっ　Teslaとかでもそうなのか？？ 364 名前：デフォルトの名無しさん mailto:sage [2009/12/11(金) 13:49:03 ] Teslaは選別品 365 名前：デフォルトの名無しさん mailto:sage [2009/12/11(金) 14:52:34 ] 低価格帯で一番安定してるGPUってなに？ 366 名前：デフォルトの名無しさん mailto:sage [2009/12/11(金) 15:53:13 ] slidesha.re/5FtABc のP.26 長崎大の人に言って選別プログラムを貰うことだなｗ 367 名前：デフォルトの名無しさん mailto:sage [2009/12/11(金) 16:12:17 ] CUDAは保障外だからCUDAやりたい奴は動くのが当たるまで買えってことなんだろう 368 名前：デフォルトの名無しさん mailto:sage [2009/12/11(金) 16:17:50 ] 仮に2/3超の確率で完動しないとしても、自前で問題無く選別できるなら Tesla買うよりコンシューマGPU買うほうがまだ安い位だから、たくさん 買う人は自前で選別できるように頑張るのが正解だよなぁ。 369 名前：デフォルトの名無しさん mailto:sage [2009/12/11(金) 16:40:06 ] 多少クロック落とせば動くんじゃないのかな。CPUのオーバークロックと一緒で。 370 名前：デフォルトの名無しさん mailto:sage [2009/12/11(金) 20:54:29 ] >>361 自分は11/27楽天でぽちって 12/01に到着。 ぼちったあとすぐ納期の表示が延びたから焦ったけど滑り込みセーフだったらしい。 書いてあることはわかり易い。しかし意外なことが触れられていなかったりする。 371 名前：デフォルトの名無しさん mailto:sage [2009/12/11(金) 22:49:23 ] >>360 その状態が続く限り、コンシューマ用パッケージソフトじゃGPGPUを使えないね。 「動かない」サポートの爆発でたいへんなことになる。 372 名前：デフォルトの名無しさん mailto:sage [2009/12/11(金) 23:34:51 ] てか>>356の作ってるソフトが選別ソフト代わりになるんじゃね？ 373 名前：デフォルトの名無しさん mailto:sage [2009/12/12(土) 13:58:47 ] おれらまだまだ、「主メモリ側が主役」って固定観念なくね？ いかんな。どかどかデバイスメモリにロード、大量のブロック×スレッドを駆使、 CPUはたまにお手伝いをさせていただく、位に思わないとな。俺まだまだ 374 名前：デフォルトの名無しさん mailto:sage [2009/12/12(土) 15:52:58 ] メモリ転送と計算を非同期で出来るのかね？ 375 名前：デフォルトの名無しさん mailto:sage [2009/12/12(土) 15:58:25 ] >>374 マルチスレッドで、CPU側は「GPU様が計算し終わるまで寝とけや」くらいの扱いにするわけよ 376 名前：デフォルトの名無しさん mailto:sage [2009/12/12(土) 16:43:09 ] >>374 できる。 基本GPUとCPUでは非同期。 ストリームを設定すれば、GPUで演算中にデータ転送させることも可能。 ちゃんと同期させるための関数もある。 377 名前：デフォルトの名無しさん mailto:sage [2009/12/13(日) 15:51:31 ] > NEXUS いつ出るんだっけ VisualStudio2008でいいんだっけ 378 名前：デフォルトの名無しさん [2009/12/15(火) 06:00:58 ] カーネルで計算した結果をCPU側に返すにはどうしたらいいの？ return aのように簡単にはできないのですか？ 379 名前：デフォルトの名無しさん mailto:sage [2009/12/15(火) 06:22:49 ] >>378 ちょｗ カーネルfunctionの引数で　float * outdata とかやってカーネルからこれに書き込む　outdata[i] = result; ホスト側でcudaMemCopyのdeviceToHostで、ホスト側でcudamallocしたfloat * に戻す みたいになる。ややこしいぞ？ 380 名前：デフォルトの名無しさん mailto:sage [2009/12/15(火) 19:31:04 ] thrust便利だなしかし 381 名前：デフォルトの名無しさん mailto:sage [2009/12/16(水) 02:52:17 ] thrust、自分で書いたカーネル関数はどう使うのっすか 382 名前：デフォルトの名無しさん mailto:sage [2009/12/16(水) 15:26:45 ] __shared__ int sh[num*2]; /* numはスレッド数 */ という共有メモリの配列をソートし、最大or最小のデータのみグローバルメモリに書き込みたいのですが これをカーネルの中でやるいい方法を教えてください。 383 名前：デフォルトの名無しさん mailto:sage [2009/12/16(水) 17:36:07 ] ソートは不要で、reductionで半分を繰り返せばいいのでは 384 名前：デフォルトの名無しさん mailto:sage [2009/12/16(水) 23:18:35 ] >>383 言われてみればそうですね。 解決しました。thx 385 名前：デフォルトの名無しさん mailto:sage [2009/12/17(木) 03:27:35 ] どこかにfortranでの使用例って ありませんか 386 名前：デフォルトの名無しさん mailto:sage [2009/12/18(金) 16:42:19 ] そういえば見かけたことねえな 387 名前：デフォルトの名無しさん mailto:sage [2009/12/18(金) 17:17:27 ] CUDAはそもそも実行速度がシビアで 高級言語向けではないと思うけどな 388 名前：デフォルトの名無しさん mailto:sage [2009/12/20(日) 02:21:29 ] 共有メモリは16個のバンクに分割されて、半ワープのスレッドが同じバンクにアクセスしなければ競合は起こらない。 みたいな事の意味がイマイチ分かりません。 例えばブロックの共有メモリが8KBなら0.5KBづつに分けられて、半ワープのスレッドが0.5KBの中ならどこにアクセスしても良いって意味なの？ それとも、先頭から4ﾊﾞｲﾄとか8ﾊﾞｲﾄとかの固定の領域に連続して分けられてるの？ でもそれだと分けられる領域が4ﾊﾞｲﾄなのか8ﾊﾞｲﾄなのか分からないじゃまいかヽ(`Д´)ﾉ カーネルの引数にvoid*とデータサイズを渡して任意のデータ長で処理させたいのですが、 こういう事をしようとすると元からバンク競合は避けられないのですか？教えて、エロイ人(´・ω・`) 389 名前：デフォルトの名無しさん mailto:sage [2009/12/20(日) 02:33:00 ] 今あるMPIのプログラムをGPUに対応させようと考えているんだけど、 CUDAでMPIを使うとき、データの転送は一端ホスト側のメモリに転送してから 他のノードへ転送することになるの？ だとするとレイテンシがすごく大きそうだね。 それとも専用のライブラリなんかあるのかな？ GPU側のグローバルメモリがホスト側のメモリにマップされていれば、 GPUを意識しないで転送ができそうなんだが。 390 名前：デフォルトの名無しさん mailto:sage [2009/12/20(日) 08:02:33 ] >>388 おれの解釈だと、共有メモリには16個のアドレス入力と16個のデータ入出力が有る と思ってるんだけど　違うかも 391 名前：デフォルトの名無しさん mailto:sage [2009/12/20(日) 09:55:19 ] >>388 連続する64バイトの領域（先頭アドレスが４の倍数）について、4バイト(32bit)のバンク１６個に分かれています。 例えばバンク０に属するアドレスは０、６４、１２８、１９２、２５６、・・・・・から始まる４バイトの領域。 任意のデータ長が４バイトより大きいのか小さいのか分かりませんが、テンプレート等を使う場合には 型に合わせて場合分けをする必要があるでしょう。 >>390 「１個のアドレス入力と１個のデータ入出力ができるバンク」が１６個あるというのが適切な表現かと。 392 名前：デフォルトの名無しさん mailto:sage [2009/12/20(日) 10:00:29 ] >>389 device memory on node 0 -> host memory on node 0 -> host memory on node 1 -> device memory on node 1 となるのでレイテンシは大きくなります。今のところCUDAではデバイスメモリをメモリマップする手段はありません。 ただし十分大きなデータを転送する場合にはパイプラン化すれば問題なくなると思います。 むしろpinnedメモリとMPIライブラリの干渉が一番の問題・・・・。 393 名前：デフォルトの名無しさん mailto:sage [2009/12/20(日) 10:47:33 ] >>392 レスサンクス やはりレイテンシは大きいのだね。 目的のアプリは数10kbyteから数100kbyteの転送を頻繁に行うから、 せっかくGPUで速くできても、転送ボトルネックで詰まりそう。 転送するサイズも小さいので、page lockさせない方がよいのかも。 394 名前：デフォルトの名無しさん mailto:sage [2009/12/20(日) 14:34:19 ] GPUの購入で悩んでいるのだが、 Tesla C1060, GTX295,GTX285のうち結局どれが 速いんですか？ちなみに流体に使います。 GTX295ってGPU2基積んでるけど並列プログラミング しないと2基機能しないとか？ 素人質問で申し訳ない。 Teslaの保証とサポートも魅力的だが。 395 名前：デフォルトの名無しさん mailto:sage [2009/12/20(日) 14:44:32 ] >>394 メモリ量でC1060になっちゃう　なんてことないかな？ 295はちゃんとマルチスレッドでプログラムしないと二基使えないです。 396 名前：デフォルトの名無しさん mailto:sage [2009/12/20(日) 15:45:47 ] >>390 >>391 なるほど…よく分かりました。ﾚｽ感謝。 処理させたい任意のデータ長は5ﾊﾞｲﾄだったり、11ﾊﾞｲﾄだったり半端な数もきます。 任意のデータ長に対してコピーやビット演算を行なうんですが、データ長が4ﾊﾞｲﾄより大きいと もうバンク競合は避けられない感じなんですね。 プログラムも汎用にさせたかったけどここは我慢するしかないのか…(´・ω・`) 397 名前：デフォルトの名無しさん mailto:sage [2009/12/20(日) 16:09:09 ] VCでプロジェクト作るのが面倒なんだけど、 なんかウィザード的なものはないのかな 398 名前：,,・´∀｀・,,）っ-○○○ mailto:sage [2009/12/20(日) 16:29:23 ] forums.nvidia.com/index.php?showtopic=65111&pid=374520&st=0&#entry374520 HTML+JavaScriptだから好きに書き換えられるだろ 399 名前：デフォルトの名無しさん mailto:sage [2009/12/20(日) 16:30:40 ] おいらも流体で使う予定だが GTX260を二つ買ってCUDAのお勉強しつつFermi待ち C1060とか中途半端に高いからねえ 400 名前：デフォルトの名無しさん mailto:sage [2009/12/20(日) 16:35:47 ] Nvidia寄りのとある企業のエンジニア曰く、 スペックはまったく同じでGTX285のチップの選別品だといっていた。 クロックがGTX285のほうが高いし、GTX285のほうが若干早いかも。 でもこのスレでもあるように、なんか計算結果がずれる可能性があるし、 メモリの多いTeslaが使えるなら使いたいよね。 401 名前：デフォルトの名無しさん [2009/12/20(日) 16:53:55 ] DodでさえIBM CellじゃなくてPS3を購入してるんだからGTX285で十分でしょう。 402 名前：デフォルトの名無しさん mailto:sage [2009/12/20(日) 17:01:10 ] >>398 まだ完成してないんじゃん 403 名前：デフォルトの名無しさん mailto:sage [2009/12/20(日) 18:24:06 ] >>396 ASCII本の絵がわかりやすいけど、 共有メモリーは[16][17]にすれば問題ないと思う。 共有メモリーは無駄飯食っている用心棒みたいなもんで、 冗長にしてでも使えるときに使わない手はない。 404 名前：394 mailto:sage [2009/12/20(日) 22:43:43 ] 皆さん、レス感謝です。 なるほど、GTX285とC1060差はさほどないんですね。 メモリと耐久性とサポート,計算結果 に関してはTeslaが有利というわけですな。 とりあえずGTX295のマルチスレッドはちょっと遠慮しようかなと思います。 もうちょい1人で悩んでみます。 405 名前：デフォルトの名無しさん mailto:sage [2009/12/21(月) 00:30:00 ] 個人的には285の2GBメモリ版がお薦めなのです。 406 名前：デフォルトの名無しさん [2009/12/21(月) 06:47:28 ] 【SIGGRAPH Asia 2009レポート】 東工大、スクウェアエニックスがCUDA実装事例を紹介 ttp://pc.watch.impress.co.jp/docs/news/event/20091221_338290.html 407 名前：デフォルトの名無しさん [2009/12/21(月) 07:04:05 ] 【SIGGRAPH Asia 2009レポート】 NVIDIAがGPUによるグラフィックスワークスタイルの変革をアピール ttp://pc.watch.impress.co.jp/docs/news/event/20091218_336837.html 408 名前：デフォルトの名無しさん mailto:sage [2009/12/21(月) 07:14:00 ] >>405 どっか、285の4GB版とか安く出してくれないもんかね。 409 名前：デフォルトの名無しさん mailto:sage [2009/12/21(月) 07:15:58 ] 　　　　 ...|￣￣ |　＜ Fermiはまだかね？ 　　　／:::|　　___|　　　　　　 ∧∧　 　 ∧∧ 　　/::::＿|＿＿_|＿　　　　（　｡_｡）.　 （　｡_｡） 　　||:::::::（　・∀・）　　 　　／<▽>　 ／<▽> 　　||::／ <ヽ∞/>＼　　 |::::::;;;;::/ 　|::::::;;;;::/ 　　||::|　　 <ヽ/>.-　|　　|:と),__｣ 　 |:と),__｣ ＿..||::|　　 o　　o　...|_ξ|:::::::::| 　 　.|::::::::| ＼　　＼＿_(久)__／＿＼::::::|　　　 |:::::::| .||.i＼　　　　　　　　､__ﾉﾌ ＼|　　 　|:::::::| .||ヽ .i＼　_ __ ＿＿＿＿ __ _.＼　　 |::::::| .||　ﾞヽ i　　　 ﾊ i　ﾊ i　ﾊ i　ﾊ　|　　し'_つ .||　　 ﾞ|i〜^~^〜^~^〜^~^〜 410 名前：デフォルトの名無しさん mailto:sage [2009/12/21(月) 11:01:34 ] めるせんぬついすた、GPU版の使い方＠Windowsがやっとわかったわー 411 名前：デフォルトの名無しさん mailto:sage [2009/12/21(月) 15:33:23 ] 　　　　　|┃三　　　　　　　　　 /::::::::ハ､＼､::::::::＼＼::::::::::::', 　　　　　|┃　　　　　　　　　 　 ｉ:::::::ｲ　　｀> ー─--ミ::::::::::::| 　　　　　|┃　　　　　 　　　　　 {::::::::|　 　 ::＼:::／:::: 　＼:::ﾘ-} 　ｶﾞﾗｯ.　|┃　　　　　　　　　 　 ',::r､:|　 <●>　　<●>　　!> ｲ 　　　　　|┃　 ﾉ//　　　　　　　　|:､`{　 ｀> .::　 ､　　　 　 __ﾉ 　　　　　|┃三　　　　　　　　　　|::∧ﾍ　　／､__r)＼　　　|:::::| 　　　　　|┃ 　　　　　　　　　　　|::::::`~',　〈　,＿ｨｪｧ 〉　　l::::::》 ＜フェニミストはまだかね　辛抱たまらん 　　　　　|┃　　　　　　 　　　　　|:::::::::::::'､　 `=＝'´　　,,ｲ::ﾉﾉ从 　　　　　|┃三　　　　　　　　　ノ从､:::::::::`ｉ､,,　...　..,,／ |::::://:从 412 名前：デフォルトの名無しさん mailto:sage [2009/12/21(月) 16:19:40 ] おれの物欲も辛抱たまらん、Core i7＋GTX260Mノートが欲しくて。 でもFermiのモバイル版が乗ったノートが出るのを松 413 名前：デフォルトの名無しさん mailto:sage [2009/12/21(月) 21:07:30 ] 故あってCUDAを使うことになったのですが自宅にはRadeonを積んだものしかありません。 コンパイルオプションでCPUでのエミュレーションができると何かで読んだのですが これを利用すればRadeonの環境でも一応の動作はするのでしょうか？ 414 名前：デフォルトの名無しさん mailto:sage [2009/12/21(月) 21:24:33 ] CPUでのエミュレーションなので radeon関係なくCPU上で動くよ 415 名前：デフォルトの名無しさん mailto:sage [2009/12/21(月) 21:44:22 ] 即レスどうもです。学校には設備があるので問題ないのですが自宅でもやりたかったので あともう一つ質問よろしいでしょうか。VineLinuxでの動作は可能でしょうか? 気になったのですがどうもサポートしてないようなので 416 名前：デフォルトの名無しさん mailto:sage [2009/12/22(火) 05:22:26 ] >>412 悪魔：「単精度なら大して変わらないんだから、買っちゃえよ！Fermi出たらまた買えよ！」 417 名前：デフォルトの名無しさん mailto:sage [2009/12/22(火) 06:31:17 ] >>415 Linuxでサポート期待しちゃいかんだろう あとエミュレーションはテクスチャ周りなど微妙に？おかしい模様なんで、あんまり期待しないほうがいい。 5000円くらいで適当なカード買うのが一番無難かもよ。 418 名前：デフォルトの名無しさん mailto:sage [2009/12/23(水) 09:13:22 ] なんか、GeForceの在庫がどんどん無くなっていってない？ まもなくFermi出るな！これは！ 419 名前：デフォルトの名無しさん mailto:sage [2009/12/23(水) 11:06:25 ] C2050/2070よりも先にFermi搭載のGeForceが出ると発表しているしね NVIDIA、“Fermi”採用第1弾GPU「Tesla 20」シリーズ発表 - ITmedia +D PC USER ttp://plusd.itmedia.co.jp/pcuser/articles/0911/17/news039.html まあ、実際に2010Q1に出るかは怪しいわけだが C2050が2499ドルってのは自宅用で用意するには結構きつい値段なので、 Fermi搭載のGeForceがどの程度の値段なのかが今から気になっている 420 名前：デフォルトの名無しさん mailto:sage [2009/12/23(水) 11:07:45 ] >>415 とりあえず玄人思考のGTX260でも買ってみよう でかいからPCケースに入るか確認してからね 421 名前：デフォルトの名無しさん mailto:sage [2009/12/26(土) 03:28:46 ] Device 0: "GeForce 8600 GTS" Total number of registers available per block: 8192 CUDAやろうと思ってますが、レジスタの領域が少なすぎませんか？ __device__ void swap(float *d1, float *d2); 例えばこのような関数呼び出すのに引数とtempで計12byte、他にもthIDや作業用でローカル変数使うから、 最大のパフォーマンス求めようとすると実質スレッドは300個くらいになるんだけど… こんなんだと何万、何千のスレッドとか無理じゃね？ みんなカーネル以外の関数は作らずにカーネルの中に処理を直書きしてるのですか？ 422 名前：デフォルトの名無しさん mailto:sage [2009/12/26(土) 04:01:38 ] >>421 レジスタとスタックとかバイトとレジスタ個数とかごっちゃになってないか色々と 423 名前：デフォルトの名無しさん mailto:sage [2009/12/26(土) 11:11:36 ] 4Gamer.net ― Fermi時代の製品展開が分かってきたNVIDIAのGPUロードマップ。DX11世代が出揃うのは2010年Q2以降に（Fermi（開発コードネーム）） www.4gamer.net/games/099/G009929/20091012001/ Fermi搭載のGeForceは3月くらいかなあ…？ この記事自体が2ヶ月前だからなんともだけどね 424 名前：デフォルトの名無しさん mailto:sage [2009/12/26(土) 17:54:34 ] あんどうさん更新 ttp://www.geocities.jp/andosprocinfo/wadai09/20091226.htm 425 名前：デフォルトの名無しさん mailto:sage [2009/12/26(土) 18:03:48 ] >>421 >Device 0: "GeForce 8600 GTS" >Total number of registers available per block: 8192 直訳すると、「ブロックあたり使用可能なレジスター数：8192」 ブロックあたりのスレッド数は数千・数万もいらない。 128〜256程度でだいたいパフォーマンスは出る。（それを複数コアで何セットも同時に動かす） あとカーネル１個でなんでもかんでもするもんでもない。（直列的な）処理ごとに分割すればいい。 426 名前：421 mailto:sage [2009/12/27(日) 23:15:38 ] >>422 >>425 うん…見直すと自分とても恥ずかしい事言いましたね(´・ω・`) 当初の疑問は理解して解決しました。 で、また別の話のGPGPUに関して質問です。 今DirectX使ってゲームを作っていて処理の重い部分(当たり判定とか)をGPUにやらせようと思っているんだけど、 ゲームって元々グラフィック描画でGPU使うので、それに加えてCUDAやプログラマブルシェーダーでGPGPUしたら相当パフォーマンス落ちますか？ 目的はCPU使用率を抑えたいんだけど、ゲームのような1フレーム毎にGPGPUをするのは実用で使えうるものなんでしょうか。 427 名前：デフォルトの名無しさん mailto:sage [2009/12/27(日) 23:28:32 ] 当たり前の話だけど、GPUをフルに使うようなレンダリングをしていたら遅くなる フレーム枚にGPGPUの処理を入れるのはそんなに問題ない ただし処理の重い部分ってのが上手く並列化できるアルゴリズムであることや、そこそこの量あることが求められる 分岐ばっかだったり、処理よりメインメモリ-VRAM間の転送のほうが時間かかると逆に遅くなる ようは、本当にGPGPUを使ったほうがいい処理なのかを吟味する必要がある 428 名前：デフォルトの名無しさん mailto:sage [2009/12/27(日) 23:54:43 ] 質問です。 C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK\C\src\template\template_vc90.vcproj を開き、ビルドした所、「入力ファイル 'cutil32D.lib' を開けません。」というエラーが出ました。「x64」では、スキップされてしまいます。 >>34 とほぼ同じ状況なので、このスレに書かれている事をしてみたのですが、変化ありません。 また、リンカの設定でcutil32D.libをcutil64D.libに変えた所、x64でビルド出来るようになったのですが、 「モジュールのコンピュータの種類 'x64' は対象コンピュータの種類 'X86' と競合しています。」と別のエラーが発生しました。 環境は windows7 64bit Visual C++ 2008 Express Edition GTS250 CUDAのドライバ、toolkit、sdkは2.3でwin7 64bitの物を入れました。 改善策がありましたら、ぜひお願いします。 429 名前：デフォルトの名無しさん mailto:sage [2009/12/28(月) 00:17:23 ] >>428 Expressは64bitでコンパイルできないかも？ www.microsoft.com/japan/msdn/vstudio/2008/product/compare.aspx 430 名前：デフォルトの名無しさん mailto:sage [2009/12/28(月) 00:25:04 ] >>429 みたいです。やはりproを買うしかないのでしょうか… レスありがとうございました。 431 名前：デフォルトの名無しさん mailto:sage [2009/12/28(月) 00:48:50 ] >>428 >>430 答えるのに十分な情報量を持った質問を久しぶりに見たような気がしたｗ proが良いけどねぇ。 今は最適化コンパイラもついてるし、standardで良い気もする。 （standard使ったことはないので参考まで） 432 名前：デフォルトの名無しさん mailto:sage [2009/12/28(月) 02:01:56 ] PlatformSDKにx64コンパイラあったような 433 名前：デフォルトの名無しさん mailto:sage [2009/12/28(月) 20:43:04 ] 落ち着いて、なぜこういうエラーがでるのか？ と、考えるしかないよ。 434 名前：デフォルトの名無しさん mailto:sage [2009/12/29(火) 11:29:33 ] CUDAが使えるGPUが載ってるか否かの判断をプログラム上で行うには どうするといいのでしょう？ 435 名前：デフォルトの名無しさん mailto:sage [2009/12/29(火) 11:57:15 ] OpenCL入門 - マルチコアCPU・GPUのための並列プログラミング が出版されるそうな アスキードットテクノロジーズの記事書いた人たちらしいが GPGPUでもっともメジャーなのはCUDAだと思うけど OpenCLがこの本の売り文句通りスタンダードになるんだろうか…？ よくわからないけど、無関心でいるわけにも行かないのでとりあえずポチってみる 感想あったらアマゾンのレビューに書く >>434 プログラム上でその判断をやる必要というのがよくわからない コマンドラインでやるのでは駄目なのか 436 名前：デフォルトの名無しさん mailto:sage [2009/12/29(火) 11:57:22 ] >>434 それはやはり cudaGetDeviceCountして、CUDAデバイス個数調べ cudaGetDevicePropertyをDeviceCountまわして、 .major、.minorでバージョンチェック .multiProcessorCountでコアの個数調べ して使うしかないんじゃ。 437 名前：434 mailto:sage [2009/12/29(火) 12:21:50 ] >>436 どうも >>435 CUDA依存部分を共有ライブラリの形でプログラム本体から切り出しておいて、 実行時に動的にリンク出来たらとか考えてます。 …CUDA版バイナリと、CUDA無し版バイナリを用意して、インストール時に 選ばせればいいのか。 438 名前：デフォルトの名無しさん mailto:sage [2009/12/29(火) 13:05:23 ] >>435 スタンダードはどう考えてもDirectComputeだろ。 いまのWindowsのシェアとDirectXの普及率から考えてCUDAがスタンダードになるにはポテンシャルが違いすぎる。 439 名前：438 mailto:sage [2009/12/29(火) 13:06:36 ] CUDAはスタンダードになれる可能性はあるけど、OpenCLはたぶん廃れる。 OpenCLって所詮はCUDAのラッパでしょ？ ラッパライブラリがスタンダードになった事例ってあんまり知らないんだけどさぁ。 440 名前：デフォルトの名無しさん mailto:sage [2009/12/29(火) 13:16:26 ] Appleはラッパ好きだよな。 ObjCもラッパではないにしろマクロ駆使してるだけで、骨はCに変わりはないでしょ？ Appleは声は大きいけど技術が無いから、OpenCLもあんまり期待してないよ。 所詮ラッパライブラリだから、コンパイルしたらどのGPUでも使えるってわけではないし。 OpenCLのDLLロードしたらどのGPUでも同じバイナリでおｋみたいになったら使いやすいけどね。 441 名前：デフォルトの名無しさん mailto:sage [2009/12/29(火) 13:19:23 ] てか俺が作っちゃう？ｗｗｗ 442 名前：デフォルトの名無しさん mailto:sage [2009/12/29(火) 13:27:41 ] >>437 436でCUDA使えるかどうか確認して、 プログラムとしてはCPU用ルーチンとGPU用ルーチンを両方持っておく でいいんじゃね？(やはり、動作確認用にCPU/GPU双方で試すのは必要と思うし) 443 名前：デフォルトの名無しさん mailto:sage [2009/12/29(火) 13:30:38 ] >>438 おれ、 ttp://openvidia.sourceforge.net/index.php/DirectCompute 見て「うざっ！！」と思って挫折した。　ダメな俺 444 名前：デフォルトの名無しさん mailto:sage [2009/12/31(木) 03:50:21 ] pyCUDAってどうかな DriverAPIを使ってモジュールを実行時に作成、ロードするみたいですね。 445 名前：デフォルトの名無しさん mailto:sage [2010/01/03(日) 16:34:54 ] 3次元のデータをデバイス側に送って計算したいんですけど、 >>328の ＞固めてまとめて渡しちゃったほうがオーバーヘッドが少ないと思います。 ＞トータル何列あるよ、は別にパラメータで渡す。 というのはどこにどのように記述すればいいのでしょう？ 446 名前：デフォルトの名無しさん mailto:sage [2010/01/03(日) 16:57:35 ] >>445 もちろん、少ないならデバイス関数の引数として渡す。パラメータ複数個まとめた配列で渡しても良いし。 計算の対象、計算の結果、に加えて、計算のパラメータ、も同様にホスト側で確保→cudaMemCopyするわけ。 447 名前：デフォルトの名無しさん mailto:sage [2010/01/03(日) 17:38:23 ] >>446 ありがとうございます。 物分かりが悪くて申し訳ないのですが、 cudamemcpyは確保した領域にデータをコピーする使い方しか思いつかないので、 ＞計算の対象、計算の結果、に加えて、計算のパラメータ、も同様にホスト側で確保→cudaMemCopyするわけ。 を簡単に例で説明していただけないでしょうか？ 448 名前：デフォルトの名無しさん mailto:sage [2010/01/04(月) 00:23:45 ] えっ　なぜ分からないんだ？　たとえばこんな感じだよ __device__ calculate(float * input, float * output, int * params){ : if (params[i] == 0) 　　　output[i] = func_A(input[i]); if (params[i] == 1) 　　　output[i] = func_B(input[i]); 　__syncthreads(); 449 名前：デフォルトの名無しさん mailto:sage [2010/01/04(月) 11:06:25 ] >>448 ありがとうございます。 なんだか変な方向に考えすぎていました。 450 名前：デフォルトの名無しさん mailto:sage [2010/01/04(月) 12:52:14 ] CUDAってGPUカーネルの中で関数読んで値をreturnさせることできたっけ？ あと__global__にしないとホストから関数呼べなくね？ 451 名前：448 mailto:sage [2010/01/04(月) 13:31:34 ] calculateが__global__で、func_A()、func_B()が__device__ですたorz 452 名前：デフォルトの名無しさん mailto:sage [2010/01/04(月) 14:00:48 ] Fermi搭載GeForceの発売がまた不透明になったとか

---

---

[ 続きを読む ] / [ 携帯版 ]

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