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