【GPGPU】NVIDIA CUDA質問スレッド

1 名前：デフォルトの名無しさん mailto:sage [2007/09/17(月) 14:54:28 ] 本家 developer.nvidia.com/object/cuda.html 960 名前：954 mailto:sage [2009/01/17(土) 08:44:25 ] >>955=956 えーと、都合数MB程度ってことか。それだったら>959の言うようにカーネル関数のタイムアウトじゃないかな。 DEBUGつきでCUT_SAFE_CALL(間違ってたら失敬)を使うか、自前でcudaThreadsSynchronize()を呼んで戻り値を見てみて。 カーネル関数起動後は、復帰を待たずにホストの処理に戻るからカーネル関数自体のエラーはカーネル関数起動では捕捉できないのよ。 で、まさかとは思うけど同一データを100回コピーならバス越しにやってはダメ。別のデータならホスト側の用件次第かな。 # でも多分、一旦ホスト側に転送してから分配した方が速い気がする。CPUとメモリ周りがよっぽど遅くなければ。 >>957 下手な突っ込みはお郷が知れるよ。この場合、CPUが遅いんじゃないかとかバスがx1なんじゃないかとかGPUがモニタ表示か何かで よっぽど負荷が掛かっているんじゃないかとかGPU側メモリが巨大なSTELAなのかとか、組み合わせ次第で色々条件が変わってくるのよ。 961 名前：デフォルトの名無しさん [2009/01/17(土) 09:52:40 ] >>960 ＞ 下手な突っ込みはお郷が知れるよ。 質問者が答えてないのを指摘しただけなんだが… 誤解させてすまんかった 962 名前：デフォルトの名無しさん mailto:sage [2009/01/17(土) 11:09:08 ] そういう意味か。あの書き方じゃぁ、誤解されるよ。 まぁ、大量転送じゃないと言う意味で、今回は関係なさそうだね。 963 名前：デフォルトの名無しさん [2009/01/17(土) 15:38:06 ] >>952の話 >>957 ATIの統合チップセットでオンボードのATIビデオカードを無効 PCI-E 1.0 x16バスにGeforce9600GT 512MB CPUはAMDデュアルプロセッサ(結構遅いやつ) 964 名前：デフォルトの名無しさん [2009/01/17(土) 15:55:29 ] >>952の話 >>959 カーネル実行 a:cudaのエラーチェック(カーネルのエラーのありなし) cudaThreadsSynchronize for(100回 結果データをデバイスからホストにコピー b:cudaのエラーチェック(直前のcudaMemcpyのエラーのありなし) ) bのとこでthe launch timed out and was terminatedが出てるんです。 先週は時間がきてあきらめて帰ったのでここまでです。 言われたようにcudaThreadsSynchronizeの戻り値見てみたほうがよさそう ですね。 965 名前：デフォルトの名無しさん mailto:sage [2009/01/17(土) 18:45:09 ] >>951 共有メモリはカーネル起動時に動的に確保されるメモリ領域だから カーネルが終了したら消えるし、１つの実行スレッドから全部コピーなんて論理的におかしいし 共有メモリはあくまで高速動作と同期が取れる作業用のメモリとして考えたほうがいい 共有メモリを使って計算した結果は１個づつ普通のグローバルメモリに入れてやるのが正しいやり方 >>952 10000を一度に転送して実行しても 1を10000回繰り返して転送しても 実行時間は大差ないんですよ CUDAで実行する部分は出来るだけコンパクトにまとめて 呼び出す本体の方で特殊な演算関数くらいの感覚でループさせて使うのが正解 966 名前：デフォルトの名無しさん mailto:sage [2009/01/18(日) 23:59:40 ] JCublasについて おしえてエロイ人 967 名前：デフォルトの名無しさん [2009/01/19(月) 11:04:50 ] >>925の話 cudaThreadsSynchronizeの戻り値をチェックしたら the launch timed out and was terminatedが出ていました。結果コピーで エラーで落ちてたのだけどエラーデータは、その前に実行していた cudaThreadsSynchronizeの問題だったようです。 (cudaThreadsSynchronizeが正常になるまで待つとしても配列を 100x1000回、100x10000回、回すと1分待っても同期とれないようです。) つまりカーネルに同期できない処理が書かれていたのが原因だと思います。 NVIDIA_CUDA_Programming_Guide_2.0.pdfのp34-35のようなコードで かつそこのカーネルのdevPtrを参照するだけでなく書き戻す処理をや っているのでそれがだめなのだと思います。恐らくこういうピッチ処理 の場合は、参照のみが許されるのでは?と思います。 問題のコードをこのpdfの変数に直して下記に書いておきます。 __global__ void myKernel(float* devPtr, int pitch) { for (int r = 0; r < height - 1; r++) { float* row1 = (float*)((char*)devPtr + r * pitch); float* row2 = (float*)((char*)devPtr + (r + 1) * pitch); for (int c = 1; c < width - 1; c++) { row2[c] = row1[c]とrow1[c + 1]とrow1[c -1]やらの計算 } } } このp34-35を見て思うのはindex使わずにこんな処理書いてGPUでやる意味 あるの?と思いますが。自分で書いてなんなのですが 968 名前：デフォルトの名無しさん mailto:sage [2009/01/19(月) 17:28:20 ] お前はあほかｗ 969 名前：デフォルトの名無しさん mailto:sage [2009/01/19(月) 17:49:31 ] タイムアウトになる原因はそのでかいループのせい せいぜいミリ秒単位でタイムアウトを判断してるから ミリ秒以下で応答しないようなカーネルは絶対に書いては駄目 cudaThreadsSynchronizeは発行したカーネルがすべて終了するのを待つだけ グローバルメモリは読み書きは出来るが前後は保障されないので １スレッドが書き込みする箇所は限定する必要がある 共有メモリを使って他のスレッドが書き込んだ内容を参照して利用する場合に カーネル内部で___syncthreadを使う これが本来の同期の意味 970 名前：デフォルトの名無しさん [2009/01/19(月) 19:44:51 ] >>952の話 >>968のように言われるのは分かって書いてみたんだけど NVIDIA_CUDA_Programming_Guide_2.0.pdfのp34-35のようなコードが マニュアルに書いてあるのがおかしいと思う。 __global__ void myKernel(float* devPtr, int pitch){} そもそもこんな書き方じたいが書けるけど間違えな使い方だと。 この書き方しているとこにやらないようにこの部分に×印つけてほしい。 あとはコンパイラがえらくなったらfor多重ループをうまく処理する アセンブラぐらい作ってほしい。(OpenMPでパラレルfor指定すると#のタグだけ でプロセッサ使ってきって高速化なるんだけどね。そこまでぐらい将来 的にはしてほしい。) 971 名前：デフォルトの名無しさん mailto:sage [2009/01/19(月) 20:53:12 ] 文句言っても何にもならない。 そう悟る時がくるまで気長に待ちましょうよ。 そうすれば成長しまっせ。 972 名前：デフォルトの名無しさん mailto:sage [2009/01/19(月) 23:38:15 ] それよりも先ず、日本語を何とかしてくれ。 ＞間違えな使い方 ＞しているとこにやらないように ＞プロセッサ使ってきって 973 名前：デフォルトの名無しさん mailto:sage [2009/01/19(月) 23:41:25 ] そのサンプルはあくまでもcudaMallocPitch()がどういう風にメモリを確保するかの説明だからなぁ。 きっとライターは、まさかそのまま動かそうとされるは思っちゃいまいよ。 974 名前：デフォルトの名無しさん mailto:sage [2009/01/20(火) 02:02:49 ] ところでドライバ180なんだが2Dモードだかなんだかで消費電力抑える為に ドライバが自動的にクロックを下げて3Dゲーム動作させると上がる仕様なんだけど CUDA実行しても上がらないぞｗ 975 名前：デフォルトの名無しさん mailto:sage [2009/01/20(火) 02:08:30 ] ドライバを替えてくださいw つーか、それって169.09のときに指摘されて直したんじゃなかったのか?＞NVIDIA 976 名前：デフォルトの名無しさん [2009/01/20(火) 03:40:32 ] d_a[0][blockIdx.x][blockIdx.y][threadIdx.x] って使おうとしたら"expression must have arithmetic or enum type" ってでたんだけど何がいけないんすか？ 教えてください。 977 名前：デフォルトの名無しさん mailto:sage [2009/01/20(火) 07:58:55 ] >>976 d_aがポインタなら、途中で解決できない状況があるのかも。 978 名前：デフォルトの名無しさん mailto:sage [2009/01/20(火) 08:41:39 ] 共有メモリで構造体を使いたいのだけど 例えば struct data { int a; char b; } func<<<dim3(), dim3(), SIZE * sizeof(data)>>>(); __global__ kernel(){ __shared__ data d[SIZE]; } こんな感じでやるとCUDAとC側のパック構造が違うせいだと思うけどおかしくなる どうやれば出来るの？ 979 名前：デフォルトの名無しさん mailto:sage [2009/01/20(火) 08:59:56 ] パックが違うのだけが理由なら、 struct data {int a; char b; char dummy[3];}とでもしておけばいいじゃん。 980 名前：デフォルトの名無しさん mailto:sage [2009/01/20(火) 13:53:14 ] ただでさえ少ない共有メモリをそんな無駄に使えない 981 名前：デフォルトの名無しさん mailto:sage [2009/01/20(火) 20:44:37 ] >>977 世の中夜勤帰りで朝から寝てる人だっているんだよ？ 引っ越しの時ちゃんと挨拶行った？ 顔合わせたら軽く会話するとかしてちゃんとコンタクト取り続けてる？ 日頃からそういうコミニュケーションが取れてればいつ洗濯機を回していいのか いつ静かにしなければならないのか 迷惑を掛けないように生活出来るはずなんだが 982 名前：デフォルトの名無しさん mailto:sage [2009/01/20(火) 23:36:58 ] >>980 マジで言っているのなら、設計が悪い。 どうしてもパディングしたくないくらい逼迫しているなら、int一個を切り分けて3バイトと1バイトで使え。 983 名前：デフォルトの名無しさん mailto:sage [2009/01/21(水) 00:22:25 ] 共有メモリが制限されてるのに無駄な領域作って ほとんどをグローバルメモリに追いやる方がよっぽど設計が悪いでしょｗ 984 名前：デフォルトの名無しさん mailto:sage [2009/01/21(水) 00:25:51 ] GPUのメモリレイテンシって１２とかの世界だぞ CPU用のDDR2で５だからな intを内部でcharとして扱うプロセスを考慮しても共有メモリ使った方が早いんだよ 985 名前：デフォルトの名無しさん mailto:sage [2009/01/21(水) 00:35:27 ] CUDAで共有メモリを使うこと自体、無駄な同期待ちが発生するから避けたいところだと思うが。 coaxschedな読み書きができるなら、共有メモリより遅くないぞw 986 名前：デフォルトの名無しさん mailto:sage [2009/01/21(水) 02:24:30 ] >>975 (エンバグの)歴史は繰り返す。 987 名前：デフォルトの名無しさん mailto:sage [2009/01/21(水) 12:18:47 ] >>CUDAで共有メモリを使うこと自体、無駄な同期待ちが発生するから避けたいところだと思うが。 え？ 共有メモリってCPUで言う所のただのレジスタみたいなもんで同期は指定しないと発生しないと思うけど ところでローカル変数をすべて共有メモリに置いた場合のが早いんだけど ローカル変数ってデバイス上のメモリに確保されるだけで共有メモリより遅いレベルのメモリを使ってるっぽいね 988 名前：デフォルトの名無しさん mailto:sage [2009/01/21(水) 18:11:15 ] あー、表現が悪かった。共有メモリを共有目的に使うと同期を取る必要が出てくるから同期待ちが発生するということ。 レジスタみたいに使うのなら確かに関係なかったね。 で、レジスタよりも速いかどうかについてはptx見てみたいところ。 989 名前：,,・´∀｀・,,）っ-○◎● mailto:sage [2009/01/21(水) 18:18:05 ] 見た感じリードレイテンシはこれくらい レジスタ＞＞Const Memory＞Shared Memory＞＞＞＞DRAM

---

---

[ 新着レスの取得/表示 (agate) ] / [ 携帯版 ]

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