- 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/
- 110 名前：デフォルトの名無しさん mailto:sage [2011/10/19(水) 13:25:43.68 ]
- >>108
＞cudaMalloc((void**)&hoge_d,sizeof(hoge_t));
意味のおさらい。
& hoge_dをcudaMalloc()に渡すことで、ローカル変数hoge_dにデバイス上のアドレスが格納される。
＞cudaMalloc((void**)&hoge_d->Pt,sizeof(int)*100);
& hoge_d->Ptは、意味としては & (* hoge_d).Pt に等しい。
つまり、hoge_dがデバイス上のアドレスを保持しているにも拘らず
ホスト上で* hoge_d するからsegmentation faultを引き起こす。
データ効率を考えると固定長配列にできないか検討するなどして見直すべきだが、
どうしてもこのままの構造でやりたいなら sizeof(int) * 100 確保したデバイス上の
アドレスをcudaMemcpy()でhoge_d->ptにコピーする必要がある。
って、この下りは>109が書いていることと同じだね。
- 111 名前：108 mailto:sage [2011/10/19(水) 17:14:09.68 ]
- >>109,>>110
お忙しい中、レス有難うございます。
>>110
データ構造については検討しようと思いますが、
いただいたアドバイスを参考に以下のようにしてみたところsegmentation faultがなくなりました。
109,110さんの意図していたものとなっておりますでしょうか？
int main(){
hoge_t *hoge:
hoge=(hoge_t*)malloc(sizeof(hoge_t));
cudaMalloc((void**)&(hoge->Pt),sizeof(int)*100);
hoge_t *hoge_d;
cudaMalloc((void**)&(hoge_d),sizeof(hoge_t));
cudaMemcpy(hoge_d,hoge,sizeof(hoge_t),cudaMemcpyHostToDevice);
//終了処理は省略
}
あるいはこういう感じでしょうか・・・
int main(){
int *tmp_pt;
cudaMalloc((void**)&(tmp_pt),sizeof(int)*100);
hoge_t *hoge_d;
cudaMalloc((void**)&(hoge_d),sizeof(hoge_t));
cudaMemcpy(hoge_d->Pt,tmp_pt,sizeof(int),cudaMemcpyHostToDevice);
//終了処理は省略
}
- 112 名前：デフォルトの名無しさん mailto:sage [2011/10/19(水) 21:10:14.17 ]
- 概ねいいんでない?
前者なら、私はhoge_t hoge;で定義しちゃうけど。
- 113 名前：109 mailto:sage [2011/10/19(水) 22:04:08.49 ]
- 俺が考えてたのは後者だな。
合ってると思うよ。
でもmemcpyのsizeof(int)はsizeof(int *)にした方が良いと思う。
- 114 名前：108 mailto:sage [2011/10/19(水) 23:57:45.47 ]
- 未だにポインタが使いこなせてないと痛感しました。
このたびはアドバイス有難うございました。
- 115 名前：デフォルトの名無しさん mailto:sage [2011/10/20(木) 00:12:23.07 ]
- ポインタ使える男の人ってモテそうだよね
- 116 名前：デフォルトの名無しさん mailto:sage [2011/10/20(木) 17:52:20.01 ]
- ポインタでいろんな所を指して、
覗いたりいじくり回したりするんですね。
- 117 名前：やんやん ◆yanyan72E. mailto:sage [2011/10/21(金) 01:37:52.34 ]
- トラックポイントではない
- 118 名前：デフォルトの名無しさん mailto:sage [2011/11/03(木) 20:05:27.91 ]
- デバイスから出る「the launch timed out and was terminated.」とかのメッセージを
CPU側で受信する方法ってどうやるのかな？コンソール出力されるんだからできると思うけど
- 119 名前：デフォルトの名無しさん mailto:sage [2011/11/03(木) 21:13:42.56 ]
- >>118
デバイス側といっても、cudart辺りがドライバからのステータスを見て出していると思われ。
従って、当然リダイレクトできる。
- 120 名前：デフォルトの名無しさん [2011/11/07(月) 04:01:37.49 ]
- CUBLASとCUSPARSEを使って倍精度で計算を行うCG法のプログラムを作成したのですが、
CUSPARSEのcsrmv関数だけが異様に遅く速度が出ません。
この原因としては何が考えられますでしょうか？
自分としてはCUBLASとCUSPARSEでそれぞれ別のhandleを使っているせいかと思っているのですが
それを確かめる方法や、CUBLASとCUSPARSEでhandleを共有する方法がわかりません...
- 121 名前：デフォルトの名無しさん mailto:sage [2011/11/07(月) 09:05:53.46 ]
- CG法と書いてあるので問題行列は対称疎行列でCUSPARSEからはcsrmv関数だけ、CUBLASからは
Level-1 BLASルーチンしか使っていないのだと思いますが、
実際にどの程度遅いのでしょうか？
ランダムアクセスを含むcsrmvはLevel-1 BLASルーチンとくらべてメモリアクセス効率が極端に下がるのは
仕様というか当然の結果です。
また古いバージョンのライブラリを使っていたりしませんか？
handle云々は基本的に関係ありません。関係するとすればよほど小さな行列を扱っている場合に限られます。
- 122 名前：デフォルトの名無しさん mailto:sage [2011/11/07(月) 15:42:01.27 ]
- 仰る通り対称疎行列でCUSPARSEからはcsrmv関数のみです。
行列のサイズはn×n n=122187で、ライブラリのバージョンは4.0です。
BLASでの総計算時間が12.66sに対し、
CUSPARSEでの総計算時間が80.36sという結果です。
このCUSPARSEでの計算時間はCPUで計算を行った場合よりも遅く、
nVidiaのCUDA toolkit 4.0 Performance Reportの結果と比較しても
非常に遅いと思います。
- 123 名前：デフォルトの名無しさん mailto:sage [2011/11/07(月) 18:49:04.47 ]
- 構造体の中にある配列はどのように確保してmemコピーすればいいんでしょうか？
#define N (32)
struct var_box{
int boxi;
float boxf;
};
struct var_str{
int *vari; // ~N*N
float var_g;
struct var_box *vb; // ~N
};
のような構造体がある時
struct var_str *vvv,*vvv_d;
vvv = (struct var_str*)malloc(sizeof(struct var_str));
vvv->vari = (int*)malloc((N*N)*sizeof(int));
vvv->vb = (struct var_box*)malloc(sizeof(struct var_box)*(N));
値代入
cudaMallocHost((void**)&vvv_d,sizeof(struct var_str));
cudaMallocHost((void**)&vvv_d->vari,sizeof(int)*(N*N));
cudaMallocHost((void**)&vvv_d->vb,sizeof(struct var_box)*(N));
cudaMemcpy(vvv_d, vvv, sizeof(struct var_str), cudaMemcpyHostToDevice);
cudaMemcpy(vvv_d->vari, vvv->vari, sizeof(int)*(N*N), cudaMemcpyHostToDevice);
cudaMemcpy(vvv_d->vb, vvv->vb, sizeof(struct var_box)*(N), cudaMemcpyHostToDevice);
GPUに送ってCPUに戻しただけですが float var_g; に関しては問題なくできていますが配列にした部分が送れていないみたいです。
そもそも確保の部分だけで0.5sかかっているのでちゃんとできてるかどうか怪しいです。
- 124 名前：デフォルトの名無しさん mailto:sage [2011/11/07(月) 18:52:46.57 ]
- 構造体にポインタがあるのはダメ
理屈で考えればわかるよな
- 125 名前：デフォルトの名無しさん mailto:sage [2011/11/07(月) 19:13:06.14 ]
- >>124
そうでしたか
わかりました
- 126 名前：デフォルトの名無しさん mailto:sage [2011/11/07(月) 19:34:23.39 ]
- 馬鹿でプログラミング初心者には難しいれす（^ρ^）
- 127 名前：やんやん ◆yanyan72E. mailto:sage [2011/11/07(月) 19:39:20.38 ]
- 無駄にでっかくメモリ確保された
ただのポインタがコピーされてるだけに見えるのは気のせい？
実体をコピーしようとしている意図は読み取れるけれど、
コードはその意図を反映してないような。。。
- 128 名前：やんやん ◆yanyan72E. mailto:sage [2011/11/07(月) 19:50:21.46 ]
- あ、考え違いか、すまん
- 129 名前：デフォルトの名無しさん mailto:sage [2011/11/07(月) 19:52:32.74 ]
- >>124
わからないのでその理屈を教えていただけないでしょうか
- 130 名前：デフォルトの名無しさん mailto:sage [2011/11/07(月) 20:15:14.16 ]
- cudaMallocHostはあるけどcudaMallocはしてないの?書き洩れ？
- 131 名前：デフォルトの名無しさん mailto:sage [2011/11/07(月) 20:28:54.09 ]
- デバイス側のcudamallocが無いよね？
cudaMallocHost((void**)&vvv_d->vari,sizeof(int)*(N*N));
ってvvv_d->variをホスト側にメモリ確保しちゃってんるんだけど
必要なことは例えば
cudaMallo(reinterpret_cast<void**>(&vvv_d->vari,sizeof(int)*(N*N));
じゃないの？
- 132 名前：デフォルトの名無しさん mailto:sage [2011/11/08(火) 09:53:28.63 ]
- 取り敢えず、ホスト側とデバイス側の二種類のメモリ空間を扱っていることを肝に銘じよう。
「構造体にポインタがある」ある場合、そのポインタがホスト側のメモリを指しているのか
デバイス側のメモリを指しているのか常に意識しないといけない。
それがいやだったら、ポインタではなく配列にしてしまえ。
cf.
struct var_str {
int var[N * N];
float var_g;
struct var_box vb[N];
};
- 133 名前：デフォルトの名無しさん mailto:sage [2011/11/08(火) 09:54:43.06 ]
- それ以前に、そんなデータ構造を渡しても素直に並列化できなくて速度が出ない悪寒。
- 134 名前：デフォルトの名無しさん mailto:sage [2011/11/08(火) 10:50:48.01 ]
- >>129
>>108-114の流れを読め．
- 135 名前：デフォルトの名無しさん [2011/11/08(火) 13:04:40.64 ]
- C# (.NET Framework 4.0)で動作するCUDAのライブラリってありますか？
CUDA.NETがそうだとは思うのですが、.NET Framework 2.0 or newerとしか書いていないので、
結局のところ.NET 4.0で動作するのかどうかわかりません。
- 136 名前：デフォルトの名無しさん mailto:sage [2011/11/08(火) 15:31:03.58 ]
- おまえ、そのレベルでよくこのスレに辿り着けたな。。
- 137 名前：デフォルトの名無しさん mailto:sage [2011/11/08(火) 18:02:17.87 ]
- とりあえず動かしてみろｗ
- 138 名前：デフォルトの名無しさん mailto:sage [2011/11/09(水) 04:38:11.56 ]
- 4.1RC出てるんだな
デベロッパプログラムに登録したいけど審査あるから、
他人のソースコンパイルして遊ぶだけの俺には無理
- 139 名前：デフォルトの名無しさん mailto:sage [2011/11/09(水) 14:48:59.26 ]
- 試しにデタラメ並べ立てて申請したら通った
んでSDKのサンプルを上から順番に走らせて遊んでたらブルスク食らった
危ねえ・・・
- 140 名前：デフォルトの名無しさん mailto:sage [2011/11/09(水) 22:34:20.87 ]
- >>138
審査あんの？俺普通に登録してOKでたけど
- 141 名前：デフォルトの名無しさん mailto:sage [2011/11/09(水) 22:45:40.52 ]
- 審査ってなにを調べるの？
- 142 名前：デフォルトの名無しさん mailto:sage [2011/11/09(水) 23:51:20.06 ]
- ウンコの重さとか
- 143 名前：デフォルトの名無しさん mailto:sage [2011/11/10(木) 00:31:38.44 ]
- Parallel Nsightのアカウントとは別なのか。
面倒くさすぎだろ。
- 144 名前：デフォルトの名無しさん mailto:sage [2011/11/10(木) 06:20:27.98 ]
- CUDA.NETってCUDA 3.0までで止まってるよね？
- 145 名前：デフォルトの名無しさん mailto:sage [2011/11/10(木) 12:49:08.28 ]
- >>144
くだんねーと（CUDANET）思ったんだろ作ってる方も
- 146 名前：デフォルトの名無しさん mailto:sage [2011/11/11(金) 00:24:14.22 ]
- 【審議中】
∧,,∧ ∧,,∧
∧ (´・ω・) (・ω・｀) ∧∧
( ´・ω) Ｕ) ( つと ﾉ(ω・｀ )
｜ Ｕ ( ´・) (・｀ ) と ﾉ
ｕ-ｕ （l ) ( ﾉu-ｕ
`ｕ-ｕ'. `ｕ-u'
- 147 名前：デフォルトの名無しさん mailto:sage [2011/11/14(月) 02:32:16.94 ]
- kernelの中でどの部分でレジスタがMAXになっているか、どの値を保持しているかっていうのは簡単にわかるものですか？
- 148 名前：デフォルトの名無しさん mailto:sage [2011/11/14(月) 11:50:00.02 ]
- 細かいことは兎も角、取り敢えずptx出力を見てみたら医院で内科医。
- 149 名前：デフォルトの名無しさん mailto:sage [2011/11/14(月) 18:06:51.94 ]
- shrQATest.hって4.0のtoolkitとSDK入れたらどこにあるはずですか？
- 150 名前：デフォルトの名無しさん mailto:sage [2011/11/15(火) 02:43:04.47 ]
- %ProgramData%\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.0\shared
- 151 名前：デフォルトの名無しさん mailto:sage [2011/11/15(火) 02:44:06.11 ]
- \inc
- 152 名前：デフォルトの名無しさん mailto:sage [2011/11/15(火) 16:34:29.61 ]
- 見つかりました
ありがとうございます
- 153 名前：デフォルトの名無しさん mailto:sage [2011/11/15(火) 21:47:31.35 ]
- linux kernelからGPGPU使いたいんだけどKGPU以外に何かいい方法ある？
- 154 名前：デフォルトの名無しさん mailto:sage [2011/11/15(火) 22:55:00.82 ]
- CUDA developer用としてCUDAサイトで配布されてるドライバと普通のドライバってどんな違いがあるの？
CUDAに必須と思ったら、普通のドライバインストールしてもCUDA開発できると書いている所があった。
- 155 名前：デフォルトの名無しさん mailto:sage [2011/11/15(火) 23:18:16.01 ]
- 一時期、developer用と一般用のバージョンが同じ時があったけど
md5sumは一致した。
おそらく、一般用ドライバのバージョンで対応が間に合っていないときに
beta版を出してるだけだと思う。
だから、両者でバージョンの新しい方を使えば問題ない。
- 156 名前：デフォルトの名無しさん mailto:sage [2011/11/16(水) 00:34:28.85 ]
- SC2011 OpenACC Joint Press Release - main
www.openacc-standard.org/announcements-1/nvidiacraypgicapsunveil%E2%80%98openacc%E2%80%99programmingstandardforparallelcomputing
- 157 名前：デフォルトの名無しさん mailto:sage [2011/11/16(水) 05:00:22.99 ]
- >>155
了解。レスありがと
- 158 名前：デフォルトの名無しさん [2011/11/22(火) 17:29:12.79 ]
-
すいません、経験者の方教えてください。
ホストメモリをページロックメモリで確保して、
ストリームを300個くらい起動して
MemcpyAnsyncでh->d ,kenrel, h->dで非同期実行させてるんですけど、
どうやらkernelでセグって停止してしまいます。
しかし、cuda-gdbで見てみるとどうもセグるようなポジションじゃないんです。
そこで質問なんですが、MemcpyAsyncでコピーリクエストをだしてコピーが始まり、
制御がカーネルに入り、コピーが完了したストリームを実行してる途中に、
まだコピー中の領域でセグった場合、どこでセグったように見えるのでしょうか？
やはりその時実行しているカーネルの中でセグっているように見えるのでしょうか？
- 159 名前：デフォルトの名無しさん mailto:sage [2011/11/27(日) 22:05:36.51 ]
- 全スレッドの計算結果を1コアに集約して1スレッドで処理したいんだけどいい方法あるかな？
リダクションのサンプルだと最後はCPUで合計だしこれでは都合悪いんだよね
全てGPUでやりたいのよね
- 160 名前：デフォルトの名無しさん mailto:sage [2011/11/27(日) 23:20:19.00 ]
- しかたないじゃん、CPUで合計出した方が速いんだから。
- 161 名前：デフォルトの名無しさん mailto:sage [2011/11/27(日) 23:28:38.68 ]
- >>159
バイナリツリーで足していくしか。
Ex. thread数が64のとき
if (threadIdx < 32) data[threadIdx] += data[threadIdx + 32];
__syncthreads();
if (threadIdx < 16) data[threadIdx] += data[threadIdx + 16];
__syncthreads();
if (threadIdx < 8) data[threadIdx] += data[threadIdx + 8];
__syncthreads();
if (threadIdx < 4) data[threadIdx] += data[threadIdx + 4];
__syncthreads();
if (threadIdx < 2) data[threadIdx] += data[threadIdx + 2];
__syncthreads();
if (threadIdx < 1) result = data[threadIdx] + data[threadIdx + 1];
>>160
データ量によってはそうでもないよ。
- 162 名前：デフォルトの名無しさん mailto:sage [2011/11/27(日) 23:35:14.30 ]
- >>161
＞データ量によってはそうでもないよ。
>>159は「1スレッドで処理したい」と言っている。
それに、「計算結果を集約」と言っているだけで足し算とは言っていない。
この条件の下では、データ量が多くなれば多くなるほどCPUで処理した方が速い。
データ数が極端に少なければ、転送時間の関係でGPUでしたほうが速いかもしれないが。
- 163 名前：デフォルトの名無しさん mailto:sage [2011/11/27(日) 23:40:05.87 ]
- あーそうか、まさか1スレッドだけ動かす方法が判らんと言う話とは思わなかった。
単純に、「CPUで合計」の代わりをやりたいのかと思ったよ。
- 164 名前：159 mailto:sage [2011/11/27(日) 23:55:11.26 ]
- >>160-163
レスサンクス
一度集約して最少二乗法を適用してから再度マルチスレッドで処理しようとしてるんだけど
一度CPUに戻すしかないかな・・・
- 165 名前：デフォルトの名無しさん mailto:sage [2011/11/28(月) 00:02:09.10 ]
- 最小二乗法ならマルチスレッドでできるじゃん。
- 166 名前：デフォルトの名無しさん mailto:sage [2011/11/29(火) 03:07:20.82 ]
- リダクションしろよ
1ブロック分しかできないってこと？
- 167 名前：デフォルトの名無しさん mailto:sage [2011/11/29(火) 10:48:39.03 ]
- サイズが決まっている3次元配列どうしの演算で演算回数も決まっているんだけど、
ブロック数とスレッド数を調整して演算回数ぴったりにするとコアレッシングしなくて効率悪いんだ。
こーいう時ってどゆアプローチ？
- 168 名前：デフォルトの名無しさん mailto:sage [2011/11/29(火) 14:42:48.05 ]
- CPUで計算させてみる
- 169 名前：デフォルトの名無しさん mailto:sage [2011/11/29(火) 15:26:04.48 ]
- もちろんCPUでは計算させてみてるよぉぉん(><)
1つ目の演算対象配列はギリギリL2に乗らないサイズ、2つ目の配列はVRAMにも乗らないサイズなので悩んでる。
- 170 名前：159 mailto:sage [2011/11/29(火) 16:52:31.44 ]
- もっと良いCPUかGPU買え
- 171 名前：デフォルトの名無しさん mailto:sage [2011/11/30(水) 02:20:28.60 ]
- >>159
同じブロック内の複数スレッドは共有メモリで集約できるけど、
ブロックをまたがっての集約ができないってこと？
手元のコードではブロック数の一時保存のための領域を確保し
そこにブロックごとの結果を放り込むカーネルを起動し、
次にブロック数が1、スレッド数が前のカーネルのブロック数の
カーネルを起動。その中でさらに集約して結果を1つにまとめている。
カーネル2回起動よりもCPUでやらせた方が速いかもしれないが。
- 172 名前：デフォルトの名無しさん mailto:sage [2011/11/30(水) 04:29:51.70 ]
- code.google.com/p/thrust/
のリダクション使ってみれば？
中身で何やってるかはよくわからんが
- 173 名前：デフォルトの名無しさん mailto:sage [2011/11/30(水) 08:40:15.94 ]
- カーネルの中で関数作って足せばいいんじゃない
threadが終わったら自作関数に投げて足すとか
バリアがいるけど
- 174 名前：デフォルトの名無しさん mailto:sage [2011/11/30(水) 08:52:34.48 ]
- 微妙に話が噛み合っていない悪寒。
つーか、>159が混乱に輪を掛けているし。
- 175 名前：159 mailto:sage [2011/11/30(水) 10:04:46.33 ]
- 各スレッドでサンプリングされたXiとYiを使って
Σ（Ai*Xi-Yi）^2を計算して最小化するパラメータのAiを求めたいんだよ
求まったAiを使ってまたGPUでマルチスレッドで処理をするから
CPUに戻したくなかっただけ
戻せばできるのは分かるけど、戻さないでできるならその方がいいと思ったわけね
- 176 名前：やんやん ◆yanyan72E. mailto:sage [2011/11/30(水) 10:24:24.37 ]
- なんか式がおかしい、Ai=Yi/Xi なんて答えを求めたい訳じゃなかろう？
- 177 名前：159 mailto:sage [2011/11/30(水) 11:22:19.91 ]
- >>176
うん
正確には多項式近似の最小二乗法
めんどいので略した
- 178 名前：171 mailto:sage [2011/11/30(水) 14:43:43.94 ]
- >>177
非線形最小二乗法ってこと？非線形最適化法しか知らないけど、目的関数が
特殊な場合にそれを上手に利用したのが非線形最小二乗って理解。
手元では対数尤度（二乗和みたいなもの）を171の手順で計算する関数を
CUDAで書いて、MCMCの中で使って推定値を求めているけど、同じことは
最小二乗法でもできると思うんだが…。
- 179 名前：159 mailto:sage [2011/11/30(水) 14:50:41.39 ]
- >>178
カーネルの二度起動だよね？
その方法でもできると思うけど、一度で済ませたいというのがニーズ
それよかMCMCをマルチスレッドでやるってのが分からんのだが…
あれはマルコフ連鎖だからほとんどマルチスレッドで効果上がらんだろw
- 180 名前：178 mailto:sage [2011/11/30(水) 15:21:59.61 ]
- >>179
こちらも一度の起動で済ませたかったけど方法が分からなかったので、
分かったら是非報告よろしく。
尤度の計算対象とするデータ数があまりにばかでかくて…。マルコフ連鎖は
順次計算だけど、その中の並列化は相当のメリットがあってね。
- 181 名前：やんやん ◆yanyan72E. mailto:sage [2011/11/30(水) 16:50:24.20 ]
- >>179
二度起動がいやなら、ここの
ttp://neuralgorithm.org/documents/global-thread-synchronization-on-cuda-ja/
globalsyncみたいなのを使って計算→集計→計算って
やらせればいいんだろうけれど、
なんか怖いなw
- 182 名前：デフォルトの名無しさん mailto:sage [2011/11/30(水) 16:50:29.44 ]
- CUDA Visual Plofiler への入力ファイル(cpjファイルとcsvファイル)をコンソールを使って作成したいのですが、可能でしょうか。
GUI環境ではメモリが足りず実行できないプログラムのビジュアルプロファイルを取りたいのですが…
- 183 名前：デフォルトの名無しさん mailto:sage [2011/12/02(金) 18:08:55.43 ]
- linuxで4.0のtoolkitを導入した時に古いものをリムーブするかと出てきたのでyesを選んだのですが、もしかして同じ階層のものを全部削除とかしています？
X windowすら上手く開かなくなったのですが…
- 184 名前：デフォルトの名無しさん mailto:sage [2011/12/03(土) 18:12:31.51 ]
- unixは知らんけど、windowsだとインストールしたファイルしか消されないよ
新しいの入れる度に毎回cl.hppを手動で追加してるけど、
アンインストールしてもそれだけ残ってる
- 185 名前：デフォルトの名無しさん mailto:sage [2011/12/03(土) 18:46:52.04 ]
- Toolkitを入れたからじゃなくて、Driverを入れたからじゃね?
Driverを入れるときにroot権限がないとXのconfigを更新できないから。
- 186 名前：デフォルトの名無しさん [2011/12/06(火) 06:03:27.50 ]
- 4.1 RC2
| NVIDIA Developer Zone
developer.nvidia.com/cuda-toolkit-41
- 187 名前：デフォルトの名無しさん mailto:sage [2011/12/06(火) 10:46:45.36 ]
- 蔵人にメル栓抜きツイスターきたのか
- 188 名前：デフォルトの名無しさん mailto:sage [2011/12/06(火) 11:57:45.78 ]
- コンパイルどれくらい早くなってんのかな
VCでコンパイルしてて小さなcuファイルでも数秒待たされるのがガチでイライラするんだけど
- 189 名前：デフォルトの名無しさん mailto:sage [2011/12/06(火) 12:05:12.77 ]
- vcでコンパイルしているなら、変わるわけないだろ。
vsからnvccを起動しているなら兎も角も。
- 190 名前：デフォルトの名無しさん mailto:sage [2011/12/06(火) 12:18:58.00 ]
- cuファイルとcppファイルに分けるのはどういう意図があるんでしょうか？thrustって何に使えるんですか？
- 191 名前：デフォルトの名無しさん mailto:sage [2011/12/06(火) 12:26:24.32 ]
- >>190
・nvccに任せておきたくない部分は、.cppに書く。
・抽象的に取り敢えず書くのに便利。
- 192 名前：デフォルトの名無しさん mailto:sage [2011/12/06(火) 16:39:18.49 ]
- cuにkernel_foo()を書いて
cppに書いたfoo()からkernel_foo()を呼ぶのが定石かと思ってた
そうすればopenclに移行したりしてもocl_kernel_foo()を呼び出すようにすれば変更すくないし。
- 193 名前：デフォルトの名無しさん mailto:sage [2011/12/06(火) 17:11:07.38 ]
- >>192
その定石に則れば、>191の前半を満たせるから問題ないよ。
ただ、それだとCPU側とGPU側で共通のロジックを更に別のソースに書かないといけなくなるから適材適所だと思う。
- 194 名前：デフォルトの名無しさん [2011/12/07(水) 00:20:58.16 ]
- 青木氏の本読んでも、ガイド読んでも、くすだれcudaを見ても全くわからなかったので質問します。
おそらくwarpの話で、cudaをまったく分かっていないゆえの質問だと思います。
cudaのデータの取扱いを調べるために、以下のような構造のデータをGPU側に送り、
typedef struct data{
int i,j; //初期化でi=1,j=0
}DATA;
Grid1コ,ThreadBlock１コ,総Thread数512コと指定して、以下のようなコードを実行させました。
__global__ void test(DATA *A){
int i = blockDim.x*blockIdx.x + threadIdx.x;
if(i%2==0){//threadIdの奇遇でiの値及び加算値を変更
A->i=2; A->j+=1;
}else{
A->i=3; A->j+=2;
}}
5,6回実行して、iの値はthreadの総数を奇遇どちらにしても3で不変でした。
jの値は実行するたび値が異なり、j=3,5,7,9のいずれかになりました。
iの結果は各warpの32コのThreadが順次if文を実行してて、
idが奇数のときの場合が後に実行されるから、結果がi=3となるのか？という理解でよろしいのでしょうか。
また、jの結果は青木氏の言う「加算命令を実行した結果は有効なスレッドに対してのみ反映される」
の理解がいると思うのですが、そもそも有効なスレッドがどう判定されているのかわかりません。
また512コのthreadがあるのに、jの値の結果が10以下になるのはどうも腑に落ちません。
i,jの値を決めているものは何か、ご教示願います。
- 195 名前：デフォルトの名無しさん mailto:sage [2011/12/07(水) 00:45:36.69 ]
- +=で同じアドレスに同時書き込みしてるから
- 196 名前：デフォルトの名無しさん [2011/12/07(水) 01:09:59.92 ]
- >>195
+=で同じアドレスに同時書き込みすると、内部で何が起こるんですか？
- 197 名前：デフォルトの名無しさん mailto:sage [2011/12/07(水) 01:27:31.21 ]
- 競合状態が発生してんじゃないの？
atomic演算とか同期が必要だと思うよ。
512スレッドで同一アドレスの変数の読み書きしてんだから。
まず５１２個の要素の配列作って、添え字にスレッド番号（上のi）を指定して確認してみたら？
書籍ならcuda exampleも買って読むといいかもね
- 198 名前：デフォルトの名無しさん [2011/12/07(水) 02:05:18.21 ]
- >>197
>競合状態が発生してんじゃないの？
>>１９５のコメと合わせて考えるに、なんとなく予想はしてましたけど、取り合いになってるんですね...
>atomic演算とか同期が必要だと思うよ。
まだザックリとしか勉強してないので、atomic演算は知らなかったです。あとで試してみます。
同期を行う場合だったら、どうすればいいのだろう。
>まず５１２個の要素の配列作って、添え字にスレッド番号（上のi）を指定して確認してみたら？
それは分岐条件がきちんと実行しているのか見るためのテストをしたとき確認しました。
- 199 名前：デフォルトの名無しさん mailto:sage [2011/12/07(水) 07:39:11.25 ]
- 根本的にプログラミングの基礎が抜けている悪寒。
- 200 名前：デフォルトの名無しさん mailto:sage [2011/12/07(水) 09:07:22.99 ]
- >>198
atomicという排他的処理を知らないのなら並列化プログラムをやるのは早すぎる。
増してCUDAという特殊な並列化をやるのはもっと早すぎる。
- 201 名前：デフォルトの名無しさん [2011/12/07(水) 12:02:55.64 ]
- >>199
どこまで戻ればいいんでしょうか...
>>200
...これ使って課題提出しなきゃならんので、早すぎると言われても後に引けないです...
- 202 名前：デフォルトの名無しさん [2011/12/07(水) 13:13:17.16 ]
- >>198
>>201
ですけど、
>>200の"atomicという排他的処理を知らないのなら並列化プログラムをやるのは早すぎる。"
では並列化プログラムをやるにあたり、
どういったことを勉強して、どういった手順でやればいいんでしょうか？
そこがよくわからず、私はいきなりcudaに突っ込んだので四苦八苦してるわけですが...
- 203 名前：デフォルトの名無しさん mailto:sage [2011/12/07(水) 14:10:39.88 ]
- >>202
同期とロック
- 204 名前：デフォルトの名無しさん mailto:sage [2011/12/07(水) 15:51:27.47 ]
- 青木本読んだんじゃないのん？
x+=1ってのは
1.xを読み出す
2.読み出した値に1を加算する
3.結果をxに格納する
って手順なんだけど以下略
- 205 名前：デフォルトの名無しさん [2011/12/07(水) 16:45:28.63 ]
- >>203
同期とロックですか。勉強します。
>>204
青木本は読みましたけど、ものにしたっていう状態じゃないです...
>x+=1ってのは...
普段、なんとなく使ってるので... 勉強不足ですみません。
- 206 名前：デフォルトの名無しさん [2011/12/07(水) 17:29:38.47 ]
- ここの一番下に載ってる資料の4-7ページにatomicの説明が少しある。
ttp://accc.riken.jp/HPC/training.html
- 207 名前：デフォルトの名無しさん [2011/12/07(水) 17:58:07.38 ]
- >>206
ありがとうございます。参照します。
- 208 名前：デフォルトの名無しさん mailto:sage [2011/12/07(水) 23:52:25.72 ]
- この本でも読んでみると良い。
日本語でわかりやすい。
www.amazon.co.jp/gp/product/4798014621/
- 209 名前：デフォルトの名無しさん [2011/12/08(木) 00:42:53.72 ]
- >>208
ありがとうございます。明日早速本屋行ってきます。
- 210 名前：デフォルトの名無しさん [2011/12/08(木) 17:37:10.02 ]
- 久しぶりに新しいCUDAの本が出たようだ。
www.amazon.co.jp/dp/4906608000/ref=cm_sw_r_fa_dp_454Zob03VNZ25
|
 |
