- 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/
- 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の発売がまた不透明になったとか
- 453 名前：デフォルトの名無しさん mailto:sage [2010/01/04(月) 16:05:01 ]
- ⇒を何らかの演算として 0⇒256 256⇒512 512⇒0 といった感じにしたいのですが
a = (a+256)%513;
とする以外思いつきません。けど剰余は遅いので違う計算にしたいんです。
いっそのことテーブル作ったりif文にしたほうが早いのでしょうか？
馬鹿すぎてわかりません。誰か助けて
- 454 名前：デフォルトの名無しさん mailto:sage [2010/01/04(月) 16:10:45 ]
- >>453
Cudaだと四則演算はほとんど処理時間には影響しないし、その式をつかうといいんじゃない？
っていおうとおもったけど、a=512のとき0にならなくね？
- 455 名前：デフォルトの名無しさん mailto:sage [2010/01/04(月) 16:13:31 ]
- 257⇒0 258⇒1 … でいいの？
でもそうすると512⇒0がヘンだけどな。
int val;
val = a + 256;
if (val > 512) val = val - 513;
でいいんじゃないの？
というかC言語の疑問をCUDAスレで聞くのはまちがっとる、ここは人口が少ないぞ
- 456 名前：デフォルトの名無しさん mailto:sage [2010/01/04(月) 16:26:48 ]
- じゃあCUDAっぽく
const int TABLE[512] = { 256, ... , 512, ... , 0 }
__device__ int f(int a){
return TABLE[a];
}
- 457 名前：デフォルトの名無しさん mailto:sage [2010/01/04(月) 17:09:36 ]
- レスありがとうございます
わけわかんないこと書いてすみません
テーブルとif文試して出直してきます
- 458 名前：デフォルトの名無しさん mailto:sage [2010/01/04(月) 17:37:31 ]
- >>453
b⇒aとして、
a = (b+256)%768;
でいいのかな？
0⇒256
256⇒512
511⇒767
512⇒0
- 459 名前：デフォルトの名無しさん mailto:sage [2010/01/04(月) 21:45:29 ]
- 釣りだよな？
- 460 名前：デフォルトの名無しさん mailto:sage [2010/01/05(火) 01:06:37 ]
- 申し訳ないっす
- 461 名前：デフォルトの名無しさん mailto:sage [2010/01/07(木) 00:21:36 ]
- ランダム関数の作成で悩んでます。
・CPUで作成した乱数リストをGPUに転送しない
・1スレッドの中で何度も呼ばれる
・グローバルメモリは使わない。スレッドIDやclock()をseedとして使う
この制約で良い乱数生成アルゴリズムを教えて下さい。
MTGPっていうのを参考にしようとしましたがムズ過ぎて挫折…まずコンパイルが通らず(´・ω・`)
- 462 名前：デフォルトの名無しさん mailto:sage [2010/01/07(木) 01:01:35 ]
- おれMTGPコンパイルできたよ。GPU側で作成してそのままGPUで使えるし。
（シンボルコピー）
- 463 名前：デフォルトの名無しさん mailto:sage [2010/01/07(木) 08:15:53 ]
- 使ったファイルはこれだけ
mtgp32-fast.h
mtgp32-cuda.cu
mtgp32-fast.c
mtgp32-param-fast.c
make_single_random()とかの末尾を適当に書き換える。
- 464 名前：デフォルトの名無しさん [2010/01/07(木) 12:09:10 ]
- 誰かＯｐｔｉＸやってる人いませんかー？
- 465 名前：デフォルトの名無しさん mailto:sage [2010/01/08(金) 09:15:11 ]
- CES 2010: NVIDIA Shows GeForce GF100 Fermi Video Card - Legit Reviews
www.legitreviews.com/news/7122/
3/2に本当に出るのかしらん
- 466 名前：デフォルトの名無しさん mailto:sage [2010/01/08(金) 09:22:19 ]
- 気にするな、出たら買う、出るまではCUDAの修行をする。
単純でｲｲじゃないか。
- 467 名前：デフォルトの名無しさん mailto:sage [2010/01/08(金) 09:41:48 ]
- >>461
Perlinノイズやってみてよ
ttp://mrl.nyu.edu/~perlin/doc/oscar.html
Cのソースが有る
- 468 名前：デフォルトの名無しさん mailto:sage [2010/01/08(金) 11:02:05 ]
- Fermiもダメそうってどこかのブログで見た
NVIDIAまじやばい
- 469 名前：デフォルトの名無しさん mailto:sage [2010/01/08(金) 11:05:40 ]
- ようやくCUDAに慣れてきたところなのに、StreamSDKとかOpenCLとか
DirectComputeとかまた勉強しないといけないの？
辛いのうｗｗプログラマは辛いのうｗｗｗ orz
- 470 名前：デフォルトの名無しさん mailto:sage [2010/01/08(金) 13:36:24 ]
- windows7 64bit、Visual C++ 2008 Pro環境に、CUDA toolkit、sdk 2.3をインストールして、
C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK\C\srcにある、
_vc90.slnファイルを開いて、問題なくビルドできます。
exeファイルの出力先は、デフォルトでC:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK\C\bin\win64\ですが、
これ以外の出力先に変更する方法を教えてください。
- 471 名前：デフォルトの名無しさん mailto:sage [2010/01/08(金) 13:42:23 ]
- プロジェクトのプロパティで変更できるよ
$(OutDir)\$(ProjectName).exe
これを丸ごとフルパスで指定してしまいなさい。
- 472 名前：デフォルトの名無しさん mailto:sage [2010/01/08(金) 17:14:12 ]
- XP,32bitでコンパイルしたプログラムを7,64bitで動かそうとしたら動かなかった･･･
こういうものなのか
- 473 名前：デフォルトの名無しさん [2010/01/08(金) 17:23:41 ]
- 32/64ビットのcudart.dllが必要
- 474 名前：デフォルトの名無しさん mailto:sage [2010/01/08(金) 19:20:38 ]
- >>471
thanks
- 475 名前：デフォルトの名無しさん mailto:sage [2010/01/10(日) 12:33:02 ]
- 現在CUDA SDKを導入しVisual C++ 2005 Express Editionを使用しCPUでのエミュレーションを行っています。
いくつかのサンプルは動かせたのですがparticlesなど複数のサンプルでコンパイルはできているようですが実行すると以下のようなエラーが出ます
device emulation mode and device execution mode cannot be mixed.
エラーが出ている行は
cutilCheckMsg("cudaMalloc failed");
cutilSafeCall(cudaGLRegisterBufferObject(vbo));
cutilSafeCall( cudaMalloc3DArray(&noiseArray, &channelDesc, size) );
のように「cutil〜」になってます。オプションから設定しているのでデバイスエミュレーション自体はうまく行ってるようなのですが
エミュレーションを行う際はソースを書き換える必要があるのでしょうか?
- 476 名前：デフォルトの名無しさん mailto:sage [2010/01/10(日) 12:55:52 ]
- メッセージのとおりエミュONOFFまぜてる
またはリビルドしてないんじゃ？
プロジェクト設定じゃなくてビルド構成のほうでエミュレーションにかえられないっけ？
- 477 名前：デフォルトの名無しさん mailto:sage [2010/01/10(日) 13:05:00 ]
- メッセージの内容は把握できています。
リビルドも多分できてると思います。
エミュレーションの設定は
プロジェクトのプロパティ→CUDA Build Rule→Emulation Mode はい
にしてます。ほかの方法はよくわからないのですがほかの方法があるのでしょうか?
- 478 名前：デフォルトの名無しさん mailto:sage [2010/01/10(日) 22:13:13 ]
- ビルド構成のほうからエミュレーションの設定ができました。問題なくビルド&実行できました。ありがとうございます。
ただFPSが0.1なんでほぼ無意味でした。
- 479 名前：デフォルトの名無しさん [2010/01/10(日) 23:24:17 ]
- エミュレーションなんてHello world以上のことをやるもんじゃないな
- 480 名前：デフォルトの名無しさん mailto:sage [2010/01/11(月) 00:20:59 ]
- 処理する画面領域を１０００分の１ぐらいに限定して、バグを調査するくらいのことはできる。
- 481 名前：デフォルトの名無しさん mailto:sage [2010/01/11(月) 01:26:19 ]
- なんだっけ、エミュじゃなくて実機で実デバッグできるやつ
早く来てほしいな
- 482 名前：デフォルトの名無しさん mailto:sage [2010/01/12(火) 14:25:34 ]
- CPU版アプリの移植の際に、エミュはデータ検証に使える。
- 483 名前：デフォルトの名無しさん [2010/01/12(火) 20:21:38 ]
- -G オプションって使える？WinXPで
「'assembler' は、内部コマンドまたは外部コマンド、
操作可能なプログラムまたはバッチ ファイルとして認識されていません。」
と出てビルドできないけど。
--deviceemuもつけないとだめなの？
- 484 名前：デフォルトの名無しさん [2010/01/13(水) 09:38:31 ]
- CUDAのバージョンを書き出しておきたいのですが、
CUDAのバージョンを調べる関数か何かありませんか？
- 485 名前：デフォルトの名無しさん mailto:sage [2010/01/13(水) 09:41:29 ]
- >>484
>>436 でどう？
- 486 名前：デフォルトの名無しさん mailto:sage [2010/01/13(水) 09:50:27 ]
- こっちを要求している気がする。
cudaRuntimeGetVersion()
cudaDriverGetVersion()
- 487 名前：デフォルトの名無しさん [2010/01/13(水) 09:56:08 ]
- WindowsでCUDAを使っています。
ループが一定回数超える度に、coutで経過を表示しています。
Linuxに比べて遅かったので、いろいろ試してみたところ、
DOSウインドウの前で、マウスを動かすと、速くなります。
垂直同期か何かで引っかかっているのかと思い、設定で切ってみましたが
状況が変わりません。
どなたか似たような状況になって解決された方はいないでしょうか？
- 488 名前：487 [2010/01/13(水) 11:05:28 ]
- 追加報告です。
cudaMemcpyでcudaMemcpyDeviceToHostでコピーするところで異常に時間がかかっているようです。
この1行を消すだけで、100秒ほどかかっていたのが2秒にまでなりました。
逆方向のcudaMemcpyHostToDeviceは問題ありません。
- 489 名前：デフォルトの名無しさん mailto:sage [2010/01/13(水) 11:13:17 ]
- CUDAプログラムがGPUで計算等をしている間は画面表示の更新が遅くなります。
１回のループにどれくらいの時間がかかるのか、一定回数がどれくらいかによりますが、
経過の表示される間隔が短いとその影響を受ける可能性があります。
あるいは、マウスを動かすことでCPUに負荷をかけないとパワーセーブモードに入ってしまうとか・・・・。
- 490 名前：デフォルトの名無しさん mailto:sage [2010/01/13(水) 11:17:06 ]
- >>488
転送に１００秒かかるのは現実的にありえないので、恐らくcudaThreadSynchronize()等でカーネルの実行完了待ちを
していないのではないかと思います。カーネル呼び出しそのものは一瞬で終わりますが、
後続のcudaMemcpy(.....,...,...., cudaMemcpyDeviceToHost)時に暗黙的にそれまでに呼び出したカーネルの実行完了を待っています。
・・・・・ということかどうかは分かりませんが。
- 491 名前：デフォルトの名無しさん mailto:sage [2010/01/13(水) 11:29:27 ]
- >>489,490さん
早速の返信ありがとうございます。
cudaMemcpyでcudaMemcpyDeviceToHostを消したプログラムで、
プログラムの一番最後にだけ結果を取り出してみたところ、
正しく計算されていました。
次に、カーネルの実行部分のみをコメントアウトした場合、
やはり非常に時間がかかってしまいました。
やはり、GPUからCPUへのデータ転送に（というよりなにか同期の部分の様な気もしますが）
時間がかかっているようです。
計算用マシンなので省電力の設定はしていません。
- 492 名前：デフォルトの名無しさん mailto:sage [2010/01/13(水) 11:48:15 ]
- >>491
> cudaMemcpyでcudaMemcpyDeviceToHostを消したプログラムで、
> プログラムの一番最後にだけ結果を取り出してみたところ、
> 正しく計算されていました。
これを読むと、そもそも何度もDeviceToHostの転送をする必要がない処理という解釈で
よいのでしょうか？
DeviceToHostが遅くなる理由はハードウェア的なものから色々あります。
マザーボードを交換したという人もいました。
SDKに含まれているbandwidthTestの結果はどうなっていますか？
"--memory=pinned" を付けた場合と付けなかった場合をそれぞれ調べてみてください。
- 493 名前：デフォルトの名無しさん [2010/01/13(水) 11:59:12 ]
- Device 0: GeForce GTX 285
Quick Mode
Host to Device Bandwidth, 1 Device(s), Paged memory
Transfer Size (Bytes) Bandwidth(MB/s)
33554432 4340.2
Device to Host Bandwidth, 1 Device(s), Paged memory
Transfer Size (Bytes) Bandwidth(MB/s)
33554432 4193.7
Device to Device Bandwidth, 1 Device(s)
Transfer Size (Bytes) Bandwidth(MB/s)
33554432 119962.3
--memory=pinned
Device 0: GeForce GTX 285
Quick Mode
Host to Device Bandwidth, 1 Device(s), Pinned memory, Write-Combined Memory Enabled
Transfer Size (Bytes) Bandwidth(MB/s)
33554432 5758.1
Device to Host Bandwidth, 1 Device(s), Pinned memory, Write-Combined Memory Enabled
Transfer Size (Bytes) Bandwidth(MB/s)
33554432 5471.9
Device to Device Bandwidth, 1 Device(s)
Transfer Size (Bytes) Bandwidth(MB/s)
33554432 119852.3
- 494 名前：デフォルトの名無しさん mailto:sage [2010/01/13(水) 12:00:32 ]
- 結果はこんな感じです。
特に片方だけ遅いというわけでもないようです。
Linuxで起動した場合、遅くなるという問題は起きませんでした。
DeviceToHostが必要なのは結果を取り出すときだけで、
計算自体には必要ありません。
- 495 名前：デフォルトの名無しさん [2010/01/13(水) 12:03:47 ]
- 似たような話をnvidiaのforumで見つけました
forums.nvidia.com/lofiversion/index.php?t68705.html
forums.nvidia.com/lofiversion/index.php?t73047.html
返信は結構あるのですが、解決策はよくわかりません。
- 496 名前：デフォルトの名無しさん mailto:sage [2010/01/13(水) 14:46:26 ]
- ・CUDA内での計算処理、ループ数とループ内のステップ数、スレッド数ブロック数、だいたいどんなもんすか
・HostToDevice、DeviceToHostで転送するメモリのサイズはどのくらいですか
- 497 名前：デフォルトの名無しさん mailto:sage [2010/01/13(水) 15:02:38 ]
- とりあえず
タイマー1,2作成
cudaMemCopy(HostToDevice)
タイマー1,2起動
<<<>>>カーネル呼ぶ
cudaThreadSyncronize()
タイマー1停止
cudaMemCopy(DeviceToHost)
タイマー2停止
してほんとにMemCopyなのかカーネルなのか確かめてみる必要は。
- 498 名前：デフォルトの名無しさん mailto:sage [2010/01/13(水) 19:23:55 ]
- コピーが細切れでforループでやたら回数呼んで転送してたりしない？
１発で全部転送するのと1/1000ずつ1000回コピーするのとでは、
かかる時間に雲泥の差があるけど。
- 499 名前：487 mailto:sage [2010/01/14(木) 04:10:12 ]
- PCを再起動したら、上記の問題は出なくなりました。
お騒がせしました。
- 500 名前：デフォルトの名無しさん mailto:sage [2010/01/14(木) 09:16:24 ]
- おいｗｗ 終わりかよｗｗ
- 501 名前：デフォルトの名無しさん [2010/01/14(木) 11:44:30 ]
- cutil64.dllってどこに置けばいいの？
- 502 名前：デフォルトの名無しさん [2010/01/15(金) 05:05:43 ]
- GPU側で配列の合計を求める方法を教えてください
- 503 名前：デフォルトの名無しさん mailto:sage [2010/01/15(金) 07:59:04 ]
- サンプルのreduction見れ
- 504 名前：デフォルトの名無しさん [2010/01/15(金) 09:52:56 ]
- syncthreadsって異なるブロック間のスレッドも同期されるのでありましょうか？
- 505 名前：デフォルトの名無しさん mailto:sage [2010/01/15(金) 20:54:27 ]
- >>504
programing guide よめ
- 506 名前：デフォルトの名無しさん [2010/01/15(金) 23:10:51 ]
- Best Practices Guide
もいいこと盛りだくさん書いてあるよ
- 507 名前：デフォルトの名無しさん mailto:sage [2010/01/16(土) 19:28:55 ]
- cudaでは多次元配列の使用は推奨されてないの？
- 508 名前：デフォルトの名無しさん mailto:sage [2010/01/17(日) 12:17:08 ]
- pthreadとの組み合わせでうまく動作せずに困っています。
引数として渡す構造体に、float * device_Xを用意して
for(int k=0;k<NUM_THREADS;k++){
cudaMalloc((void**)&targ[k].device_X,
sizeof(float)*(hostX.nRow)*(hostX.nCol));
cudaMemcpy((void*)targ[k].device_X,
host_X,
sizeof(float)*(hostX.nRow)*(hostX.nCol),
cudaMemcpyHostToDevice);
pthread_create(&handle[k], NULL,
compute_thread_func,
(void *)&targ[k]);
}
と渡すと、pthread内で
cudaMemcpy(temp,argument->device_X, //argumentは渡されたvoid* argをキャストしたもの
sizeof(float)*(nRow)*(nCol),
cudaMemcpyDeviceToHost);
としてもうまく取り出せません。（中身がぐちゃぐちゃになる）
同様の処理をpthreadで呼び出す関数内で行うとうまく動作します。
スレッドごとにメモリを取るのだから、中でやってもいいのですが気になります。
何か原因になっていそうなことは無いでしょうか。
- 509 名前：デフォルトの名無しさん mailto:sage [2010/01/17(日) 12:59:47 ]
- >>508
＃1個のGPUに複数のスレッドからアクセスしたいアプリケーションなのかよく分かりませんが・・・・。
CUDAはスレッド毎にGPUのリソース（コンテキスト）を管理しています。
従ってこの例では
子スレッドのcompute_thread_func内の処理は
親スレッドがcudaMallocした領域にアクセスできません。
- 510 名前：デフォルトの名無しさん mailto:sage [2010/01/17(日) 13:32:14 ]
- >>509
ありがとうございます。
複数のスレッドから一台のGPUへのアクセスです。
（後々、スレッドごとに一個のGPUにしたいのですが）
スレッドごとに管理しているのははじめて知りました。
解決しました。
- 511 名前：デフォルトの名無しさん mailto:sage [2010/01/18(月) 20:51:33 ]
- 「Picture Motion Browser VAIO Edition」を追加。
超解像HDアップコンバートやノイズ除去、手ぶれ補正などの編集や動画作成を、NVIDIA CUDAによるGPU演算を使って行なえる。
- 512 名前：デフォルトの名無しさん mailto:sage [2010/01/18(月) 22:00:36 ]
- GT200 Fermi
トランジスタ数 14億 30億
倍精度浮動小数点 30FMA Oｐｓ/clock 256FMA Oｐｓ/clock
単精度浮動小数点 240 512
シェアドメモリ(SM) 16KB 64KB
L1メモリ 無し 48KB or 16KB
L2メモリ 無し 768KB
アドレス幅 32bit 64bit
複数のカーネル実行
無し 16個まで（GigaThredエンジン)
＊L1キャッシュ搭載
GT200では16KBのSMが32個あり、それぞれCUDAコアを8個割り当てられ、L1キャッシュが無し。
Fermiでは64KBのSMが16個。それぞれにCUDAコアが32個割り当てられ、SMから16KBか48KBのどちらかをL1キャッシュとして使用可能。
GT200に対して、3倍のSMと16KBのL1が使用可もしくは同じサイズのSMと48KBのL1が使用できるようになった。これにより、今までCUDA化できなかったプログラムの対応を増やし、さらに高速化もできる。
各CUDAコアに含まれるFPユニットは倍精度浮動少数演算を強化し、GT200に対し8倍の能力。
＊L2メモリの搭載 グローバルメモリに対する高速化。
＊C++をサポート
＊複数のカーネルの動作をサポート
SM内部のパイプラインを強化。SFUが複数に分けられたのでタスクとデータをより効率化。スレッドスケジューラを2個。
＊双方向PCIE通信
GT200ではPCIEバスの送受信をどちらか片方しか一度に実行できず、理論値8GB/s・実測4〜5GB/s程度だが
Fermiでは双方向通信が可能になり12GB/sを実現
＊新しいメモリコントローラ FermiよりGDDR5まで対応し、ECCにも対応する。
＊コア内部を各部でモジュール化
設定された価格帯や用途ごとにコアを設計しなおさず、機能をカットオフしたり追加したりできる。
SM単位でCUDAのコアを減らしたり、D3DやOpenGLなどの固定ハードウェアが不要なTeslaなどでオフになったりする可能性もある。
- 513 名前：,,・´∀｀・,,）っ-○○○ mailto:sage [2010/01/19(火) 02:32:35 ]
- GT200 理想のFermi 現実のFermi
トランジスタ数 14億 30億 30億
倍精度浮動小数点 30FMA Oｐｓ/clock 256FMA Oｐｓ/clock 224FMA Oｐｓ/clock
単精度浮動小数点 240 512 448FMA Oｐｓ/clock
シェアドメモリ(SM) 16KB 64KB 〃
L1メモリ 無し 48KB or 16KB 〃
L2メモリ 無し 768KB 〃
アドレス幅 32bit 64bit 〃
複数のカーネル実行 無し 16個まで（GigaThredエンジン) 14個まで（GigaThredエンジン)
- 514 名前：デフォルトの名無しさん mailto:sage [2010/01/19(火) 04:04:04 ]
- OpenCL待ちしてたけど、GPUの対応はCUDAだけっていう開発ばっかな
Fermiで本気出すから
- 515 名前：デフォルトの名無しさん mailto:sage [2010/01/20(水) 02:23:16 ]
- GT100ってなんかメモリアクセスのレイテンシがすごくでかくない？
それを隠蔽するために、L2キャッシュが128Kになっているけど、
コヒーレンシとかどうなのかな？
やっぱりGPUはデータ並列なアプリしか向いていないのかね。
- 516 名前：デフォルトの名無しさん mailto:sage [2010/01/20(水) 16:04:40 ]
- Fermiのグローバルメモリのレイテンシが遅いってのは何処から来た情報？GDDR5に対応するんだから、帯域は大きくなりそうだけど。
- 517 名前：デフォルトの名無しさん mailto:sage [2010/01/20(水) 17:08:55 ]
- >>512
SFUは今のGT200でも複数ある。
ただ、SFUが外部のパイプラインに分かれたからSFUを使っているときに違う種類の演算ができるようになったってことじゃね・。
>>515
最適化しなければいけない
- 518 名前：デフォルトの名無しさん mailto:sage [2010/01/21(木) 03:41:00 ]
- >>516
これをみて思った。
pc.watch.impress.co.jp/img/pcw/docs/343/352/html/kaigai-01.jpg.html
でもGT200と一緒なんだね。
pc.watch.impress.co.jp/docs/2008/0617/kaigai_10l.gif
いずれにしても、何でGPUがあんなにレイテンシが大きいのかがわかったよ。
だから、coalesced accessが必要なのね。
- 519 名前：デフォルトの名無しさん mailto:sage [2010/01/21(木) 04:03:50 ]
- >>518
>いずれにしても、何でGPUがあんなにレイテンシが大きいのかがわかったよ。
俺のために説明してくれ！
- 520 名前：デフォルトの名無しさん mailto:sage [2010/01/21(木) 22:43:04 ]
- GPUのメモリのレイテンシが大きいというよりも、CPUの場合と同じぐらいはかかるといった方がいい
- 521 名前：デフォルトの名無しさん mailto:sage [2010/01/22(金) 01:50:01 ]
- CPUだとメモリチャンネルが精々2〜3チャネル位しか無いが、GT100だと12チャネルもあるから、DRAMのチャネルのスイッチングに
それ以上かかると思っていい。あと、仮にGPC1?がDRAM0チャネル上にあるデータと、
DRAM11チャネル上にあるデータを同時にアクセスする場合、当然レイテンシがかかる。
もし、GPC1がDRAM0、DRAM1だけのデータアクセスですむなら、CPUと変わらない。
最適化が必要というのはそういったところだと思う。
- 522 名前：デフォルトの名無しさん mailto:sage [2010/01/22(金) 05:43:22 ]
- MTGPがコンパイルできません。
>>463の言う、4つのファイル使って試したところ100個近くコンパイルエラーが出てしまいます。
あと「末尾を適当に書き換える。」の意味がよく分からんどす…
エラー内容は
1>\略)\mtgp32-fast.h(117) : error C2054: 'inline' の後に '(' が必要です。
1>\略)\mtgp32-fast.h(120) : error C2057: 定数式が必要です。
1>\略)\mtgp32-fast.h(120) : error C2466: サイズが 0 の配列を割り当てまたは宣言しようとしました。
1>\略)\mtgp32-fast.h(120) : error C2085: 'mtgp32_do_recursion' : 仮パラメータ リスト内にありません。
1>\略)\mtgp32-fast.h(121) : error C2061: 構文エラー : 識別子 'inline'
1>\略)\mtgp32-fast.h(122) : error C2054: 'inline' の後に '(' が必要です。
...
WinXP32bit、VC++2005、CUDA SDK2.3
該当箇所のコード見てもどこが悪いのか分からない…助けて…。
|
 |
