12/01/24 19:53:55.48
コンパイル時に以下のエラーメッセージが出ているのですが、
ptxas error : Entry function '...' uses too much local data (0x16054 bytes, 0x4000 max)
これはローカルメモリーと何か関係あるような気がするのですけど、
ちょっとわからないので教えて頂けないでしょうか。
宜しくお願いします。
356:デフォルトの名無しさん
12/01/24 21:03:52.65
それともう1つ。
externを使って、グローバル変数を共有させたいのですが、
ホスト・デバイスともに共有させたい場合はどうしたらいいのでしょうか?
357:デフォルトの名無しさん
12/01/24 21:56:32.01
> ptxas error : Entry function '...' uses too much local data (0x16054 bytes, 0x4000 max)
単に配列の次元が大きすぎるのでは？
最大 0x4000バイトのところを0x16054 bytes使おうとしている？
> ホスト・デバイスともに共有させたい場合はどうしたらいいのでしょうか?
物理的に無理でないの？
ホスト←→デバイス間でコピーしないといけないのでは？？
358:デフォルトの名無しさん
12/01/24 21:56:54.10
>>355
メッセージそのままの意味だろ。
359:デフォルトの名無しさん
12/01/25 01:32:34.89
CUDAの思想はいいと思うけど、ターゲットデバイスを完全には
隠蔽しきれないのが残念だな。
360:やんやん ◆yanyan72E.
12/01/25 01:41:55.08
スピード勝負の世界で、しかもかなり複雑なアーキテクチャなのに、
ハードウェアを完全に隠蔽する訳にいかないだろ。
361:357
12/01/25 12:34:12.36
> ホスト・デバイスともに共有させたい場合はどうしたらいいのでしょうか?
ヘッダーファイルに
#DEFINE
などで定義して、ホストとデバイスのプログラムに そのヘッダーファイルをinclude する手があった♪
362:355,356
12/01/25 15:21:00.98
>>358
青木氏の「はじめてのCUDA」に似たようなエラーメッセージが載ってるのですが、
local dataがどの部分の話なのか、ちょっとわからない状態でして...
>>357
>単に配列の次元が大きすぎるのでは？
予定だと1×256スレッドしか使うので、足りないはずないような... 意味が違うか。
ちょっとわからないので詳細お願いします。
>ヘッダーファイルに
>#DEFINE
>などで定義して、ホストとデバイスのプログラムに そのヘッダーファイルをinclude する手があった♪
ちょっと試してみます。それができなかったら、コピーするしかないかもしれません。
__device__ extern TEST test;
のような宣言してて、実行ファイルができたときがあったのですが、何か違うよな、と引っかかってたので、
ここで質問しました。ありがとうございます。
363:デフォルトの名無しさん
12/01/25 17:40:16.56
差し支えなければプログラムをアップしてくださいませ♪
364:デフォルトの名無しさん
12/01/25 18:01:06.82
>>362
このサイトの一番下にある資料の3-41ページにそのエラーメッセージが載っている模様。
URLﾘﾝｸ(accc.riken.jp)
365:355,356
12/01/25 18:04:09.31
>>362です。
>>363
申し訳ないのですが、プログラムのアップは出来ないです。すみません。
なんかもう分からなくなったので、ホストとデバイスの住み分けを行ったところ、
「ptxas error...」云々のメッセージが消え、コンパイルできました。
（さっきまでhost,device,global宣言関数がごちゃ混ぜな状態だった。）
何が原因だったのか分からず仕舞いです。
グローバル変数の共有はとりあえず、コンスタントメモリにデータコピーで様子見することにしました。
元コード(.c)をもう一度読み直したところ、デバイス側の方はReadOnlyで十分なようでしたから。
皆様回答ありがとうございました。またよろしくお願いします。
366:365
12/01/25 21:05:50.75
>>364
資料ありがとうございます。
うまくいったと思った途端、また同じエラーが出てきてしまったので、確認します。
.cファイルに.cuファイルを組み込ませるようにしたら、
__host__ __device__で修飾した関数が定義されてないと.cファイル側に言われ、
.cファイルと.cuファイルそれぞれ単独で動かせるようにしたら、
（同じ内容の関数を.cと.cuそれぞれ別名で実体作った。.c:test .cu:ktestみたいな）
今度はさっきと同じエラー。
実行ファイル作るだけなのに難しい...
367:デフォルトの名無しさん
12/01/26 08:42:18.26
Host側とDevice側で型を別にすれば開発もっと楽になると思うんだけど
368:デフォルトの名無しさん
12/01/26 08:53:13.45
取り敢えずCudaのサンプルは捨てて、インクルード関係とオブジェクトの生成手順を確認するんだ。
369:デフォルトの名無しさん
12/01/26 10:14:19.95
見せられないところは削除して、問題の部分だけ残して、アップすれば？？
370:366
12/01/26 16:19:08.01
>>368
>取り敢えずCudaのサンプルは捨てて、インクルード関係とオブジェクトの生成手順を確認するんだ。
のように手順を踏んでコンパイルしたところ、実行ファイルができました。
皆様ありがとうございました。
371:370
12/01/26 20:33:24.55
お恥ずかしながら、また戻ってきました。
あれからセグメンテーション違反が出てきてしまったので、あれこれ探していた結果、
どうやらデバイスのメモリ確保＆送信に失敗していたようです。
しかし解せないことがあって、
構造体A a,構造体B b[100],構造体C c[100],...(以下略 をデバイス側に送るのですが、
（1つ１つのサイズは結構大きめ。グローバルは1GBあって余裕で確保出きるハズ...）
void main(){
・・・
test(&a,b,c....);
}
void test(A *a,B *b,C *c...){
A *d_a; B *d_b; C *d_c;
CUT_SAFE_CALL(cudaMalloc((void**)&d_a,sizeof(A)));
CUT_SAFE_CALL(cudaMemcpy(d_a,a,sizeof(A),cudaMemcpyHtoD));
CUT_SAFE_CALL(cudaMalloc((void**)&d_b,sizeof(B)*100));
CUT_SAFE_CALL(cudaMemcpy(d_a,b,sizeof(B)*100,cudaMemcpyHtoD));
・・・
｝
はじめてcuda,cuda_by_exampleで確認したところ、
文法的ミスはないはずなのに確保ミスしてるらしく、中断してます。
この原因は一体全体なんなんでしょうか。
372:デフォルトの名無しさん
12/01/26 21:28:56.32
4.1正式版
CUDA Toolkit 4.1 | NVIDIA Developer Zone
URLﾘﾝｸ(developer.nvidia.com)
373:デフォルトの名無しさん
12/01/26 21:55:28.42
構造体の中身がよく分からんが，allocateのところで *100は必要なのかな？
構造体Bのメモリーを１００個分用意しようとしている？？
構造体Bの中にすでに [100]個の配列を取っているのに？？？
数は合っておるのか？？？？
374:デフォルトの名無しさん
12/01/26 22:17:44.08
ｒｈｅｌ6.0のがリンクミスかでダウンロドできん
375:デフォルトの名無しさん
12/01/27 00:03:32.49
>>372
遂にRC取れたか！4.0からどう変わってるか知らんけど。
Kepler向けのプログラムを作れるのは5.0とかになるんだろうか。
そしてそれはいつ出るんだろう。
>>371
1GBのVRAM積んでるカードでも400MBぐらいの確保で失敗したことがあったような覚えがある。
確保するサイズを小さくしてもエラーが出るってんならこの話は忘れてくれ。
376:デフォルトの名無しさん
12/01/27 01:34:46.65
100x100の行列の積を1万回くらい実行したいのですが、
CUBLASだと小さすぎてCPUより遅くなってしまいます。
この1万回は相互に依存関係は無いので、並列化が可能なのですが、
このプログラムは自分で書くしかないのでしょうか。
TESLA C1070があるので、丸ごとVRAMに載ります。
377:デフォルトの名無しさん
12/01/27 03:17:55.02
4.1にcublas{S,D,C,Z}gemmBatchedってのが追加されてました。
especially for smaller matricesに効くってあったので、使ってみます。
378:デフォルトの名無しさん
12/01/27 07:52:34.05
一つのブロック内で100x100の行列の積を１回行う？
379:デフォルトの名無しさん
12/01/27 07:52:51.42
間違った
380:デフォルトの名無しさん
12/01/27 07:53:00.44
あれ？
381:デフォルトの名無しさん
12/01/27 07:53:07.78
一つのブロック内で100x100の行列の積を１回行う？
382:デフォルトの名無しさん
12/01/27 07:55:11.45
たびたび，失礼
一つのブロック内で100x100の行列の積を１回行う？
そのブロックを１万個用意する？
それらを並列に計算する？
と言うコンセプトかな？？？
自分でプログラムを作った方が，よほど勉強になると思う
383:371
12/01/27 18:15:48.81
>>373
CUT_SAFE_CALL(cudaMalloc((void**)&d_b,sizeof(B)*100));
CUT_SAFE_CALL(cudaMemcpy(d_a,b,sizeof(B)*100,cudaMemcpyHtoD));
これはミスですね....
正しくは
CUT_SAFE_CALL(cudaMalloc((void**)&d_b,sizeof(B)));
CUT_SAFE_CALL(cudaMemcpy(d_a,b,sizeof(B),cudaMemcpyHtoD));
です。すみません。
>>375
4000MBの領域確保でも失敗することがあったんですか...
ちょっと試してみます。
384:デフォルトの名無しさん
12/01/27 20:31:39.53
d_aに送信してるあたりが
385:デフォルトの名無しさん
12/01/27 21:52:59.41
d_b のサイズは構造体Bと同じ？？
d_aのサイズはbと同じ？すなわち構造体Bと同じ？？
違っていたらエラーが出て当たり前では？？？
386:デフォルトの名無しさん
12/01/28 13:33:36.75
それから，ホスト側で も メモリーは確保しているのか？？？？
387:383
12/01/28 17:54:35.37
>>371
なんかサンプルがめちゃくちゃなんで書き直します。
388:387
12/01/28 20:55:21.49
確認なのですけど、カーネル関数の引数はポインタ限定ですか?
389:デフォルトの名無しさん
12/01/28 22:58:27.22
配列はポインタで，１変数はそのままではなかったかい？
なぜか知らんけど
390:デフォルトの名無しさん
12/01/28 23:07:13.34
夜の埼玉は
391:デフォルトの名無しさん
12/01/29 01:49:05.91
>388
ホスト側のポインタ以外なら何でもOKのはず。少なくともintは渡せる。
クラスを渡せるかどうかは知らないけど。
392:デフォルトの名無しさん
12/01/29 06:58:36.12
ホストのポインタ送って、GPU側のload/store時に
メインメモリから自動転送ってこともできるから、
その説明は正しくない。
393:388
12/01/30 15:13:56.68
>>391
セグメンテーション違反してた部分の引数に（デバイスで使用する）double（実体）を指定してたのですけど、
（デバイスで使用する）ポインタに変えたら、問題なく実行できました。
dtempはcudaMalloc/cudaMemcpy使ってます。
そもそもな話cudaはポインタ宣言したdtempの実体は*dtempではないのかな。
394:デフォルトの名無しさん
12/01/30 15:16:54.53
>>393
何を言っているのか判らない。問題が再現する最低限のコードでも貼ってみていただきたい。
395:デフォルトの名無しさん
12/01/30 16:19:41.78
>>393
__global__ void kernel(double dtemp){
}
int main(){
double temp=0.0;
double *dtemp;
cudaMalloc((void**)&dtemp,sizeof(double));
cudaMemcpy(dtemp,&temp,sizeof(double),cudaMemcpyHostToDevice);
kernel<<<1,1>>>(*dtemp);//(←※1)
cudaMemcpy(&temp,dtemp,sizeof(double),cudaMemcpyDeviceToHost);
return 0;
}
※1の部分でsegmentation fault。
__global__ kernel(double *dtemp){
}として
kernel<<<1,1>>>(dtemp);
とすると問題無し。
dtempのアドレスを引数にするのはいいけど、dtempの中身（0.0）を引数にしようとするのはダメ？
そもそも※1の宣言自体間違ってる?どうなんだろう。
396:デフォルトの名無しさん
12/01/30 16:29:17.80
393じゃないけど試してみた
__global__ void kernel(double dtemp){
}
int main(){
double temp=0.0;
double dtemp;
// cudaMalloc((void**)&dtemp,sizeof(double));
// cudaMemcpy(dtemp,&temp,sizeof(double),cudaMemcpyHostToDevice);
kernel<<<1,1>>>(dtemp);//(1)
// cudaMemcpy(&temp,dtemp,sizeof(double),cudaMemcpyDeviceToHost);
return 0;
}
これでコンパイルも実行もおkだったけど
397:デフォルトの名無しさん
12/01/30 16:57:38.16
>>395
dtempはdevicePointerなんだから、Hostコード中でデリファレンスしちゃダメ。
double値を渡したいだけなら引き数はdoubleで充分。
それと、※1は宣言じゃなくて呼び出しだから構文としてはあっている。
そもそも何がしたいのだろう。値を渡したいなら値を渡せばいいし、
ポインタを渡したいならポインタをそのまま渡せばいい。
devicePointerをHostでデリファレンスしちゃダメだし、
(その後どうせdevice側でデリファレンスできないから)HostPointerをdeviceにコピーしてもダメ。
398:デフォルトの名無しさん
12/01/30 17:04:05.78
>>396
>>397
色々ごちゃまぜだったようです。スッキリしました。
ありがとうございました。
399:デフォルトの名無しさん
12/01/30 20:58:23.57
これからCUDAの勉強を始めようと思っています。
GPGPUプログラミングはシェアードメモリを有効活用できるかがポイントであるように思います。
そこで、あえて最新のグラボに買い換えず、キャッシュのないTesla世代のGTX285を所持し続けています。
キャッシュがあると、シェアードメモリをうまく使えてなくても、それをカバーしてしまって、
うまくプログラムできているかどうかを把握し辛いと考えたからです。
ただ、心配なのが、Fermi世代以降に発売された書籍だと、内容がTesla世代のGPUにそぐわない
といった状況が起きてしまうことです。
GTX285のまま勉強するに際して気をつけること等あれば教えてください。
問題がなければ関連書籍（日本語）を全部購入して読もうと思います。
よろしくお願いします。
400:デフォルトの名無しさん
12/01/30 21:08:43.55
＞＞GPGPUプログラミングはシェアードメモリを有効活用できるかがポイントであるように思います。
そんなの計算する内容によるわ
活用しようがないのもあるからな
計算（使用）目的は明確なの？
401:399
12/01/30 21:39:43.45
>>400
信号処理に使う予定です。
FIRフィルタリングやFFTなどを大きな1次元、あるいは2次元配列に適用したいと思っています。
402:デフォルトの名無しさん
12/01/30 21:45:58.41
俺は400じゃないが、気をつけることは、勉強がひと段落していざ使うときになって知識が陳腐化してても泣かないことかな。
ハードウェアはどんどん変わるし、TeslaやFermiで良しとされたことが数年後に通用するかはわからないからね。
せっかく今から始めるならFermi・Keplerを使ったほうがいいと俺は思うけど、
10年後にはどれも等しく陳腐化してるだろうから長期的に考えると確かにTeslaでもいいのかもしれない。
ただ長期的に考えること自体がGPGPUでは…と1行目にループする。
403:デフォルトの名無しさん
12/01/30 22:12:50.10
古い方がいろいろ制限があって，勉強になるとは思う．
404:デフォルトの名無しさん
12/01/30 22:17:32.23
ｆｆｔはcpu側からはインターフェースあるけどディバイス内では自前で作るしかない
405:デフォルトの名無しさん
12/01/30 22:18:26.89
> キャッシュがあると、シェアードメモリをうまく使えてなくても、それをカバーしてしまって、
> うまくプログラムできているかどうかを把握し辛いと考えたからです。
シェアードメモリーを使うようにプログラミングすればよいかと
どのメモリーを使っているのか分からない，と言うことはないと思う
あれば，そこをしっかり勉強すべきかと
406:デフォルトの名無しさん
12/01/30 22:19:12.36
そう言えば古いのは倍精度の計算ができないのでは？？
407:デフォルトの名無しさん
12/01/30 22:22:30.66
まあ，デバイスメモリーだけで計算してみるとか，
シェアードメモリーも使って計算してみるとか，
いろんなプログラムを作って比較するのが良いかと
408:デフォルトの名無しさん
12/01/30 22:26:01.98
最初はシェアドするスレッドはどれだってだけでも頭が痛くなった
409:402
12/01/30 23:28:19.90
あとTeslaでやるなら、FermiとTeslaでバンクコンフリクトの発生条件が違ってるから
Teslaでのこれの回避には深入りしないほうがいいかと。
他にシェアードメモリの大きさなどの量的な違いがあるけど、そういった量的な
制約のために質的なというかアルゴリズムの種類を変えるようなことまで深入りするのも
避けたほうが。だってFermiやそれ以降で量が変わると一から見直しになるから。
これはFermi使ってる人もいずれ同じことなんだけど、勉強じゃなくて性能が欲しいんだからしょうがない。
410:デフォルトの名無しさん
12/01/30 23:57:48.21
sharedがregister並に速いなんて嘘だから。
occupancyなんて上げるな。
warpあたりのregisterを増やして
ILPを利用してレイテンシを隠蔽すべきだと。
URLﾘﾝｸ(www.cs.berkeley.edu)
411:399
12/01/31 00:31:29.61
>>403
やはり、そうですか！
>>404
参考になります。
FFTは重要になるのでしっかりと取り組みたいです。
>>405
>>407
>>408
なるほど、どのメモリかをよく意識してプログラミングします。
CPUの最適化に取り組んでいて思ったのですが、
キャッシュは便利な反面、こちらから意図的に挙動を制御できないことが
パフォーマンスの考察を難しくしてしまう側面もあると感じました。
キャッシュなしのTeslaで鍛えたいです。
>>406
単精度しか使わないので問題なしです。
もし半精度でパフォーマンスが上がるなら、そのほうがイイくらいです。
>>402
>>409
世代を跨ぐにあたって非常に参考になるご助言をありがとうございます。
各種メモリの特性や帯域を意識して取り組むことで、
固有のデバイスに限定されない定性的なところを理解したいと思います。
Keplerではこんな組み方がイイかもしれないな、と思索できたら楽しそうです。
>>410
興味深い資料をありがとうございます。
各部のレイテンシ、スループットを頭に叩き込みます。
412:デフォルトの名無しさん
12/01/31 08:31:17.17
>>411
最終的に動かす機器で最もよい性能が出るように最適化するべきだし
自動でやってくれることを手でやるってのは勉強にはなるかもしれないけど意味ないのではって思うけど。
コンパイラの最適化を切ってアセンブリ書くようなものでしょ。
413:デフォルトの名無しさん
12/01/31 12:08:28.56
時間がいくらでもあるorそういう研究ならともかく
短期間で組めてそこそこのパフォーマンスが出せればいいなら
キャッシュのあるFermiで適当に書いてもいいと思うんだけどね
414:399
12/01/31 21:32:30.15
>>412
おっしゃる通りです
そのあたりは目的をよく考えて判断したいです。
今回はどちらかというと、GPUのポテンシャルを見極めたいというのが主です。
コードの可読性を優先して、性能はコンパイラやキャッシュ等にできるだけ委ねるというスタイルも
今後検討していきたいと思っています。
>>413
今回は研究用途ですが、
将来的にはFermi以降のGPUで、ちゃっちゃと書いたコードが快適に動くというところを目指したいです。
ただ、あまりにも下手なプログラムにならないよう、最初に少し深くやっておきたいです。
415:デフォルトの名無しさん
12/02/02 04:23:36.19
レジスタといえば、GK104では強化されるのか、それともさらにピーキーになるのか気になる。
GF104はグラフィック特化で、レジスタ数やワープ数の上限に悩まされたし。
とりあえずアーキテクチャが一般公開されるまで待つしかないかな。
416:デフォルトの名無しさん
12/02/04 00:16:58.91
また出てるな、これで最後らしいが
URLﾘﾝｸ(page4.auctions.yahoo.co.jp)
417:デフォルトの名無しさん
12/02/08 09:16:57.43
誰かいるかな…
URLﾘﾝｸ(includehoge.blog.fc2.com)
このブログのソースをコンパイルして動かしてみようと思ったのですが
1>cudamatrix.obj : error LNK2019: 未解決の外部シンボル _cufftDestroy@4 が関数 _main で参照されました。
1>cudamatrix.obj : error LNK2019: 未解決の外部シンボル _cufftExecC2C@16 が関数 _main で参照されました。
1>cudamatrix.obj : error LNK2019: 未解決の外部シンボル _cufftPlan1d@16 が関数 _main で参照されました。
1>../../bin/win32/Debug/template.exe : fatal error LNK1120: 外部参照 3 が未解決です。
というエラーが出ます。
cufft.h自体は読み込んでいて、引数の過少があるときはそれに対してエラーを返すので
定義されていないという扱いになっているわけではないみたいなんですけど
解決法があったら教えて下さい
418:デフォルトの名無しさん
12/02/08 09:17:57.58
誰かいるかな…
URLﾘﾝｸ(includehoge.blog.fc2.com)
このブログのソースをコンパイルして動かしてみようと思ったのですが
1>cudamatrix.obj : error LNK2019: 未解決の外部シンボル _cufftDestroy@4 が関数 _main で参照されました。
1>cudamatrix.obj : error LNK2019: 未解決の外部シンボル _cufftExecC2C@16 が関数 _main で参照されました。
1>cudamatrix.obj : error LNK2019: 未解決の外部シンボル _cufftPlan1d@16 が関数 _main で参照されました。
1>../../bin/win32/Debug/template.exe : fatal error LNK1120: 外部参照 3 が未解決です。
というエラーが出ます。
cufft.h自体は読み込んでいて、引数の過少があるときはそれに対してエラーを返すので
定義されていないという扱いになっているわけではないみたいなんですけど
解決法があったら教えて下さい
419:デフォルトの名無しさん
12/02/08 09:22:20.66
誰か居るかな・・・いたら助けてください
URLﾘﾝｸ(includehoge.blog.fc2.com)
このブログのソースをコンパイルして動かしてみようと思ったのですが
1>cudamatrix.obj : error LNK2019: 未解決の外部シンボル _cufftDestroy@4 が関数 _main で参照されました。
1>cudamatrix.obj : error LNK2019: 未解決の外部シンボル _cufftExecC2C@16 が関数 _main で参照されました。
1>cudamatrix.obj : error LNK2019: 未解決の外部シンボル _cufftPlan1d@16 が関数 _main で参照されました。
1>../../bin/win32/Debug/template.exe : fatal error LNK1120: 外部参照 3 が未解決です。
以上のようなエラーが出てコンパイルが通りません
構造体の生成や関数の過少には反応するので、cufft.hが読み込めてないってことはないと思うんですが…
解決法があったら教えて下さい
420:デフォルトの名無しさん
12/02/08 09:22:49.74
誰か居るかな・・・いたら助けてください
URLﾘﾝｸ(includehoge.blog.fc2.com)
このブログのソースをコンパイルして動かしてみようと思ったのですが
1>cudamatrix.obj : error LNK2019: 未解決の外部シンボル _cufftDestroy@4 が関数 _main で参照されました。
1>cudamatrix.obj : error LNK2019: 未解決の外部シンボル _cufftExecC2C@16 が関数 _main で参照されました。
1>cudamatrix.obj : error LNK2019: 未解決の外部シンボル _cufftPlan1d@16 が関数 _main で参照されました。
1>../../bin/win32/Debug/template.exe : fatal error LNK1120: 外部参照 3 が未解決です。
以上のようなエラーが出てコンパイルが通りません
構造体の生成や関数の過少には反応するので、cufft.hが読み込めてないってことはないと思うんですが…
解決法があったら教えて下さい
421:デフォルトの名無しさん
12/02/08 09:23:33.66
うわなんか四回も書き込んじゃったごめん
422:デフォルトの名無しさん
12/02/08 09:25:46.81
コンパイル詳しくないけどlibcufft.so（linuxの場合）のリンクしてる？
windowsならlibcufft.dllだとおもう
423:デフォルトの名無しさん
12/02/08 09:27:16.52
コンパイルじゃなくてリンクの問題でしょ
424:デフォルトの名無しさん
12/02/08 09:58:44.53
nvcc -lcufft sample.cu
425:デフォルトの名無しさん
12/02/08 10:20:43.62
このスレでときどき思うのはC/C++をほとんど知らないのにもかかわらず
CUDAに挑戦しようとする勇壮な挑戦者が多いな、ということ
彼らはその後うまくいってるのだろうか？
426:デフォルトの名無しさん
12/02/08 13:41:41.97
CUDA.NETも3.0で止まってるしなぁ・・・
427:デフォルトの名無しさん
12/02/08 21:44:42.85
コンピュート・ケイパビリティによって何が変わったかを俯瞰できるような一覧を掲載したサイトってご存知ありませんか？？
428:デフォルトの名無しさん
12/02/08 22:16:51.83
>>417
追加のライブラリファイルにナントカ.libを追加すればいけそう。
ナントカは>>422さんのレスから察するにlibcufft.libかなぁ。
このへんは>>425さんの言うとおりcuda云々じゃなくてVisual Studio(C++)を使う上で頻繁に設定するところだねー。
>>427
俺も知りたいなぁそれ。その手の情報はSDKpdfのProgramming GuideとBest Practiceくらいしか知らない（片方だけかも）
429:デフォルトの名無しさん
12/02/08 22:26:54.49
cufft.libなのだが。
430:デフォルトの名無しさん
12/02/09 07:17:56.91
てすてす
431:デフォルトの名無しさん
12/02/09 07:18:12.24
コマンドなら
432:デフォルトの名無しさん
12/02/09 07:18:18.71
あれ？
433:デフォルトの名無しさん
12/02/09 07:18:34.70
以下のよう
434:デフォルトの名無しさん
12/02/09 07:18:39.54
nvcc fft.cu --compiler-bindir="C:\Program Files\Microsoft Visual Studio 10.0\VC\bin" -arch=sm_10 -L"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.1\lib\Win32" -lcufft
435:デフォルトの名無しさん
12/02/09 07:19:01.92
ファイル名をfft.cuとしてみた
436:デフォルトの名無しさん
12/02/09 07:20:23.21
結果の一部
1014 6.347414 0.587785
1015 6.318954 -0.000000
1016 6.293659 -0.587785
1017 6.271449 -0.951057
1018 6.252275 -0.951057
1019 6.236131 -0.587785
1020 6.222946 0.000000
1021 6.212736 0.587785
1022 6.205431 0.951057
1023 6.201046 0.951057
437:デフォルトの名無しさん
12/02/09 07:21:26.93
-arch=sm_10 は適宜，変更のこと
438:デフォルトの名無しさん
12/02/09 07:25:00.70
これでも行けた♪
nvcc fft.cu -lcufft
439:デフォルトの名無しさん
12/02/09 07:26:45.21
既出でした．失礼 >> 424
440:デフォルトの名無しさん
12/02/09 11:41:33.64
GPUの勉強を始めたばかりなのですが、スパース行列の格納形式種類で
ELLの格納方法とメリットが分かりません
調べたらlisのマニュアルに配列方法はあるのですが、説明と図とが噛み合ない点があり
理解できずにいます。
だれか教えてください
441:デフォルトの名無しさん
12/02/09 13:59:04.44
>>426
CUDA#とかCUDAfy.NETとか代替のライブラリはあるっぽいよ
使ったことないから詳細は知らないけど
442:デフォルトの名無しさん
12/02/09 14:23:54.18
>>440
LisはGPU上で動くの？
443:デフォルトの名無しさん
12/02/09 14:39:02.94
sparse matrixの格納形式はCUDAの問題じゃね～だろ
444:デフォルトの名無しさん
12/02/09 15:45:25.65
>>440
コアレッシング状態で転送できるようになる。
445:デフォルトの名無しさん
12/02/09 17:19:58.52
CUDA.NETは使えなくはないぞ♪
446:デフォルトの名無しさん
12/02/09 18:44:06.72
>>440
スパース行列の格納方法がのっているのがLisのオンラインマニュアルくらいだったので
>>443
ELLはGPUに適した格納方法と、ネットのどこかにあったので関連があるかと思い書き込みました
たしかにCUDA以前の問題ですよね
>>444
なるほど、見えてきました！ありがとうございます
447:デフォルトの名無しさん
12/02/09 19:02:29.63
>>422->>439
こんなにたくさん答えてくれるとは・・・
おかげで動きましたーありがとう
448:デフォルトの名無しさん
12/02/09 19:08:48.94
おお、よかった^^
449:デフォルトの名無しさん
12/02/10 10:39:58.03
URLﾘﾝｸ(detail.chiebukuro.yahoo.co.jp)
450:デフォルトの名無しさん
12/02/10 13:35:49.60
>>441
CUDAfy.NETは知らんかった
最近出てきたのね、ありがとう
451:デフォルトの名無しさん
12/02/10 21:23:36.65
2次元配列にアクセスするときのインデックス計算で
gy = blockDim.y * blockIdx.y + threadIdx.y;
gx = blockDim.x * blockIdx.x + threadIdx.x;
というのを頻繁にやると思うのですが、
この積和演算ってCUDAコアが使われるのでしょうか？
それとも、CUDAコアとは別にアドレス演算器（？）的なハードウェアがあって、
それで計算してくれるのでしょうか？
後者だとうれしいです。
452:デフォルトの名無しさん
12/02/10 21:37:17.01
アドレスもプログラムもデータ
453:デフォルトの名無しさん
12/02/10 23:52:12.69
>>451
CUDAコア上に展開された（定義された）ブロックの中で演算をしていると思う
454:451
12/02/11 01:18:08.05
>>453
CUDAコアのサイクルを消費して計算しているということですね？
ありがとうございました。
455:デフォルトの名無しさん
12/02/11 01:24:48.08
インデックス計算かどうかをGPUがどう見分けるというんだ
456:451
12/02/11 10:05:14.94
>>455
単にCUDAの予約語であるblockDimやthreadIdだけで構成されていて、
○○ * ○○ + ○○ の積和の形になっていれば、
専用の回路に計算させる、とかできそうなものですが。
CPUではメモリアドレス式のカタチを解析してアドレス演算器に割り当てています。
457:デフォルトの名無しさん
12/02/11 11:29:23.45
そもそもグリッド，ブロック，スレッドの概念を理解しているのかな？
458:デフォルトの名無しさん
12/02/11 11:29:54.96
○○ * ○○ + ○○ だけではプログラムは作れないと思うのだが
459:デフォルトの名無しさん
12/02/11 11:55:45.33
>>456
それ、どのCPUでの話ですか？
460:デフォルトの名無しさん
12/02/11 19:16:02.78
環境： CUDA 4.1 + Visual Studio 2010 Pro. + Windows 7 Pro. 64bit
C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.1\C\common\cutil_vs2010.sln
をバッチビルドして、
cutil64.lib, cutil64.dll,
cutil64D.lib, cutil64D.dll
を作ろうとしたところエラーがでてしまいます。
エラーは出るのですが、cutil64.lib 以外は生成でき、
そのライブラリを使った行列積のサンプルプログラムもコンパイル＆実行できました。
＃ちょっと気持ち悪いですが。
しかし、cutil64.lib はどうしても生成できません。
何が悪いのでしょうか。
cutil64.lib, cutil64.dll,
cutil64D.lib, cutil64D.dll
をビルドしたときのエラーメッセージは以下の通りです。
461:デフォルトの名無しさん
12/02/11 19:16:34.66
------ すべてのリビルド開始: プロジェクト: cutil, 構成: Release x64 ------
C:\Program Files (x86)\MSBuild\Microsoft.Cpp\v4.0\Microsoft.CppClean.targets(74,5): warning : パス 'C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.1\C\common\lib\x64\cutil64.dll' へのアクセスが拒否されました。
bank_checker.cpp
cmd_arg_reader.cpp
cutil.cpp
stopwatch.cpp
stopwatch_win.cpp
Generating Code...
LINK : warning LNK4075: /INCREMENTAL は /OPT:ICF の指定によって無視されます。
LINK : fatal error LNK1104: ファイル 'lib\x64\\cutil64.dll' を開くことができません。
------ すべてのリビルド開始: プロジェクト: cutil, 構成: Debug x64 ------
bank_checker.cpp
cmd_arg_reader.cpp
cutil.cpp
stopwatch.cpp
stopwatch_win.cpp
Generating Code...
ライブラリ C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.1\C\common\lib\x64\\cutil64D.lib とオブジェクト C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.1\C\common\lib\x64\\cutil64D.exp を作成中
LINK : fatal error LNK1158: 'cvtres.exe' を実行できません。
========== すべてリビルド: 0 正常終了、2 失敗、0 スキップ ==========
462:デフォルトの名無しさん
12/02/11 19:21:14.69
追記です。
cutil64.lib がないので 行列積サンプルプログラムは、Releace モードではビルドできていません。
ビルド＆実行でてきているのは、debug モードです。
463:デフォルトの名無しさん
12/02/11 20:12:55.27
Releace + x64モードでcutil64.dllはできるが，cutil64.libはできないと言うことですか？
464:デフォルトの名無しさん
12/02/11 20:14:45.64
エラーが出てもライブラリはできている，と言うことはないですか？
この中に↓
C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.1\C\common\lib\x64
465:デフォルトの名無しさん
12/02/11 20:16:27.24
失礼 Releace -> Release
466:デフォルトの名無しさん
12/02/11 20:26:05.18
まあ、CUDAのインデックス計算ほど
アホ臭いものは無いと思うがな。
テクスチャアドレスではもっと複雑な計算しているのに。
467:デフォルトの名無しさん
12/02/11 20:50:55.57
CUDAを数値計算に使おうという話だよ
468:デフォルトの名無しさん
12/02/11 22:18:37.30
>>463
>>464
お返事ありがとうございます。
C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.1\C\common\cutil_vs2010.sln
のバッチビルドの前後で以下のようにライブラリが増えます。
ビルド前
common\lib\x64\ には cutil64.dll のみ存在
ビルド後
common\lib\x64\ には cutil64.dll cutil64D.dll cutil64D.lib の３つになる
cutil64.libは生成されません。ビルド時に >>461のエラーのためと思われます。
469:デフォルトの名無しさん
12/02/12 00:49:32.40
変だね．こちらも同じようなエラーが出るが（xx.dllを開始できません），xx.libはできている．
試しに，ソルーションエクスプローラーのcutilを右クリックして，「追加」，「既存の項目」で，cutil64.dll を追加してみる．．．．とか
470:デフォルトの名無しさん
12/02/12 00:51:08.01
で，何でビルド前に「cutil64.dll のみ存在 」するのかな？
一度，全部クリーンしてみる．．．．とか
471:デフォルトの名無しさん
12/02/12 01:08:37.66
>>469
お返事ありがとうございます。
エラーは出るんですね。貴重な情報です。
エラーを気にしないですみます。
>cutil64.dll を追加してみる．．．．とか
無理を承知でやってみました。当然ですがダメでした。
>>470
そういうものだと思っていましたが、確かにそうですね。
クリーンインストールしてみます。
472:デフォルトの名無しさん
12/02/12 02:45:15.54
クリーンインストールしてみました。
C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.1\C\common\lib\x64
に最初から cutil64.dll が作られるみたいです。
なお、インストールたCUDAファイルは、
devdriver_4.1_winvista-win7_64_286.16_notebook.exe
cudatoolkit_4.1.28_win_64.msi
gpucomputingsdk_4.1.28_win_64.exe
の３つです。
473:デフォルトの名無しさん
12/02/12 07:49:26.70
cutilのクリーンを実行すれば，すべてのdll，libが消えるはずなのに，もし， cutil64.dllが残っているとすれば，
書き込み，読み込みの権限がないと言うことでは？
cutil64.dll はマニュアルで消せますか？
474:デフォルトの名無しさん
12/02/12 15:39:25.13
>>437
お返事ありがとうございます。
cutil64.dll をマニュアルで削除して、再度
C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.1\C\common\cutil_vs2010.sln
でビルドしたら、すべてのライブラリが生成されました。
ご推察の通り、cutil64.dll にアクセスできなかったのが原因だったようです。
本当にありがとうございました。
＜追記＞
adminユーザーで環境構築作業をしていたので、
cutil64.dll にアクセスできなかった原因は謎のままです。
自分で構築してしまえば、その後は、上書きビルドは問題なくできます。
同じく、CUDA-SDKの中に、simpleTemplate というソースコードがあるのですが、
このソースコードを書き換えて、Releaseモードでビルドすると、
インストール初期からある、simpleTemplate.exe にコンフリクトするらしく、
ビルドエラーが出ました。この場合はも、初期からある、simpleTemplate.exe を削除すれば
問題なく繰り返しビルドができるようになりました。
同じ原因だと思います。
475:デフォルトの名無しさん
12/02/12 15:40:34.77
↑ 番号ミスです。ごめんなさい。
>>473 が正しいです。
476:デフォルトの名無しさん
12/02/12 15:49:09.03
よかった，よかった♪
477:デフォルトの名無しさん
12/02/12 22:34:46.87
OS：32bit CUDA:32bit VisualStudio:32bit
でCUDAプログラムをしている人が、速度を上げることを目的に
OS：64bit CUDA:64bit VisualStudio:64bit
に変更することって意味あります？
478:デフォルトの名無しさん
12/02/12 23:58:34.00
ないと思う。ホスト側で2GB以上のメモリを使いたいわけでもないのだろうし。
64bitビルドして動かすとポインタがレジスタを2つ使ってしまうという話があるので、逆に遅くなる可能性も。
本当に64bitビルド時にデバイス側でもポインタが64bitで扱われるかどうかを
自分の目で確かめたわけではないけど、少なくともここではそういうことを言っている。
高速演算記 第3回 「チューニング技法その１ ＣＵＤＡプログラミングガイドからピックアップ」 | G-DEP
URLﾘﾝｸ(www.gdep.jp)
＞ここでレジスタ数は32bitレジスタの数で、long long, double値や64ビット環境を対象としたCUDAプログラムでは
＞ポインタなどが64bitとして扱われますのでレジスタを2つ使用することになります。
479:デフォルトの名無しさん
12/02/13 10:45:03.24
>>478
なるほど。ありがとうございます。
これと少し環境が違う場合ですが、
OS：64bit CUDA:32bit VisualStudio:32bit
と、
OS：64bit CUDA:64bit VisualStudio:64bit
では、速度差はどうなると思われますか？
ネット上でよく目にするmatrix積のサンプルコードで実験したところ、
同程度か、どうかすると前者の方が速い場合がありました。
前者は、WOW上で動いていると思われるので、
遅いだろうと予測したのですが、それに反する結果でした。
これをどう解釈すればよいか謎なんです。
480:デフォルトの名無しさん
12/02/13 10:46:42.16
479です。OSは、Windows7 64bit Pro です。
481:デフォルトの名無しさん
12/02/13 11:17:41.23
プログラムのどの部分を計測したのであろうか？
CUDAの部分のみ？？
482:デフォルトの名無しさん
12/02/13 12:26:28.94
CUDAの部分のみです。
483:デフォルトの名無しさん
12/02/13 12:35:43.06
CUDAはOSに依存せずに動くとかで、OSのbit数には無関係ということですか？
そのため、
>>478で説明して頂いたように、同程度もしくは32bitが若干速くなる場合もあるということでしょうか。
484:デフォルトの名無しさん
12/02/13 13:17:49.82
ちなみに測定時間はいくらでしょうか？
演算量を増やすとどうなるでしょうか？
485:デフォルトの名無しさん
12/02/13 16:09:51.62
ソースは、以下のHPの最後の「Matrix_GPU」です。
URLﾘﾝｸ(www.gdep.jp)
計算は、スレッドサイズ一杯で組まれているようですので、
#define MATRIX_SIZE 1024/*行列１辺の数*/
をこれ以上上げることができませんでした。
ちなみに、ノートパソコン搭載のGPUチップは、GT-540M。
Releaseでビルドして 650ms付近を中心に数10msばらつきます。
win32とx64で大差はなく、
win32の方が平均10ms程度速いような気がするという程度です。
486:デフォルトの名無しさん
12/02/13 17:03:00.91
CUDAの計算のところを１０００回程度，繰り返すとかはできるでしょうか？
そうすれば有意な差が出るかも知れません．
487:デフォルトの名無しさん
12/02/13 19:18:58.94
200回繰り返してみました。
OS：64bit CUDA:32bit VisualStudio:32bit → 131,671 ms
OS：64bit CUDA:64bit VisualStudio:64bit → 132,146 ms
0.3%の差で、前者が速かったです。
100回繰り返しでも、0.3%差でした。
488:デフォルトの名無しさん
12/02/13 19:39:39.79
なるほど，有り難うございます．
有意な差が出たと言うことですね．
では，後は分かる人，よろしく♪
ε=ε=ε=ε=ε=┌(； ・＿・)┘
489:デフォルトの名無しさん
12/02/14 10:35:58.19
GPU側では64bitのメリットは使用可能なメモリが増えることだけで、その代わりにアドレス計算の命令数や使用レジスタ数が増えるデメリットがある。
前はGPU側だけ32bitモードで動かすという恐ろしい手法があったが、今使うのはなかなか難しい。
490:デフォルトの名無しさん
12/02/14 16:21:21.98
ptxオプションつけてptx出力を見れば見当がつくね。
64ビットアドレッシングだとアドレス計算の為にint32→int64の変換がしばしば行なわれる。
491:デフォルトの名無しさん
12/02/14 18:19:16.13
ざっくり捉えると、32bitと64bitでは、
CUDA自体はさほど変わらない。
CUDAと切り離されてコンパイルされるC++言語部分は、64bitの方がいくらか速い。
という感じでしょうか。
492:デフォルトの名無しさん
12/02/14 18:24:53.01
>>491
アドレス計算多用したり、レジスタ沢山使うようなカーネルだと
CUDA自体32bit版の方が無視できないほど
速くなる可能性はある。
493:デフォルトの名無しさん
12/02/14 18:53:10.54
doubleが速くなるっていう噂はどうなったの？
494:デフォルトの名無しさん
12/02/14 22:53:20.23
GPU上にあるコンスタントキャッシュメモリって、CUDAではどのようにして使われるのでしょうか？
計算の中にリテラルがあった場合（CPUでは命令の即値として埋め込まれますが）や、
constantを付けて宣言した定数が割り当てられたりするのでしょうか？？
495:デフォルトの名無しさん
12/02/16 05:51:50.91
llvmで使えるようになったのなら
boostも使える？
496:デフォルトの名無しさん
12/02/16 08:10:03.40
>>494
適当に数値を入れて，確かめてみるとか
497:デフォルトの名無しさん
12/02/16 08:17:55.24
>>495
２年前にパッチが出ているようです．boostが何か知らんけど
URLﾘﾝｸ(svn.boost.org)
498:494
12/02/17 00:11:19.70
>>496
試してみました。
const変数もリテラルもレジスタにマッピングされました。
ひょっとして、DirectXAPIなどのグラフィックスAPIを利用していた場合に
使われるだけ（要は旧来通りの使われ方）かも･･･
プロファイラを見ると、テクスチャキャッシュのヒット/ミスヒットが見れるようになっていますが
これなんかはグラフィックスAPI専用のような気が。
しかし同じようにコンスタントキャッシュの項目もあって然りと思うが在らず。
う～ん、気にしないことにしますｗ
499:デフォルトの名無しさん
12/02/17 06:31:17.52
貴重な報告、ありがとうございます♪
500:デフォルトの名無しさん
12/02/17 07:54:24.51
GPUを演算器としてフルに使いたい場合って
ビデオ出力をオンボードのものになるようBIOSで切り替えた方がいいのでしょうか
linuxだったらxconfigいじるだけでいけるとか？
501:デフォルトの名無しさん
12/02/17 07:57:38.45
フルの定義は？
502:デフォルトの名無しさん
12/02/17 07:58:52.97
>>494
CPUでも即値として埋め込めない配列や文字列リテラルは
.rdataセクションのデータとして書き込まれるのと同様に、
GPUでも即値として使えないconstデータはコンスタント領域に置かれる。
で、コンスタント領域の値を直接オペランド指定できない
命令を使う場合は当然レジスタを介して使用することになる。
逆に一部の変数にデフォルトで定数が入っているなんてのは
ABIで決まったもの以外は普通無理だから、コンスタント境域などの
メモリに置いて読み込む以外やり様が無い。
503:デフォルトの名無しさん
12/02/17 08:33:23.95
>>500
普通はGPUカードに何もモニタを繋がなければ(≒出力する設定をしなければ)
GVRAMをほぼ全部使い切れるよ。
Linuxなら、Xを立ち上げなければGPUカードから出力していてもかなり使えるよ。
ついでに、3秒だか5秒だからの制限もなくなるからGPUカーネルで長時間計算させられるよ。
504:デフォルトの名無しさん
12/02/17 20:22:12.27
VRAM 4GBのボードが六千円台
URLﾘﾝｸ(akiba-pc.watch.impress.co.jp)
505:494
12/02/18 16:32:22.87
>>502
ありがとうございます。
コンスタント変数の扱われ方、レジスタが必要な理由、理解しました。
書籍を読んでいたら、コンスタントメモリおよびテクスチャメモリについて書かれているのを見つけました。
コンスタントメモリは__constant__修飾子を付けて宣言することで割り当てられるようです。
物理的にはVRAMに置かれ、GPU内部のコンスタントキャッシュでキャッシュが効くものと思われます。
空間フィルタリングなどでフィルタ係数を格納しておくのに適しているのではと思いました。
（繰り返し参照されるのでキャッシュが効きそう）
テクスチャメモリはテンプレートクラス型のtexture<･･･>で宣言することで割り当てられるようです。
物理的にはVRAMに置かれ、GPU内部のテクスチャキャッシュでキャッシュが効くものと思われます。
さらに、グラフィック用途時の同様、テクスチャフィルタリングやサンプリング方式などのハードワイヤード回路機能を
利用できるよう、それらを指定するプロパティを保持しているようです。
506:デフォルトの名無しさん
12/02/18 21:02:09.38
複数カーネルを呼び出すんだけどカーネルのソースはそれぞれ別ファイル
それぞれのカーネルから同じDevice関数を呼び出したいんだけどファイルまたがることできないようだし
ソースごとに同じ内容の別名関数作るしかないのかな？
507:デフォルトの名無しさん
12/02/18 22:43:38.30
ファイルが別でも，プログラムとしてつながっていれば問題はないと思うが．
関数のプロトタイプ宣言をしっかりしておけば良いかと．．．
508:デフォルトの名無しさん
12/02/18 22:46:25.76
共通のdevice関数を別のヘッダに纏めて
カーネルソースからincludeして使う。
device関数はどうせinline展開されるのだから
通常のプログラムでのinline関数と同様に扱えばいい。
509:デフォルトの名無しさん
12/02/18 22:54:19.36
>>508
サンクス
でも今やってみたんだけど全ファイルで展開されるんで
multiple definition of 関数名
ってエラーになっちゃう
実際やってみました？
510:デフォルトの名無しさん
12/02/18 23:22:56.53
カーネル起動時、1スレッドの起動につき1クロックかかり、128スレッドなら128クロックかかると書籍に書いてあるのですが、
これはCPU側でもかかるのでしょうか？
それとも、CPU側はカーネル起動の合図を送ってすぐ制御を戻し、
GPU側の制御ブロックがクロックをかけてスレッドを起動するというかたちでしょうか？
511:デフォルトの名無しさん
12/02/19 01:11:11.40
教えて欲しいことがあります．
CPU側で複数スレッドを立てて，互いに独立した処理を行っています．
そして，それぞれスレッドにおいて，CUDAを使って画像処理をさせたいのですが，（たとえば，グレースケール化とか）
この時，CUDAでの処理にテクスチャメモリを使いたい場合は，どのようなコードを書けばいいのでしょうか？
テクスチャメモリを使う場合，グローバルで宣言しなけらばならないですよね？
たとえば，5枚のRGB画像をグレースケール化するときに，
CPU側で5スレッド立てて，各スレッドでCUDAを使ってグレースケールへの変換処理をしたいのですが，
テクスチャメモリをグローバルで宣言するとおかしなことになるきがするんですが．
どなたか教えて頂けないでしょうか？
宜しくお願い致します．
512:デフォルトの名無しさん
12/02/19 01:21:41.45
CUDA対応のGPUを5枚挿す。
またはGTX590などを3枚挿す。
513:デフォルトの名無しさん
12/02/19 01:26:39.83
>>512
回答ありがとうございます．
テクスチャメモリを使った処理の場合，マルチGPUじゃないと出来ない
とういことでしょうか？
514:デフォルトの名無しさん
12/02/19 01:44:29.68
公式マニュアル見れば分かること
試してみれば分かること
これを聞く人が多すぎる
515:デフォルトの名無しさん
12/02/19 01:53:02.78
4-way SLI でも、ホスト上の4スレッドで利用する場合、
各々1GPUを独立して割り当てるため、4GPU間で通信して仕事をこなす本来的なSLIにはならない、って認識でイイんですよね？
（SLIコネクタがなくても動く？）
516:デフォルトの名無しさん
12/02/19 02:32:56.76
>>511
テクスチャを5個宣言じゃダメなの？
517:デフォルトの名無しさん
12/02/19 02:47:03.16
>>516
回答ありがとうございます．
CPU側でスレッドを立てる段階で決めたいので，
できることなら，可変にしたいと考えています．
514さんは「これを聞く人が多すぎる 」と仰っていますが，
そもそも，このようなことをシングルGPUでやろうとするのが間違いなのでしょうか？
518:デフォルトの名無しさん
12/02/19 02:59:57.92
いやいや、まず試せば分かることじゃん。
その労力を2chに押しつけるのは、今この瞬間は楽では良いかもしれないけど、
自分の成長には繋がらないよ。
519:デフォルトの名無しさん
12/02/19 05:31:26.29
>>513
OS側のスレッド何本あろうとGPUが1つしかないんだから
順番に処理されるだけだろってことでしょ
520:デフォルトの名無しさん
12/02/19 05:43:18.09
concurrent kernel execution…
521:509
12/02/19 07:12:00.07
>>508
失礼しました
インライン関数化で無事できました
ありがとうございました
522:デフォルトの名無しさん
12/02/19 11:42:52.58
>>518
心遣い感謝します．
>>519
>>520
回答ありがとうございます．
やっぱり，そうなんですね．
「concurrent kernel execution」ついて調べて確信しました．
回答・アドバイスして下さった皆様，ありがとうございました．
おかげで解決しました．
523:デフォルトの名無しさん
12/02/19 11:49:26.29
次に同じ疑問を持たれた方のために，
参考資料のアドレスを貼っておきます．
URLﾘﾝｸ(www.nvidia.co.jp)
上記アドレスにあるpdfの「コンカレントカーネル実行」に書かれていることが，参考になるかと思います．
524:デフォルトの名無しさん
12/02/19 16:33:50.20
誰が質問者なのかわからないので謎な流れだが、とりあえずグレースケール化にテクスチャは不要だと思う今日この頃。
525:デフォルトの名無しさん
12/02/19 17:00:58.89
NPPで一発解決だよな
526:デフォルトの名無しさん
12/02/20 22:07:59.44
atomicCASのCASってどういう意味ですか？？
527:デフォルトの名無しさん
12/02/20 22:24:51.56
Compare And Swap
528:526
12/02/20 22:28:38.87
>>527
ありがとうございました！
529:デフォルトの名無しさん
12/03/07 03:16:50.68
ｃｕｄａの本には2次元配列より1次元の方が速いってあったけど実際どのくらい違うんですか？
また、コピーとgpu上での演算の両方で遅くなるんですか？
530:デフォルトの名無しさん
12/03/07 03:39:04.89
なんでこのスレ突然止まってたの？
卒論修論シーズンが終わったから？
531:デフォルトの名無しさん
12/03/07 06:34:38.99
>>529
一行目: 速い
二行目: 遅くなる
馬鹿?
532:デフォルトの名無しさん
12/03/07 16:56:25.77
学会はこれからだというのに
533:デフォルトの名無しさん
12/03/07 17:15:29.75
> ｃｕｄａの本には2次元配列より1次元の方が速いってあったけど実際どのくらい違うんですか？
メモリーのアクセス時間の早さのこと？
ちなみに２次元であっても，例えば f[i, j] で，先にｊを変化させるか，i を変化させるのかで
アクセス時間が違ったと思う．
要はメモリーが並んでいる順番にアクセスするのが早いはず．どの位，早くなるかは知りませぬ♪
> また、コピーとgpu上での演算の両方で遅くなるんですか？
コピーとは？ CPUとGPU間のメモリーの転送のことかな？？
534:デフォルトの名無しさん
12/03/08 10:13:29.60
人いなくなったな
春になってGPUネタで研究する学生が増えるのを待つしかないか
535:デフォルトの名無しさん
12/03/08 11:44:54.65
二次元配列なんて存在しねぇ
って考えると楽なのに
536:デフォルトの名無しさん
12/03/08 19:50:20.17
昔のオカルト本では「謎の四次元」「四次元失踪」とかあって、
「四次元」とか謎めいた雰囲気を感じた。
今なら四次元配列とか当たり前だｗ
537:デフォルトの名無しさん
12/03/09 01:08:22.32
>>536
そうそうｗ
四次元の神秘性が薄らいじゃうｗ
538: ◆QZaw55cn4c
12/03/09 22:10:56.52
ハミルトンの四元数というのがあって近年計算機工学に応用されるようになって云々かんぬん
539:デフォルトの名無しさん
12/03/09 22:22:34.47
発見自体は1800年代だったっけ？
たしか橋の上かなんかで思いついたとかｗ
540:デフォルトの名無しさん
12/03/11 11:39:47.33
HLSLとどう違うの？
541:デフォルトの名無しさん
12/03/11 11:49:19.59
>>540
C言語っぽくなってる。
542:デフォルトの名無しさん
12/03/11 14:43:17.16
四次元配列と四次元ベクトルは別物だろ
後者は要素数4の一次元配列
543:デフォルトの名無しさん
12/03/11 17:04:08.60
>>540
HLSLはShader Languageなんで整数とか扱えなかったと思う。
あとステップ数も制限があったような。
544:デフォルトの名無しさん
12/03/11 17:26:35.76
Compute Shaderを記述する言語も
HLSLじゃ無かったっけ？
リソースベースのHLSLと、ポインタ・配列ベースのCUDA
545:デフォルトの名無しさん
12/03/11 17:38:23.47
リソースベースって言い方、分かり易いね。
HLSLはまさにそんな感じだ。
546:デフォルトの名無しさん
12/03/12 13:07:18.89
>>541>>543
サンクス
自由度がまして打ちやすくなってるんだな
547:デフォルトの名無しさん
12/03/12 13:35:35.95
>>544>>545
サンクス
548:デフォルトの名無しさん
12/03/15 21:09:26.09
誰かいますかね、三つほど質問が。
■__syncthreads()だけでは共有メモリへの書き込みを保証できない？
１つのスレッドがグローバルメモリから共有メモリにデータを書き込み、
その後全てのスレッドがそのデータを使用して計算を行うような場合、
書き込み後に__syncthreads()だけではなく__threadfence_block()も必要なのでしょうか?
青木本には__threadfence_block()について特に言及ありませんでしたが・・・。
■ブロック内の全スレッドからの同一グローバルメモリへのアクセス
ブロック内で共通で使用する構造体などをグローバル→共有メモリに移す場合
全スレッドで行うよりもやはり
if(threadIdx.x==0)・・・
のようにした方が良いでしょうか?
■カーネル内でのreturn文の使用悪影響あるか
スレッドごとに計算を行うか判定をする場合、if文で囲っている例をよく見ますが
これは
if(条件)return;
と書いてはいけないのでしょうか？
上のように書いてもとりあえず計算は流れたのですが何か悪影響はあるでしょうか？
549:デフォルトの名無しさん
12/03/15 21:25:56.04
>>548
・取り敢えずsyncしか使ってないけど問題になったことはない。
・全スレッドから共有メモリへの書き込みを行なうのは多分遅くなるんじゃないかな?
・どちらで書いても同じこと。普通のCPUのような分岐とは違うことを判っていればOK。
550:548
12/03/15 22:39:42.94
>>549
ありがとうございます!!
一個目の_syncthreads()と__threadfence_block()の件ですが、
syncだけだと今日うまくいかなかったもので。
ただ他のバグの影響なども考えられるのでもうちょっと調べてみます。
551:デフォルトの名無しさん
12/03/16 04:06:36.41
>>548
・__syncthreads()は__threadfence_block()相当の処理を
含んでいた気がするけど気のせいかも。
・全スレッドで同じメモリにアクセスするのはたとえfermiでも遅くなるはず。
・カーネル内部で_syncthreads()使う必要があるなら
returnは使っちゃ駄目だろう。
552:デフォルトの名無しさん
12/03/16 20:10:45.37
いまだにアプリ開発環境すらまともに構築できてない・・・
visual studio 2008でやろうと思って一応ビルドは通ったけど
実行するとまずcutil32.dllがありませんって出た。
次にcutil.dllをデバッグ.exeと同じフォルダに置き実行！！
CUDA version is insufficient for CUDART version.
ってなる・・・orz
まずなにからはじめるべきですか？
553:デフォルトの名無しさん
12/03/16 20:17:42.17
ちなみに.cuの中身は拾ってきたちょっと複雑なコード
#include <stdio.h>
#include <cutil.h>
int main( int argc, char** argv )
{
CUT_DEVICE_INIT(argc, argv);
CUT_EXIT(argc, argv);
return 0;
}
・・・・・orz
#include <stdio.h>
#include <cutil.h>
int main( int argc, char** argv )
{
return 0;
}
これに書き換えると
プログラムが完成し、エラーもなく実行もできる
554:デフォルトの名無しさん
12/03/16 20:51:20.83
GPUドライバのアップデート
555:デフォルトの名無しさん
12/03/17 01:39:06.58
>>554
ありがとうございました！！！！！！！！！！
動いた！！！！！！！！
556:デフォルトの名無しさん
12/03/17 08:54:13.85
おおｗ
よかったな！
557:デフォルトの名無しさん
12/03/17 12:49:24.04
>LINK : /LTCG が指定されましたが、コードの生成は必要ありません。リンク コマンド ラインから /LTCG を削除し、リンカの性能を改善してください。
と表示されるのですがリンク コマンド ラインは固定されて編集できません。
解決方法はありますか？
558:デフォルトの名無しさん
12/03/17 13:37:04.85
>>577
補足：
開発環境はVisualStudio2008
cuda ver 2.3
559:デフォルトの名無しさん
12/03/17 14:59:10.69
windowsを窓から投げ捨てろ
560:509
12/03/17 15:16:23.02
そんなことして道歩いてる人の頭に当たっちゃったら大変ですよ
561:デフォルトの名無しさん
12/03/17 15:44:28.84
角に当たったら痛そうだもんね・・・
562:デフォルトの名無しさん
12/03/17 18:01:09.23
>>557
リンカ -> 最適化 -> リンク時のコード生成 (/LTGG)
C/C++ -> 最適化 -> プログラム全体の最適化 (/GL)
Visual Studio 2008 の使い方なのでスレが違うかも。
563:デフォルトの名無しさん
12/03/18 15:04:39.12
>>562
ありがとうございます。
CUDA-Zの実行結果はどのように見たらいいですか？
日本のサイトが全然ないです。
564:デフォルトの名無しさん
12/03/18 15:07:18.71
>>563です
すみません。解決しました。
565:デフォルトの名無しさん
12/03/21 21:06:39.97
コンスタントメモリキャッシュへのアクセスはバンクコンフリクトとかないんでしょうか？？
566:デフォルトの名無しさん
12/03/21 22:20:52.26
>>565
そりゃキャッシュはバンクになってないからねー
567:565
12/03/21 22:44:39.32
>>566
おお、やっぱり。
できるだけコンスタントメモリ使うようにしまつ。
568:デフォルトの名無しさん
12/03/21 23:02:34.40
アドレスが静的に解決できないというのが前提だけど
16ポートのSRAMなんてコスト的に不可能だからマルチバンク以外無いんじゃないの？
569:デフォルトの名無しさん
12/03/22 00:27:08.40
Fermi以前はコンスタントメモリ使う意味あったけど、
Fermi以降はL2キャッシュとあんまり変わらない印象
570:デフォルトの名無しさん
12/03/22 22:48:01.36
GTX680が発表されたけど、CUDA的には好ましくない方向の進化が多い。。
571:デフォルトの名無しさん
12/03/22 23:03:25.80
チップ名 GTX 680 GTX 580
GPC*1 4 4
SM*2 8 16
CUDAコア 1536 512
テクスチャーユニット 128 64
ROPユニット 32 48
ベースクロック*3 1.006GHz 772M/1.544GHz
ブーストクロック 1.058GHz －
メモリー転送レート 6.008Gbps 4.008Gbps
メモリー容量 GDDR5 2048MB GDDR5 1536MB
メモリーバス幅 256ビット 384ビット
メモリー転送速度 192.26GB／秒 192.4GB／秒
製造プロセス 28nm 40nm
補助電源端子 6ピン×2 8ピン＋6ピン
推奨電源ユニット出力 550W 600W
TDP*4 195W 244W
572:デフォルトの名無しさん
12/03/22 23:06:16.47
GK104はミドルレンジだからGK110は全体的に上回ってくるでしょ
573:デフォルトの名無しさん
12/03/22 23:16:01.94
kepler誕生おめ！
.o゜*。o
／⌒ヽ*゜*
∧_∧ ／ヽ ）｡*o ｯﾊﾟ
(・ω・)丿゛￣￣' ゜
. ノ/ /
ノ￣ゝ
574:デフォルトの名無しさん
12/03/22 23:21:31.67
Keplerｷﾀ━━━(ﾟ∀ﾟ)━━━ !!!!!
575:デフォルトの名無しさん
12/03/22 23:22:01.12
gen3じゃないんだっけ？
576:デフォルトの名無しさん
12/03/22 23:42:00.14
>>570
もともとはミドルレンジでグラフィック向けだったから仕方ない気もする。
予想以上にグラフィック方面に舵を切ったという感はあるけど。
このままグラフィック向けとGPGPU向けで大きく分かれていくのではないかという心配はあるかな。
577:デフォルトの名無しさん
12/03/22 23:53:32.45
1SM = 192コアか。おっそろしいなあ。
578:デフォルトの名無しさん
12/03/23 00:00:34.24
nVidia始まったな。
579:デフォルトの名無しさん
12/03/23 00:17:06.29
>>577
warp の扱いどうなるんかな。。。
580:デフォルトの名無しさん
12/03/23 00:54:42.91
>>579
URLﾘﾝｸ(pc.watch.impress.co.jp)
> 32スレッドのWARPに同じ命令を実行する、この基本は、Keplerでも変わっていない。
らしいから、変わらないんじゃないかな。
GF104/114のSMには48コアと2ワープスケジューラ、4ワープディスパッチャで
GK104のSMXには192コアと4ワープスケジューラ、8ワープディスパッチャになっている。
その上レジスタ数は倍、L1キャッシュ/シェアードメモリはそのままってことは
GF104/114よりさらにピーキーになっているのかな？
581:デフォルトの名無しさん
12/03/23 01:01:32.84
>>580
あれ？
コンスタントキャッシュって無くなった？
L1/L2キャッシュがその役割を担ってる？
ということはFermiからか･･･
582:デフォルトの名無しさん
12/03/23 01:06:15.11
48コアが192コアになったのに
レジスタは2倍、
共有メモリは据え置き。
どーすんだこれ。。
583:デフォルトの名無しさん
12/03/23 01:54:37.82
レジスタ足りんくなりそうな。
584:デフォルトの名無しさん
12/03/23 04:09:37.63
Keplerはクロック落としてパイプラインを浅くする設計
演算器のレイテンシが小さくなるならレジスタの消費量は変わらない
Fermiの18cycleは頭おかしすぎた
これが例えば6cycleにになればレイテンシ隠蔽に必要なスレッド数が1/3になるから問題ない
585:デフォルトの名無しさん
12/03/23 04:14:42.78
x86 CPUと同じ道を辿ってるのか
586:デフォルトの名無しさん
12/03/23 15:07:43.87
誰か２６次元計算してくれ、1000ｺｱくらいじゃマジに足らんぞｗ
587:デフォルトの名無しさん
12/03/23 15:59:10.59
float a[1][1][1][1][1][1][1][1][1][1][1][1][1][1][1][1][1][1][1][1][1][1][1][1][1][1] = NULL;
*a += 1;
588:デフォルトの名無しさん
12/03/23 17:47:27.55
こんにちは
国際暗号学会のプレプリントサーバにこんな論文があがってました
Usable assembly language for GPUs: a success story
Daniel J. Bernstein, et. al.
URLﾘﾝｸ(eprint.iacr.org)
GPUのことはさっぱりわかりませんが、なにかこのスレの足しにでもなれば幸いです
それでは
589:デフォルトの名無しさん
12/03/23 20:49:13.93
>584
グローバルメモリアクセスのレイテンシ隠匿とか、ループが遅いとかの情報が頭にあったんで
今まで深く考えず1024スレッド突っ込んでたんだけど、
スレッド減らしてループ回すような構造にしたほうがいい、って解釈でいいんだろうか？
590:デフォルトの名無しさん
12/03/25 02:00:59.95
>>588
PTXよりもっとネイティブ寄りのアセンブラ言語qhasm-cudasmを使って
パフォーマンスクリティカルな場面で力を発揮(nvccの148％)するよ！って話かな？
暗号学会で発表されるんだね。
591:デフォルトの名無しさん
12/03/25 20:09:00.66
メモリアクセスに対する演算の比率を上げないと、性能をフルに発揮できないことは分かったんですが、
具体的にどれくらいの比まで高めるべきかの目標はどうやって決めればイイでしょうか？？
592:デフォルトの名無しさん
12/03/26 10:43:34.08
理論性能（カタログ値）がでるまで頑張れば良いのでは？
それからグローバルメモリーのアクセス速度が，カタログ値の何％になっているのかも
チェックすべきだと思う．
593:591
12/03/26 18:59:54.70
>>592
ありがとうございます！
やはり、
理論性能が出ない → ボトルネックを割り出して改善 → 先頭に戻る
のループで追い込んでいくやり方ですね。
594:デフォルトの名無しさん
12/03/27 00:14:52.50
本を読んで、Visual Profilerを知ったのですが、
ひょっとして今はParallel Nsightで同じことができるでしょうか？
595:デフォルトの名無しさん
12/03/28 05:31:51.00
GTX 680駄目すぎるわ
死んだ
596:デフォルトの名無しさん
12/03/28 06:40:03.41
まだ未公開？
URLﾘﾝｸ(developer.download.nvidia.com)
URLﾘﾝｸ(www.abload.de)
PTX ISA3.0
597:デフォルトの名無しさん
12/03/28 08:00:17.14
warp shuffle
598:デフォルトの名無しさん
12/03/28 10:36:38.09
VLIWの腐ったようなアーキテクチャになったくさいな
599:デフォルトの名無しさん
12/03/28 10:51:03.77
どのへんが？
600:デフォルトの名無しさん
12/03/28 10:54:22.36
斜め読みしたうえで完全にESPだが
命令は各スロットごとに別という点でVLIWでデータパスはSIMDみたいに独立とみた
ソフトウェア的にはもちろん別スレッドとして書けるみたいな
全然違ったらごめんね
601:デフォルトの名無しさん
12/03/28 13:35:57.98
いやPTXレベルの命令だから全然関係ないね
>>598-600は忘れてくれ
602:デフォルトの名無しさん
12/03/28 21:20:01.15
やっぱりGCNと同じでshuffle入れてきたな。
603:デフォルトの名無しさん
12/03/28 22:40:28.29
CUDA C Programming Guide Version 4.2
74p.
Table 5-1. Throughput of Native Arithmetic Instructions (Operations per Clock Cycle per Multiprocessor)
いろいろやばすぎるな。
604:営利利用に関するLR審議中＠詳細は自治スレへ
12/03/29 11:33:59.98
>>589
いまさらだが
物理レジスタが足りてるなら同時に多数スレッドを保持しておけるが
複雑なカーネルだとレジスタは不足しがち
605:営利利用に関するLR審議中＠詳細は自治スレへ
12/03/29 12:27:07.17
>604があるから時として64ビットアドレッシングより32ビットアドレッシングの方が有利なんだよね。
606:営利利用に関するLR審議中＠詳細は自治スレへ
12/03/29 18:48:17.20
Visual studio2010、CUDA4.1、Windows7 64bitではじめようと思ったんだけど、ネットで拾ったプログラムとか動かすとcutil_inline.hが見つからないって出る。
これってどうすればいいの？
607:営利利用に関するLR審議中＠詳細は自治スレへ
12/03/29 19:47:51.42
>>606
GPU Computing SDKをインストールしてパスを通す
608:営利利用に関するLR審議中＠詳細は自治スレへ
12/03/29 22:36:38.83
スレタイがイイね。
くだすれくーだすれｗ
609:営利利用に関するLR審議中＠詳細は自治スレへ
12/04/06 21:14:51.29
なんでGPGPUはみんな同じようなアプリしかつくらないん？
想像力が欠如してるから他人の猿真似ばかりしてんの？
GTX680の性能を生かすアプリ教えろ
610:営利利用に関するLR審議中＠詳細は自治スレへ
12/04/06 21:37:23.28
>>609
邪魔
611:営利利用に関するLR審議中＠詳細は自治スレへ
12/04/06 22:02:22.54
わりぃなぁ。そんじょそこらにはないアプリ作っているんだが公表できないんだわ。
612:営利利用に関するLR審議中＠詳細は自治スレへ
12/04/06 22:56:59.64
恥ずかしいやつが湧いたな
613:営利利用に関するLR審議中＠詳細は自治スレへ
12/04/07 01:09:23.67
>>609
BOINCでもやってな
なんぼぶん回しても次から次へ宿題出してくれるから
614:営利利用に関するLR審議中＠詳細は自治スレへ
12/04/07 10:44:11.21
数値流体力学シンポジウムにでも参加すれば良いぞ♪
615:営利利用に関するLR審議中＠詳細は自治スレへ
12/04/07 19:20:32.56
公表できないと言ってるけど、どうせCUDA ZONEに登録されてるようなものだろ？
外人の真似と負け惜しみしかできないの？
素人だからよくわからないけど、欧米に対して技術面で遅れているから
この分野で日本に有名な人がいないでしょ？
自称一流の教授に俺の書き込み見せてあげて
616:営利利用に関するLR審議中＠詳細は自治スレへ
12/04/07 23:56:42.42
だから数値流体力学シンポジウムにでも参加すれば？
この分野では例えば東京工大の青木先生が有名だが
617:営利利用に関するLR審議中＠詳細は自治スレへ
12/04/07 23:57:49.78
ちなみに 615
618:営利利用に関するLR審議中＠詳細は自治スレへ
12/04/07 23:58:19.73
あれ？
ちなみに６１５は大学生か？？
619:営利利用に関するLR審議中＠詳細は自治スレへ
12/04/08 00:08:49.21
大学生なら，物理や力学関係の学会に参加すれば，GPGPUを使った
シミュレーションの研究結果が報告されていることがわかるはず♪
高卒なら縁はないが．．．
620:営利利用に関するLR審議中＠詳細は自治スレへ
12/04/08 00:11:11.59
大学生がこんな幼稚な文章書いてたら日本終っちまうぞｗ
621:営利利用に関するLR審議中＠詳細は自治スレへ
12/04/08 09:03:46.81
GTX480である程度大きなサイズをホストからデバイスに転送するのに3回に2回ぐらいセグメンテーションエラーで落ちる。
うまく行くとなんの問題もなく実行できる。
デバイス側でメモリ確保ができてないみたいなんだが、こんなもんなのかね？
622:営利利用に関するLR審議中＠詳細は自治スレへ
12/04/08 09:17:09.30
そのカードでモニターを表示させているとか？
623:営利利用に関するLR審議中＠詳細は自治スレへ
12/04/08 12:09:14.45
ケプラーの倍精度計算は速くなったの？
それとも以前と同じ？
624:営利利用に関するLR審議中＠詳細は自治スレへ
12/04/08 12:27:33.01
HPC向けがでてみないと分からんけど
GTX680じゃSPに対して1/24だよ
625:営利利用に関するLR審議中＠詳細は自治スレへ
12/04/08 15:32:40.75
URLﾘﾝｸ(www.tml.tkk.fi)
レイトレは最適化するとスペックなりのパフォーマンスだな680
626:営利利用に関するLR審議中＠詳細は自治スレへ
12/04/08 20:11:08.50
>>622
表示させている。
それがダメなのかな？
627:営利利用に関するLR審議中＠詳細は自治スレへ
12/04/08 21:17:10.19
２枚さして，一枚はディスプレイ用
もう一枚は演算用にしないと，一枚では負荷に耐えられないのでは？
628:営利利用に関するLR審議中＠詳細は自治スレへ
12/04/10 00:35:51.46
みんなすごいね
せいぜい準備して
九九の掛け算一括処理くらいしかできないよ
629:営利利用に関するLR審議中＠詳細は自治スレへ
12/04/10 00:59:53.96
CUDAを学ぼうとするからそうなるんだと思うよ。
何か問題があって、それをCUDAで解こう！って始めた方が早く習得できる。
630:営利利用に関するLR審議中＠詳細は自治スレへ
12/04/10 16:05:07.35
>>624
それって今までよりも遅いってこと？
どこかに倍精度のベンチマークの比較はありませんか？
631:営利利用に関するLR審議中＠詳細は自治スレへ
12/04/10 16:47:44.33
今研究室でGTX680か580のどっちかを買おうって話になってるんだけど
CUDA的にはどっちがいいと思う？
一任されて困ってる・・・
632:営利利用に関するLR審議中＠詳細は自治スレへ
12/04/10 17:43:14.82
floatかdoubleか。
あと整数演算も遅くなったらしい。
テスラ売るためとはいえ、なんかいやーんな感じ
633:営利利用に関するLR審議中＠詳細は自治スレへ
12/04/10 18:28:46.60
> 631
両方買う．
が，テスラの方が安定していると聞いているよ（速度は若干GTXより落ちるが）
634:営利利用に関するLR審議中＠詳細は自治スレへ
12/04/10 18:32:30.11
>>631
迷わず580かと。
680はレジスタとか整数演算・論理演算のスループットとか色々と問題になりそう。
それに最新の正式版Toolkit 4.1のプログラミングガイド見てもKepler載っていないし・・・
>>632
グラフィック性能のワットパフォーマンスを上げるためというのが一番じゃないかな。
635:営利利用に関するLR審議中＠詳細は自治スレへ
12/04/10 19:03:34.66
>>631
勉強用に1枚2枚買うって話ならGTX580だろうね。
１．資料がぜんぜん揃ってないKeplerを今買ってもしょうがない。
２．GPGPUとしての性能がGTX580のほうが「上」
（完全上位互換というわけではないし処理にもよるしスペック上FLOPSでは負けてはいるが）
GTX680のグラフィック・GPGPU性能を調べる ≪ dokumaru
URLﾘﾝｸ(dokumaru.wordpress.com)
10数枚買って研究室全体に大量導入…なら先生が決めるよね。
両方買ってもいい。でもそれならKeplerは対応するToolkitが出てからでも遅くはないかと。
あるいは一刻一秒を争うならなおさらKeplerは冒険かと。
636:営利利用に関するLR審議中＠詳細は自治スレへ
12/04/10 20:08:49.43
今はまだ早いGK110をまて
5月にイベントあるから、そこでなんかあるかも
637:営利利用に関するLR審議中＠詳細は自治スレへ
12/04/11 16:27:30.42
Teslaはメモリ容量が良いよね
GTXだと計算領域が足りないよ・・・
638:デフォルトの名無しさん
12/04/11 23:21:32.32
FLOPS/MBで見ると、Teslaでも全然足りない。
639:デフォルトの名無しさん
12/04/12 07:21:34.36
何の計算？？
640:デフォルトの名無しさん
12/04/13 04:19:22.20
倍精度計算が主なのですが、Ivyと680と580、どれがコストパフォーマンス的にお薦めですか？
641:デフォルトの名無しさん
12/04/13 09:48:41.62
IvyはCUDA動かないよ。
642:デフォルトの名無しさん
12/04/13 10:12:31.27
言語の問題じゃなくて、プログラムはこれから作るから倍精度計算をわんさかやろうと思うんだけどどれがいいかなあ？
程度の話じゃないかと。
643:デフォルトの名無しさん
12/04/13 10:17:18.06
そういうことか。
超並列に対応できるのであればGPUのほうがイイね。
644:デフォルトの名無しさん
12/04/13 10:27:10.68
えっ！？
ここCUDAのスレだよね？
てか、Ocelotとかどうなのかな。
645:デフォルトの名無しさん
12/04/13 12:45:42.72
倍精度計算が中心なら、CPUで最適化するのが一番。
例えば近似計算のようにGPUの単精度で近づけてから、
CPUの倍精度で収束させるとかならありだけど。
646:デフォルトの名無しさん
12/04/13 16:13:12.61
cuda zoneがメンテナス中・・・
647:デフォルトの名無しさん
12/04/13 17:11:45.77
toolkit4.2が来るのかな
648:デフォルトの名無しさん
12/04/13 18:28:30.79
>>640
GPU用の倍精度プログラムを書く気があるならTeslaにしとけ。
コストが厳しいならRadeonの最上位にしとけ。
649:デフォルトの名無しさん
12/04/14 01:50:49.98
sdk 4.1とtoolkit 4.1インストールしたんだけど
アンインストールせずにそのまま
sdk 2.3とtoolkit２．３をインストールしたらコンパイルやリンクの挙動とかおかしくなりますか？
650:デフォルトの名無しさん
12/04/14 12:46:31.14
自分でpathやMakefileなどを管理できるのなら無問題。
651:デフォルトの名無しさん
12/04/14 22:18:52.58
parallel nsightは甘え
652:デフォルトの名無しさん
12/04/14 22:34:41.85
>>651
甘えと言えるほどすごいのか。
今度使ってみようｗ
653:デフォルトの名無しさん
12/04/16 19:24:07.51
シングルGPUでもデバッグできるようになったのが凄くうれしい
654:デフォルトの名無しさん
12/04/17 03:52:59.70
680を手に入れたんだけど、ガッカリ性能だった
ゲーム系は速くなってるんだけどね
655:デフォルトの名無しさん
12/04/17 08:51:17.06
何をやったの
656:デフォルトの名無しさん
12/04/17 15:39:00.54
PTX直接書いてプログラミングする人とかいるの？
657:デフォルトの名無しさん
12/04/17 19:51:20.08
一応いるけど。
658:デフォルトの名無しさん
12/04/17 21:23:06.59
GTX580が生産終了なんだとか
659:デフォルトの名無しさん
12/04/17 23:38:03.55
>658
在庫はけたからか。
しかしIntelのCPUはGPU載せてんのにGPGPUにはさっぱり対応しないからあんま意味ないな。
CUDA対応とかにしないのは戦略的判断なんだろうけど、なんとももったいない。
660:デフォルトの名無しさん
12/04/17 23:51:03.54
大丈夫
Intelも来週からGPGPU対応する
661:デフォルトの名無しさん
12/04/18 16:31:10.22
倍精度計算じゃまだインテルに分があるしね
662:デフォルトの名無しさん
12/04/18 16:57:28.76
ivyのeuどうなってのかね
663:デフォルトの名無しさん
12/04/20 18:04:01.73
> 654
今，それに触手を伸ばしているところだけど，どこがダメだった？？
664:デフォルトの名無しさん
12/04/21 20:08:12.47
>>660
ivyってGPGPUとして使えんの？
665:デフォルトの名無しさん
12/04/21 20:18:04.15
OpenCLならできなくもないのかな？
666:デフォルトの名無しさん
12/04/24 13:30:34.56
誰かHMPP使ったことある人いる？
667:デフォルトの名無しさん
12/04/24 20:58:18.61
CUDA4.2きたな
URLﾘﾝｸ(developer.nvidia.com)
668:デフォルトの名無しさん
12/04/25 20:18:30.00
Adobe Creative Suite 6: Bye bye CUDA, Hello OpenCL!
URLﾘﾝｸ(www.geeks3d.com)
669:デフォルトの名無しさん
12/04/25 20:24:16.57
ivyの影響だな
670:デフォルトの名無しさん
12/04/26 23:10:43.77
CPU、GPUを利用（プログラム）するには？
URLﾘﾝｸ(togetter.com)
671:デフォルトの名無しさん
12/04/29 14:44:53.67
GPUのデメリットは同じ変数計算を毎回糞真面目に超高速で行うところ
672:デフォルトの名無しさん
12/04/29 18:55:51.68
メモリ読むより速いからな
673:デフォルトの名無しさん
12/04/30 07:47:23.73
GTX690
URLﾘﾝｸ(pc.watch.impress.co.jp)
674:デフォルトの名無しさん
12/04/30 08:35:43.64
>> 671
一つ一つの計算は超高速でもなんでもない
並列で行うので早くなるだけ
超高速になるか否かはプログラミングの問題
>>672
演算にはメモリーの読み書きを伴うので，演算が「メモリ読むより速い 」とはならないのでは？
675:デフォルトの名無しさん
12/04/30 21:06:45.57
>>671
意味が分からん。
アーキの概念の理解ができていないじゃねーか？
676:デフォルトの名無しさん
12/05/01 00:39:15.46
>>675
明日短いわかりやすいソースアップするからコンパイルして実行してみて
言いたいことがわかると思う。
CPUにはあってGPUにはない機能を使うことになる、まぁホントしょうもないことだけど・・・
677:デフォルトの名無しさん
12/05/01 00:53:32.34
？？？
678:デフォルトの名無しさん
12/05/01 00:59:29.65
言ったら悪いかも知れんけど単にアルゴリズムが悪いんじゃないのか。
679:デフォルトの名無しさん
12/05/01 08:09:04.59
>>671
アドレス計算とかまさにそれだよね。
普通のループなら+4で済むところが、
ptr + threadIdx.x*4 + threadIdx.y*hoge
とかになっちゃう。
680:デフォルトの名無しさん
12/05/01 08:24:43.67
それはGPUのデメリットじゃないな。
GPU(nvcc)でもループなら普通に書いたら普通に最適化してくれる。
681:デフォルトの名無しさん
12/05/01 10:38:58.63
>>679
これはデメリットと違う。
CPUでマルチスレッドでやれば同じように明示的にアドレス計算を行う必要がある。
682:デフォルトの名無しさん
12/05/01 13:31:24.76
シングルスレッドでの最適化が、そのままマルチスレッドに使えると思ってるなら、並列で組むのに向いてないな。
ひとつの処理として見たとき無駄でも、それで大多数の演算を同時に走らせることができるなら、
並列処理においてはそれこそが効率的なんだよ。
683:デフォルトの名無しさん
12/05/02 01:02:11.09
>>669
残念、AMDとアドビのコラボでした
URLﾘﾝｸ(www.4gamer.net)
Ivyは端から相手にされてません
684:デフォルトの名無しさん
12/05/02 01:48:54.23
URLﾘﾝｸ(ascii.jp)
このカード使ってる人居ませんか？
メモリがいっぱい欲しいけど、高いカードは買えないので
試しに買ってみようかと思うのですが。
685:デフォルトの名無しさん
12/05/02 02:05:34.87
同じ世代のGPUでも生産地の違いで演算速度は全く違うからね
もっと言うと転送速度が全く違う
まあフラッシュメモリでも同じこと言えるけど
686:デフォルトの名無しさん
12/05/02 04:31:00.56
>683
ずいぶんニッチなところだな
687:デフォルトの名無しさん
12/05/02 09:30:37.50
>>684
マジレスすると、CUDAでやるメリットはない。
Sandyやivyの方がはるかに高速。
まあ、CUDA勉強するだけならいいが、もっと別のカードのほうがいいだろ。
688:デフォルトの名無しさん
12/05/02 11:01:38.16
>>684
メモリ転送が遅過ぎて4GBのメモリを活かしきれない悪寒。
689:デフォルトの名無しさん
12/05/02 12:41:42.31
SRAMを４GBつんでるカードはないのか？
690:デフォルトの名無しさん
12/05/02 16:10:34.75
>>667
>New Features
>Support for GK10x Kepler GPUs.
とりあえず、GK104対応にしました的か。
691:デフォルトの名無しさん
12/05/02 18:28:26.54
VRAM 4GB以上のカードって、ほとんどないんだね。
TeslaかQuadroしか見つからなかった。
お値段10万円越
692:デフォルトの名無しさん
12/05/02 18:42:59.54
4年後にはVRAM16GBが普通にでまわるんだよ
693:679
12/05/02 22:18:41.81
>>680 >>681
あれ、そういう話じゃないのか…
そうだとすると、 >>671 が何を言いたかったのか思い付かないな…
694:デフォルトの名無しさん
12/05/03 15:04:28.21
>671は皮肉だろ。
695:デフォルトの名無しさん
12/05/03 16:07:25.13
ということにしたいのですね。
696:デフォルトの名無しさん
12/05/04 01:03:14.33
デメリットに感じる境地まで辿りついたんだよ、きっと
俺にはまだメリットにしか思えないんだけど・・・
697:デフォルトの名無しさん
12/05/05 21:18:13.51
>>692
今から4年前の2008年頃はG80世代で大体1GBだったが、4年後のGTX680でまだ2GBだから、
4年後はせいぜい4GBなんじゃないの？Tesla系で16GBにはなっていそうだけど。
698:デフォルトの名無しさん
12/05/05 21:24:38.22
そういやPCのDRAM搭載量に比べて、あんまり伸びないよね＞ビデオカードのメモリ
699:デフォルトの名無しさん
12/05/05 21:26:52.65
GDDRは数が出ないからね。
DRAMメーカーがあんな状態だから尚更でしょう。
700:デフォルトの名無しさん
12/05/05 21:33:02.38
プロセスシュリンクが汎用DRAMと同じように進めば同じようにでかくなると思うんだけど。
だんだん引き離されてるってこと？
701:デフォルトの名無しさん
12/05/10 01:35:01.75
日本のメモリの会社が潰れたのはかなり痛いな・・・
702:デフォルトの名無しさん
12/05/10 15:36:00.11
ptxコード読まなきゃいけなくなったんだけど、typeの.predって何なのかいまいちわかってない
703:デフォルトの名無しさん
12/05/10 16:46:59.88
述部(predicate)だね。
ptxの場合は単に、比較などの結果を保持するだけのような希ガス。
で、そのレジスタの結果に依存してインストラクションの実行する、と。
例えば、
--
setp.gt.s32 %p1, %r5, %r7;
@%p1 bra $Lt_0_12802;
--
なら r5 > r7のときに分岐するし、
--
setp.lt.s32 %p2, %r9, %r11;
@%p2 sub.s32 %r14, %r11, %r14;
--
なら r9 < r11のときに引き算を行なう。
704:デフォルトの名無しさん
12/05/10 21:33:54.24
分岐マスクのためのレジスタは何本あるんだろ
それとも汎用レジスタと共用なのか
705:702
12/05/11 16:21:50.31
>>703
thx
そういう意味だったのか……
CUDAはC言語の延長だから大丈夫とか考えた三月の俺を叩きのめしたい
PTXコードの読み方って英語のやつしかないよねたぶん
706:デフォルトの名無しさん
12/05/11 17:11:50.74
>>705
私が書いたメモならあるよw
>>704
実験コードで見たところ、汎用レジスタと述語レジスタの合計で制限されてたかと。
述語レジスタだけでどこまで増やせるかは実験してない。
707:702
12/05/11 17:26:43.70
>>706
恵んでください。
割と切実に。卒業したいので。
708:デフォルトの名無しさん
12/05/11 18:02:12.84
ISA的にはwarpあたり7本か6本じゃね。
3bitのどれかが常にalways扱いだったような。
709:デフォルトの名無しさん
12/05/14 15:19:34.27
初心者質問です。
お願いします。
cufftってcuda3.2でも使えるのでしょうか？
cufftdestroyが未解決の外部シンボルだと言われてしまうのですが？
ただ単に、リンクできてないだけなのでしょうか？
710:デフォルトの名無しさん
12/05/14 18:14:14.93
>>709
使えたと思うよ。
つーか、cufftdestroy()が未解決って、あんたの間違いだろ。
711:デフォルトの名無しさん
12/05/15 14:03:46.41
>>710
返信ありがとうございます
他の関数はコンパイルが通る(通っているように見えるだけ？)のに
cufftdestroy()
cufftExecZ2Z()
cufftPlan1d()
だけが未解決となっているのですが、
この関数だけ、他のライブラリが必要だなんてことがあるのでしょうか？
712:デフォルトの名無しさん
12/05/15 15:52:04.02
destroyはDestroy。
z2zは未実装。
Plan1dはしらね。
警告レベル引き上げれば?
713:デフォルトの名無しさん
12/05/17 11:36:49.94
cufft.hはインクルードしているのかな？
714:デフォルトの名無しさん
12/05/17 13:08:31.44
■後藤弘茂のWeekly海外ニュース■
NVIDIAが世界最多トランジスタ数のチップ「GK110」を公開
URLﾘﾝｸ(pc.watch.impress.co.jp)
715:デフォルトの名無しさん
12/05/17 15:21:47.32
警告レベルって、デフォルトは最大なんですよね？
Destroyに関しては、タイプミスです。
z2zは未実装っていうのが、よくわからないんですけど。。。。
716:デフォルトの名無しさん
12/05/17 17:10:10.69
GPGPU上でソケット通信とかって出来るかな
717:デフォルトの名無しさん
12/05/17 17:20:54.19
GPGPUの仮想マシン同士のn対n通信をシミュレートとかそういうのをイメージした
718:デフォルトの名無しさん
12/05/17 18:16:22.69
>>715
未実装: 実装されていないこと。
cufftのライブラリの中にz2zの関数そのものが存在していないのよ。
で、あんたがどんな環境で開発しているか判らんのに警告レベルがどうなっているかなんか判るかい。
そんなことは自分で調べなさいよ。
>>716
cuda5でLAN接続されているGPU同士で連携させる機能がつくらしいよ。
719:デフォルトの名無しさん
12/05/17 21:52:07.76
多次元配列を扱えないのは何でなんだろう．
ブロックとスレッドインデックスで一次元化するの面倒なんだけど．
720:デフォルトの名無しさん
12/05/17 23:46:21.89
ピンメモリを確保すると、スワップによる退避を防げるのは分かったのですが、
実際はスワップ以外にも、メモリフラグメンテーション解消のためのコンパクションでも
メモリアドレスの変化って起こり得ますよね？
それもないようにするのがピンメモリですよね？
721:デフォルトの名無しさん
12/05/18 01:25:35.91
>>719
別に扱えなくはないぞ。普通にdata[blockIdx.x][thiredIdx.x]ってできると思う。
スレッド数を定数にしなくちゃならなくなるから却って煩わしいと思うけど。
つーか、面倒ったってオフセット計算する関数を作るだけじゃん。