【GPGPU】くだすれCUDAスレ part5【NVIDIA】

【GPGPU】くだすれCUDAスレ part5【NVIDIA】 at TECH

[前50を表示]
100:デフォルトの名無しさん
11/10/10 20:49:05.67
カーネル側のソース分割ってできないんだっけ？
環境はLinux　cuda4.0

101:デフォルトの名無しさん
11/10/10 22:00:17.43
できるけどコンスタントメモリが変な事になったような気がする

102:デフォルトの名無しさん
11/10/10 22:04:05.94
このままだと5万行くらいになっちゃいそう・・・

103:デフォルトの名無しさん
11/10/10 22:39:08.05
>>102
正直処理対象がCUDAに向いてないような。。。

104:デフォルトの名無しさん
11/10/14 02:20:30.66
とある文献でレイテンシを表す際, 種類によってclock, cycleが使い分けられていました。

clock : texture cache, constant cache
cycle : global memory, shared memory

なんとなくメモリ, キャッシュで使い分けられているのはわかりますが, このままでは速度の比較ができません。
二つの違いは何なのか, またcycle-clock間で単位の変換が可能なのか, もしご存知の方がいたらご教授ください。

105:デフォルトの名無しさん
11/10/17 13:50:45.89
カーネル関数内で立てたフラグをホスト側で判断する事って出来る？

106:デフォルトの名無しさん
11/10/17 14:13:41.38
グローバルメモリ介してできるだろ

107:デフォルトの名無しさん
11/10/18 02:35:47.28
はじめてのCUDAプログラミングの9章に載ってるcublas使うプログラム実行したんだけど、
cublasInit()あたりで「～0x7c812afb で初回の例外が発生しました: Microsoft C++ の
例外: cudaError_enum～」って出てうまく動かないんだけどなして？

108:デフォルトの名無しさん
11/10/19 00:53:08.52
かなり初歩的な質問で申し訳ないのですが
構造体hoge_tにバッファを持たせPtにバッファのアドレスが入るようにしたいのですが
以下のようにするとsegmentation faultになってしまいます。
このようなことは可能なのでしょうか。

struct hoge_t{
int* Pt;
int a,b,c,d;
}

int main(){
hoge_t *hoge_d;
cudaMalloc((void**)&hoge_d,sizeof(hoge_t));
cudaMalloc((void**)&hoge_d->Pt,sizeof(int)*100);
//終了処理は省略
}

109:デフォルトの名無しさん
11/10/19 11:56:08.38
適当なポインタで確保した後に，
そのアドレス転送してやればいいんじゃない？

110:デフォルトの名無しさん
11/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
11/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:デフォルトの名無しさん
11/10/19 21:10:14.17
概ねいいんでない?
前者なら、私はhoge_t hoge;で定義しちゃうけど。

113:109
11/10/19 22:04:08.49
俺が考えてたのは後者だな。
合ってると思うよ。
でもmemcpyのsizeof(int)はsizeof(int *)にした方が良いと思う。

114:108
11/10/19 23:57:45.47
未だにポインタが使いこなせてないと痛感しました。
このたびはアドバイス有難うございました。

115:デフォルトの名無しさん
11/10/20 00:12:23.07
ポインタ使える男の人ってモテそうだよね

116:デフォルトの名無しさん
11/10/20 17:52:20.01
ポインタでいろんな所を指して、
覗いたりいじくり回したりするんですね。

117:やんやん ◆yanyan72E.
11/10/21 01:37:52.34
トラックポイントではない

118:デフォルトの名無しさん
11/11/03 20:05:27.91
デバイスから出る「the launch timed out and was terminated.」とかのメッセージを
CPU側で受信する方法ってどうやるのかな？コンソール出力されるんだからできると思うけど

119:デフォルトの名無しさん
11/11/03 21:13:42.56
>>118
デバイス側といっても、cudart辺りがドライバからのステータスを見て出していると思われ。
従って、当然リダイレクトできる。

120:デフォルトの名無しさん
11/11/07 04:01:37.49
CUBLASとCUSPARSEを使って倍精度で計算を行うCG法のプログラムを作成したのですが、
CUSPARSEのcsrmv関数だけが異様に遅く速度が出ません。
この原因としては何が考えられますでしょうか？

自分としてはCUBLASとCUSPARSEでそれぞれ別のhandleを使っているせいかと思っているのですが
それを確かめる方法や、CUBLASとCUSPARSEでhandleを共有する方法がわかりません...

121:デフォルトの名無しさん
11/11/07 09:05:53.46
CG法と書いてあるので問題行列は対称疎行列でCUSPARSEからはcsrmv関数だけ、CUBLASからは
Level-1 BLASルーチンしか使っていないのだと思いますが、

実際にどの程度遅いのでしょうか？

ランダムアクセスを含むcsrmvはLevel-1 BLASルーチンとくらべてメモリアクセス効率が極端に下がるのは
仕様というか当然の結果です。

また古いバージョンのライブラリを使っていたりしませんか？

handle云々は基本的に関係ありません。関係するとすればよほど小さな行列を扱っている場合に限られます。

122:デフォルトの名無しさん
11/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:デフォルトの名無しさん
11/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:デフォルトの名無しさん
11/11/07 18:52:46.57
構造体にポインタがあるのはダメ
理屈で考えればわかるよな

125:デフォルトの名無しさん
11/11/07 19:13:06.14
>>124
そうでしたか
わかりました

126:デフォルトの名無しさん
11/11/07 19:34:23.39
馬鹿でプログラミング初心者には難しいれす（^ρ^）

127:やんやん ◆yanyan72E.
11/11/07 19:39:20.38
無駄にでっかくメモリ確保された
ただのポインタがコピーされてるだけに見えるのは気のせい？
実体をコピーしようとしている意図は読み取れるけれど、
コードはその意図を反映してないような。。。

128:やんやん ◆yanyan72E.
11/11/07 19:50:21.46
あ、考え違いか、すまん

129:デフォルトの名無しさん
11/11/07 19:52:32.74
>>124
わからないのでその理屈を教えていただけないでしょうか

130:デフォルトの名無しさん
11/11/07 20:15:14.16
cudaMallocHostはあるけどcudaMallocはしてないの?書き洩れ？

131:デフォルトの名無しさん
11/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:デフォルトの名無しさん
11/11/08 09:53:28.63
取り敢えず、ホスト側とデバイス側の二種類のメモリ空間を扱っていることを肝に銘じよう。
「構造体にポインタがある」ある場合、そのポインタがホスト側のメモリを指しているのか
デバイス側のメモリを指しているのか常に意識しないといけない。

それがいやだったら、ポインタではなく配列にしてしまえ。
cf.
struct var_str {
int var[N * N];
float var_g;
struct var_box vb[N];
};

133:デフォルトの名無しさん
11/11/08 09:54:43.06
それ以前に、そんなデータ構造を渡しても素直に並列化できなくて速度が出ない悪寒。

134:デフォルトの名無しさん
11/11/08 10:50:48.01
>>129
>>108-114の流れを読め．

135:デフォルトの名無しさん
11/11/08 13:04:40.64
C# (.NET Framework 4.0)で動作するCUDAのライブラリってありますか？
CUDA.NETがそうだとは思うのですが、.NET Framework 2.0 or newerとしか書いていないので、
結局のところ.NET 4.0で動作するのかどうかわかりません。

136:デフォルトの名無しさん
11/11/08 15:31:03.58
おまえ、そのレベルでよくこのスレに辿り着けたな。。

137:デフォルトの名無しさん
11/11/08 18:02:17.87
とりあえず動かしてみろｗ

138:デフォルトの名無しさん
11/11/09 04:38:11.56
4.1RC出てるんだな
デベロッパプログラムに登録したいけど審査あるから、
他人のソースコンパイルして遊ぶだけの俺には無理

139:デフォルトの名無しさん
11/11/09 14:48:59.26
試しにデタラメ並べ立てて申請したら通った
んでSDKのサンプルを上から順番に走らせて遊んでたらブルスク食らった
危ねえ・・・

140:デフォルトの名無しさん
11/11/09 22:34:20.87
>>138
審査あんの？俺普通に登録してOKでたけど

141:デフォルトの名無しさん
11/11/09 22:45:40.52
審査ってなにを調べるの？

142:デフォルトの名無しさん
11/11/09 23:51:20.06
ウンコの重さとか

143:デフォルトの名無しさん
11/11/10 00:31:38.44
Parallel Nsightのアカウントとは別なのか。
面倒くさすぎだろ。

144:デフォルトの名無しさん
11/11/10 06:20:27.98
CUDA.NETってCUDA 3.0までで止まってるよね？

145:デフォルトの名無しさん
11/11/10 12:49:08.28
>>144
くだんねーと（CUDANET）思ったんだろ作ってる方も

146:デフォルトの名無しさん
11/11/11 00:24:14.22
【審議中】
　　　 ∧,,∧　 ∧,,∧
　∧ (´・ω・) (・ω・｀) ∧∧
(　´・ω)　Ｕ) ( つとﾉ(ω・｀ )
｜　Ｕ (　　´・) (・｀　　)　とﾉ
　ｕ-ｕ（l 　　 ) (　　　ﾉu-ｕ
　　　　 `ｕ-ｕ'.　`ｕ-u'

147:デフォルトの名無しさん
11/11/14 02:32:16.94
kernelの中でどの部分でレジスタがMAXになっているか、どの値を保持しているかっていうのは簡単にわかるものですか？

148:デフォルトの名無しさん
11/11/14 11:50:00.02
細かいことは兎も角、取り敢えずptx出力を見てみたら医院で内科医。

149:デフォルトの名無しさん
11/11/14 18:06:51.94
shrQATest.hって4.0のtoolkitとSDK入れたらどこにあるはずですか？

150:デフォルトの名無しさん
11/11/15 02:43:04.47
%ProgramData%\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.0\shared

151:デフォルトの名無しさん
11/11/15 02:44:06.11
\inc

152:デフォルトの名無しさん
11/11/15 16:34:29.61
見つかりました
ありがとうございます

153:デフォルトの名無しさん
11/11/15 21:47:31.35
linux kernelからGPGPU使いたいんだけどKGPU以外に何かいい方法ある？

154:デフォルトの名無しさん
11/11/15 22:55:00.82
CUDA developer用としてCUDAサイトで配布されてるドライバと普通のドライバってどんな違いがあるの？
CUDAに必須と思ったら、普通のドライバインストールしてもCUDA開発できると書いている所があった。

155:デフォルトの名無しさん
11/11/15 23:18:16.01
一時期、developer用と一般用のバージョンが同じ時があったけど
md5sumは一致した。
おそらく、一般用ドライバのバージョンで対応が間に合っていないときに
beta版を出してるだけだと思う。

だから、両者でバージョンの新しい方を使えば問題ない。

156:デフォルトの名無しさん
11/11/16 00:34:28.85
SC2011 OpenACC Joint Press Release - main
URLﾘﾝｸ(www.openacc-standard.org)

157:デフォルトの名無しさん
11/11/16 05:00:22.99
>>155
了解。レスありがと

158:デフォルトの名無しさん
11/11/22 17:29:12.79

すいません、経験者の方教えてください。

ホストメモリをページロックメモリで確保して、
ストリームを300個くらい起動して
MemcpyAnsyncでh->d ,kenrel, h->dで非同期実行させてるんですけど、
どうやらkernelでセグって停止してしまいます。
しかし、cuda-gdbで見てみるとどうもセグるようなポジションじゃないんです。

そこで質問なんですが、MemcpyAsyncでコピーリクエストをだしてコピーが始まり、
制御がカーネルに入り、コピーが完了したストリームを実行してる途中に、
まだコピー中の領域でセグった場合、どこでセグったように見えるのでしょうか？

やはりその時実行しているカーネルの中でセグっているように見えるのでしょうか？

159:デフォルトの名無しさん
11/11/27 22:05:36.51
全スレッドの計算結果を1コアに集約して1スレッドで処理したいんだけどいい方法あるかな？
リダクションのサンプルだと最後はCPUで合計だしこれでは都合悪いんだよね
全てGPUでやりたいのよね

160:デフォルトの名無しさん
11/11/27 23:20:19.00
しかたないじゃん、CPUで合計出した方が速いんだから。

161:デフォルトの名無しさん
11/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:デフォルトの名無しさん
11/11/27 23:35:14.30
>>161
＞データ量によってはそうでもないよ。

>>159は「1スレッドで処理したい」と言っている。
それに、「計算結果を集約」と言っているだけで足し算とは言っていない。
この条件の下では、データ量が多くなれば多くなるほどCPUで処理した方が速い。
データ数が極端に少なければ、転送時間の関係でGPUでしたほうが速いかもしれないが。

163:デフォルトの名無しさん
11/11/27 23:40:05.87
あーそうか、まさか1スレッドだけ動かす方法が判らんと言う話とは思わなかった。
単純に、「CPUで合計」の代わりをやりたいのかと思ったよ。

164:159
11/11/27 23:55:11.26
>>160-163
レスサンクス
一度集約して最少二乗法を適用してから再度マルチスレッドで処理しようとしてるんだけど
一度CPUに戻すしかないかな・・・

165:デフォルトの名無しさん
11/11/28 00:02:09.10
最小二乗法ならマルチスレッドでできるじゃん。

166:デフォルトの名無しさん
11/11/29 03:07:20.82
リダクションしろよ
1ブロック分しかできないってこと？

167:デフォルトの名無しさん
11/11/29 10:48:39.03
サイズが決まっている3次元配列どうしの演算で演算回数も決まっているんだけど、
ブロック数とスレッド数を調整して演算回数ぴったりにするとコアレッシングしなくて効率悪いんだ。
こーいう時ってどゆアプローチ？

168:デフォルトの名無しさん
11/11/29 14:42:48.05
CPUで計算させてみる

169:デフォルトの名無しさん
11/11/29 15:26:04.48
もちろんCPUでは計算させてみてるよぉぉん(><)
1つ目の演算対象配列はギリギリL2に乗らないサイズ、2つ目の配列はVRAMにも乗らないサイズなので悩んでる。

170:159
11/11/29 16:52:31.44
もっと良いCPUかGPU買え

171:デフォルトの名無しさん
11/11/30 02:20:28.60
>>159
同じブロック内の複数スレッドは共有メモリで集約できるけど、
ブロックをまたがっての集約ができないってこと？
手元のコードではブロック数の一時保存のための領域を確保し
そこにブロックごとの結果を放り込むカーネルを起動し、
次にブロック数が1、スレッド数が前のカーネルのブロック数の
カーネルを起動。その中でさらに集約して結果を1つにまとめている。
カーネル2回起動よりもCPUでやらせた方が速いかもしれないが。

172:デフォルトの名無しさん
11/11/30 04:29:51.70
URLﾘﾝｸ(code.google.com)
のリダクション使ってみれば？

中身で何やってるかはよくわからんが

173:デフォルトの名無しさん
11/11/30 08:40:15.94
カーネルの中で関数作って足せばいいんじゃない
threadが終わったら自作関数に投げて足すとか
バリアがいるけど

174:デフォルトの名無しさん
11/11/30 08:52:34.48
微妙に話が噛み合っていない悪寒。
つーか、>159が混乱に輪を掛けているし。

175:159
11/11/30 10:04:46.33
各スレッドでサンプリングされたXiとYiを使って
Σ（Ai*Xi-Yi）^2を計算して最小化するパラメータのAiを求めたいんだよ

求まったAiを使ってまたGPUでマルチスレッドで処理をするから
CPUに戻したくなかっただけ

戻せばできるのは分かるけど、戻さないでできるならその方がいいと思ったわけね

176:やんやん ◆yanyan72E.
11/11/30 10:24:24.37
なんか式がおかしい、Ai=Yi/Xi なんて答えを求めたい訳じゃなかろう？

177:159
11/11/30 11:22:19.91
>>176
うん
正確には多項式近似の最小二乗法
めんどいので略した

178:171
11/11/30 14:43:43.94
>>177
非線形最小二乗法ってこと？非線形最適化法しか知らないけど、目的関数が
特殊な場合にそれを上手に利用したのが非線形最小二乗って理解。
手元では対数尤度（二乗和みたいなもの）を171の手順で計算する関数を
CUDAで書いて、MCMCの中で使って推定値を求めているけど、同じことは
最小二乗法でもできると思うんだが…。

179:159
11/11/30 14:50:41.39
>>178
カーネルの二度起動だよね？
その方法でもできると思うけど、一度で済ませたいというのがニーズ

それよかMCMCをマルチスレッドでやるってのが分からんのだが…
あれはマルコフ連鎖だからほとんどマルチスレッドで効果上がらんだろw

180:178
11/11/30 15:21:59.61
>>179
こちらも一度の起動で済ませたかったけど方法が分からなかったので、
分かったら是非報告よろしく。
尤度の計算対象とするデータ数があまりにばかでかくて…。マルコフ連鎖は
順次計算だけど、その中の並列化は相当のメリットがあってね。

181:やんやん ◆yanyan72E.
11/11/30 16:50:24.20
>>179
二度起動がいやなら、ここの
URLﾘﾝｸ(neuralgorithm.org)
globalsyncみたいなのを使って計算→集計→計算って
やらせればいいんだろうけれど、
なんか怖いなw

182:デフォルトの名無しさん
11/11/30 16:50:29.44
CUDA Visual Plofiler への入力ファイル(cpjファイルとcsvファイル)をコンソールを使って作成したいのですが、可能でしょうか。
GUI環境ではメモリが足りず実行できないプログラムのビジュアルプロファイルを取りたいのですが…

183:デフォルトの名無しさん
11/12/02 18:08:55.43
linuxで4.0のtoolkitを導入した時に古いものをリムーブするかと出てきたのでyesを選んだのですが、もしかして同じ階層のものを全部削除とかしています？
X windowすら上手く開かなくなったのですが…

184:デフォルトの名無しさん
11/12/03 18:12:31.51
unixは知らんけど、windowsだとインストールしたファイルしか消されないよ
新しいの入れる度に毎回cl.hppを手動で追加してるけど、
アンインストールしてもそれだけ残ってる

185:デフォルトの名無しさん
11/12/03 18:46:52.04
Toolkitを入れたからじゃなくて、Driverを入れたからじゃね?
Driverを入れるときにroot権限がないとXのconfigを更新できないから。

186:デフォルトの名無しさん
11/12/06 06:03:27.50
4.1 RC2
| NVIDIA Developer Zone
URLﾘﾝｸ(developer.nvidia.com)

187:デフォルトの名無しさん
11/12/06 10:46:45.36
蔵人にメル栓抜きツイスターきたのか

188:デフォルトの名無しさん
11/12/06 11:57:45.78
コンパイルどれくらい早くなってんのかな
VCでコンパイルしてて小さなcuファイルでも数秒待たされるのがガチでイライラするんだけど

189:デフォルトの名無しさん
11/12/06 12:05:12.77
vcでコンパイルしているなら、変わるわけないだろ。
vsからnvccを起動しているなら兎も角も。

190:デフォルトの名無しさん
11/12/06 12:18:58.00
cuファイルとcppファイルに分けるのはどういう意図があるんでしょうか？thrustって何に使えるんですか？

191:デフォルトの名無しさん
11/12/06 12:26:24.32
>>190
・nvccに任せておきたくない部分は、.cppに書く。
・抽象的に取り敢えず書くのに便利。

192:デフォルトの名無しさん
11/12/06 16:39:18.49
cuにkernel_foo()を書いて
cppに書いたfoo()からkernel_foo()を呼ぶのが定石かと思ってた
そうすればopenclに移行したりしてもocl_kernel_foo()を呼び出すようにすれば変更すくないし。

193:デフォルトの名無しさん
11/12/06 17:11:07.38
>>192
その定石に則れば、>191の前半を満たせるから問題ないよ。
ただ、それだとCPU側とGPU側で共通のロジックを更に別のソースに書かないといけなくなるから適材適所だと思う。

194:デフォルトの名無しさん
11/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:デフォルトの名無しさん
11/12/07 00:45:36.69
+=で同じアドレスに同時書き込みしてるから

196:デフォルトの名無しさん
11/12/07 01:09:59.92
>>195
+=で同じアドレスに同時書き込みすると、内部で何が起こるんですか？

197:デフォルトの名無しさん
11/12/07 01:27:31.21
競合状態が発生してんじゃないの？
atomic演算とか同期が必要だと思うよ。
512スレッドで同一アドレスの変数の読み書きしてんだから。

まず５１２個の要素の配列作って、添え字にスレッド番号（上のi）を指定して確認してみたら？
書籍ならcuda exampleも買って読むといいかもね

198:デフォルトの名無しさん
11/12/07 02:05:18.21
>>197
>競合状態が発生してんじゃないの？
>>１９５のコメと合わせて考えるに、なんとなく予想はしてましたけど、取り合いになってるんですね...

>atomic演算とか同期が必要だと思うよ。
まだザックリとしか勉強してないので、atomic演算は知らなかったです。あとで試してみます。
同期を行う場合だったら、どうすればいいのだろう。

>まず５１２個の要素の配列作って、添え字にスレッド番号（上のi）を指定して確認してみたら？
それは分岐条件がきちんと実行しているのか見るためのテストをしたとき確認しました。

199:デフォルトの名無しさん
11/12/07 07:39:11.25
根本的にプログラミングの基礎が抜けている悪寒。

200:デフォルトの名無しさん
11/12/07 09:07:22.99
>>198
atomicという排他的処理を知らないのなら並列化プログラムをやるのは早すぎる。
増してCUDAという特殊な並列化をやるのはもっと早すぎる。

201:デフォルトの名無しさん
11/12/07 12:02:55.64
>>199
どこまで戻ればいいんでしょうか...

>>200
...これ使って課題提出しなきゃならんので、早すぎると言われても後に引けないです...

202:デフォルトの名無しさん
11/12/07 13:13:17.16
>>198
>>201
ですけど、
>>200の"atomicという排他的処理を知らないのなら並列化プログラムをやるのは早すぎる。"
では並列化プログラムをやるにあたり、
どういったことを勉強して、どういった手順でやればいいんでしょうか？
そこがよくわからず、私はいきなりcudaに突っ込んだので四苦八苦してるわけですが...

203:デフォルトの名無しさん
11/12/07 14:10:39.88
>>202
同期とロック

204:デフォルトの名無しさん
11/12/07 15:51:27.47
青木本読んだんじゃないのん？

x+=1ってのは
1.xを読み出す
2.読み出した値に1を加算する
3.結果をxに格納する
って手順なんだけど以下略

205:デフォルトの名無しさん
11/12/07 16:45:28.63
>>203
同期とロックですか。勉強します。

>>204
青木本は読みましたけど、ものにしたっていう状態じゃないです...
>x+=1ってのは...
普段、なんとなく使ってるので... 勉強不足ですみません。

206:デフォルトの名無しさん
11/12/07 17:29:38.47
ここの一番下に載ってる資料の4-7ページにatomicの説明が少しある。
URLﾘﾝｸ(accc.riken.jp)

207:デフォルトの名無しさん
11/12/07 17:58:07.38
>>206
ありがとうございます。参照します。

208:デフォルトの名無しさん
11/12/07 23:52:25.72
この本でも読んでみると良い。
日本語でわかりやすい。

URLﾘﾝｸ(www.amazon.co.jp)

209:デフォルトの名無しさん
11/12/08 00:42:53.72
>>208
ありがとうございます。明日早速本屋行ってきます。

210:デフォルトの名無しさん
11/12/08 17:37:10.02
久しぶりに新しいCUDAの本が出たようだ。
URLﾘﾝｸ(www.amazon.co.jp)

211:デフォルトの名無しさん
11/12/08 18:59:25.20
>>210
グラフィックをメインにしてCUDAを道具としてこれから使おうとする人にはいいかも？
3章第3項がCUDA入門。第4項の「応用プログラム」でおもしろい話が読めたらいいね。

こっち（2011/11/14発売）も気になったけど、内容説明読んだだけじゃどの程度の本なのかわからなかった。
値段とタイトルからはちょっと期待させられる。
Amazon.co.jp： CUDA Application Design and Development: Rob Farber: 洋書
URLﾘﾝｸ(www.amazon.co.jp)

212:デフォルトの名無しさん
11/12/09 05:54:08.89
CUDAというかVisualStudioが分からない。鬱だ死のう。

213:デフォルトの名無しさん
11/12/09 09:40:45.47
VSがわからないんじゃなく、C++がわからないんだろ

214:デフォルトの名無しさん
11/12/11 23:39:15.74
fortranもここでいいかな？
pgiのコンパイラ買って、三重ループの前後に指示行追加してやったけどまったく速くならない。憂鬱。。

215:デフォルトの名無しさん
11/12/11 23:56:05.05
fortran使う人って量子論とかやってる人？
どの世界の人が使うのかしら？

216:デフォルトの名無しさん
11/12/12 08:31:48.58
fortranを使うのは過去の遺産のためだろう。
fortranを使えば速くなるわけでもないしなあ。

217:デフォルトの名無しさん
11/12/12 09:39:06.64
fortran懐かし過ぎる
先日、二度と使うことは無いと思って本を捨てたばっかりだわ

218:デフォルトの名無しさん
11/12/12 20:09:34.35
高レベル言語使うのは時代の趨勢じゃないかな。

LAPACKはfortranで書かれている。
fortranはnvccよりもアプリケーションに近い。
CUDAを使うのにnvccのレベルでなきゃダメということは無いと思う。
逆に、nvcc使わない人をバカにするヤツは、本物のバカだと思う。

219:デフォルトの名無しさん
11/12/12 20:20:14.55
道具にこだわることが目的となって、肝心の結果を出せない人間を馬鹿というでは？

220:デフォルトの名無しさん
11/12/12 21:18:03.01
TSUBAME2.0でnvccを使うのと京でfortran使うのと比べるとどうなんだろう。
京はSPARCだけで計算して、TSUBAME2.0の10倍くらいの速度みたいだ。

東工大は教育機関なのに対して、理化学研究所は結果を求められる独立行政法人。

221:デフォルトの名無しさん
11/12/12 22:25:59.20
既存のコードを使う限りでは京のほうが速いだろう。
ただ、富士通のコンパイラがダメダメだと聞いたけど、
それは解消されたのかな？

222:デフォルトの名無しさん
11/12/12 22:29:03.43
富士通のコンパイラは、8コアCPUへの割り当ては自動らしい。
これ、すごいことだと思うんだけど、
このコンパイラは完成しているのだろうか。

223:デフォルトの名無しさん
11/12/12 22:46:19.38
>>222
openMPに対応してたら各CPUのコアへの割り当ては普通にコンパイラ側でやると思うけど？

224:デフォルトの名無しさん
11/12/13 11:47:52.86
>>223
OpenMPがそんな事やるかよ。やるのはOSだぞ。
>>222が言っているのはNUMAのことだ。
現時点でも普通のLinuxならnumactlを使えばかなりのケースでノードの割り当てができるようになる。
CUDAを使う場合でも、複数GPUを使う場合に、MPIを使う場合に有効だ。

225:やんやん ◆yanyan72E.
11/12/13 14:16:05.38
なんだ？OpenMPやNUMAを何のことだと思ってるんだ？
というか釣られたか？

226:デフォルトの名無しさん
11/12/13 17:55:43.92
>>222
> これ、すごいことだと思うんだけど、
スパコンの世界では、ずうううっと前からやってるんじゃない？

227:デフォルトの名無しさん
11/12/13 20:56:29.90
じゃあ、そろそろ俺らも本気出すか

228:デフォルトの名無しさん
11/12/14 09:03:43.27
>>227
本気だしたら、負けだと思うんだ。

229:デフォルトの名無しさん
11/12/14 20:16:13.92
4Gamer.net ― NVIDIA，CUDA 4.1をリリース。CUDAコンパイラのソースコード公開も
URLﾘﾝｸ(www.4gamer.net)
中国時間2011年12月14日，NVIDIAは「CUDA 4.1」をリリースした。
最大の注目点は, LLVMベースとなるCUDAコンパイラが搭載されそのソースコードが公開された点だ。
LLVM（Low Level Virtual Machine）はAppleなどが参加する“言語非依存”のコンパイラ環境だ。
NVIDIAは同時に「CUDAプラットフォームのオープン化」を明言し，
LLVMベースとなるCUDAコンパイラのソースコードを研究者やツールベンダーに公開している。
要するに，NVIDIAのGPUでしか利用できなかったCUDA環境がほかのCPUやGPUに広がっていく可能性が出てきたのだ。
たとえばCPUや，それこそRadeonなどの他社製GPUでもCUDAを利用できる可能性が出てきたわけで，
CUDAの標準化をさらに推し進める起爆剤となり得るのである。

230:デフォルトの名無しさん
11/12/14 22:38:03.27
これって、個人でもソースもらえるのかな

231:デフォルトの名無しさん
11/12/14 23:11:41.54
へえ、ソースを公開したんだ。
でも思いっきりOpenCLとかぶるな。
OpenCLが死亡か。

232:デフォルトの名無しさん
11/12/15 08:04:39.43
なんで、OpneCLが死亡するんだ？

233:デフォルトの名無しさん
11/12/15 11:54:57.42
まだ来てなくない？

234:デフォルトの名無しさん
11/12/15 11:59:38.42
Radeon用CUDA作っても、特定の機種用になりそうだな

235:デフォルトの名無しさん
11/12/18 02:35:32.30
そもそもRADEON上でCUDAってマトモに動くのか？
一応HD3000世代のRV670からはFP64対応らしいが

236:デフォルトの名無しさん
11/12/19 02:14:51.68
動かないよ
動かそうとする人も現時点ではいないだろ
まだまだGPGPU向けには問題ありそうだし

237:デフォルトの名無しさん
11/12/20 01:48:20.74
cudaスレだが、
研究室や自分の周りでけで高速化すればいいんだったらcudaでいいんだろうけど
製品に組みこむならintelの普及度やなるべく多くの人に使ってもらえることを
考えるとOpenCLのほうがいいんじゃないかと思ってる。

photoshopのGPU対応もOpenCLでやってるみたいだし。

238:デフォルトの名無しさん
11/12/20 03:10:58.69
Photoshopは現状CUDAとATi Streamです
高速化しようとすると結局はデバイス毎に最適化したプログラムを書くことになるので、
デバイスに特化したライブラリの方が融通が利きます
Parallel NsightでのデバッグはOpenCLではできません

239:デフォルトの名無しさん
11/12/20 03:38:50.36
NVIDIA Opens Up CUDA, Compiler Source Code is Out - Bright Side Of News*
URLﾘﾝｸ(www.brightsideofnews.com)

240:デフォルトの名無しさん
11/12/20 18:56:37.44
CUDAのハードウェア種別を問わない展開が可能になったという事は、
CUDAの上に乗っかる形で実装されているPhysXやOptiX等もまた
ハードウェア種別を問わず動作可能になる可能性があるという事だよな
期待したい

241:デフォルトの名無しさん
11/12/20 22:23:45.55
別にソースコードなくたって、ある程度の知識を持った人なら
CUDAを移植できたと思うけど

242:デフォルトの名無しさん
11/12/20 22:49:46.93
RADEONでPhysXはマーケティングの理由で使えないんだろ？

243:デフォルトの名無しさん
11/12/20 23:19:03.65
AMDがCUDAを実装されるの嫌がってPhysX蹴ったんだっけ

244:デフォルトの名無しさん
11/12/21 19:57:27.52
質問：
GPU側に2次元配列を送るには、
１.2次元を1次元に直してからcudaMallocでメモリ確保→cudaMemcpyで送信
２.cudaMallocPitchでメモリ確保→cudaMemcpy2Dで送信
という2つの方法がありますけど、
構造体のメンバーに2次元配列があるものをGPU側に送る場合はどのようにしたらいいんでしょうか？

245:デフォルトの名無しさん
11/12/21 20:20:10.20
ポインタじゃなく配列ならそのままでうまくいかんか？

246:デフォルトの名無しさん
11/12/21 21:23:30.55
>>244
再度検証してたけど、サジ投げた。
最終的には以下の構造体をGPU側に送信して処理した後、CPU側に戻したい。
typedef struct test{
int a; int b[5];
}TEST;

typedef struct data{
int c; TEST d[1][10];
}DATA;
上をcudaMallocでメモリ確保＆cudaMemcpy送信しようとしてるけど、違うのかな。
教示お願いします。

247:デフォルトの名無しさん
11/12/21 21:36:40.66
そのまま転送すればいいじゃん

248:デフォルトの名無しさん
11/12/21 21:49:19.70
>>246です。
すいません。できました。お騒がせしました。

249:デフォルトの名無しさん
11/12/22 08:21:40.56
>>229
このコンパイラって、謎のCをPTXに変換する部分だけじゃないの？
今のフロントエンド(cudafe)は、他社製品のEDGベースだから、出てこないと思うが。
　cudafe -> PTXへのコンパイラ(今だとnvopencc) -> ptxas
のまんなかが出来るだけなので、この部分のソースがあったところで、CUDAコンパイラを自分で作れるとか、
CUDAを他のプロセッサで動かせるようになるとかいうものではない。

250:デフォルトの名無しさん
11/12/22 13:07:41.69
.cファイルがコンパイルできません
たすけて

251:デフォルトの名無しさん
11/12/22 13:13:00.92
んｖｃｃじゃできんだよ

252:デフォルトの名無しさん
11/12/26 15:17:58.63
あの、OpenCLってオワコンなの？IntelとAMDが支持してるのに全然盛り上がってないんだが。

253:デフォルトの名無しさん
11/12/26 21:25:42.32
科学技術計算向けがtesla一択だから他を使う理由もないってっところなのかね。
わざわざcudaで組んだものを他に移植する理由もないだろうし。

254:デフォルトの名無しさん
11/12/26 21:42:01.58
OpenCLで書くと長くなってしまう。
それだけのこと

255:デフォルトの名無しさん
11/12/27 02:04:51.35
C++ Bindings使えよ。

256:デフォルトの名無しさん
11/12/27 10:05:27.99
ちょっとしたラッパーかくのも嫌なのか

257:デフォルトの名無しさん
11/12/27 11:26:29.65
超素人の質問です。
コマンドラインでnvccを使いコンパイルを行っているのですが、OpenGLのサンプルコードをコンパイルできません。
freeglutを入れてます。
gcc -lglut program.cではコンパイル可能でした。
しかしnvcc -lglut program.cuとするとコンパイル不可となります。
-Lや-lによってライブラリやヘッダファイル先を指定したところで、コンパイルできませんでした。
阿呆な質問かもしれませんが、自己解決できそうにないのでよろしくおねがいします。

258:デフォルトの名無しさん
11/12/27 11:33:58.46
凄いアフォだな
エラーメッセージすべてコピペし

259:帝徒＝繪璃奈＝啓北商業の野島えり
11/12/27 14:53:52.03
主犯少頭劣一族＝蔗冽一族とは

中国華喃の山の梺の村八分の家。

二間位の横長で玄関の右側がお勝手。

大正に猿のままで生まれたのが

鈴木あゆみ(网(アミ) 范襤の子)

フィリピン人の范襤と同じくフィリピン人のモンゴルに逃げた『シバ』との間の男児。

日本名鈴木ひろしと聞いた。

その後、親戚の鈴木大樹を殺し戸籍を使用。

今は一文字雉。

260:デフォルトの名無しさん
11/12/29 21:08:02.00
母体になるプログラムなしで、
CUDAのコード単体で動かせるツールって無いかねぇ。
外出先でCUDAのサンプル見せるのがメンドイ。
例えば ./cuda-run < source.cuda みたいな感じで実験できるといい感じ。
codepadみたいな感じで動かせればなおよしだけど。

261:デフォルトの名無しさん
11/12/29 22:01:29.67
そういうスクリプトかけばよい

262:デフォルトの名無しさん
12/01/05 10:51:17.15
>>260
＞母体になるプログラムなしで、
それってなんのこと? 外出先のPCで動かしたいってこと?
コンパイラのことなら、事前にコンパイルしておくかリモートでアクセスすればいい。
ドライバのことなら、外出先の端末にインストールさせて貰うしかない。
外出先のPCなんか頼らず、自前のPCを持っていけばいいじゃん。

263:デフォルトの名無しさん
12/01/05 13:36:51.62
>>257
手元に開発環境がなく、エラーメッセージをコピー忘れしてました。すみません。

# cc Drawtest.c -lglut
# ./a.out

（空のウィンドウが製作される）
（Drawtest.cをDrawtest.cuに中身は同じで拡張子だけ変えて）

# nvcc Drawtest.cu -lglut
/usr/bin/ld: skipping incompatible /usr/local/cuda/bin/../lib64/libglut.a when searching for -lglut

（-lglutを見付かられなかった？）

# nvcc -L /usr/local/cuda/lib64/ -l /usr/local/cuda/include/ Drawtest.cu -lglut
/usr/bin/ld: cannot find -l/usr/local/cuda/include/
collect2: ld はステータス 1 で終了しました

となりました。

264:デフォルトの名無しさん
12/01/05 14:25:08.28
インクルードパスの指定は -l じゃなくて -I な

265:デフォルトの名無しさん
12/01/05 15:05:37.06
/usr/bin/ld: skipping incompatible～って64/32ビット混在時のエラーメッセージじゃなかったかな

266:デフォルトの名無しさん
12/01/05 15:09:35.86
>>264
なんというミス。
目で見て違いが分からない。
L（小文字）→i（大文字）に変えて同じようにしてみました。
# nvcc -L /usr/local/cuda/lib64/ -I /usr/local/cuda/include/ Drawtest.cu -lglut
/usr/bin/ld: skipping incompatible /usr/local/cuda/lib64//libglut.a when searching for -lglut
/usr/bin/ld: skipping incompatible /usr/local/cuda/bin/../lib64/libglut.a when searching for -lglut

結局パス指定なしと同じ見たいですね。

267:デフォルトの名無しさん
12/01/05 15:37:13.27
おれも64/32の違いかと思ったが、パス指定から指摘してみた
実行している環境は？ -m64 オプション付けてみ

268:デフォルトの名無しさん
12/01/06 13:53:42.06
# nvcc -L /usr/local/cuda/lib64/ -I /usr/local/cuda/include/ Drawtest.cu -lglut -m64
/usr/bin/ld: skipping incompatible /usr/local/cuda/lib64//libglut.a when searching for -lglut
/usr/bin/ld: skipping incompatible /usr/local/cuda/bin/../lib64/libglut.a when searching for -lglut
ううむ・・・・。

centOS 64bit
Device 0: "Quadro 4000"
Device 1: "GeForce 9500 GT"
パスは
LD_LIBRARY_PATH = /usr/local/cuda/lib64
です。
openGLのないコードはコンパイル及び実行まで問題ありません。

openGLを使うにあたって
/usr/local/cuda/include/にglut.h freeglut.h freeglut_ext.h freeglut_std.h
/usr/local/cuda/lib64/にlibglut.a
をコピペしましたが、まだ足りないのでしょうか？

269:デフォルトの名無しさん
12/01/06 13:59:40.26
libglutは64ビット？
32ビットなんじゃねか？

270:デフォルトの名無しさん
12/01/06 15:03:44.65
glut.h freeglut.h freeglut_ext.h freeglut_std.h libglut.aをccが見ているであろう先のものと置き換えてしてみました。

# nvcc -L /usr/local/cuda/include/ -I /usr/local/cuda/include/ Drawtest.cu -lglut -m64
/tmp/tmpxft_000030a4_00000000-12_Drawtest.o: In function `disp()':
tmpxft_000030a4_00000000-1_Drawtest.cudafe1.cpp:(.text+0x14e6): undefined reference to `glClearColor'
tmpxft_000030a4_00000000-1_Drawtest.cudafe1.cpp:(.text+0x14f0): undefined reference to `glClear'
/usr/local/cuda/bin/../lib64/libglut.a(libglut_la-freeglut_init.o): In function `glutInit':
…
(中略)
…
(.text+0x898): undefined reference to `glPopClientAttrib'
collect2: ld はステータス 1 で終了しました

openGL系の関数が読まれていないみたいです。

271:デフォルトの名無しさん
12/01/06 15:37:17.18
いや、だからfreeglutは64ビットを使用してるのかと聞いてるんだが
質問に対しきちんと回答できないようでは教えようがない罠

272:デフォルトの名無しさん
12/01/06 18:39:39.36
>>271
freeglutは64ビットのものを入れ直してやってみましたが、同じ結果になりました。

>>257
そこでCUDAのヘッダフォルダとライブラリフォルダに入れていたものを消し、-Lと-Iで直接freeglutのフォルダを指定したら
解決しました。
ライブラリが足りなかったのか、それともCUDAと一緒くたにしたのがダメだったのか解かりませんが、ちゃんとPATH通せってことみたいです。
お騒がせしました。

273:デフォルトの名無しさん
12/01/06 19:45:07.40
>>262
だよねぇ・・・

274:デフォルトの名無しさん
12/01/06 20:14:34.67
環境
Windows7 Professional 64bit
Microsoft Visual C++ 2010 Express Version 10.0.40219.1
Microsoft .NET Framwork Version 4.0.30319
GeForce GTX 580
CUDA Toolkit 4.0.17
SDK 4.0.19
devdriver_4.0_winvista-win7_64_270.81_general

このサイトを参考に環境を構築しました。
URLﾘﾝｸ(feather.cocolog-nifty.com)
そして以下のサイトのサンプルプログラムを実行してみました。
URLﾘﾝｸ(www.gdep.jp)
Hello,Worldと99 bottles of beerはcpu、gpu共に実行できました。
しかし、Matrixのプログラムはcpuの方は実行できるのですがgpuの方が実行できません。
以下のエラーを吐きます。
matrix_gpu.cu(5) : fatal error C1083: Cannot open include file: 'cutil_inline.h'
: No such file or directory
どうやらVisualStudioのパスがうまく通ってないということまでわかり、以下のサイトなどを参考にCUDA_INC_PATHなどを変えてみましたが、一向に変わりません。
URLﾘﾝｸ(d.hatena.ne.jp)
SDK内のcutil_inline.h自体をtoolkitのincフォルダにコピペすると、他の.hファイルもいくつか同じエラーが出たのでエラーになったものをすべてコピペしたところ、LNK2019"link.exe"というエラーで先に進めませんでした。
一度VisualStudioを再インストールしてみましたが、状況は変わりません。
Nvidia GPU computing SDK Browserではサンプルプログラムを実行できているので、CUDAの環境は整っていると思われます。
どうすれば解決できますでしょうか・・・。かれこれ１週間以上格闘しています。
VisualStudioは2010よりも2008にした方がいいでしょうか？

275:デフォルトの名無しさん
12/01/06 20:23:04.24
C初心者にはきついと思うんだが…
とりあえず 'cutil_inline.h'のある場所を見つけて
そこを-I　/'cutil_inline.hのある場所'と指定する

意味わからなければCのコンパイルを勉強すること

276:デフォルトの名無しさん
12/01/06 20:39:53.25
274ではありませんが､「Cのコンパイル」の勉強にオススメの書籍とかありましたら紹介していただけるとうれしいです

277:デフォルトの名無しさん
12/01/06 20:55:31.81
275だけど、お薦めは他の人に任せるけど
一つ言えるのはCUDAはC/C++より数段難しい
C/C++の質問をしてるようでは先が思いやられる、と正直思うんだが…

超低レベルな質問もOKなくだすれだけど、その程度は自力で調べる能力が必要ではないだろうか？

278:デフォルトの名無しさん
12/01/06 21:22:02.03
>>275
>>274です。ありがとうございます。
今実行できる環境にないので後で試してみます。
また結果報告しにきます。

279:デフォルトの名無しさん
12/01/07 12:14:43.10
一時期は同価格帯のCPUとGPUを比較してGPUは５０倍の性能があるって言われてたけど
今１０～１５倍程度しか無くない？

3930k
５万円
158GFlops

GTX590
６．３万円
2480GFlops

GTX 580
３．８万円
1500GFlops

次に出ると言われてるGTX700系は500系の２倍の性能らしいけどそれでも５０倍には遠く及ばない。
この「同価格帯でGPUはCPUの１０～１５倍程度の性能しかない」という評価は正しいだろうか？

280:デフォルトの名無しさん
12/01/07 12:43:15.62
>>279
それは単にCPUは1コアとして比較しているからGPUを使えばうん十倍と言われるんだよ。
単純にメモリ帯域で3～4倍、演算器としては、単純にAVXの8並列x6コアxクロック2倍で比較すると5倍にしかならない。
CPUは大きなキャッシュがあるから局所性の高い処理の場合はメモリ帯域の差が縮まるので、10倍差がつくならいい方。
もちろんアプリによるけど、CPU向けのコードがカリカリに最適化されていた場合は思ったほど差がつかない印象。
それでも3～5倍位は早くなるからいいんだけどね。
問題はメンテナンスだな。。。CUDAのみと割り切れればいいんだけど、なかなかそうは行かないからなあ。

281:デフォルトの名無しさん
12/01/07 16:09:01.15
>>279
GTX580で500GFLOPS程度。そのデータはおかしい
CPUなんてE8500で25GFLOPSくらいだから20倍程度だな。
もちろん最近のコア多いやつならもっと差は縮まるだろうな。

それに理論値の話だから多数のスレッドでまわしたらコンテキストスイッチの
オーバーヘッドとかGPUプログラミングは難しくて効率上げるの難しいとか
そういうので実効は変わってくるんじゃね？

282:デフォルトの名無しさん
12/01/07 16:14:51.24
あ、同一価格帯か、E8500が2万しなかったことを考えると
580にそろえて2倍したとして10倍程度だから、まあおおざっぱには
「同価格帯でGPUはCPUの１０～１５倍程度の性能しかない」は正しいんじゃないかと思う。

580と同時期の4万弱のCPUならもっと差は縮まるだろうし。

次ページ