【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/ 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側で型を別にすれば開発もっと楽になると思うんだけど 368 名前：デフォルトの名無しさん mailto:sage [2012/01/26(木) 08:53:13.45 ] 取り敢えずCudaのサンプルは捨てて、インクルード関係とオブジェクトの生成手順を確認するんだ。 369 名前：デフォルトの名無しさん mailto:sage [2012/01/26(木) 10:14:19.95 ] 見せられないところは削除して、問題の部分だけ残して、アップすれば？？ 370 名前：366 [2012/01/26(木) 16:19:08.01 ] >>368 >取り敢えずCudaのサンプルは捨てて、インクルード関係とオブジェクトの生成手順を確認するんだ。 のように手順を踏んでコンパイルしたところ、実行ファイルができました。 皆様ありがとうございました。 371 名前：370 [2012/01/26(木) 20:33:24.55 ] お恥ずかしながら、また戻ってきました。 あれからセグメンテーション違反が出てきてしまったので、あれこれ探していた結果、 どうやらデバイスのメモリ確保＆送信に失敗していたようです。 しかし解せないことがあって、 構造体A a,構造体B b[100],構造体C　c[100],...(以下略　をデバイス側に送るのですが、 （1つ１つのサイズは結構大きめ。グローバルは1GBあって余裕で確保出きるハズ...） void main(){ ・・・ test(&a,b,c....); } void test(A *a,B *b,C *c...){ A *d_a; B *d_b; C *d_c; CUT_SAFE_CALL(cudaMalloc((void**)&d_a,sizeof(A))); CUT_SAFE_CALL(cudaMemcpy(d_a,a,sizeof(A),cudaMemcpyHtoD)); CUT_SAFE_CALL(cudaMalloc((void**)&d_b,sizeof(B)*100)); CUT_SAFE_CALL(cudaMemcpy(d_a,b,sizeof(B)*100,cudaMemcpyHtoD)); ・・・ ｝ はじめてcuda,cuda_by_exampleで確認したところ、 文法的ミスはないはずなのに確保ミスしてるらしく、中断してます。 この原因は一体全体なんなんでしょうか。 372 名前：デフォルトの名無しさん [2012/01/26(木) 21:28:56.32 ] 4.1正式版 CUDA Toolkit 4.1 | NVIDIA Developer Zone developer.nvidia.com/cuda-toolkit-41 373 名前：デフォルトの名無しさん mailto:sage [2012/01/26(木) 21:55:28.42 ] 構造体の中身がよく分からんが，allocateのところで　*100は必要なのかな？ 構造体Bのメモリーを１００個分用意しようとしている？？ 構造体Bの中にすでに [100]個の配列を取っているのに？？？ 数は合っておるのか？？？？ 374 名前：デフォルトの名無しさん mailto:sage [2012/01/26(木) 22:17:44.08 ] ｒｈｅｌ6.0のがリンクミスかでダウンロドできん 375 名前：デフォルトの名無しさん mailto:sage [2012/01/27(金) 00:03:32.49 ] >>372 遂にRC取れたか！4.0からどう変わってるか知らんけど。 Kepler向けのプログラムを作れるのは5.0とかになるんだろうか。 そしてそれはいつ出るんだろう。 >>371 1GBのVRAM積んでるカードでも400MBぐらいの確保で失敗したことがあったような覚えがある。 確保するサイズを小さくしてもエラーが出るってんならこの話は忘れてくれ。 376 名前：デフォルトの名無しさん mailto:sage [2012/01/27(金) 01:34:46.65 ] 100x100の行列の積を1万回くらい実行したいのですが、 CUBLASだと小さすぎてCPUより遅くなってしまいます。 この1万回は相互に依存関係は無いので、並列化が可能なのですが、 このプログラムは自分で書くしかないのでしょうか。 TESLA C1070があるので、丸ごとVRAMに載ります。 377 名前：デフォルトの名無しさん mailto:sage [2012/01/27(金) 03:17:55.02 ] 4.1にcublas{S,D,C,Z}gemmBatchedってのが追加されてました。 especially for smaller matricesに効くってあったので、使ってみます。 378 名前：デフォルトの名無しさん mailto:sage [2012/01/27(金) 07:52:34.05 ] 一つのブロック内で100x100の行列の積を１回行う？ 379 名前：デフォルトの名無しさん mailto:sage [2012/01/27(金) 07:52:51.42 ] 間違った 380 名前：デフォルトの名無しさん mailto:sage [2012/01/27(金) 07:53:00.44 ] あれ？ 381 名前：デフォルトの名無しさん mailto:sage [2012/01/27(金) 07:53:07.78 ] 一つのブロック内で100x100の行列の積を１回行う？ 382 名前：デフォルトの名無しさん mailto:sage [2012/01/27(金) 07:55:11.45 ] たびたび，失礼 一つのブロック内で100x100の行列の積を１回行う？ そのブロックを１万個用意する？ それらを並列に計算する？ と言うコンセプトかな？？？ 自分でプログラムを作った方が，よほど勉強になると思う 383 名前：371 [2012/01/27(金) 18:15:48.81 ] >>373 CUT_SAFE_CALL(cudaMalloc((void**)&d_b,sizeof(B)*100)); CUT_SAFE_CALL(cudaMemcpy(d_a,b,sizeof(B)*100,cudaMemcpyHtoD)); これはミスですね.... 正しくは CUT_SAFE_CALL(cudaMalloc((void**)&d_b,sizeof(B))); CUT_SAFE_CALL(cudaMemcpy(d_a,b,sizeof(B),cudaMemcpyHtoD)); です。すみません。 >>375 4000MBの領域確保でも失敗することがあったんですか... ちょっと試してみます。 384 名前：デフォルトの名無しさん mailto:sage [2012/01/27(金) 20:31:39.53 ] d_aに送信してるあたりが 385 名前：デフォルトの名無しさん mailto:sage [2012/01/27(金) 21:52:59.41 ] d_b のサイズは構造体Bと同じ？？ d_aのサイズはbと同じ？すなわち構造体Bと同じ？？ 違っていたらエラーが出て当たり前では？？？ 386 名前：デフォルトの名無しさん mailto:sage [2012/01/28(土) 13:33:36.75 ] それから，ホスト側で　も　メモリーは確保しているのか？？？？ 387 名前：383 [2012/01/28(土) 17:54:35.37 ] >>371 なんかサンプルがめちゃくちゃなんで書き直します。 388 名前：387 [2012/01/28(土) 20:55:21.49 ] 確認なのですけど、カーネル関数の引数はポインタ限定ですか? 389 名前：デフォルトの名無しさん mailto:sage [2012/01/28(土) 22:58:27.22 ] 配列はポインタで，１変数はそのままではなかったかい？ なぜか知らんけど 390 名前：デフォルトの名無しさん mailto:sage [2012/01/28(土) 23:07:13.34 ] 夜の埼玉は 391 名前：デフォルトの名無しさん mailto:sage [2012/01/29(日) 01:49:05.91 ] >388 ホスト側のポインタ以外なら何でもOKのはず。少なくともintは渡せる。 クラスを渡せるかどうかは知らないけど。 392 名前：デフォルトの名無しさん mailto:sage [2012/01/29(日) 06:58:36.12 ] ホストのポインタ送って、GPU側のload/store時に メインメモリから自動転送ってこともできるから、 その説明は正しくない。 393 名前：388 [2012/01/30(月) 15:13:56.68 ] >>391 セグメンテーション違反してた部分の引数に（デバイスで使用する）double（実体）を指定してたのですけど、 （デバイスで使用する）ポインタに変えたら、問題なく実行できました。 dtempはcudaMalloc/cudaMemcpy使ってます。 そもそもな話cudaはポインタ宣言したdtempの実体は*dtempではないのかな。 394 名前：デフォルトの名無しさん mailto:sage [2012/01/30(月) 15:16:54.53 ] >>393 何を言っているのか判らない。問題が再現する最低限のコードでも貼ってみていただきたい。 395 名前：デフォルトの名無しさん [2012/01/30(月) 16:19:41.78 ] >>393 __global__ void kernel(double dtemp){ } int main(){ double temp=0.0; double *dtemp; cudaMalloc((void**)&dtemp,sizeof(double)); cudaMemcpy(dtemp,&temp,sizeof(double),cudaMemcpyHostToDevice); kernel<<<1,1>>>(*dtemp);//(←※1) cudaMemcpy(&temp,dtemp,sizeof(double),cudaMemcpyDeviceToHost); return 0; } ※1の部分でsegmentation fault。 __global__ kernel(double *dtemp){ }として kernel<<<1,1>>>(dtemp); とすると問題無し。 dtempのアドレスを引数にするのはいいけど、dtempの中身（0.0）を引数にしようとするのはダメ？ そもそも※1の宣言自体間違ってる?どうなんだろう。 396 名前：デフォルトの名無しさん mailto:sage [2012/01/30(月) 16:29:17.80 ] 393じゃないけど試してみた __global__ void kernel(double dtemp){ } int main(){ double temp=0.0; double dtemp; // cudaMalloc((void**)&dtemp,sizeof(double)); // cudaMemcpy(dtemp,&temp,sizeof(double),cudaMemcpyHostToDevice); kernel<<<1,1>>>(dtemp);//(1) // cudaMemcpy(&temp,dtemp,sizeof(double),cudaMemcpyDeviceToHost); return 0; } これでコンパイルも実行もおkだったけど 397 名前：デフォルトの名無しさん mailto:sage [2012/01/30(月) 16:57:38.16 ] >>395 dtempはdevicePointerなんだから、Hostコード中でデリファレンスしちゃダメ。 double値を渡したいだけなら引き数はdoubleで充分。 それと、※1は宣言じゃなくて呼び出しだから構文としてはあっている。 そもそも何がしたいのだろう。値を渡したいなら値を渡せばいいし、 ポインタを渡したいならポインタをそのまま渡せばいい。 devicePointerをHostでデリファレンスしちゃダメだし、 (その後どうせdevice側でデリファレンスできないから)HostPointerをdeviceにコピーしてもダメ。 398 名前：デフォルトの名無しさん [2012/01/30(月) 17:04:05.78 ] >>396 >>397 色々ごちゃまぜだったようです。スッキリしました。 ありがとうございました。 399 名前：デフォルトの名無しさん mailto:sage [2012/01/30(月) 20:58:23.57 ] これからCUDAの勉強を始めようと思っています。 GPGPUプログラミングはシェアードメモリを有効活用できるかがポイントであるように思います。 そこで、あえて最新のグラボに買い換えず、キャッシュのないTesla世代のGTX285を所持し続けています。 キャッシュがあると、シェアードメモリをうまく使えてなくても、それをカバーしてしまって、 うまくプログラムできているかどうかを把握し辛いと考えたからです。 ただ、心配なのが、Fermi世代以降に発売された書籍だと、内容がTesla世代のGPUにそぐわない といった状況が起きてしまうことです。 GTX285のまま勉強するに際して気をつけること等あれば教えてください。 問題がなければ関連書籍（日本語）を全部購入して読もうと思います。 よろしくお願いします。 400 名前：デフォルトの名無しさん mailto:sage [2012/01/30(月) 21:08:43.55 ] ＞＞GPGPUプログラミングはシェアードメモリを有効活用できるかがポイントであるように思います。 そんなの計算する内容によるわ 活用しようがないのもあるからな 計算（使用）目的は明確なの？ 401 名前：399 mailto:sage [2012/01/30(月) 21:39:43.45 ] >>400 信号処理に使う予定です。 FIRフィルタリングやFFTなどを大きな1次元、あるいは2次元配列に適用したいと思っています。 402 名前：デフォルトの名無しさん mailto:sage [2012/01/30(月) 21:45:58.41 ] 俺は400じゃないが、気をつけることは、勉強がひと段落していざ使うときになって知識が陳腐化してても泣かないことかな。 ハードウェアはどんどん変わるし、TeslaやFermiで良しとされたことが数年後に通用するかはわからないからね。 せっかく今から始めるならFermi・Keplerを使ったほうがいいと俺は思うけど、 10年後にはどれも等しく陳腐化してるだろうから長期的に考えると確かにTeslaでもいいのかもしれない。 ただ長期的に考えること自体がGPGPUでは…と1行目にループする。 403 名前：デフォルトの名無しさん mailto:sage [2012/01/30(月) 22:12:50.10 ] 古い方がいろいろ制限があって，勉強になるとは思う． 404 名前：デフォルトの名無しさん mailto:sage [2012/01/30(月) 22:17:32.23 ] ｆｆｔはcpu側からはインターフェースあるけどディバイス内では自前で作るしかない 405 名前：デフォルトの名無しさん mailto:sage [2012/01/30(月) 22:18:26.89 ] > キャッシュがあると、シェアードメモリをうまく使えてなくても、それをカバーしてしまって、 > うまくプログラムできているかどうかを把握し辛いと考えたからです。 シェアードメモリーを使うようにプログラミングすればよいかと どのメモリーを使っているのか分からない，と言うことはないと思う あれば，そこをしっかり勉強すべきかと 406 名前：デフォルトの名無しさん mailto:sage [2012/01/30(月) 22:19:12.36 ] そう言えば古いのは倍精度の計算ができないのでは？？ 407 名前：デフォルトの名無しさん mailto:sage [2012/01/30(月) 22:22:30.66 ] まあ，デバイスメモリーだけで計算してみるとか， シェアードメモリーも使って計算してみるとか， いろんなプログラムを作って比較するのが良いかと 408 名前：デフォルトの名無しさん mailto:sage [2012/01/30(月) 22:26:01.98 ] 最初はシェアドするスレッドはどれだってだけでも頭が痛くなった 409 名前：402 mailto:sage [2012/01/30(月) 23:28:19.90 ] あとTeslaでやるなら、FermiとTeslaでバンクコンフリクトの発生条件が違ってるから Teslaでのこれの回避には深入りしないほうがいいかと。 他にシェアードメモリの大きさなどの量的な違いがあるけど、そういった量的な 制約のために質的なというかアルゴリズムの種類を変えるようなことまで深入りするのも 避けたほうが。だってFermiやそれ以降で量が変わると一から見直しになるから。 これはFermi使ってる人もいずれ同じことなんだけど、勉強じゃなくて性能が欲しいんだからしょうがない。 410 名前：デフォルトの名無しさん mailto:sage [2012/01/30(月) 23:57:48.21 ] sharedがregister並に速いなんて嘘だから。 occupancyなんて上げるな。 warpあたりのregisterを増やして ILPを利用してレイテンシを隠蔽すべきだと。 www.cs.berkeley.edu/~volkov/volkov10-GTC.pdf 411 名前：399 mailto:sage [2012/01/31(火) 00:31:29.61 ] >>403 やはり、そうですか！ >>404 参考になります。 FFTは重要になるのでしっかりと取り組みたいです。 >>405 >>407 >>408 なるほど、どのメモリかをよく意識してプログラミングします。 CPUの最適化に取り組んでいて思ったのですが、 キャッシュは便利な反面、こちらから意図的に挙動を制御できないことが パフォーマンスの考察を難しくしてしまう側面もあると感じました。 キャッシュなしのTeslaで鍛えたいです。 >>406 単精度しか使わないので問題なしです。 もし半精度でパフォーマンスが上がるなら、そのほうがイイくらいです。 >>402 >>409 世代を跨ぐにあたって非常に参考になるご助言をありがとうございます。 各種メモリの特性や帯域を意識して取り組むことで、 固有のデバイスに限定されない定性的なところを理解したいと思います。 Keplerではこんな組み方がイイかもしれないな、と思索できたら楽しそうです。 >>410 興味深い資料をありがとうございます。 各部のレイテンシ、スループットを頭に叩き込みます。 412 名前：デフォルトの名無しさん mailto:sage [2012/01/31(火) 08:31:17.17 ] >>411 最終的に動かす機器で最もよい性能が出るように最適化するべきだし 自動でやってくれることを手でやるってのは勉強にはなるかもしれないけど意味ないのではって思うけど。 コンパイラの最適化を切ってアセンブリ書くようなものでしょ。 413 名前：デフォルトの名無しさん mailto:sage [2012/01/31(火) 12:08:28.56 ] 時間がいくらでもあるorそういう研究ならともかく 短期間で組めてそこそこのパフォーマンスが出せればいいなら キャッシュのあるFermiで適当に書いてもいいと思うんだけどね

---

---

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

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