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

1 名前：デフォルトの名無しさん [2011/08/23(火) 22:08:06.09 ] このスレッドは、他のスレッドでは書き込めない超低レベル、 もしくは質問者自身何が何だが分からない質問を勇気を持って書き込むスレッドです。 CUDA使いが優しくコメントを返しますが、 お礼はCUDAの布教と初心者の救済をお願いします。 CUDA・HomePage ttp://developer.nvidia.com/category/zone/cuda-zone 関連スレ GPGPU#5 ttp://hibari.2ch.net/test/read.cgi/tech/1281876470/ 前スレ 【GPGPU】くだすれCUDAスレ【NVIDIA】 ttp://pc12.2ch.net/test/read.cgi/tech/1206152032/ 【GPGPU】くだすれCUDAスレ pert2【NVIDIA】 ttp://pc12.2ch.net/test/read.cgi/tech/1254997777/ 【GPGPU】くだすれCUDAスレ pert3【NVIDIA】 ttp://hibari.2ch.net/test/read.cgi/tech/1271587710/ 【GPGPU】くだすれCUDAスレ pert4【NVIDIA】 ttp://hibari.2ch.net/test/read.cgi/tech/1291467433/ 267 名前：デフォルトの名無しさん mailto:sage [2012/01/05(木) 15:37:13.27 ] おれも64/32の違いかと思ったが、パス指定から指摘してみた 実行している環境は？ -m64 オプション付けてみ 268 名前：デフォルトの名無しさん [2012/01/06(金) 13:53:42.06 ] # nvcc -L /usr/local/cuda/lib64/ -I /usr/local/cuda/include/ Drawtest.cu -lglut -m64 /usr/bin/ld: skipping incompatible /usr/local/cuda/lib64//libglut.a when searching for -lglut /usr/bin/ld: skipping incompatible /usr/local/cuda/bin/../lib64/libglut.a when searching for -lglut ううむ・・・・。 centOS 64bit Device 0: "Quadro 4000" Device 1: "GeForce 9500 GT" パスは LD_LIBRARY_PATH = /usr/local/cuda/lib64 です。 openGLのないコードはコンパイル及び実行まで問題ありません。 openGLを使うにあたって /usr/local/cuda/include/にglut.h freeglut.h freeglut_ext.h freeglut_std.h /usr/local/cuda/lib64/にlibglut.a をコピペしましたが、まだ足りないのでしょうか？ 269 名前：デフォルトの名無しさん mailto:sage [2012/01/06(金) 13:59:40.26 ] libglutは64ビット？ 32ビットなんじゃねか？ 270 名前：デフォルトの名無しさん [2012/01/06(金) 15:03:44.65 ] glut.h freeglut.h freeglut_ext.h freeglut_std.h libglut.aをccが見ているであろう先のものと置き換えてしてみました。 # nvcc -L /usr/local/cuda/include/ -I /usr/local/cuda/include/ Drawtest.cu -lglut -m64 /tmp/tmpxft_000030a4_00000000-12_Drawtest.o: In function `disp()': tmpxft_000030a4_00000000-1_Drawtest.cudafe1.cpp:(.text+0x14e6): undefined reference to `glClearColor' tmpxft_000030a4_00000000-1_Drawtest.cudafe1.cpp:(.text+0x14f0): undefined reference to `glClear' /usr/local/cuda/bin/../lib64/libglut.a(libglut_la-freeglut_init.o): In function `glutInit': … (中略) … (.text+0x898): undefined reference to `glPopClientAttrib' collect2: ld はステータス 1 で終了しました openGL系の関数が読まれていないみたいです。 271 名前：デフォルトの名無しさん mailto:sage [2012/01/06(金) 15:37:17.18 ] いや、だからfreeglutは64ビットを使用してるのかと聞いてるんだが 質問に対しきちんと回答できないようでは教えようがない罠 272 名前：デフォルトの名無しさん [2012/01/06(金) 18:39:39.36 ] >>271 freeglutは64ビットのものを入れ直してやってみましたが、同じ結果になりました。 >>257 そこでCUDAのヘッダフォルダとライブラリフォルダに入れていたものを消し、-Lと-Iで直接freeglutのフォルダを指定したら 解決しました。 ライブラリが足りなかったのか、それともCUDAと一緒くたにしたのがダメだったのか解かりませんが、ちゃんとPATH通せってことみたいです。 お騒がせしました。 273 名前：デフォルトの名無しさん mailto:sage [2012/01/06(金) 19:45:07.40 ] >>262 だよねぇ・・・ 274 名前：デフォルトの名無しさん [2012/01/06(金) 20:14:34.67 ] 環境 Windows7 Professional 64bit Microsoft Visual C++ 2010 Express Version 10.0.40219.1 Microsoft .NET Framwork Version 4.0.30319 GeForce GTX 580 CUDA Toolkit 4.0.17 SDK 4.0.19 devdriver_4.0_winvista-win7_64_270.81_general このサイトを参考に環境を構築しました。 feather.cocolog-nifty.com/weblog/2011/07/visual-studio-2.html そして以下のサイトのサンプルプログラムを実行してみました。 www.gdep.jp/page/view/218 Hello,Worldと99 bottles of beerはcpu、gpu共に実行できました。 しかし、Matrixのプログラムはcpuの方は実行できるのですがgpuの方が実行できません。 以下のエラーを吐きます。 matrix_gpu.cu(5) : fatal error C1083: Cannot open include file: 'cutil_inline.h' : No such file or directory どうやらVisualStudioのパスがうまく通ってないということまでわかり、以下のサイトなどを参考にCUDA_INC_PATHなどを変えてみましたが、一向に変わりません。 d.hatena.ne.jp/Levi/20090921/1253535802#c SDK内のcutil_inline.h自体をtoolkitのincフォルダにコピペすると、他の.hファイルもいくつか同じエラーが出たのでエラーになったものをすべてコピペしたところ、LNK2019"link.exe"というエラーで先に進めませんでした。 一度VisualStudioを再インストールしてみましたが、状況は変わりません。 Nvidia GPU computing SDK Browserではサンプルプログラムを実行できているので、CUDAの環境は整っていると思われます。 どうすれば解決できますでしょうか・・・。かれこれ１週間以上格闘しています。 VisualStudioは2010よりも2008にした方がいいでしょうか？ 275 名前：デフォルトの名無しさん mailto:sage [2012/01/06(金) 20:23:04.24 ] C初心者にはきついと思うんだが… とりあえず 'cutil_inline.h'のある場所を見つけて そこを-I　/'cutil_inline.hのある場所'と指定する 意味わからなければCのコンパイルを勉強すること 276 名前：デフォルトの名無しさん mailto:sage [2012/01/06(金) 20:39:53.25 ] 274ではありませんが､「Cのコンパイル」の勉強にオススメの書籍とかありましたら紹介していただけるとうれしいです 277 名前：デフォルトの名無しさん mailto:sage [2012/01/06(金) 20:55:31.81 ] 275だけど、お薦めは他の人に任せるけど 一つ言えるのはCUDAはC/C++より数段難しい C/C++の質問をしてるようでは先が思いやられる、と正直思うんだが… 超低レベルな質問もOKなくだすれだけど、その程度は自力で調べる能力が必要ではないだろうか？ 278 名前：デフォルトの名無しさん mailto:sage [2012/01/06(金) 21:22:02.03 ] >>275 >>274です。ありがとうございます。 今実行できる環境にないので後で試してみます。 また結果報告しにきます。 279 名前：デフォルトの名無しさん [2012/01/07(土) 12:14:43.10 ] 一時期は同価格帯のCPUとGPUを比較してGPUは５０倍の性能があるって言われてたけど 今１０〜１５倍程度しか無くない？ 3930k ５万円 158GFlops GTX590 ６．３万円 2480GFlops GTX 580 ３．８万円 1500GFlops 次に出ると言われてるGTX700系は500系の２倍の性能らしいけどそれでも５０倍には遠く及ばない。 この「同価格帯でGPUはCPUの１０〜１５倍程度の性能しかない」という評価は正しいだろうか？ 280 名前：デフォルトの名無しさん mailto:sage [2012/01/07(土) 12:43:15.62 ] >>279 それは単にCPUは1コアとして比較しているからGPUを使えばうん十倍と言われるんだよ。 単純にメモリ帯域で3〜4倍、演算器としては、単純にAVXの8並列x6コアxクロック2倍で比較すると5倍にしかならない。 CPUは大きなキャッシュがあるから局所性の高い処理の場合はメモリ帯域の差が縮まるので、10倍差がつくならいい方。 もちろんアプリによるけど、CPU向けのコードがカリカリに最適化されていた場合は思ったほど差がつかない印象。 それでも3〜5倍位は早くなるからいいんだけどね。 問題はメンテナンスだな。。。CUDAのみと割り切れればいいんだけど、なかなかそうは行かないからなあ。 281 名前：デフォルトの名無しさん mailto:sage [2012/01/07(土) 16:09:01.15 ] >>279 GTX580で500GFLOPS程度。そのデータはおかしい CPUなんてE8500で25GFLOPSくらいだから20倍程度だな。 もちろん最近のコア多いやつならもっと差は縮まるだろうな。 それに理論値の話だから多数のスレッドでまわしたらコンテキストスイッチの オーバーヘッドとかGPUプログラミングは難しくて効率上げるの難しいとか そういうので実効は変わってくるんじゃね？ 282 名前：デフォルトの名無しさん mailto:sage [2012/01/07(土) 16:14:51.24 ] あ、同一価格帯か、E8500が2万しなかったことを考えると 580にそろえて2倍したとして10倍程度だから、まあおおざっぱには 「同価格帯でGPUはCPUの１０〜１５倍程度の性能しかない」は正しいんじゃないかと思う。 580と同時期の4万弱のCPUならもっと差は縮まるだろうし。 283 名前：デフォルトの名無しさん mailto:sage [2012/01/07(土) 16:22:20.40 ] 32nmと40nmで比較とか 284 名前：デフォルトの名無しさん mailto:sage [2012/01/07(土) 21:34:45.46 ] void *p; CUDA_SAFE_CALL(cudaSetDevice(0)); CUDA_SAFE_CALL(cudaMalloc(&p, 16)); でcudaMallocからcudaError_enum例外が飛ぶのですが 原因はなにが考えられますか？ ちなみにSDKのサンプルは動きます。 285 名前：284 mailto:sage [2012/01/07(土) 22:23:52.14 ] 異なるディレクトリに複数のCUDAのライブラリが入ってて 古いバージョンが参照されてるのが原因だった・・・ 286 名前：デフォルトの名無しさん mailto:sage [2012/01/08(日) 14:34:35.46 ] CUDA4.1でcomputeProfはnvvcになったんですか？ だいぶ代わってしまって驚きますた 287 名前：デフォルトの名無しさん mailto:sage [2012/01/08(日) 14:35:50.30 ] nvvp 288 名前：デフォルトの名無しさん mailto:sage [2012/01/08(日) 20:50:11.30 ] >>282 GPUの論理性能ってFMAとかMADとかの積和で倍になっているから、 GTX 580 1500GFlopsとかって、それらを使わなければ750GFlopsとかになるんだよな。 だから実効れべるではもっと縮まるんだろうな。それでも5倍以上差がでるとかなり有効なんだけど。 289 名前：デフォルトの名無しさん mailto:sage [2012/01/09(月) 10:37:28.98 ] 個人で買うやつも居るのかね？ ttp://page16.auctions.yahoo.co.jp/jp/auction/u38857709 290 名前：デフォルトの名無しさん mailto:sage [2012/01/09(月) 12:04:44.39 ] いる。というか俺が買う。 291 名前：デフォルトの名無しさん mailto:sage [2012/01/09(月) 13:09:16.68 ] こういう需要が限られていてかつ堅くて高い製品はamazonに出した方が高く売れるんじゃないかねえ。 292 名前：デフォルトの名無しさん mailto:sage [2012/01/09(月) 18:01:50.03 ] >>275 274です。 -Iオプションを付けてヘッダファイルのある場所を指定してみましたが、相変わらず同じエラーです。 C:\cuda_practice\test>nvcc -o matrix_gpu.exe matrix_gpu.cu -I C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.0\C\common\inc matrix_gpu.cu matrix_gpu.cu(5) : fatal error C1083: Cannot open include file: 'cutil_inline.h': No such file or directory 自分でも試行錯誤してみますが周りにcudaがわかる人がいなくて行き詰ってます・・・。 293 名前：デフォルトの名無しさん mailto:sage [2012/01/09(月) 18:20:34.55 ] もっとコンパイル早くなんねーかなぁ… thrust::sortとか使うようにしただけでコンパイルに十数秒かかってイライラする 294 名前：デフォルトの名無しさん mailto:sage [2012/01/09(月) 20:40:26.69 ] 速いマシンでやる。 295 名前：デフォルトの名無しさん mailto:sage [2012/01/09(月) 20:45:40.42 ] nvccコンパイラをGPGPUで実装すればいい！ CUDAがオープンになるどさくさにまぎれてそういうのを作る人が出るかもしれないこともないかも 296 名前：デフォルトの名無しさん mailto:sage [2012/01/09(月) 21:57:43.85 ] コンパイラの動作は現状GPU向きとは言えないんじゃね？ GPU自体が単純な性能向上だけではない、思いもよぬ進化でもしない限り 297 名前：デフォルトの名無しさん mailto:sage [2012/01/09(月) 22:54:12.52 ] コンパイルなんて依存関係ありまくりで、並列処理が苦手な最たるものだからまずやる意味がないだろう。 CUDAにすれば何でも速くなると思ってんじゃね？ 298 名前：デフォルトの名無しさん mailto:sage [2012/01/10(火) 17:15:02.61 ] sharedメモリって同一WARP内でのやり取りだったら同期なしで大丈夫なのでしょうか？ たとえば全部で32スレッドで次のkernelを実行した場合、 WARP内で一斉にsharedメモリに書きに行くので同期しないで大丈夫かと 思ったのですが、実際にはうまくいきません。 globalメモリに読みに行く段階でコアレッシングが発生していないので それが原因なのでしょうか？ どなたか教えてください。 __global__ void kernel(float *g_v, float *g_x){ float x = 0.0f; int i = blockDim.x * blockIdx.x + threadIdx.x; __shared__ float s_v[32]; s_v[i] = g_v[i+i%3]; __syncthreads(); // これが必要かどうか？ x = s_v[(i+3)%32]; g_x[i] = x; } 299 名前：デフォルトの名無しさん mailto:sage [2012/01/10(火) 17:35:07.81 ] >>292 参照パスを ""で囲む♪ 300 名前：デフォルトの名無しさん mailto:sage [2012/01/10(火) 17:39:16.65 ] >>298 iが３２以上になることはないかい？？ ブロックの数は？？ 301 名前：274 mailto:sage [2012/01/10(火) 17:52:39.67 ] >>299 書き込みしてからそのミスに気付きました・・・。 それを直してもうまくPathが通らなかったので、CUDA4.0、VisualStudio2008でやることにしました。 302 名前：デフォルトの名無しさん mailto:sage [2012/01/10(火) 17:56:37.24 ] >>298 fence入れてないからじゃね。 303 名前：298 mailto:sage [2012/01/10(火) 17:59:16.44 ] >>300 ごめんなさい、書いたコードが一部誤ってました。 iが32以上になることがあります。 ブロック数は数100程度になります。 このときは、下のソースのようになると思うのですが、 やはりうまくいきません。 __global__ void kernel(float *g_v, float *g_x){ 　float x = 0.0f; 　int i = blockDim.x * blockIdx.x + threadIdx.x; 　__shared__ float s_v[blockDim.x]; 　s_v[threadIdx.x] = g_v[i+i%3]; 　__syncthreads();　　　// これが必要かどうか？ 　x = s_v[(threadIdx.x+3)%32]; 　g_x[i] = x; } 304 名前：デフォルトの名無しさん mailto:sage [2012/01/10(火) 18:12:32.43 ] >>303 コンパイル後のバイナリを見たわけじゃないから予測だが、 s_v[threadIdx.x] = g_v[i+i%3]; x = s_v[(threadIdx.x+3)%32]; g_x[i] = x; は 1.1:tmp=g_v[i+i%3]; 　　//グローバルアクセスなので遅い 1.2:s_v[threadIdx.x]=tmp;　// 2.1:x = s_v[(threadIdx.x+3)%32]; 3.1:g_x[i] = x; みたいに解釈される。 で、単純な式だからコンパイラは s_v[threadIdx.x]とs_v[(threadIdx.x+3)%32]; が必ず別のアドレスになる事を検知して 1.1 2.1 3.1 1.2のように命令を入れ替えてしまう。 だから__threadfence_block()がいるはず。 305 名前：298 mailto:sage [2012/01/10(火) 18:17:13.87 ] >>302 >>304 ありがとうございました。 試してみたいと思います。 306 名前：デフォルトの名無しさん mailto:sage [2012/01/10(火) 19:45:54.97 ] >> 305 だいたい　　x = s_v[(threadIdx.x+3)%32]; 　の前にすべてのs_v[??}が定義されていないといけないだろ？？ スレッドが３２に設定されているのならば　(threadIdx.x+3)%32　は０か１にしかならないが．．． 307 名前：デフォルトの名無しさん mailto:sage [2012/01/10(火) 21:10:42.27 ] >>301 -lcutil32 とその参照パスを追加すれば Visual Studio 2010 でもコンパイル，実行できる♪ ちなみに私の場合は nvcc matrix.cu -I "C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.0\CUDALibraries\common\inc" -L "C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.0\CUDALibraries\common\lib\Win32" -lcutil32 tmpxft_000024b8_00000000-3_matrix.cudafe1.gpu tmpxft_000024b8_00000000-8_matrix.cudafe2.gpu matrix.cu tmpxft_000024b8_00000000-3_matrix.cudafe1.cpp tmpxft_000024b8_00000000-14_matrix.ii C:\cuda> a Processing time: 1936.354858 (msec) Press ENTER to exit... 308 名前：デフォルトの名無しさん mailto:sage [2012/01/10(火) 21:12:58.21 ] ちなみに家のはGeForce GT 220なので非力．．．． 309 名前：306 mailto:sage [2012/01/10(火) 22:21:29.79 ] 間違えた 　↓ (threadIdx.x+3)%32　は０か１にしかならないが．．． ３２で割った余りであった．．． 310 名前：298 mailto:sage [2012/01/11(水) 09:16:02.50 ] >>306 >>309 いろいろ試してみたところ、とりあえず上記の問題は解決しました。 回答ありがとうございました。 311 名前：デフォルトの名無しさん mailto:sage [2012/01/11(水) 11:25:42.40 ] すごい。 見た感じ、ちょっと不安だったから 参考になったみたいでよかった。 312 名前：274 mailto:sage [2012/01/11(水) 19:30:44.93 ] >>307 VS2008でもうまくできませんでした。 再度VS2010を入れてみるも手順通りいかず、当然実行もできませんでした。 リカバリしたのでこれからもう一度2010でやってみます。もう心が折れそうです・・・。 ちなみに2010expressとprofessionalの違いで実行できないこととかあり得ますか？ 別のPCでprofessionalで試してみてくれた方がいて、うまくできたようなのですが・・・。 313 名前：デフォルトの名無しさん mailto:sage [2012/01/11(水) 20:43:20.68 ] >>312 どんなエラーがでますか？ そもそもnvccでコンパイルするのだから，Visual Studioは関係ないでしょう？ VisualStudioのないLinux上でもコンパイル，実行できるはず． もう一度やるのならばライブラリ lcutil32を付ければ良いかと（３２ビットならば） 314 名前：274 mailto:sage [2012/01/12(木) 01:51:21.33 ] >>313 VisualStudioであれば多少使ったことがあるので導入しました。 少なくともwindowsでは「nvccがVisualStudioのcl.exeを必要とする」と複数のサイトに書いてあるようです。 gpu.fixstars.com/index.php/Windows_Vista_%E3%81%ABCUDA%E9%96%8B%E7%99%BA%E7%92%B0%E5%A2%83%E3%82%92%E3%82%A4%E3%83%B3%E3%82%B9%E3%83%88%E3%83%BC%E3%83%AB%E3%81%99%E3%82%8B exth.net/~ohshima/wordpress/category/windows/ エラーはプロジェクトをビルドしたときに「CL.exeがないからビルドできない」というものでした。 そもそもcl.exeは32ビットのコンパイラらしいのですが、64ビットPCでも使うのでしょうか？ また以前、いい感じのところまでいったときのエラーは C:\cuda_practice\test>nvcc -o matrix_gpu.exe matrix_gpu.cu matrix_gpu.cu tmpxft_00000bc0_00000000-3_matrix_gpu.cudafe1.gpu （省略） symbol __imp_cutCheckCmdLineFlag referenced in function "void __cdecl __cutilExi t(int,char * *)" (?__cutilExit@@YAXHPEAPEAD@Z) matrix_gpu.exe : fatal error LNK1120: 6 unresolved externals でした。 どうやらリンカーの設定を見直せとのことだったのでwww.scribd.com/doc/66757447/を参考に VisualStudioでプロジェクトのプロパティから「リンカ > 入力 > 追加の依存ファイル」にcutil32.lib, cutil64.lib を追加。 しかし「cutil64.libが見つからない」とのエラー。 実際に探してみてもcutil64.libがどこにも見つかりませんでした。 今まで3回もリカバリしてやってみましたが、どうもPathがうまく通せていない気がします。 もうすでにリカバリしてしまったので、明日また環境構築をはじめからやり直します。。 Linuxの環境はないんです、すみません・・・。 今の環境はWindows7 64bitです。 詳しくはこちら>>274に書きました。 315 名前：デフォルトの名無しさん mailto:sage [2012/01/12(木) 07:17:27.30 ] >>314 はじめはVSのプロジェクトを構築するのではなく、VSの（32bit用の）コマンドプロンプトを立ち上げて、そこで コマンドを打ち込み、コンパイル、実行した方がよいのでは。 これがうまく行ってから、プロジェクトの構築をした方がよいでしょう。 cutil64.libがないと言うことは64bit用のSDKがインストールされていないのでは？ ここにあるはず。 C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.0\CUDALibraries\common\lib\x64 プロジェクトを構築するときに32bit, 64bit、あるいは両方の実行ファイルを作るかどうか設定できますよ。 316 名前：デフォルトの名無しさん mailto:sage [2012/01/12(木) 07:18:06.91 ] >>314 64bit版のcl.exeがある。 見つからないなら-ccbinオプションでcl.exeの場所を指定。 cutil32.libとcutil64.libを両方指定してはいけない。どちらか必要なものを。 で、cutil64は確か自分でビルドする必要がある。 SDKのディレクトリの中にプロジェクトファイルがあるからそれをビルド。 317 名前：デフォルトの名無しさん mailto:sage [2012/01/12(木) 08:36:26.90 ] codepad.org/Kyns70af 上記のプログラムが期待と違う結果を出力するのだが原因がわかる方がいれば教えていただけないだろうか． （CUDA 3.2 + GeForce 320M + Fedora 14 環境と CUDA 3.2 + Tesla C1060 + CentOS 5.5 環境でテスト） -- 期待した出力 -- 1.00 0.00 1.00 0.00 1.00 0.00 1.00 0.00 -- 実際の出力 -- 0.00 0.00 0.00 0.00 0.00 0.00 0.00 0.00 ちなみに，11行目の右辺を 2.0f に変更すると期待通りの結果を出力してくれる． -- 出力 -- 1.00 2.00 1.00 2.00 1.00 2.00 1.00 2.00 318 名前：317 mailto:sage [2012/01/12(木) 08:37:35.44 ] すみません． >>317のコードのURLを間違えました． 誤: codepad.org/Kyns70af 正: codepad.org/cfbTSsdF 319 名前：デフォルトの名無しさん mailto:sage [2012/01/12(木) 09:55:39.89 ] やってみたけど、問題は再現されなかった cuda_funcの中に printf(" %d %d %f\n",threadIdx.x,threadIdx.x & 1,a[threadIdx.x]); を入れて、チェックしてみれば？？ ちなみにこんな結果となる 0 0 1.000000　　　<- cuda_funcの中では問題はなさそう 1 1 0.000000 2 0 1.000000 3 1 0.000000 4 0 1.000000 5 1 0.000000 6 0 1.000000 7 1 0.000000 1.00 <--　外でも問題はなかった 0.00 1.00 0.00 1.00 0.00 1.00 0.00 320 名前：デフォルトの名無しさん mailto:sage [2012/01/12(木) 10:24:08.53 ] >>317 同じく問題再現されず（CUDA4.1 + GF560Ti） そのデバイス両方とも compute capability 1.x だし 2.x 環境で試してみたらどうだ 321 名前：317 mailto:sage [2012/01/12(木) 10:56:51.02 ] >>319 >>320 コメントありがとうございます 現在の私の環境ではカーネル関数内で printf は使えないですが 変なことになるのは私の環境だけのようなので Compute capability 2.x 環境で後日試してみます 322 名前：デフォルトの名無しさん mailto:sage [2012/01/12(木) 11:18:04.11 ] ちなみにcuPrintf.cuh　とcuPrintf.cu　をネット上からひらってくれば、printfと同等の関数（cuPrintf）が使えますよ♪ 323 名前：デフォルトの名無しさん mailto:sage [2012/01/12(木) 11:57:42.71 ] CUDA 3.2のコード出力に問題がある気がします。 カーネル部分をどうも a[threadIdx.x] = (float)(FUNC(-(threadIdx.x & 1)); に変換しようとしている気配が。 FUNCはISET.S32命令による処理だけどそれがミスっているとしか思えない。 自分で以下のように書くか、あるいはFermi以降またはCUDA4.0以降を使うべし。 a[threadIdx.x] = (float)(1 - (1 & threadIdx.x)); ちなみに実行はしてないので同じ問題が起こるかどうかは知りません。 324 名前：274 mailto:sage [2012/01/12(木) 13:42:05.81 ] >>315 >>316 ありがとうございます。 とりあえず315さんの通りコマンドライン上で実行できるところを目指そうと思います。 cl.exeを指定するオプションがあるんですね。 自分はコンパイラオプションの勉強不足のようです・・・。 cutil64は自分でビルドするんですか、それは今までやってなかったと思います。 恐縮ですが、これから環境を整えるので何か参考になるサイトがあれば教えていただけますか？ いくつかサイトを見て設定したのですが、どれもうまくいかなかったので・・・。 何度もすみません。よろしくお願いします。 325 名前：デフォルトの名無しさん mailto:sage [2012/01/12(木) 16:35:09.29 ] 質問です。 同じkernel関数をfor文で何度も実行するのですが、 回数によって1回あたりのkernel関数の実行時間が異なるみたいです。 具体的には for(i=0; i<10000; i++){ 　kernel_foo <<<grid_size, block_size >>>(...); } とすると300,000msかかるのが for(i=0; i<100; i++){ 　kernel_foo <<<grid_size, block_size >>>(...); } では1,000msくらいになります。同じ関数を実行しているので 単純に考えれば10000回ループを回したときは100回の100倍時間がかかるはずですが、 実際には300倍程度になりました。 もう少し細かく見てみると、最初60回程度はkernel関数を1回実行するのに0.04ms程で済んでいたのですが それ以降は1回あたり25ms程になっていました。 似た事例として ttp://d.hatena.ne.jp/ymizushi/20091214/p1 という記事を見つけましたが、この記事内でも原因はよく分かっていません。 どなたか詳細をご存じないでしょうか？ 326 名前：デフォルトの名無しさん mailto:sage [2012/01/12(木) 17:29:28.93 ] >> 324 参考になるやら、ならないやら www.fxfrog.com/?p=3660 www.ademiller.com/blogs/tech/2011/05/visual-studio-2010-and-cuda-easier-with-rc2/ cudasample.net/sample/1/001.html tech.ckme.co.jp/cuda_first.shtml cudasample.net/ others2.blog.so-net.ne.jp/2010-09-12 >> 325 1． カーネルの中の演算数がiによって異なる 2. 60回目で解が発散し NaNになっているが、お構いなしに計算は続けられている 3. カーネルからCPUにメモリーコピーしている時間が紛れ込んでいる かな？ 327 名前：274 mailto:sage [2012/01/12(木) 17:37:04.49 ] >>326 ありがとうございます。 全て拝見しましたが既読のものでした。やはり普通に設定すればできるようですね・・・。 今格闘しているので、また報告します。 328 名前：325 mailto:sage [2012/01/12(木) 17:50:39.27 ] >>326 ありがとうございます。 全項目について確認してみましたが、 1. 演算数はどれも一定 2. ループ回数によらず正しい値が得られている 3. 測定時間中にcudaMemcpy等の関数は呼ばれていない という状態です。 329 名前：デフォルトの名無しさん mailto:sage [2012/01/12(木) 18:22:29.85 ] >>328 1プロセスで2つのカーネル実行してるなら順番入れ替えて試してみた？ 最初の実行は起動コストだかで時間かかるのが普通なんだが 330 名前：デフォルトの名無しさん mailto:sage [2012/01/12(木) 19:29:37.96 ] 時間の計測が間違っている。 cudaThreadSynchronize()を読んだあとに終了時刻を取得する必要がある。 1024個までのカーネル呼び出しはキューに積むだけですぐに次に進んでしまう。 30000回ループの方はキューが詰まるので(30000-1024)回くらいの実行完了を待つことになり、比較的妥当な計算時間に見えるはず。 もちろん正確ではないのだが。 331 名前：325 mailto:sage [2012/01/12(木) 19:35:41.01 ] >>329-330 回答ありがとうございました。 cudaThreadSynchronize()を入れるのを失念しておりました。 332 名前：デフォルトの名無しさん mailto:sage [2012/01/12(木) 20:59:37.22 ] cudaDeviceSynchronize() がナウいらしい 333 名前：sage [2012/01/13(金) 08:22:49.98 ] >>330 便乗して質問ですが、 キューに一度に入るカーネル数が1024個で、 それ以上はカーネル呼び出しが待ちになるというのは、 CUDAのマニュアルかどこかに書いてありますか？ 334 名前：デフォルトの名無しさん mailto:sage [2012/01/13(金) 09:57:50.47 ] >>333 実験したらそうだったけど、文書化されているかどうかは知らない。 これだけ多ければ普通は問題ないし、バージョンアップで変わるかも。 335 名前：デフォルトの名無しさん mailto:sage [2012/01/13(金) 11:09:53.09 ] >>334 了解しました。 どうも有難うございました。 336 名前：デフォルトの名無しさん mailto:sage [2012/01/13(金) 16:58:47.32 ] それにしても　きゅー　にカーネルが溜まっているからと言って，GPUの処理能力が落ちるとは思えないぞ 337 名前：デフォルトの名無しさん mailto:sage [2012/01/13(金) 17:17:25.93 ] きゅー　にカーネル田丸？ それってstream使ってる時のこと？ 文脈呼んでないからわからんけど 338 名前：デフォルトの名無しさん mailto:sage [2012/01/13(金) 20:07:11.66 ] キューに空きがあれば積んで即制御が戻る。 空きが無ければ空くまで待ってる。 詰まってるから遅くなるわけではなく、遅い方が正しい実行時間なだけ。 339 名前：デフォルトの名無しさん mailto:sage [2012/01/18(水) 22:19:57.72 ] CUDA4.0からデバイスエミュレーションがなくなったらしいけど 実機で実行するとOSごと固まることがある みんなどうしているのでしょうか 340 名前：デフォルトの名無しさん mailto:sage [2012/01/19(木) 00:43:55.87 ] デバイスを物理的に叩き割る 341 名前：デフォルトの名無しさん mailto:sage [2012/01/19(木) 04:33:29.60 ] 2.3くらいからなかったような気が 342 名前：デフォルトの名無しさん mailto:sage [2012/01/19(木) 15:25:40.77 ] >>339 具体的には？？ 343 名前：デフォルトの名無しさん mailto:sage [2012/01/20(金) 18:49:16.04 ] >>342 メモリアクセスのバグがあるとWindowsが固まって動かなくなったり ディスプレイにノイズが出てくる OSが画面表示に使っている領域を壊しているようだけど 普通は起こらないことなんでしょうか？ 344 名前：デフォルトの名無しさん mailto:sage [2012/01/21(土) 10:42:26.35 ] 私の経験では、CUDAのプログラムをチェックするために printf を使って変数を書き出したところ、 その量が多すぎて、画面が真っ暗になったり、PCがシャットダウンしたことがあった。 書き出し量を減らしたところ、問題はなくなった。 こんなことかな？？ 345 名前：344 mailto:sage [2012/01/21(土) 11:00:22.89 ] 追加です。 グラフィックカードを2枚刺して、一枚をCUDA用、もう一枚をディスプレイ用にすればよいとのことを聞いたことがあるが、 私のはGTX５９０でカードが2枚、入っているとのこと。 CUDAの計算に Device(0)にしても(1)にしても、前述の問題が起きたので、枚数には関係ないこともあるらしい？？ 346 名前：デフォルトの名無しさん mailto:sage [2012/01/22(日) 05:18:00.51 ] printf使えるとか便利になったわ 347 名前：デフォルトの名無しさん mailto:sage [2012/01/22(日) 16:18:23.06 ] >>344 fermiではないのでprintfは使えないです。 GTX260なんですけど、それが悪いのかも。 計算用GPUを別に刺すってのは考えていました。 ただ、そうするべきという常識みたいなものがあるのかなっていう疑問があったので。 あとこういう状態でデバイスエミュレーション外したならひどい話だなと思ったのですが みなさんあまり起きていないようで。 348 名前：デフォルトの名無しさん mailto:sage [2012/01/22(日) 16:23:00.85 ] グローバルメモリの添え字計算計算が間違っていたりで確保した範囲外にアクセスすると Windowsが固まったりディスプレイにノイズが出てくることがあるって状態です。 349 名前：やんやん ◆yanyan72E. mailto:sage [2012/01/22(日) 16:48:55.32 ] グラフィックカード二枚差しにすると、 nSghtでハードウェアデバッグができるとかじゃなかったっけ？ 350 名前：デフォルトの名無しさん [2012/01/22(日) 17:02:34.79 ] YES もしくはPC2台で片方をデバッグ用にしても良い。 実際に動かすGPUがCUDA対応なら、開発側は非nVIDIAでもOK 351 名前：デフォルトの名無しさん mailto:sage [2012/01/23(月) 09:39:54.36 ] streamについて分からない点があるので どなたか教えてください。 以下で「カーネル」はカーネル関数の呼び出し、 「転送」はcudaMemcpyAsync、()内はstreamだとします。 (質問1)以下の状態のとき、転送(1)は、カーネル(1)と同じstreamなので 実行できませんが、キュー内でそれより後にある転送(2)は実行できますか？ カーネル(1) 実行中 転送(1) キュー内の最初のタスク 転送(2) キュー内の2番目のタスク (質問2)以下の状態のとき、転送(2)は実行できますか？ (キュー内で転送(2)の前にstream(0)のタスクがあります） カーネル(1) 実行中 カーネル(0) キュー内の最初のタスク 　　転送(2) キュー内の2番目のタスク (質問3)以下の状態のとき、転送(1)は実行できますか？ (実行中のタスクはstream(0)です) カーネル(0) 実行中 転送(1) キュー内の最初のタスク 「CUDA C Programming Guide」を読んでもよく分かりませんでした。 宜しくお願いします。 352 名前：デフォルトの名無しさん mailto:sage [2012/01/23(月) 10:54:29.21 ] 1はFermiならOK。古い世代ではダメ。 2と3は絶対ダメ。stream0は前後をぶった切る役目がある。 353 名前：デフォルトの名無しさん mailto:sage [2012/01/24(火) 08:17:05.28 ] >>352 1の動作がFermiとそれ以前で異なるのは知りませんでした。 回答どうも有難うございました。 354 名前：デフォルトの名無しさん mailto:sage [2012/01/24(火) 10:48:47.33 ] Fermiでプロファイラのtimestampを使ってテストしたところ カーネル(1) 実行中 転送(HostToDevice)(1) キュー内の最初のタスク 転送(HostToDevice)(2) キュー内の2番目のタスク だと転送(2)はカーネル(1)と同時に動かず、 カーネル(1) 実行中 転送(HostToDevice)(1) キュー内の最初のタスク 転送(DeviceToHost)(2) キュー内の2番目のタスク だと転送(2)はカーネル(1)と同時に動きました。 355 名前：デフォルトの名無しさん [2012/01/24(火) 19:53:55.48 ] コンパイル時に以下のエラーメッセージが出ているのですが、 ptxas error : Entry function '...' uses too much local data (0x16054 bytes, 0x4000 max) これはローカルメモリーと何か関係あるような気がするのですけど、 ちょっとわからないので教えて頂けないでしょうか。 宜しくお願いします。 356 名前：デフォルトの名無しさん [2012/01/24(火) 21:03:52.65 ] それともう1つ。 externを使って、グローバル変数を共有させたいのですが、 ホスト・デバイスともに共有させたい場合はどうしたらいいのでしょうか? 357 名前：デフォルトの名無しさん mailto:sage [2012/01/24(火) 21:56:32.01 ] >　ptxas error : Entry function '...' uses too much local data (0x16054 bytes, 0x4000 max) 単に配列の次元が大きすぎるのでは？ 最大　0x4000バイトのところを0x16054 bytes使おうとしている？ > ホスト・デバイスともに共有させたい場合はどうしたらいいのでしょうか? 物理的に無理でないの？ ホスト←→デバイス間でコピーしないといけないのでは？？ 358 名前：デフォルトの名無しさん mailto:sage [2012/01/24(火) 21:56:54.10 ] >>355 メッセージそのままの意味だろ。 359 名前：デフォルトの名無しさん mailto:sage [2012/01/25(水) 01:32:34.89 ] CUDAの思想はいいと思うけど、ターゲットデバイスを完全には 隠蔽しきれないのが残念だな。 360 名前：やんやん ◆yanyan72E. mailto:sage [2012/01/25(水) 01:41:55.08 ] スピード勝負の世界で、しかもかなり複雑なアーキテクチャなのに、 ハードウェアを完全に隠蔽する訳にいかないだろ。 361 名前：357 mailto:sage [2012/01/25(水) 12:34:12.36 ] > ホスト・デバイスともに共有させたい場合はどうしたらいいのでしょうか? ヘッダーファイルに #DEFINE などで定義して、ホストとデバイスのプログラムに そのヘッダーファイルをinclude する手があった♪ 362 名前：355,356 [2012/01/25(水) 15:21:00.98 ] >>358 青木氏の「はじめてのCUDA」に似たようなエラーメッセージが載ってるのですが、 local dataがどの部分の話なのか、ちょっとわからない状態でして... >>357 >単に配列の次元が大きすぎるのでは？ 予定だと1×256スレッドしか使うので、足りないはずないような... 意味が違うか。 ちょっとわからないので詳細お願いします。 >ヘッダーファイルに >#DEFINE >などで定義して、ホストとデバイスのプログラムに そのヘッダーファイルをinclude する手があった♪ ちょっと試してみます。それができなかったら、コピーするしかないかもしれません。 __device__ extern TEST test; のような宣言してて、実行ファイルができたときがあったのですが、何か違うよな、と引っかかってたので、 ここで質問しました。ありがとうございます。 363 名前：デフォルトの名無しさん mailto:sage [2012/01/25(水) 17:40:16.56 ] 差し支えなければプログラムをアップしてくださいませ♪ 364 名前：デフォルトの名無しさん mailto:sage [2012/01/25(水) 18:01:06.82 ] >>362 このサイトの一番下にある資料の3-41ページにそのエラーメッセージが載っている模様。 ttp://accc.riken.jp/HPC/training.html 365 名前：355,356 [2012/01/25(水) 18:04:09.31 ] >>362です。 >>363 申し訳ないのですが、プログラムのアップは出来ないです。すみません。 なんかもう分からなくなったので、ホストとデバイスの住み分けを行ったところ、 「ptxas error...」云々のメッセージが消え、コンパイルできました。 （さっきまでhost,device,global宣言関数がごちゃ混ぜな状態だった。） 何が原因だったのか分からず仕舞いです。 グローバル変数の共有はとりあえず、コンスタントメモリにデータコピーで様子見することにしました。 元コード(.c)をもう一度読み直したところ、デバイス側の方はReadOnlyで十分なようでしたから。 皆様回答ありがとうございました。またよろしくお願いします。 366 名前：365 [2012/01/25(水) 21:05:50.75 ] >>364 資料ありがとうございます。 うまくいったと思った途端、また同じエラーが出てきてしまったので、確認します。 .cファイルに.cuファイルを組み込ませるようにしたら、 __host__ __device__で修飾した関数が定義されてないと.cファイル側に言われ、 .cファイルと.cuファイルそれぞれ単独で動かせるようにしたら、 （同じ内容の関数を.cと.cuそれぞれ別名で実体作った。.c:test .cu:ktestみたいな） 今度はさっきと同じエラー。 実行ファイル作るだけなのに難しい... 367 名前：デフォルトの名無しさん mailto:sage [2012/01/26(木) 08:42:18.26 ] Host側とDevice側で型を別にすれば開発もっと楽になると思うんだけど

---

---

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

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