【GPGPU】NVIDIA CUDA質問スレッド

【GPGPU】NVIDIA CUDA質問スレッド at TECH

[前50を表示]
950:デフォルトの名無しさん
09/01/16 14:10:31
2.1をつついた人に質問
2.1ではカーネルに渡す引数でポインタ配列は使えますか?

951:デフォルトの名無しさん
09/01/16 18:48:04
共有メモリからグローバルメモリに連続領域をコピーするにはどうすればいいでしょうか
カーネルから呼び出せるmemcpyみたいなのがあれば教えてください

952:デフォルトの名無しさん
09/01/16 18:58:33
cudaMemcpyが長いようで、the launch timed out and was terminated.が
でてしまいす。これなんとかするほうほうありましたっけ?

953:デフォルトの名無しさん
09/01/16 21:53:12
>>952
転送量を減らす。

>>951
cudaMemcpy

954:デフォルトの名無しさん
09/01/17 00:16:47
>>951
スレッド内で連続的にメモリにアクセスするようなコードは望ましくない。
coalscedにアクセスするとなると、自分でコードを書くしかない。
共有メモリから転送を始める前に、同期を取ることも忘れずに。

>>952
参考までに、どんなハードウェアの組み合わせでどれくらい転送しようとしている?
8800GT辺りだと転送できるだけ全部転送してもタイムアウトしないのだけれど。

955:デフォルトの名無しさん
09/01/17 02:52:44
>>954
カーネル実行後の結果データコピーでエラーになっているのですが
sizeof(float)*1000バイトを100回連続でコピーしています。
100個とも別のアドレスに結果をコピーしたいためそうしてます。やはり
いっぺんにメインメモリにいったんコピーしてからやったほうがいいので
しょうか?

956:デフォルトの名無しさん
09/01/17 03:03:49
>>955
sizeof(float)*10000バイトを100回連続コピーでした。

957:デフォルトの名無しさん
09/01/17 07:16:40
>>954
＞ >>952
＞どんなハードウェアの組み合わせ

958:デフォルトの名無しさん
09/01/17 07:19:52
>>956
同一データ部分を100箇所にコピー？
独立の100箇所をコピー？

959:デフォルトの名無しさん
09/01/17 07:55:29
>>956
それ、本当にメモリのコピーでタイムアウトしてる？
カーネル実行後に同期取らずにすぐにメモリのコピー始めて
100回コピーのどこかでカーネル実行がタイムアウトしてるとか。

cudaMemcpyの戻り値は
Note that this function may also return error codes from previous, asynchronous launches.
ということだし。

960:954
09/01/17 08:44:25
>>955=956
えーと、都合数MB程度ってことか。それだったら>959の言うようにカーネル関数のタイムアウトじゃないかな。
DEBUGつきでCUT_SAFE_CALL(間違ってたら失敬)を使うか、自前でcudaThreadsSynchronize()を呼んで戻り値を見てみて。
カーネル関数起動後は、復帰を待たずにホストの処理に戻るからカーネル関数自体のエラーはカーネル関数起動では捕捉できないのよ。
で、まさかとは思うけど同一データを100回コピーならバス越しにやってはダメ。別のデータならホスト側の用件次第かな。
# でも多分、一旦ホスト側に転送してから分配した方が速い気がする。CPUとメモリ周りがよっぽど遅くなければ。

>>957
下手な突っ込みはお郷が知れるよ。この場合、CPUが遅いんじゃないかとかバスがx1なんじゃないかとかGPUがモニタ表示か何かで
よっぽど負荷が掛かっているんじゃないかとかGPU側メモリが巨大なSTELAなのかとか、組み合わせ次第で色々条件が変わってくるのよ。

961:デフォルトの名無しさん
09/01/17 09:52:40
>>960
＞下手な突っ込みはお郷が知れるよ。

質問者が答えてないのを指摘しただけなんだが…
誤解させてすまんかった

962:デフォルトの名無しさん
09/01/17 11:09:08
そういう意味か。あの書き方じゃぁ、誤解されるよ。
まぁ、大量転送じゃないと言う意味で、今回は関係なさそうだね。

963:デフォルトの名無しさん
09/01/17 15:38:06
>>952の話
>>957
ATIの統合チップセットでオンボードのATIビデオカードを無効
PCI-E 1.0 x16バスにGeforce9600GT 512MB
CPUはAMDデュアルプロセッサ(結構遅いやつ)

964:デフォルトの名無しさん
09/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:デフォルトの名無しさん
09/01/17 18:45:09
>>951
共有メモリはカーネル起動時に動的に確保されるメモリ領域だから
カーネルが終了したら消えるし、１つの実行スレッドから全部コピーなんて論理的におかしいし
共有メモリはあくまで高速動作と同期が取れる作業用のメモリとして考えたほうがいい
共有メモリを使って計算した結果は１個づつ普通のグローバルメモリに入れてやるのが正しいやり方

>>952
10000を一度に転送して実行しても
1を10000回繰り返して転送しても
実行時間は大差ないんですよ
CUDAで実行する部分は出来るだけコンパクトにまとめて
呼び出す本体の方で特殊な演算関数くらいの感覚でループさせて使うのが正解

966:デフォルトの名無しさん
09/01/18 23:59:40
JCublasについて
おしえてエロイ人

967:デフォルトの名無しさん
09/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:デフォルトの名無しさん
09/01/19 17:28:20
お前はあほかｗ

969:デフォルトの名無しさん
09/01/19 17:49:31
タイムアウトになる原因はそのでかいループのせい
せいぜいミリ秒単位でタイムアウトを判断してるから
ミリ秒以下で応答しないようなカーネルは絶対に書いては駄目

cudaThreadsSynchronizeは発行したカーネルがすべて終了するのを待つだけ

グローバルメモリは読み書きは出来るが前後は保障されないので
１スレッドが書き込みする箇所は限定する必要がある

共有メモリを使って他のスレッドが書き込んだ内容を参照して利用する場合に
カーネル内部で___syncthreadを使う
これが本来の同期の意味

970:デフォルトの名無しさん
09/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:デフォルトの名無しさん
09/01/19 20:53:12
文句言っても何にもならない。
そう悟る時がくるまで気長に待ちましょうよ。
そうすれば成長しまっせ。

972:デフォルトの名無しさん
09/01/19 23:38:15
それよりも先ず、日本語を何とかしてくれ。
＞間違えな使い方
＞しているとこにやらないように
＞プロセッサ使ってきって

973:デフォルトの名無しさん
09/01/19 23:41:25
そのサンプルはあくまでもcudaMallocPitch()がどういう風にメモリを確保するかの説明だからなぁ。
きっとライターは、まさかそのまま動かそうとされるは思っちゃいまいよ。

974:デフォルトの名無しさん
09/01/20 02:02:49
ところでドライバ180なんだが2Dモードだかなんだかで消費電力抑える為に
ドライバが自動的にクロックを下げて3Dゲーム動作させると上がる仕様なんだけど
CUDA実行しても上がらないぞｗ

975:デフォルトの名無しさん
09/01/20 02:08:30
ドライバを替えてくださいw
つーか、それって169.09のときに指摘されて直したんじゃなかったのか?＞NVIDIA

976:デフォルトの名無しさん
09/01/20 03:40:32
d_a[0][blockIdx.x][blockIdx.y][threadIdx.x]
って使おうとしたら"expression must have arithmetic or enum type"
ってでたんだけど何がいけないんすか？
教えてください。

977:デフォルトの名無しさん
09/01/20 07:58:55
>>976
d_aがポインタなら、途中で解決できない状況があるのかも。

978:デフォルトの名無しさん
09/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:デフォルトの名無しさん
09/01/20 08:59:56
パックが違うのだけが理由なら、
struct data {int a; char b; char dummy[3];}とでもしておけばいいじゃん。

980:デフォルトの名無しさん
09/01/20 13:53:14
ただでさえ少ない共有メモリをそんな無駄に使えない

981:デフォルトの名無しさん
09/01/20 20:44:37
>>977
世の中夜勤帰りで朝から寝てる人だっているんだよ？
引っ越しの時ちゃんと挨拶行った？
顔合わせたら軽く会話するとかしてちゃんとコンタクト取り続けてる？
日頃からそういうコミニュケーションが取れてればいつ洗濯機を回していいのか
いつ静かにしなければならないのか
迷惑を掛けないように生活出来るはずなんだが

982:デフォルトの名無しさん
09/01/20 23:36:58
>>980
マジで言っているのなら、設計が悪い。
どうしてもパディングしたくないくらい逼迫しているなら、int一個を切り分けて3バイトと1バイトで使え。

983:デフォルトの名無しさん
09/01/21 00:22:25
共有メモリが制限されてるのに無駄な領域作って
ほとんどをグローバルメモリに追いやる方がよっぽど設計が悪いでしょｗ

984:デフォルトの名無しさん
09/01/21 00:25:51
GPUのメモリレイテンシって１２とかの世界だぞ
CPU用のDDR2で５だからな
intを内部でcharとして扱うプロセスを考慮しても共有メモリ使った方が早いんだよ

985:デフォルトの名無しさん
09/01/21 00:35:27
CUDAで共有メモリを使うこと自体、無駄な同期待ちが発生するから避けたいところだと思うが。
coaxschedな読み書きができるなら、共有メモリより遅くないぞw

986:デフォルトの名無しさん
09/01/21 02:24:30
>>975
(エンバグの)歴史は繰り返す。

987:デフォルトの名無しさん
09/01/21 12:18:47
>>CUDAで共有メモリを使うこと自体、無駄な同期待ちが発生するから避けたいところだと思うが。
え？
共有メモリってCPUで言う所のただのレジスタみたいなもんで同期は指定しないと発生しないと思うけど

ところでローカル変数をすべて共有メモリに置いた場合のが早いんだけど
ローカル変数ってデバイス上のメモリに確保されるだけで共有メモリより遅いレベルのメモリを使ってるっぽいね

988:デフォルトの名無しさん
09/01/21 18:11:15
あー、表現が悪かった。共有メモリを共有目的に使うと同期を取る必要が出てくるから同期待ちが発生するということ。
レジスタみたいに使うのなら確かに関係なかったね。

で、レジスタよりも速いかどうかについてはptx見てみたいところ。

989:,,・´∀｀・,,）っ-○◎●
09/01/21 18:18:05
見た感じリードレイテンシはこれくらい

レジスタ＞＞Const Memory＞Shared Memory＞＞＞＞DRAM