【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/ 320 名前：デフォルトの名無しさん mailto:sage [2009/12/07(月) 13:11:16 ] ベンチ見てほくそえむ 321 名前：デフォルトの名無しさん mailto:sage [2009/12/07(月) 16:15:39 ] >>317 ptx出力を眺めれば判るけど、普通のCならできる最適化もGPU向けにはできないことが多いよ。 敢えて言えば、ポインタ変数に尽くvolatileがついているかのように振る舞うみたい。 例えば、 int function(int * foo) {int bar = 3; foo[0] = bar; return foo[0];} みたいなコードはCPU向けには {foo[0] = 3; return 3;} のように最適化されるのにGPU向けには {foo[0] = 3; return foo[0];} のように律儀に解釈される。 なので、ptx出力を読めるかどうかは割りと重要かも知れず。 # つーか、メモリアクセスの個数を数えるくらいのことは普通にやってる。 322 名前：デフォルトの名無しさん mailto:sage [2009/12/08(火) 01:19:36 ] すみません、かなり初心者です。 行ごとに要素数の違う配列をデバイス側に渡したいんですけど、 a=1; //ホスト側 float** mat=(float**)malloc(size); for(i=0;i<num;i++){ mat[i]=(float*)malloc(size / a); a*=2; } //デバイス側 float** mat_d; CUDA_SAFE_CALL(cudaMalloc((void**)&mat_d,size)); for(i=0;i<num;i++){ CUDA_SAFE_CALL(cudaMalloc((void**)&mat_d[i],size/a)); a*=2; } CUDA_SAFE_CALL(cudaMemcpy(mat_d,mat,size,cudaMemcpyHostToDevice)); for(i=0;i<M;i++){ CUDA_SAFE_CALL(cudaMemcpy(mat_d[i],mat[i],size/a,cudaMemcpyHostToDevice)); a*=2; } というようにはできないんでしょうか？ 323 名前：デフォルトの名無しさん mailto:sage [2009/12/08(火) 06:39:08 ] >>322 >CUDA_SAFE_CALL(cudaMemcpy(mat_d,mat,size,cudaMemcpyHostToDevice)); これがいらないんじゃない？ ホスト側のポインタ列をデバイスに渡しても使いようがない。メモリ空間が違う。 324 名前：デフォルトの名無しさん [2009/12/08(火) 09:12:09 ] 開発環境の入っていないマシンで動かすには、どのファイルを持って行けばよいのでしょうか？ 325 名前：デフォルトの名無しさん mailto:sage [2009/12/08(火) 10:25:10 ] >>291 個人的な興味なんですけど、CUDA使ってるcodecってなんで速いんでしょうね？ BYTEアクセスじゃ無いとしても、どの部分をGPUでやらせれば爆速になるのか不思議ﾁｬﾝﾃﾞｽ 326 名前：デフォルトの名無しさん mailto:sage [2009/12/08(火) 14:53:55 ] SDKのドキュメントの通りに，VS2005で サンプルプロジェクト template を元に使っているのですが このプロジェクトに新しい.cuファイルを追加した場合に そのcuファイルへの更新がビルド時に察知されません 毎回リビルドすれば一応反映できるのですが 通常のビルドで反映させるにはどうすればよいのでしょうか？ 327 名前：322 mailto:sage [2009/12/08(火) 15:52:18 ] >>323 すみません。 抜けてましたが、 >CUDA_SAFE_CALL(cudaMemcpy(mat_d,mat,size,cudaMemcpyHostToDevice)); は、 ホスト側であらかじめ計算を行い、 その結果を初期値としてポインタ列、matに与えた上で、 その値をデバイス側のポインタ列、mat_dにコピーし、 その値を使ってデバイス側で計算をする。 というつもりで書きました。 このようにホスト側で計算した値をデバイス側に渡す時は どのように記述するのがよいのでしょうか？ 328 名前：デフォルトの名無しさん mailto:sage [2009/12/08(火) 16:15:43 ] {1個、2個、4個、8個、16個、…}　みたいなデータを渡したいのかな?? 固めてまとめて渡しちゃったほうがオーバーヘッドが少ないと思います。 トータル何列あるよ、は別にパラメータで渡す。 (実際にCUDAのルーチン書く前に、コピー/戻しの時間を色々計ってみるといいです) cudaMallocしたデータにはホストからは触れないので、 ホストでmallocしたデータ(mat)に計算結果格納 →同じサイズでcudaMalloc(mat_d) →cudaMemcpyHostToDeviceで渡す なのでそれでいいです 329 名前：323 mailto:sage [2009/12/08(火) 17:35:58 ] >>327 >CUDA_SAFE_CALL(cudaMalloc((void**)&mat_d,size)); > >for(i=0;i<num;i++){ >CUDA_SAFE_CALL(cudaMalloc((void**)&mat_d[i],size/a)); >a*=2; >} 考えてみるとここも問題があって、cudaMallocということはデバイス側でポインタ列を確保しているんだけど、 そうすると&mat_d[i]というアドレスは、デバイス側にはバッファがあるが、 ホスト側には存在しないから、ここで例外になりそう。 やるならこんな感じかな？↓（未検証） float** mat_d; // GPU側に確保する（ポインタ列用）バッファ CUDA_SAFE_CALL(cudaMalloc((void**)&mat_d, count * sizeof(float*)); float** mat_d_tmp=(float**)malloc(count * sizeof(float*)); // ホスト側に確保する。内容はGPU側ポインタ列 for(i=0;i<num;i++){ // GPU側データバッファのポインタを格納 CUDA_SAFE_CALL(cudaMalloc((void**)&mat_d_tmp[i],size/a)); a*=2; } // GPU側ポインタ列をGPUに転送 CUDA_SAFE_CALL(cudaMemcpy(mat_d,mat_d_tmp,count * sizeof(float*),cudaMemcpyHostToDevice)); free(mat_d_tmp); てやっておいて、mat_dをカーネル呼び出しの引数にしてやるとか。 ここまで、バッファ作りしかしてないので、データ転送は別途必要。 たぶん>>322の最後４行のとおりで構わない。 ただデータ転送回数多いと多少とも時間かかるから、>>328の言うように 固めたほうがベターではある。 330 名前：323 mailto:sage [2009/12/08(火) 17:48:04 ] >>329 ちなみに、ポインタサイズが32bitか64bitかはホスト側に依存してることを確認した。 forums.nvidia.com/index.php?s=&showtopic=149501&view=findpost&p=943899 331 名前：デフォルトの名無しさん mailto:sage [2009/12/08(火) 17:59:16 ] >>325 あいまいな質問だけど答えてみる。 CPUよりGPUのほうが計算能力の総量はでかいから、本来GPUが速いほうが当然なんだけど、 以下のようなものは遅くなる場合もある。 ・計算負荷よりもメモリ転送がボトルネックになるもの ・処理を細かく分解して並列化することが難しいもの codecなどは、メモリ転送がボトルネックになりやすい傾向はあるものの、 だいたい画面の領域ごとに分解できる処理が多いから、多少とも速くはできるでしょう。 あと速くならない処理はCPU側でやればいい。 332 名前：デフォルトの名無しさん mailto:sage [2009/12/08(火) 20:37:12 ] >>328>>329 参考になりました。ありがとうございます！ まだ不慣れなので色々試行錯誤しながらやってみます！ 333 名前：デフォルトの名無しさん mailto:sage [2009/12/08(火) 21:48:05 ] >>325 CUDAだけ使って早いと思ってる？ｗ 334 名前：,,・´∀｀・,,）っ-○○○ mailto:sage [2009/12/08(火) 23:12:01 ] 動き検出なんかは当然CPUにやらせたほうがいいぞ SpursEngineなりCellなりがあるならそれでもいいが 335 名前：デフォルトの名無しさん mailto:sage [2009/12/09(水) 00:11:13 ] >>333-334 流石にGPUだけとは思ってないけど、atomでh264なんかエンコすると日が暮れる勢いだけど、ionプラットフォームだと実用範囲になる。 エンコ中でも完全にCPUを使い切ってないところを見ると、やっぱGPUをかなり使ってるんだなと勝手な想像 336 名前：デフォルトの名無しさん mailto:sage [2009/12/09(水) 06:34:40 ] IONはメインメモリをグローバルメモリとして使う、つまりマッピングしておけば CPUからもGPUからもフルスピードでアクセスできるメモリ空間を作ることができる。 ION2は買おうかな・・・。 337 名前：デフォルトの名無しさん mailto:sage [2009/12/09(水) 08:29:47 ] >>336 実際はそんなに甘くはなく、GPUからフルスピードでアクセスするためにはWriteCombinedモードに する必要があって、それを付けるとCPUからのアクセスが死ぬほど遅くなります。 通常は素直に転送（実際にはメモリコピーですが）したほうがマシです。 低価格で4GB近いデバイスメモリが使えるというメリットはありますが。 338 名前：デフォルトの名無しさん mailto:sage [2009/12/09(水) 09:47:42 ] そのためのMOVNTDQA/MOVTDQだろう。 まあ前者はSSE4.1でないと使えないがな。Atomは仕様不可。あしからず。 339 名前：デフォルトの名無しさん mailto:sage [2009/12/09(水) 10:11:54 ] それらを使っても十分に遅いと思うのですが。 340 名前： ◆0uxK91AxII mailto:sage [2009/12/09(水) 10:22:13 ] WCにしつつ、CPU->GPUはmovnt*でcopyすれば良さそうだね。 CUDAとか知らないけど:b 341 名前：デフォルトの名無しさん mailto:sage [2009/12/09(水) 10:26:15 ] >>339 だからCPU専用のワーク領域と分けて読み書きを最小限にする。 342 名前：デフォルトの名無しさん mailto:sage [2009/12/09(水) 19:25:43 ] >>337 そか、CPUのキャッシュが使えなかったね。 死ぬほどまで遅くなるとは思ってなかった。 343 名前：デフォルトの名無しさん mailto:sage [2009/12/10(木) 10:12:49 ] GTX295 Vista x64 CUDA2.3 の環境で、CUDA VS WIZARDで作成したプロジェクトを使っています。 質問1. ブロックあたりのスレッド数を512にすると計算結果が正しくなりません。256なら正しいです。なぜでしょう？ shared memory はグローバル変数として宣言して、各スレッドで共用（リードオンリー）しています。スレッド内で__shared__の宣言はありません。 質問2. 完成したので、SDKのsimpleMultiGPUを参考にGPU 2個使おうとしています。 とりあえずマルチスレッド用の空の関数を書いたところ、コンパイルはできるのですが、リンクできません。 cutStartThreadなどが外部で定義されてないと言われます。 simpleMultiGPUはビルドできるので、プロジェクトのプロパティを見比べてみましたが特に違いは見あたりません。 何が悪いんでしょうか？ どなたかお助けください。 344 名前：デフォルトの名無しさん mailto:sage [2009/12/10(木) 10:45:27 ] >>343 ・sharedにはCPUから書けませんが、各スレッドがリードオンリーならどこで書き込んでいるのでしょう。 ・cutで始まる名前の関数はcutilライブラリ内にあります。cutilライブラリもビルドし、あなたのプロジェクトからリンクできるようにしないといけません。 345 名前：デフォルトの名無しさん mailto:sage [2009/12/10(木) 12:13:48 ] >>344 上については、言葉足らずでした。カーネルの最初に if( threadIdx.x == 0 ) { /* デバイスメモリからコピー */ } __syncthreads(); でコピーしています。 将来的にはコピーも並列化する予定です。 下についてはありがとうございます。既存のlib(dll?)だけでは駄目だってことですね。 346 名前：デフォルトの名無しさん mailto:sage [2009/12/10(木) 12:40:10 ] threadIdx.xはスレッドの識別子で、実行順番とは無関係では？ 347 名前：デフォルトの名無しさん mailto:sage [2009/12/10(木) 13:15:43 ] 共有メモリのつくりに関する知識が思いっきり欠如していると思われ。 threadIdxが0の人だけが書こうとしたら、他の人はみんな何もできなくて時間が無駄。 おまけに、スレッド数が一定数を超えると(GTX295の場合は256か)、別の実行単位になってしまうから(表に見えない) 共有メモリ間転送が行なわれてしまってその点でも時間の無駄。 尤も、>345の様にthreadIdx0でしかコピーしない場合は転送は発生しない代わりに結果がご覧の通りなわけで。 どうせthreadIdx0がコピーしている間他はなにもできないんだから、一斉に同じものを書いてしまえばいいんでない? 348 名前：デフォルトの名無しさん mailto:sage [2009/12/10(木) 13:33:51 ] >>347 > threadIdxが0の人だけが書こうとしたら、他の人はみんな何もできなくて時間が無駄。 開発途中なので、とりあえずこうしているだけです。 将来的にはコアレスにします。 > スレッド数が一定数を超えると(GTX295の場合は256か)、別の実行単位になってしまうから(表に見えない) 512で上手くいかないのはこのあたりが原因のようですね。 349 名前：デフォルトの名無しさん mailto:sage [2009/12/10(木) 13:49:31 ] >>347 たびたびすみません。 「別の実行単位」の意味がよく分からないのですが、 要は1つのブロックが2つのSMに割り当てられるって意味ですか？ 350 名前：デフォルトの名無しさん mailto:sage [2009/12/10(木) 14:03:14 ] >>345 Sharedは文字通り、共有できるメモリなので、 各スレッドが１ワードずつ協力して作成→どこでも読み書き可能、 ですよ。0番スレッドだけに任せなくてもいい 351 名前：デフォルトの名無しさん mailto:sage [2009/12/10(木) 17:12:55 ] 自分は>>343じゃないけど256スレッドでOKで512スレッドで計算結果がおかしくなる理由がわからない。 >おまけに、スレッド数が一定数を超えると(GTX295の場合は256か)、別の実行単位になってしまうから(表に見えない) >共有メモリ間転送が行なわれてしまってその点でも時間の無駄。 というのは時間の無駄を指摘してるだけで結果がおかしくなる理由ではないという認識なのですがあってますか？ この時の共有メモリ間転送というのは別の実行単位に転送しているということ？ >>343 256スレッドでやればOKなのでもう興味はないかもしれないが、シェアードメモリの代わりにグローバルメモリで 共用して計算した場合は512スレッドの時も結果は正しくなるのですか？ 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ケースに入るか確認してからね

---

---

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

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