【GPGPU】NVIDIA CUDA質問スレッド

>>941
異端だな

どっちでもいいんじゃね

ちなみにnudeはニュードとヌードの両方の発音が有

キュもクも似てるし、まあ仮に間違って言ったとしても気づかれないんじゃない？

ホムペにはクーダと嫁
と書かれてたはず

"The cuda and my wife"

たしか英語版のfaqには
kuh-da
と書いてあったはず

2.1はJITサポートがミソ？

2.1をつついた人に質問
2.1ではカーネルに渡す引数でポインタ配列は使えますか?

共有メモリからグローバルメモリに連続領域をコピーするにはどうすればいいでしょうか
カーネルから呼び出せるmemcpyみたいなのがあれば教えてください

cudaMemcpyが長いようで、the launch timed out and was terminated.が
でてしまいす。これなんとかするほうほうありましたっけ?

>>952
転送量を減らす。

>>951
cudaMemcpy

>>951
スレッド内で連続的にメモリにアクセスするようなコードは望ましくない。
coalscedにアクセスするとなると、自分でコードを書くしかない。
共有メモリから転送を始める前に、同期を取ることも忘れずに。

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

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

>>955
sizeof(float)*10000バイトを100回連続コピーでした。

>>954
＞ >>952
＞ どんなハードウェアの組み合わせ

>>956
同一データ部分を100箇所にコピー？
独立の100箇所をコピー？

>>956
それ、本当にメモリのコピーでタイムアウトしてる？
カーネル実行後に同期取らずにすぐにメモリのコピー始めて
100回コピーのどこかでカーネル実行がタイムアウトしてるとか。

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

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

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

>>960
＞ 下手な突っ込みはお郷が知れるよ。


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

そういう意味か。あの書き方じゃぁ、誤解されるよ。
まぁ、大量転送じゃないと言う意味で、今回は関係なさそうだね。

>>952の話
>>957
ATIの統合チップセットでオンボードのATIビデオカードを無効
PCI-E 1.0 x16バスにGeforce9600GT 512MB
CPUはAMDデュアルプロセッサ(結構遅いやつ)

>>952の話
>>959
カーネル実行
a:cudaのエラーチェック(カーネルのエラーのありなし)
cudaThreadsSynchronize
for(100回
結果データをデバイスからホストにコピー
b:cudaのエラーチェック(直前のcudaMemcpyのエラーのありなし)
)
bのとこでthe launch timed out and was terminatedが出てるんです。

先週は時間がきてあきらめて帰ったのでここまでです。
言われたようにcudaThreadsSynchronizeの戻り値見てみたほうがよさそう
ですね。

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

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

JCublasについて
おしえてエロイ人

>>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でやる意味
あるの?と思いますが。自分で書いてなんなのですが

お前はあほかｗ

タイムアウトになる原因はそのでかいループのせい
せいぜいミリ秒単位でタイムアウトを判断してるから
ミリ秒以下で応答しないようなカーネルは絶対に書いては駄目

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

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

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

>>952の話
>>968のように言われるのは分かって書いてみたんだけど
NVIDIA_CUDA_Programming_Guide_2.0.pdfのp34-35のようなコードが
マニュアルに書いてあるのがおかしいと思う。
__global__ void myKernel(float* devPtr, int pitch){}
そもそもこんな書き方じたいが書けるけど間違えな使い方だと。
この書き方しているとこにやらないようにこの部分に×印つけてほしい。

あとはコンパイラがえらくなったらfor多重ループをうまく処理する
アセンブラぐらい作ってほしい。(OpenMPでパラレルfor指定すると#のタグだけ
でプロセッサ使ってきって高速化なるんだけどね。そこまでぐらい将来
的にはしてほしい。)

文句言っても何にもならない。
そう悟る時がくるまで気長に待ちましょうよ。
そうすれば成長しまっせ。

それよりも先ず、日本語を何とかしてくれ。
＞間違えな使い方
＞しているとこにやらないように
＞プロセッサ使ってきって

そのサンプルはあくまでもcudaMallocPitch()がどういう風にメモリを確保するかの説明だからなぁ。
きっとライターは、まさかそのまま動かそうとされるは思っちゃいまいよ。

ところでドライバ180なんだが2Dモードだかなんだかで消費電力抑える為に
ドライバが自動的にクロックを下げて3Dゲーム動作させると上がる仕様なんだけど
CUDA実行しても上がらないぞｗ

ドライバを替えてくださいw
つーか、それって169.09のときに指摘されて直したんじゃなかったのか?＞NVIDIA

d_a[0][blockIdx.x][blockIdx.y][threadIdx.x]
って使おうとしたら"expression must have arithmetic or enum type"
ってでたんだけど何がいけないんすか？
教えてください。

>>976
d_aがポインタなら、途中で解決できない状況があるのかも。

共有メモリで構造体を使いたいのだけど
例えば
struct data {
int a;
char b;
}

func<<<dim3(), dim3(), SIZE * sizeof(data)>>>();

__global__ kernel(){
__shared__ data d[SIZE];
}

こんな感じでやるとCUDAとC側のパック構造が違うせいだと思うけどおかしくなる
どうやれば出来るの？

パックが違うのだけが理由なら、
struct data {int a; char b; char dummy[3];}とでもしておけばいいじゃん。

ただでさえ少ない共有メモリをそんな無駄に使えない

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

>>980
マジで言っているのなら、設計が悪い。
どうしてもパディングしたくないくらい逼迫しているなら、int一個を切り分けて3バイトと1バイトで使え。

共有メモリが制限されてるのに無駄な領域作って
ほとんどをグローバルメモリに追いやる方がよっぽど設計が悪いでしょｗ

GPUのメモリレイテンシって１２とかの世界だぞ
CPU用のDDR2で５だからな
intを内部でcharとして扱うプロセスを考慮しても共有メモリ使った方が早いんだよ

CUDAで共有メモリを使うこと自体、無駄な同期待ちが発生するから避けたいところだと思うが。
coaxschedな読み書きができるなら、共有メモリより遅くないぞw

>>975
(エンバグの)歴史は繰り返す。

>>CUDAで共有メモリを使うこと自体、無駄な同期待ちが発生するから避けたいところだと思うが。
え？
共有メモリってCPUで言う所のただのレジスタみたいなもんで同期は指定しないと発生しないと思うけど

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

あー、表現が悪かった。共有メモリを共有目的に使うと同期を取る必要が出てくるから同期待ちが発生するということ。
レジスタみたいに使うのなら確かに関係なかったね。

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

見た感じリードレイテンシはこれくらい

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

GPGPU#2
http://pc11.2ch.net/test/read.cgi/tech/1188374938/
くだすれCUDAスレ
http://pc11.2ch.net/test/read.cgi/tech/1206152032/

次スレどうする? 必要?

CUDAスレ統合で様子見るってことでいいんでない?
分けるほど流れ速くないよ。

質問スレってより雑談スレだもんなｗｗｗｗｗ