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

1 名前：デフォルトの名無しさん mailto:sage [2008/03/22(土) 11:13:52 ] このスレッドは、他のスレッドでは書き込めない超低レベル、 もしくは質問者自身何が何だが分からない質問を勇気を持って書き込むスレッドです。 CUDA使いが優しくコメントを返しますが、 お礼はCUDAの布教と初心者の救済をお願いします。 CUDA・HomePage www.nvidia.com/cuda 関連スレ 【GPGPU】NVIDIA CUDA質問スレッド pc11.2ch.net/test/read.cgi/tech/1190008468/ GPUで汎用コンピューティングを行うスレ pc11.2ch.net/test/read.cgi/tech/1167989627/ GPGPU#2 pc11.2ch.net/test/read.cgi/tech/1188374938/ 15 名前：デフォルトの名無しさん mailto:sage [2008/04/01(火) 08:47:11 ] CUDAでトリップ検索したら早いの １２８あれば１２８倍早そうな気がするんだけど すぐもれちゃいそうで怖いんだけど 16 名前：デフォルトの名無しさん mailto:sage [2008/04/01(火) 08:48:12 ] ↑早いの？です 17 名前：デフォルトの名無しさん mailto:sage [2008/04/01(火) 09:57:13 ] >>15 (比較的)単機能の(必ずしも)早くない演算器が(現在の最上位機種でも)256個あるだけなので、 単純に256倍速くなるわけではありません。 実際、4コアXeonよりも速いプログラムを作るのに苦労しているくらいですから。 18 名前：デフォルトの名無しさん mailto:sage [2008/04/03(木) 16:05:16 ] AIも256倍賢くなるのかな？ 19 名前：デフォルトの名無しさん mailto:sage [2008/04/04(金) 15:15:42 ] コンパイルすると Parsing error near '[': syntax error というエラーが出ます。何がいけないのでしょうか？ 20 名前：デフォルトの名無しさん mailto:sage [2008/04/04(金) 15:28:37 ] [の近くの構文がいけないのでしょう 21 名前：デフォルトの名無しさん mailto:sage [2008/04/04(金) 17:26:04 ] >>17 なるほど夢見すぎでした�d 22 名前：デフォルトの名無しさん [2008/04/06(日) 09:35:45 ] CUDAのプログラムを配布する場合、何を含めればよいのですか？ 23 名前：デフォルトの名無しさん mailto:sage [2008/04/06(日) 09:52:28 ] >>22 NVIDIAの対応ドライバとCUDAToolKitをインストールしてもらう前提で、実行モジュールだけ。 ToolKitのDLLが同梱できるかは不明。 24 名前：デフォルトの名無しさん mailto:sage [2008/04/06(日) 12:39:02 ] nvccの-Oオプションて数字はいくつまであるの？gccと同じ-O3まで？ 25 名前：デフォルトの名無しさん mailto:sage [2008/04/06(日) 15:28:30 ] gccにそのまま渡すだけだと思うから、きっと3。 26 名前：デフォルトの名無しさん mailto:sage [2008/04/06(日) 16:11:04 ] CUDAコードは最適化されないのですか？ 27 名前：デフォルトの名無しさん mailto:sage [2008/04/06(日) 16:14:09 ] されますよ。もう、徹底的に。 # まぁ、それほどドラスティックではないけれど。 28 名前：デフォルトの名無しさん mailto:sage [2008/04/07(月) 15:26:08 ] cudaの実行速度を最適化したい場合、Oオプション以外にどのようなオプションがありますか？ 29 名前：デフォルトの名無しさん mailto:sage [2008/04/07(月) 16:20:09 ] 頑張って自分で最適化してください。 30 名前：デフォルトの名無しさん mailto:sage [2008/04/07(月) 16:22:20 ] -O3 -use_fast_math くらいかな。私が使っているのは。 GPU側のロジックは、実は-Oの有無では変わらないかも知れないと思う。 31 名前：デフォルトの名無しさん [2008/04/09(水) 07:43:09 ] nvccのオプションで-O3を指定した場合、-Xcompilerで改めて指定する必要はないのですか？ 32 名前：デフォルトの名無しさん [2008/04/09(水) 16:50:10 ] ブロック数はどのように決めればよいのですか？ できるだけ多くするのがよいのでしょうか？ 33 名前：デフォルトの名無しさん mailto:sage [2008/04/09(水) 22:01:54 ] >>32 多くすれば良いと言う物でもないみたい 分割数を多くすることによるオーバーヘッドの方が大きくなるのか 良く分からないんだけど 34 名前：デフォルトの名無しさん [2008/04/10(木) 17:46:35 ] いくつぐらいが適当ですか？ 35 名前：デフォルトの名無しさん mailto:sage [2008/04/11(金) 01:13:42 ] もう一つのCUDAスレにヒントがあるので参考にどうぞ。 36 名前：デフォルトの名無しさん [2008/04/15(火) 06:31:17 ] 単精度を2つ使って倍精度の計算をする様なことはできませんか？ 37 名前：デフォルトの名無しさん mailto:sage [2008/04/15(火) 08:45:23 ] 実は現行CUDAにも一部倍精度演算があります。 が、まともに使うにはCUDAv2を待ちましょう。 38 名前：デフォルトの名無しさん mailto:sage [2008/04/16(水) 00:59:13 ] 今日やっと、CUDA_Profiler_1.1.txtの存在に気付いてProfilerをまともに使えるようになった。 環境変数で設定ファイル名を指定して、設定ファイルで出力項目を指定するなんてややこし過ぎる…… しかも、timestamp以外に4つまでの制限なんて、なんて半端な。 つーことで、こんな感じ。 -- cuda_profile.setting export CUDA_PROFILE=1 export CUDA_PROFILE_LOG=cuda_profile.csv export CUDA_PROFILE_CSV=1 export CUDA_PROFILE_CONFIG=cuda_profile.config -- cuda_profile.config timestamp gld_incoherent gld_coherent # gst_incoherent # gst_coherent # local_load # local_store # branch # divergent_branch instructinos warp_serialize # cta_launched -- # もうCUDAスレのどっちがどっちだか判らなくなっているのは内緒。 39 名前：デフォルトの名無しさん [2008/04/16(水) 19:29:17 ] CUDAと心中しても大丈夫でしょうか？ 40 名前：デフォルトの名無しさん mailto:sage [2008/04/17(木) 00:46:36 ] 心中する覚悟をするのは勝手ですが、実際に心中されたら大迷惑です。 41 名前：デフォルトの名無しさん mailto:sage [2008/04/17(木) 14:43:17 ] CUDA 2.0(Beta) forums.nvidia.com/index.php?showtopic=64974&pid=363542&st=0&#entry363542 42 名前：デフォルトの名無しさん mailto:sage [2008/04/19(土) 13:25:57 ] OpenGL系がダメダメだね。Vista。 ただ、nbody.exeはなぜか動いている。 43 名前：デフォルトの名無しさん mailto:sage [2008/04/19(土) 15:53:33 ] > OpenGL系がダメダメだね。 付属のバイナリはダメだが。 ソースからビルドし直せば、動くぞ。 44 名前：デフォルトの名無しさん mailto:sage [2008/04/22(火) 00:19:42 ] GPGPU#2 pc11.2ch.net/test/read.cgi/tech/1188374938/ 188 名前：デフォルトの名無しさん[sage] 投稿日：2008/04/21(月) 19:17:35 Linuxはここだ forums.nvidia.com/index.php?showtopic=65067 どうでもいいけど明後日の仕事が増えた(謎 45 名前：デフォルトの名無しさん mailto:sage [2008/04/23(水) 23:13:45 ] 質問スレッドっていきてるの？ 46 名前：デフォルトの名無しさん mailto:sage [2008/04/23(水) 23:17:54 ] たった今回答して来ましたが何か。 47 名前：デフォルトの名無しさん [2008/04/29(火) 19:05:47 ] Visual C++のバージョンはチェックしているみたいで、合っていないとコンパイルできませんが、 gccはバージョンが合ってないとコンパイルできないのでしょうか？ gccの最新版を入れても大丈夫なのでしょうか？ 48 名前：デフォルトの名無しさん mailto:sage [2008/04/29(火) 19:26:11 ] >>47 pc11.2ch.net/test/read.cgi/tech/1190008468/386 49 名前：デフォルトの名無しさん [2008/05/04(日) 01:23:12 ] ローカル変数は、どのメモリに格納されるのでしょうか？ sharedメモリでしょうか？ 50 名前：デフォルトの名無しさん mailto:sage [2008/05/04(日) 08:40:58 ] レジスター 51 名前：デフォルトの名無しさん mailto:sage [2008/05/09(金) 16:59:41 ] 2つのGPUを同時に使いたい場合、 for (int i=0;i<2;i++) {     CUDA_SAFE_CALL(cudaSetDevice(i));     somekernel<<< 64, 512>>>( ); } のような形で書けばよいのでしょうか？ それともこれだと、一つ目のカーネルが終わらないと2つめのカーネルの実行が開始されないのでしょうか？ 上記のように書いているのですが、実行速度が2倍になりません。 52 名前：デフォルトの名無しさん mailto:sage [2008/05/09(金) 23:36:57 ] >>51 スレッドを分けるか、CUDA1.1で追加されたストリーム機能を使う必要がありそうです。 後者の場合、こんな流れでしょうかね。 -- for (int ic = 0; ic < 2; ++ic) { cudaStreamCreate(& stream[ic]); cudaSetDeviece(ic); kernel<<<64, 512, 0, stream[ic]>>>(); } for (int ic = 0; ic < 2; ++ic) { cudaStreamSynchronize(stream[ic]); postProcess(ic); } cudaThreadSynchronize(); 53 名前：51 mailto:sage [2008/05/10(土) 08:54:52 ] >>52 ありがとうございます。 上記のようにやったのですが、 for (int ic = 0; ic < 2; ++ic) { の部分を for (int ic = 0; ic < 1; ++ic) { に書き換えると、実行時間がちょうど半分になるので やはり並列で実行されていないようです。 上記コード以外になにか付け加える必要があるのでしょうか。 54 名前：デフォルトの名無しさん mailto:sage [2008/05/10(土) 09:22:20 ] >>53 CUDA_SAFE_CALLマクロは毎回cudaThreadSynchronize()を呼ぶから要注意。 55 名前：51 mailto:sage [2008/05/10(土) 09:29:14 ] >>52さんの書かれたとおり、 for (int ic = 0; ic < 2; ++ic) { cudaStreamCreate(& stream[ic]); cudaSetDevice(ic); kernel<<<64, 512, 0, stream[ic]>>>(); } for (int ic = 0; ic < 2; ++ic) { cudaStreamSynchronize(stream[ic]); } cudaThreadSynchronize(); のようにしているのですが、逐次実行になってしまいます。 56 名前：51 mailto:sage [2008/05/10(土) 09:30:35 ] どこかにサンプルコードとかはないでしょうか？ 57 名前：デフォルトの名無しさん mailto:sage [2008/05/10(土) 11:34:03 ] cudaSetDevice()するとそのCPUスレッドで使うGPUデバイスが決定されるので、 それ以降の変更はできない模様。cudaSetDevice()しないでAPIを使うと、 無条件にDevice0を使うことにされてしまう模様。 CPUスレッドを分けるしかなさそうね。 NVIDIAのForumに行けばサンプルも転がっていそうだけれど…… 58 名前：57 mailto:sage [2008/05/10(土) 11:35:55 ] あー、書くの忘れた。CudaProgramingGuide(1.1)の4.5.2.2参照で。 59 名前：51 mailto:sage [2008/05/10(土) 12:35:11 ] ストリームだけではなんともならないということでしょうか？ 60 名前：57 mailto:sage [2008/05/10(土) 13:12:58 ] ならないっぽいですよ。ドキュメントを読む限り。 2.0betaのドキュメントも書き換わってないからそのままの悪寒。 61 名前：デフォルトの名無しさん [2008/05/10(土) 18:52:45 ] geforce9以降もCUDA対応してるんでしょうか？ 62 名前：デフォルトの名無しさん mailto:sage [2008/05/10(土) 20:00:46 ] >>61 勿論。 63 名前：51 mailto:sage [2008/05/11(日) 19:48:55 ] >>60 そうですか、そうなると結構面倒ですね。 ところで、スレッド分けで2つのGPUを使う場合、pthreadなどでも良いのでしょうか？ それとも、サンプルにあるような方法のみが認められているのでしょうか？ 64 名前：デフォルトの名無しさん mailto:sage [2008/05/11(日) 20:10:22 ] >>63 ちっとはcutilの中身も見ようよ。cutStartThread()は、Linux版ではpthread_create()を呼んでいる。 65 名前：デフォルトの名無しさん [2008/06/17(火) 19:19:34 ] GeForce GTX 280で倍精度を使うと、処理速度はどの程度落ちますか？ 66 名前：デフォルトの名無しさん mailto:sage [2008/06/17(火) 22:48:50 ] 単精度の半分ぐらいでしょ。JK 67 名前：デフォルトの名無しさん mailto:sage [2008/06/18(水) 00:54:10 ] いいえ、12分の1です。 68 名前：デフォルトの名無しさん mailto:sage [2008/06/18(水) 05:03:23 ] GTX280　78GFLOPS GTX260　62GFLOPS 69 名前：デフォルトの名無しさん mailto:sage [2008/06/20(金) 13:29:07 ] うーむ。GPU2枚使うのって実はコツとかいる？ ホスト側でOpenMP使って2スレッド走らせて、それぞれのスレッドに別のGPUを割り当てて るんだが、答えとしては正しいものが返ってくるんだが速くならない。 むしろ、ホスト1スレッドで1つのGPUで計算させたほうがいくらか速いくらい。 きちんとGPU2枚使ってそうだというのは確認できたんだが速くならない理由がわからない。 なんだかGPU0の処理が終わってからGPU1の処理が始まるとかやってそうな予感がしている。 GPUの切り替えの分だけちょっとだけ遅くなるというオチじゃないかと。。。 サンプルのmultigpuとか読むとホスト側のスレッド生成法は違うけど同じようなことやって るんだよなー。 2枚使って速くなったって人いたらどういう風にやったか教えてくれませんか？ かれこれ2週間くらい悩み中だ。 70 名前：デフォルトの名無しさん mailto:sage [2008/06/20(金) 15:38:55 ] >>69 pthreadでは速くなった？ 開発環境は？ 71 名前：69 mailto:sage [2008/06/20(金) 17:34:03 ] pthreadって要するにホスト側スレッド生成にcutStartThread()を使うってことだよね？ それは実はまだやっていないんだ。これからやってみようと思う。 とりあえずスレッド走らせればいいんだと思っていたから、使用経験のあるOpenMPを使った。 もしかして、cutStartThread()で生成したスレッドを使用して各GPUを使うっていうことが 暗黙の条件だったりしないよね。 cutStartThread()で速くなった人いますかね？ 開発環境はVisualStudio2005。 72 名前：デフォルトの名無しさん mailto:sage [2008/06/20(金) 17:40:20 ] うちでは速くなってるけど・・・ こちらはLinuxだけど、もしソースが載せられるならうｐしてみて。 73 名前：69 mailto:sage [2008/06/20(金) 18:41:08 ] >>72 ソースは諸事情のため色々と削ったりしたもの。 ちなみにCore2Duo＋GPU×2なのでホスト側スレッド数＝GPU数な環境。 float *h_data = 0; h_data = (float*)malloc(mem_size); (h_dataの中身をこしらえる) cudaGetDeviceCount(&num_gpus); int num_cpus = omp_get_num_procs(); omp_set_num_threads(num_cpus); #pragma omp parallel { unsigned int cpu_thread_id = omp_get_thread_num(); unsigned int num_cpu_threads = omp_get_num_threads(); unsigned int element_per_kernel = element / num_cpu_threads; float *sub_h_data = h_data + cpu_thread_id * element_per_kernel; int gpu_id = 0; CUDA_SAFE_CALL(cudaSetDevice(cpu_thread_id % num_gpus)); CUDA_SAFE_CALL(cudaGetDevice(&gpu_id)); printf("CPU thread %d (of %d) uses CUDA device %d\n", cpu_thread_id, num_cpu_threads, gpu_id); 74 名前：69 mailto:sage [2008/06/20(金) 18:42:06 ] unsigned int mem_per_kernel = mem_size / num_cpu_threads; float *d_data = 0; CUDA_SAFE_CALL(cudaMalloc((void**) &d_data, mem_per_kernel)); dim3 dimBlock(num_thread); dim3 dimGrid(num_block/num_cpu_threads); for(i=0; i<T; i++){ #pragma omp single { total = 0; for (j=0; j<element; j++) total += h_data[j]; (h_dataの値を色々いじる) } CUDA_SAFE_CALL(cudaMemcpy(d_data, sub_h_data, mem_per_kernel, cudaMemcpyHostToDevice)); calculate<<< dimGrid, dimBlock >>>(d_data, total); CUDA_SAFE_CALL(cudaMemcpy(sub_h_data, d_data, mem_per_kernel, cudaMemcpyDeviceToHost)); } } 諸事情によりomp singleを使っている場所があるんだがこいつが悪さをしているとかあるのかな。 75 名前：デフォルトの名無しさん [2008/06/27(金) 02:51:33 ] 質問です（全く素人です） OpenGLを使えばGPGPUのプログラミングが出来ると聞いたのですが、 かたやCUDAなどの専用の言語が出てきているのはどうしてでしょうか？ 76 名前：デフォルトの名無しさん mailto:sage [2008/06/27(金) 03:13:39 ] だったらcとか使わずfortran66でも塚っテロ 77 名前：デフォルトの名無しさん mailto:sage [2008/06/27(金) 12:11:54 ] >>76 いえ、ですから、OpenGLではGPUの使えない機能もあるのかなぁ、と思いまして。 78 名前：デフォルトの名無しさん mailto:sage [2008/06/27(金) 16:55:21 ] GPUで計算中でも画面は通常通り表示されますか？ 79 名前：デフォルトの名無しさん mailto:sage [2008/06/27(金) 22:56:10 ] >>78 ハードについて勉強するよろし 80 名前：デフォルトの名無しさん mailto:sage [2008/06/28(土) 00:16:42 ] CUDA実行中に、時たま怪しげな影が飛んだりフリーズしたりします。 81 名前：デフォルトの名無しさん mailto:sage [2008/06/29(日) 11:59:05 ] Folding@HomeのGPU版をしている人居ます？ PS3の280GTXは6倍以上をこなしているようですが。 ベータ版が公開されているが時間がなくてまだ何もしていない（困った。） 82 名前：デフォルトの名無しさん mailto:sage [2008/06/29(日) 12:04:52 ] >>81 マルチすんな。つーか、PS3の280GTXってなんだよpgr 83 名前：デフォルトの名無しさん [2008/06/30(月) 07:04:03 ] >>69 OpenMPの使い方が根本的に間違っていると思う 84 名前：69 mailto:sage [2008/07/01(火) 22:08:51 ] >>83 サンプルのcudaOpenMPを参考にしてやったんだけどなぁ。 これじゃどうだめなのかぜひとも教えてくれ。 ちなみにOpenMPの使い方はサンプルにあわせている。 ただsingle構文はサンプルでは使っていないがな。 85 名前：デフォルトの名無しさん mailto:sage [2008/07/02(水) 04:05:38 ] VisualStudio2005はExpress？だったらOpenMPディレクティブは機能しないよ 86 名前：デフォルトの名無しさん [2008/07/02(水) 15:57:17 ] nvidiaドライバを削除する方法ってどこかにのってないか＠CentOS5 新しいグラボ買ったんだけどドライバの削除がわからないウンコー 87 名前：デフォルトの名無しさん [2008/07/02(水) 16:15:29 ] grid, threadの意味がわからない アニメに例えて教えろ 88 名前：デフォルトの名無しさん [2008/07/02(水) 21:30:09 ] 今amazonみたらGPUGems3の日本語版の予約を受け付けてるじゃｎ 89 名前：デフォルトの名無しさん [2008/07/03(木) 01:44:22 ] ノート（OS:Vista）で動かそうとしても動きません CUDAに対応したGPUを搭載しておりますし（8700M GT） VC2005Expressをインストールしたのですが、コンパイルすると nvcc fatal: nvcc cannot find a supported cl version. Only MSVC 7.1 or 8.0. と表示してしまいます。まったくわからなくて困ってます。 よろしければお教えください 90 名前：デフォルトの名無しさん mailto:sage [2008/07/03(木) 01:47:35 ] 2005以外のバージョンのVCも入れてる？ nvccが参照しているVCが2005じゃないように見える。 91 名前：デフォルトの名無しさん [2008/07/03(木) 10:56:29 ] いれておりません。 インストールしたのは VC2005EXとVC2005SP、VC2005SP Update for Vista だけなのですが。 92 名前：69 mailto:sage [2008/07/03(木) 13:12:53 ] >>85 VisualStudio2005Professionalを使っています。 とりあえずcuda使わないで純粋にCでOpenMP使ったら高速化されているのでOpenMP自体は動いている模様。 93 名前：デフォルトの名無しさん mailto:sage [2008/07/04(金) 21:25:01 ] ttp://vipprog.net/wiki/CUDA.html 94 名前：デフォルトの名無しさん mailto:sage [2008/07/08(火) 02:11:24 ] >>93 なにこの糞サイト 95 名前：デフォルトの名無しさん mailto:sage [2008/07/08(火) 02:42:51 ] 質問です PhysXがGeforceでも使えるということを聞いたのですが、 PhysX SDKとCUDAやOpenCLというのは同じレイヤーでの話なのでしょうか？ 96 名前：デフォルトの名無しさん mailto:sage [2008/07/08(火) 05:08:31 ] >>84 pthreadでは高速化された？ 97 名前：デフォルトの名無しさん mailto:sage [2008/07/08(火) 09:16:41 ] >>88 アノ値段じゃ学生はp-子するから売れん罠 98 名前：デフォルトの名無しさん mailto:sage [2008/07/08(火) 11:01:42 ] Gem3はGPGPUに限って言うならあんまり役に立たなかった 99 名前：デフォルトの名無しさん [2008/07/08(火) 19:02:12 ] >>86 --uninstall 100 名前：86 [2008/07/09(水) 13:42:42 ] >>99 ありがトマト ところでSobelFilter見ててShareMemの使い方がわからないんだが フォーラムに質問と返答あったが内容わからん＋英語で全然わからん 誰か全訳してくれ 101 名前：86 [2008/07/09(水) 13:44:25 ] forums.nvidia.com/index.php?showtopic=52612&pid=286191&mode=threaded&show=&st=&#entry286191 誰か誰か！ 102 名前：デフォルトの名無しさん mailto:sage [2008/07/10(木) 12:22:28 ] >>86 訳すのはめんどいからパスだけど、ポイント絞って質問してくれたら回答するよ。 103 名前：デフォルトの名無しさん mailto:sage [2008/07/10(木) 12:28:53 ] >>86 C:\WINDOWS\system32\drivers から目的のドライバを手で削除すればよい。 104 名前：デフォルトの名無しさん [2008/07/13(日) 21:12:38 ] >>102 神ktkr ポイント絞るからちっと待ってて… 105 名前：102 mailto:sage [2008/07/14(月) 11:56:25 ] 問題は、自宅が規制に巻き込まれていてなかなか書き込めない辺り。 取り敢えず、注意点を列挙しておく。 ・共有メモリを確保するのはglobalFunc<<<blocks, threads, sharedMemorySize, streamNo>>>(parameters)の 　三番目のパラメータでサイズが指定されたときだけ。 ・共有メモリは一回の<<<>>>の呼び出しの間だけしか有効じゃない。 # つまり、次の回には残っていない。 ・共有メモリをハンドリングするには、extern __shared__ anyType * nameで宣言するだけ。 # つまり、コンパイラは型のマッチングやサイズのチェックをしないので自分で管理しないといけない。 ・共有メモリはblock間で独立、block内ではthread数に関わらず共有。 # つまり、実際のデバイスにそぐわないthread数を指定した場合はCUDA側で同期処理が入るのか、遅くなってしまう。 ・あるthreadが共有メモリに書いた後、別threadが読む前には__syncthreads()で同期を取らないといけない。 # ある意味当然なんだけど、その所為で遅くなるのも事実。 あー、ついでにメモリの違いを簡単に。 ※グローバルメモリ ・読み書きできる。coalescedにアクセスできれば結構速い。消えない。広い。 ・ホスト側スレッドごとに独立している。ホスト側から見ると、毎回同じアドレスになるのでどのくらい使えるか判りにくい。 ※共有メモリ ・読み書きできる。遅くない。 ・呼び出しごとに消えてしまう。余り広く取れない。事実上同期を取る必要があって使い難い。 ※定数メモリ ・速い。消えない。そこそこ広い。 ・例えばfloat2を読み込むインストラクションがないので実はグローバルメモリからfloat2を読むより遅くなる場合もある。 ・デバイス側から書き込めない。ホスト側スレッドごとに独立している。複数スレッドからCUDAを使うと毎回転送しなおすのか? ※レジスタ ・読み書きできる。速い。厳密に型チェックされる。つーか、型ごとに違うインストラクションが使われるからptxファイルで追える。 ・呼び出しごとに消える。他のメモリに較べれば狭い。 # ローカルな配列は宣言したことないからよく判らん。 106 名前：デフォルトの名無しさん mailto:sage [2008/07/14(月) 19:10:53 ] NVIDIAの仕様書見てもcoalescedの意味がいまいちわからないのだけど、どういうこと？ どこかわかりやすく解説しているサイトない？ 107 名前：デフォルトの名無しさん mailto:sage [2008/07/14(月) 19:12:13 ] カーネル内で__shared__つけて配列を宣言するのと、つけないで配列を宣言するのでは何が違うの？ 108 名前：デフォルトの名無しさん mailto:sage [2008/07/14(月) 19:24:11 ] >>106 GPUはWarp単位で同じインストラクションが走るから、要は16人17脚みたいに考えればいい。 メモリアクセスを16人17脚によるパン喰い競争みたいに考えると、自分のパンが目の前にある状態がcoalesced。 そのとき、2,3人パンを喰う必要がなくてもスルーするだけだから問題ない。 処が、二人のパンが入れ違っていたらそこで入れ替える間、みんなが待たされることになるって感じ。 # 判っている人には判るけど、判っていない人には判らない説明だなw >>107 共有メモリを使うかどうか違うだけだと思うけど。ptx出力させて較べてみたら? 109 名前：デフォルトの名無しさん mailto:sage [2008/07/15(火) 00:26:16 ] >>108 その説明、いただいてもいいですか？ 110 名前：デフォルトの名無しさん mailto:sage [2008/07/15(火) 01:56:27 ] >>109 本にするならもっと書かせてくれw Vipのwikiに載せるなら是非やってくれ 金取って講習するのに使うのなら分け前よこせw 111 名前：デフォルトの名無しさん [2008/07/15(火) 07:52:19 ] >>105 その通りですシェアードメモリとブロック数が理解できない 1. プロック数 dim3 blocks = dim3(iw/(4*BlockWidth)+(0!=iw%(4*BlockWidth)), ih/threads.y+(0!=ih%threads.y)); なぜblocks.xはiw/threads.x+(0!=iw%threads.x)じゃなくて 上の式になるのか。 2. シェアードメモリ int SharedPitch = ~0x3f&(4*(BlockWidth+2*Radius)+0x3f); int sharedMem = SharedPitch*(threads.y+2*Radius); SharedPitchはなぜ上の計算になるのか。 0x03fの意味、4*の意味、BlockWidth+2*Radiusの意味が理解できない とりあえずSharedメモリの使い方はどこを調べればわかるんだ！ｳﾝｺ！ 112 名前：デフォルトの名無しさん mailto:sage [2008/07/15(火) 08:45:54 ] >>111 どんな数字を入れるとどんな結果になるか、Excelでも使って計算してみたら? 0x3f使うのは64の倍数にするためでしょ。 113 名前：デフォルトの名無しさん [2008/07/15(火) 16:57:48 ] >>112 うーんそのなんで６４にするのかがわからないのよ 関係ないかもしれんがシェアードメモリを使ってないSobelFIilterも SobelTex<<ih, 384>>>でなぜ384かわからないお こっちは１行１グリッドにして、１スレッド１ピクセルなんだと思うが なんでスレッド数を画像の横幅iwにしないで384にするんだぁ スレッド数が384だと計算が速い理由でもあるのかお！ 114 名前：デフォルトの名無しさん mailto:sage [2008/07/16(水) 00:40:50 ] 単純に、warp数の適当な倍数になるからってだけじゃなかろか。 115 名前：デフォルトの名無しさん mailto:sage [2008/07/16(水) 18:09:02 ] GPUコードではmemcpyは使えないの？

---

---

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

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