【GPGPU】くだすれCUDAスレ pert4【NVIDIA】
レス数が1000を超えています。これ以上書き込みはできません。
0001デフォルトの名無しさん
2010/12/04(土) 21:57:13もしくは質問者自身何が何だが分からない質問を勇気を持って書き込むスレッドです。
CUDA使いが優しくコメントを返しますが、
お礼はCUDAの布教と初心者の救済をお願いします。
CUDA・HomePage
http://www.nvidia.com/cuda
関連スレ
GPUで汎用コンピューティングを行うスレ
http://hibari.2ch.net/test/read.cgi/tech/1167989627/
GPGPU#5
http://hibari.2ch.net/test/read.cgi/tech/1281876470/
前スレ
【GPGPU】くだすれCUDAスレ【NVIDIA】
http://pc12.2ch.net/test/read.cgi/tech/1206152032/
【GPGPU】くだすれCUDAスレ pert2【NVIDIA】
http://pc12.2ch.net/test/read.cgi/tech/1254997777/
【GPGPU】くだすれCUDAスレ pert3【NVIDIA】
http://hibari.2ch.net/test/read.cgi/tech/1271587710/
0002デフォルトの名無しさん
2010/12/04(土) 21:58:52CUDA Zone
http://www.nvidia.co.jp/object/cuda_home_jp.html
フォーラム
http://forum.nvidia.co.jp/EokpControl?&event=TE0001
CUDAに触れてみる
http://chihara.naist.jp/people/STAFF/imura/computer/OpenGL/cuda1/disp_content
cudaさわってみた
http://gpgpu.jp/article/61432191.html
CUDA のインストール
http://blog.goo.ne.jp/sdpaninf/e/9533f75438b670a174af345f4a33bd51
NVIDIAの「GeForce 8800 GT(G92)」と次に控える64-bit GPUアーキテクチャ
http://pc.watch.impress.co.jp/docs/2007/1031/kaigai398.htm
CUDAを使う
http://tech.ckme.co.jp/cuda.shtml
NVIDIA CUDAを弄ってみた その2
http://dvd-r.sblo.jp/article/10422960.html
CUDAベンチ
http://wataco.air-nifty.com/syacho/2008/02/cuda_2044.html
KNOPPIX for CUDA
http://www.yasuoka.mech.keio.ac.jp/cuda/
0003デフォルトの名無しさん
2010/12/04(土) 22:27:14低レベル質問にも対応するスレだったんだな
道理でCUDAは意味が無いという意見が出て来るわけだ
0004デフォルトの名無しさん
2010/12/04(土) 23:04:150005デフォルトの名無しさん
2010/12/04(土) 23:09:17ORのPERT?
だとしてもCUDAとの関係がわからん
0006デフォルトの名無しさん
2010/12/04(土) 23:32:440007デフォルトの名無しさん
2010/12/04(土) 23:43:310008デフォルトの名無しさん
2010/12/05(日) 01:01:35買ったほうがいいですか?ちなみに行列の演算を高速化したいというのが
目的です。
0009デフォルトの名無しさん
2010/12/05(日) 01:28:17GT430にしろよ・・・
0010デフォルトの名無しさん
2010/12/05(日) 01:30:180011デフォルトの名無しさん
2010/12/05(日) 11:15:340012デフォルトの名無しさん
2010/12/06(月) 02:44:34の意味がわかりません。意味を詳しく教えていただけませんか?
0013デフォルトの名無しさん
2010/12/06(月) 17:16:54activeなwarpってのは要はメモリアクセスとかで待ってる状態じゃなく今すぐ実行できるwarpの事
fermiだと0から48まで
active cycleはactiveなwarpが一つでもあるcycleの数
例えば総cycleが10cycleだとして、
毎cycleにactiveなwarpが有ればactive cycleの値は10
10cycleの内、activeなwarpが一つもないcycleが一回あったとしたら9
warpが分からないとかなら青木先生の本でも読めばいい
0014デフォルトの名無しさん
2010/12/06(月) 19:37:01warpについては理解しています、がいまだcycleがわかりません・・・
cycleとはある命令を実行するのにかかる時間であり、ある命令はwarp単位で発行されるという認識です。
ですので1cycleで48warpがアクティブになるというのがよくわかりません。
またこれらの値は大きいほうがいいのか、小さいほうがいいのかもいまいちわからない状況です。
たとえばSummary Tableにでているactive warps/active cycle やretire ipcなども
意味と値が大きいほうがいいのか、がいくら調べてもわかっていない状況です。
よろしければご教授お願いします。
0015デフォルトの名無しさん
2010/12/06(月) 21:51:210016デフォルトの名無しさん
2010/12/06(月) 21:58:05>>14高校生?
0017デフォルトの名無しさん
2010/12/06(月) 22:08:20プロファイラのマニュアルはちゃんとあるし、
プログラミングガイド、ベストパフォーマンスガイドに基本は全て載ってる
全部公式のガイドだぞ
0018デフォルトの名無しさん
2010/12/06(月) 22:26:33初心者とゆとりが融合した素晴らしい素材ってことは分かる
0019デフォルトの名無しさん
2010/12/06(月) 23:18:431cycleで複数のwarpを処理できるという点がいまいちどういう仕組みなのかわかっていませんでした。
Type SMとありましたので、この値はあるSMで計測された値だということがわかりますが、
上記のような認識だとつじつまが合わない(1SMで1warpずつ処理すると思っていたのに、1cycleで最大48warpがアクティブになるという点)ため質問させていただきました。
ややこしく、説明がいまいちになってしまい申し訳ないです。
プロファイラのマニュアルも読みましたが、読んだ上でゆとり質問となり言い返す言葉もありません。
これ以上スレを汚すのもどうかと思いますので質問を〆させていただきます。
レスくださった方ありがとうございました。
0020デフォルトの名無しさん
2010/12/06(月) 23:50:52待ち行列が無いと、演算器に空きが出来てもwarpを割り当てる事が
出来ないから、ある程度の長さの待ち行列が常にある事が
(演算器の利用効率を最大化するという意味では)重要になる。
fermiの細かいアーキテクチャまでは分からないが、
アレだけレジスタが多いとregisterがマルチバンクになっている可能性が高いから
例えば奇数cycleは奇数IDのwarp、偶数cycleは偶数IDのwarpしか発行できない
とかそういう制限があるかもしれない。
その場合はactive warpだからといって
常に任意のcycleに実行可能とならない場合もあるから
待ち行列は長い方が良い。
0021デフォルトの名無しさん
2010/12/07(火) 01:36:55質問とはちょっとずれるかもしれないけど、
http://www.softek.co.jp/SPG/Pgi/TIPS/public/accel/gpu-accel2.html
この記事が長いけど結構分かりやすいかもね
ゆとりとかは挨拶だから気にするなよ
0022デフォルトの名無しさん
2010/12/07(火) 01:52:08GT430で探したんですが1スロットで収まる商品が2つしかなく、
評判も安定性も良くないらしいのでELSAのGT240にしようかと思っています。
何かアドバイスがあればよろしくお願いします。
0023デフォルトの名無しさん
2010/12/07(火) 02:23:5621さんのレスなんか涙がレスが読めなくなるほど。
ゆとりくんは心に刻むんだぞ。
玄人のおぢさんたちに生暖かく見守られているということを。
0024デフォルトの名無しさん
2010/12/07(火) 02:39:12430使ってるが不具合ないよ。
あと今更fermi以前でPG組む気になれん。
CUDA目的ならCCのバージョンを意識しないと駄目だと思うぞ。
それをわかってて240にするなら何も言わない。
0025デフォルトの名無しさん
2010/12/07(火) 02:40:260026デフォルトの名無しさん
2010/12/07(火) 02:53:55目的、予算、制限を明確にしてほしいな
高速処理が至上命題で、予算も制限もないなら
言うまでもなく、Teslaボードをお勧めします
やっぱりfermi世代の方が無難なチョイスなので、
fermiチップのボードだとしてアドバイス
ELSAのGT240は使ってないからよくわからないけれど、
メーカーとしては鉄板の印象。安心していいのでは?
CUDAの場合、仮想メモリは使えないので、やりたい計算で
512MBのメモリだけで、溢れないか見積もることをお勧めします
1スロット幅を譲れない場合、少々値が張るが
Quadroならば1スロット幅のモデルもあるよ.
Quadro4000までは1スロット幅のようだ
市場価格は2万程度から10万弱くらい
2スロットでもよくて,そこそこ安く高性能というなら
やっぱりGTX580あたりがいいんじゃないかな
doubleやECC機能を利用しない限り、Teslaモデルより高速
コスパ最強です
それからWindowsの場合は画面出力用のボードとCUDA用のボードを分けると
デバッグしやすい気がするので、ボード2枚構成というのも考慮の余地があるかも
Linuxはしらない
0027デフォルトの名無しさん
2010/12/07(火) 02:58:31GT440だと脳内保管されてた・・・
ワシもfermiでないGT240は推奨しません
CUDAの勉強が目的ならGT240の方が
勉強になるかもしれない
fermiはキャッシュのおかげでタコなコードでも
そこそこ動いてくれるから
0029デフォルトの名無しさん
2010/12/07(火) 09:53:37詳しいことはわかりませんが、同じコア数でもGT240とGT439ではなぜ
これだけ値段の差があるのか解りません。できれば安いほうを買いた
いのですが予算は10000円程度で考えています。
何かよきアドバイスがあればよろしくお願いします。
0030デフォルトの名無しさん
2010/12/07(火) 09:59:41fermiじゃなきゃ高速演算の手間がかかる
と思えばおk
0031デフォルトの名無しさん
2010/12/07(火) 10:11:53これなんかどうでしょう?騒音がすごいというレビューを見ましたが。
0032デフォルトの名無しさん
2010/12/07(火) 12:48:45詳しくはは自分でも調べてほしいのだが、fermiってのはアーキテクチャの世代の名前。
fermiアーキテクチャのチップは4xx(480,470,465,460,450,430,)と5xx(580)。
それ以前の2xxシリーズなどはfermiではない。
cudaは最適化とかに限らず、ハードウエアで使える資源(メモリやALU)のことを意識して
プログラムを組む必要があるのだが、fermi世代だと色々制限がなくなって、プログラムを組む
労力が少なくてすむようになってるし、パフォーマンスも上がりやすい。
それで制限に関係するのがCC(Compute Capability)の番号。
この番号によって使える機能が変わるfermiは今のところ2.0と2.1。
詳しくはcuda cプログラミングガイド3.2.2 Appendix G.参照
「詳しいことは知りませんが」とか言って下調べをしないで
プログラムしようとすると痛い目にあうぞ
予算1万ならfermi世代だと430しかないな。
0033デフォルトの名無しさん
2010/12/07(火) 13:16:57最初はIONみたいなので消費電力を気にせずつかえるのがいいと思う。
速いカードの方が当然チューニングの効果がでるからおもしろい。
プロファイラで取れる情報もカードによって違うみたい。
0034デフォルトの名無しさん
2010/12/07(火) 13:24:26添字が小数点の時の線形補完を使ってみたいのですが、
メモリはチップ外にあるということはグローバルメモリと同等の速度なのですよね
キャッシュはされるみたいですが同じデータを何度も利用しなくては意味がないということでしょうか
0035デフォルトの名無しさん
2010/12/07(火) 15:47:54EVGA GeForce GT 430で迷っています。
どちらがいいかアドバイスをよろしくお願いします。
0036デフォルトの名無しさん
2010/12/07(火) 15:54:56印象でいえばELSAは手堅い印象。
EVGAはゲーム向けのチューニングとかしてて精度の点が不安。
予算的に許すならquadro600がいいかもしれない。
ただしELSAでもgeforce,quadroはcudaでの演算結果の保証はしてないから
演算を間違うコアが含まれていても交換対象ではない。
0037デフォルトの名無しさん
2010/12/07(火) 16:01:43有限体の行列なので浮動小数点は使いません。
予算は1万円です。コアの精度ってなんですか?
コアが計算を間違えることがあるのですか?
精度を求めるならどちらがいいでしょうか?
質問ばかりで申し訳ありませんが、ご回答よろしくお願いします。
0038デフォルトの名無しさん
2010/12/07(火) 16:06:170039デフォルトの名無しさん
2010/12/07(火) 16:08:52浮動小数点の意味わかってる?
0040デフォルトの名無しさん
2010/12/07(火) 16:13:50自分が計算したいのはunsigned shortなのです。
以前雑誌の特集にCUDAの記事がありNVIDIAのコアには整数演算も
可能であるということが書かれてありました。なので整数の行列も
高速化できるかと思っていました。何か間違いがありましたら
ご指摘をお願いします。
0041デフォルトの名無しさん
2010/12/07(火) 16:14:46演算結果は保証されないけど、geforceより厳しい検査がされているから
より演算結果の値にエラーがでにくいだろうという推測にもとづくもの。
演算を間違うコアがふくまれる場合があるというのは簡単に言えば1+1が2以外の結果になる
コアが含まれる場合があるって意味。
精度を求めるというのが演算結果を保証するならという意味ならTESLAしかない。
そこまでの予算を確保できない場合、上記のようなコアが含まれていないかを確認するプログラムを自分で
組んで確認しないといけない。
単純に足し算だけでなく、四則演算、浮動小数点演算、指数関数、超越関数、L1、L2キャッシュ、
VRAMに至るまで正当性を確認する必要がでると思ってる
0042デフォルトの名無しさん
2010/12/07(火) 16:21:20簡略な確認としてもいいかもしれない。それで許されるのであれば。
0043デフォルトの名無しさん
2010/12/07(火) 16:30:24間違いがないことを実験する必要があるのですね。
同じ計算機なのにどうしてそのような確率的な動作をするのか
理解しかねます。CPUだけが特別に演算が正確なのはなぜでしょう?
GPUも同じ32ビットレジスタなので構造が同じだと思ったのですが。
0044デフォルトの名無しさん
2010/12/07(火) 16:41:06CPUは元々計算用に作ってあるでしょ?CPUで1+1=2を保証しなくてだれが買うよ?
GPUは元々描画用に作られてて座標計算とかfloatで計算することが多い。
そのときに1.000002が1.000002になってもあんまり影響ないでしょ?
もちろん正確に計算できるほうがいいが、描画だったら多少不正確でも問題はないだろうっていう判断だと思う。
それにCPUに比べてALUが膨大なこともあってすべてのコアの演算精度を保証してたらキリがないでしょ。コスト的に。
そういう用途の人のために厳しいテストをしたTESLAが用意されてる訳。でもその分料金上乗せさせてもらいますよって事だよ。
0045デフォルトの名無しさん
2010/12/07(火) 16:41:22並列化のメリットがなくなってしまう気がします。
0046デフォルトの名無しさん
2010/12/07(火) 16:42:14× そのときに1.000002が1.000002になっても
○ そのときに1.000001が1.000002になっても
0047デフォルトの名無しさん
2010/12/07(火) 16:45:03一回CPUと同等の計算をして、演算結果に謝りがないことの確認をもって
GPUのコアにエラーが発生していないことの簡略確認にすれば?
ってことだよ
0048デフォルトの名無しさん
2010/12/07(火) 16:48:03計算してもらいたいと祈るばかりですね。
0049デフォルトの名無しさん
2010/12/07(火) 16:51:100050デフォルトの名無しさん
2010/12/07(火) 16:54:15が、無理なんだろうなあ
0051デフォルトの名無しさん
2010/12/07(火) 17:07:36typedef union uni {
unsigned long long int dd[2];
unsigned int cc[4];
unsigned short ss[8];
unsigned char m[16];
} on;
void code (char str[256]){
int i,c,d;
unsigned long long int u=0;
i=0;
do { c = str[i] & 0xff; e[i]=c; if((c & 0x80) > 0)
{ ++i; d = str[i] & 0xff; e[i]=d; c^=(d<<8);
printf("'%c%c' %u\n",c,d,c);
} else { printf("'%c' %d \n",c,c); } ++i; }while(c != '\0');
printf("\n");
for(i=0;i<16;i++)
printf("%u ",e[i]);
printf("\n");
for(i=0;i<8;i++)
u^=(u<<8)^e[i];
c3[0]=u;u=0;
for(i=8;i<16;i++)
u^=(u<<8)^e[i];
c3[1]=u;
printf("%llu %llu\n",c3[0],c3[1]);
}
0052デフォルトの名無しさん
2010/12/07(火) 17:08:51int a[16]={7,3,15,11,5,1,13,9,6,2,14,10,4,0,12,8};
int i,j;
unsigned int b=0,c;
__m128i GG;
for(i=0;i<16;i++)
GG.m128i_u8[a[i]]=cc.m128i_u8[i];
for(i=0;i<4;i++)
b^=(GG.m128i_u8[i]<<(8*i));
c=b&0x1f;
b=(b>>5)^(c<<27);
for(i=0;i<4;i++){
GG.m128i_u8[i]=b&0xff;
b=(b>>8);
}
return GG;
}
0053デフォルトの名無しさん
2010/12/07(火) 17:13:42int sum=0,j=0,i=0,k=0,fd,read_size,o,p,l,ll;
FILE *fp,*fq;
unsigned short m;
char fin[80],fout[80];
__m128i pa = a; __m128i pb = b; __m128i pc = c;
__m128i *pd1= (__m128i*)d1;__m128i *pd2= (__m128i*)d2;__m128i *pbuf=(__m128i*)buf;
__m128i *pbuf2=(__m128i*)buf2;
__m128i va, vb, vc, vd1, vd2;
printf("main\n");
fp=fopen(argv[2],"wb");
fq=fopen(argv[1],"rb");
hash(argv[3]);
for(i=0;i<16;i++){
d1[i]=0;
d2[i]=0;
}
0054デフォルトの名無しさん
2010/12/07(火) 17:15:17for(k=0;k<SIZE/32;k++){
memcpy(buf,&buff[32*k],32);
for(j=0;j<16;j++){
o=FG[a.m128i_u8[j]];
p=FG[b.m128i_u8[j]];
for(i=0;i<16;i++){
d1[i]^=GF[((o+h1[p][i]-2)&0xff)+1];
d2[i]^=GF[((o+h2[p][i]-2)&0xff)+1];
}
vd1 = _mm_loadu_si128(pd1); vd2 = _mm_loadu_si128(pd2);
_mm_storeu_si128(pbuf, _mm_xor_si128(_mm_loadu_si128(pbuf), vd1));
pbuf++;
_mm_storeu_si128(pbuf, _mm_xor_si128(_mm_loadu_si128(pbuf), vd2));
pbuf--;
}
_mm_storeu_si128((__m128i *)&a,_mm_xor_si128(_mm_loadu_si128((__m128i *)&a),_mm_and_si128( _mm_add_epi8(_mm_loadu_si128((__m128i *)&d1),_mm_loadu_si128((__m128i *)&c)),_mm_set1_epi8(0x0ff))));
_mm_storeu_si128((__m128i *)&b,_mm_xor_si128(_mm_loadu_si128((__m128i *)&b),_mm_and_si128( _mm_add_epi8(_mm_loadu_si128((__m128i *)&d2),_mm_loadu_si128((__m128i *)&c)),_mm_set1_epi8(0x0ff))));
buf[0]=*pbuf->m128i_u8;
pbuf++;
buf[16]=*pbuf->m128i_u8;
pbuf--;
memcpy(&buff[32*k],buf,32);
}
fwrite(buff,1,read_size,fp);
}
fclose(fp);
fclose(fq);
return 0;
}
0055デフォルトの名無しさん
2010/12/07(火) 17:15:28レジスタ
L1
共有メモリ
L2
0056デフォルトの名無しさん
2010/12/07(火) 17:16:12誤爆
0057デフォルトの名無しさん
2010/12/07(火) 17:18:52レジスタ
| ↑
L1 |コーヒレント
| ↓
共有メモリ
| ↑
L2 |コーヒレントじゃない
| ↓
グローバルメモリ
この資料に書いてあるとかいう情報があるなら
それも教えてもらえると助かります
0058デフォルトの名無しさん
2010/12/07(火) 17:19:11使っていますがRubyの実装でC言語にはしてないです。
このプログラムは疑似乱数生成器で基本的には最大32*32の行列演算を
主に行っています。これが並列化できれば32回の計算で256ビットの
乱数を生成できるので効率的になります。
というかなればいいなあと思ってます。
0059デフォルトの名無しさん
2010/12/07(火) 18:40:08int(4)なら加算だけでもfloatの4倍かかるぞ
FMAとかも使えないし
0060デフォルトの名無しさん
2010/12/07(火) 18:42:12http://sky.geocities.jp/tcshacina/grs1.c
0061デフォルトの名無しさん
2010/12/07(火) 18:45:53じゃあグラフィックボードで並列計算するのはこの場合逆効果になる
のでしょうか?買わないほうがいいですか?
0062デフォルトの名無しさん
2010/12/07(火) 18:50:41intにこだわる必要あるのか?
0063デフォルトの名無しさん
2010/12/07(火) 18:58:39計算しておいて後から一斉にunsigned charに変換することが可能で
しょうか?暗号に使う疑似乱数なので整数で処理しないといけない
という条件があります。
0064デフォルトの名無しさん
2010/12/07(火) 19:03:47あれだけコードがかけるならいいおもちゃだから買えばいいんじゃないですか。
Cellも整数の論理演算にむいていた気がします。6個のタスクを生成できる
ので使いやすい場合もありました。
0065デフォルトの名無しさん
2010/12/07(火) 19:05:17xorshiftのこと?
であればCURANDってのがあるからそれ使えば乱数生成可?
メルセンヌもSDKについてる
0066デフォルトの名無しさん
2010/12/07(火) 19:14:48という疑問です。できなければfloatのままデータを生成して、後で
CPUにまとめてfloat->intの処理をさせながらXORを計算するという
方法ではうまくいかないでしょうか?マニュアルには論理演算も
可能と書いてありますけど、整数演算の速度が浮動小数点より遅い
というのが気がかりです。
Rubyで作った暗号化モジュールいろいろ
http://sky.geocities.jp/tcshacina/crypt.html
0067デフォルトの名無しさん
2010/12/07(火) 19:24:500068デフォルトの名無しさん
2010/12/07(火) 19:29:450069デフォルトの名無しさん
2010/12/07(火) 19:32:36CURANDのマニュアル読んで目的にかなうかどうか確認し
0070デフォルトの名無しさん
2010/12/07(火) 19:34:29どうやってそれを使えばいいのですか?ソースは晒してあるので
自由に見て感想を聞かせてください。
0071デフォルトの名無しさん
2010/12/07(火) 19:37:090072デフォルトの名無しさん
2010/12/07(火) 19:40:050073デフォルトの名無しさん
2010/12/07(火) 20:49:540074デフォルトの名無しさん
2010/12/07(火) 21:00:10それを考えるのがプログラマだろうが!1から10まで聞いてないで自分でやってみろ
0075デフォルトの名無しさん
2010/12/07(火) 22:12:310076デフォルトの名無しさん
2010/12/07(火) 23:07:110077デフォルトの名無しさん
2010/12/07(火) 23:28:12いくら超低レベルスレといっても
"ちょっと調べて分かること"を次々質問するのはどうなんだ?
答えるは簡単だがこの程度のことを人に頼る時点で
そいつにまともなプログラムが作れると思えん
0078デフォルトの名無しさん
2010/12/07(火) 23:48:44同感。
そういうやつでも数学的素養が優れていたりして特定のモデルの構築とかには
優れた才能を発揮する奴がいるから、そういう所だけは認めざるを得ないこともある
のが悔しい。
0079デフォルトの名無しさん
2010/12/07(火) 23:52:59プログラミングの方法より使い方を話題にしたいなと
GPGPUでシミュレーションを高速化してみる
http://kamome.2ch.net/test/read.cgi/sim/1291687353/
0080デフォルトの名無しさん
2010/12/08(水) 01:41:45Micro-benchmarking the GT200 GPU
http://www.eecg.toronto.edu/~moshovos/CUDA08/arx/microbenchmark_report.pdf
GT200ではint加算、減算、論理演算、シフトのレイテンシ、スループットはfloat積和算と同じ。
CPUと違ってレイテンシが24もあるが、並列度が十分にあれば普通に使える。
Fermiでの資料を探したが見つからなかった。誰か頼む。
0081デフォルトの名無しさん
2010/12/08(水) 05:30:08Int add, shift, min, max: 4 cycles (Tesla), 2 cycles (Fermi)
Int mul: 16 cycles (Tesla), 2 cycles (Fermi)
FP32 add, mul, mad, min, max: 4 cycles (Tesla), 2 cycles (Fermi)
FP64 add, mul, mad: 32 cycles (Tesla), 2 cycles (Fermi)
Int divide and modulo are expensive
Divide by 2^n, use “>> n”
Modulo 2^n, use “& (2^n ? 1)”
Avoid automatic conversion of double to float
Adding “f” to floating literals (e.g. 1.0f) because the default is double
と、あった
0082デフォルトの名無しさん
2010/12/08(水) 07:54:020083デフォルトの名無しさん
2010/12/09(木) 17:13:20SDKのサンプルtempleteが2008で開けない。
カスタム ビルド規則 ファイル 'C:\Program Files\Microsoft Visual Studio 9.0\VC\VCProjectDefaults\NvCudaRuntimeApi.rules' が見つからないか、または読み込めませんでした。
って出たからrulesファイル探してるんだけどないんだよな…
SDK再インストールしたけどやっぱりない。たすけて
0084デフォルトの名無しさん
2010/12/10(金) 11:21:46複数のGPUに並列でカーネルを実行させることが出来ますか?
0085デフォルトの名無しさん
2010/12/10(金) 12:16:29008684
2010/12/10(金) 13:49:23無理ですか…
シミュレーションを複数のGPUで実行して,
結果をOpenGLで可視化したいと思っています(出力は一つから).
ランタイムAPIとOpenMPを組み合わせて作ろうと思っていました.
しかし,その場合malloc等の前処理の段階で
複数のホストスレッドでcudaSetDeviceを使ってしまうと,
いざカーネルを実行しようとしたときに,マスターホストスレッド以外のホストスレッドの
cudaSetDeviceで"もう既にセット済"というエラーが出てしまいます.
OpenGLのイベント待ちのために,
複数スレッド→マスタスレッドのみ→複数スレッド→マスタスレッドのみ→…
という状態を繰り返すのですが,1スレッドに戻っても
マスタスレッド以外で設定したコンテキストが消えないのが原因のようです.
マニュアルのドライバAPIのコンテキストの部分を読んでいたら,
色々命令があるので出来るかもしれないと思い質問させていただきました.
ホストスレッドの数は複数でも問題ないので,
コンテキストの操作に詳しい方いらっしゃいますか?
0087デフォルトの名無しさん
2010/12/10(金) 15:36:25良く見直してみ
0088デフォルトの名無しさん
2010/12/10(金) 17:16:41GlobalからのReadキャッシュをL1キャッシュにおいて有効にしたときと無効にしたときにL1キャッシュが無効のほうが早い場合があり、
Visual Profilerで調べてみるとinstruction issueが93152(L1有効)と67666(L1無効)
で差があり、replayを含まないinstruction executed が54592(有効無効どちらも同じ値)
ですのでreplayに問題があると思ったのですが、どういう状況下でreplayが起こるのかいまいちわかりません。
replayはコアレッシングできていないreadのときにおこると思っていたのですが、間違っているなら教えてください。
0089デフォルトの名無しさん
2010/12/10(金) 23:41:08ビルドルールのemulation modeにて、はい (-deviceemu -D_DEVICEEMU)を選択しても、
下記エラー分が消えません。
calling a host function from a __device__/__global__ function is not allowed
どうすればよいのでしょう?
CUDAのバージョンは3.1です。
0090デフォルトの名無しさん
2010/12/10(金) 23:51:250091デフォルトの名無しさん
2010/12/10(金) 23:55:04calling a host function from a __device__/__global__ function is not allowed
するとやはりエラーが消えないのですが・・・
0092デフォルトの名無しさん
2010/12/11(土) 00:13:47おまえ言ってること矛盾してるし、何をやってそういうエラーがでたかの説明してなくて
アドバイスを求めても無理。
vbにprintfとかないよね?どういうことなの?
0093デフォルトの名無しさん
2010/12/11(土) 00:14:07Fermi系列なら確かそのまんまprintf使えるけど
windowsならVisual Nsightとかいうのでデバッグできたはず
あとはどうしてもエミュレーション使いたいならcuda2.3使うとか
0094デフォルトの名無しさん
2010/12/11(土) 00:16:07printfは確かfermi系だけのサポートのはず。
そして2.3はfermiサポートしてない。3.0からだったはず。
0095デフォルトの名無しさん
2010/12/11(土) 00:19:120096デフォルトの名無しさん
2010/12/11(土) 00:24:18同様の現象が起きていて困ってる
GTX480ではinstruction issueが取れなかったんだが、Tesla使い?
0098デフォルトの名無しさん
2010/12/11(土) 18:03:20探しています。
フロートの性能は欲しいがC2050は高すぎるということで。
0099デフォルトの名無しさん
2010/12/11(土) 21:19:193DVantageとかの260と580の性能差を4倍したものを求めて見積もればいいんじゃない?
0100デフォルトの名無しさん
2010/12/11(土) 22:44:360101デフォルトの名無しさん
2010/12/11(土) 23:25:140102デフォルトの名無しさん
2010/12/12(日) 00:01:33とりあえずsoftekでググると何か出てくるかもしれんよ
0103デフォルトの名無しさん
2010/12/12(日) 00:58:43アプリ?はR+CUDAかMATLABで多体シミュレーションです。
けどRって言語だよな。
0104デフォルトの名無しさん
2010/12/12(日) 02:22:330105デフォルトの名無しさん
2010/12/12(日) 13:53:06誰かfermi系で128台ぐらいのGPUクラスター作ってくんねーかな
0106デフォルトの名無しさん
2010/12/12(日) 13:57:080107デフォルトの名無しさん
2010/12/12(日) 14:12:03ttp://tsubame.gsic.titech.ac.jp/ja/hardware-architecture
東工大にでも逝けや
0108デフォルトの名無しさん
2010/12/12(日) 14:35:17その条件でfloatの性能求めるならたいして変わらないんじゃないかな
doubleなら別だけど
0109デフォルトの名無しさん
2010/12/12(日) 14:47:34メモリの転送が非同期で双方向通信できるとか
コンカレントカーネル実行ができるとか、自動キャッシュがあるとか
のおかげで最適化(上記の機能を使うように修正)すればかなり早くなると思う。
修正なしでそのままならちょっと〜けっこう早くなると思う。
でもR+CUDAを知らないので無責任な推測です。
0110デフォルトの名無しさん
2010/12/12(日) 16:01:53最適化してしまえばそんなに性能変わらないんじゃね。
0111デフォルトの名無しさん
2010/12/12(日) 16:45:01C1060に最適化したものをそのままFermiで動かしても差はそんなにないだろうって事だと思うけど
それでもコアレスの制限が緩和された影響と自動キャッシュのおかげで結構速くなると思う。
さらにFermiのストリームとかを使用してデータ転送の最適化を行えば段違いになると思う。
特に4枚挿してるということからx8*4の転送速度になって転送がボトルネックになりそうだから
そこの最適化は効いてくるんじゃないかと思ってる
ところで103のMBはガルフ対応だからPCI-e2.0対応だよね?
0112デフォルトの名無しさん
2010/12/12(日) 16:55:360113デフォルトの名無しさん
2010/12/12(日) 17:16:590114デフォルトの名無しさん
2010/12/12(日) 18:35:23ストリームはCC依存でtoolkit依存のようだが?
0115デフォルトの名無しさん
2010/12/12(日) 18:36:43ストリームはCC依存ではなくtoolkit依存のようだが?
0116デフォルトの名無しさん
2010/12/12(日) 20:48:29多体問題って演算バウンドだったんじゃないの確か
GTX580に交換するうまみはそれほど無い気がする
0117デフォルトの名無しさん
2010/12/12(日) 21:13:38floatなら計算パワーだけで考えると大して性能向上してなかったの?
floatだけならメモリ周りのボトルネックを解消して全体的にFlopsが上がりやすくなりましたって
だけだったのか・・・・
0118デフォルトの名無しさん
2010/12/12(日) 21:34:3820万+なら自分の財布でやってしまえ。
ちなみにMathematica8もCUDA対応で年始位にデリバリーとか何とか。
>>104
>MATLAB 2010bのGPUArray
R+CUDAの方に感けて,まだやってませんでした。すみません。
0119デフォルトの名無しさん
2010/12/12(日) 21:43:37C1060 443.4GFLOPS
580 447.154 single-precision GFLOP/s at 20 flops per interaction
580は過去スレからc1060はググった結果
なにこれ!?
0120デフォルトの名無しさん
2010/12/12(日) 21:46:090121デフォルトの名無しさん
2010/12/12(日) 21:46:5116384 bodies, total time for 10 iterations: 69.775 ms
= 38.472 billion interactions per second
= 769.433 single-precision GFLOP/s at 20 flops per interaction
0122デフォルトの名無しさん
2010/12/12(日) 21:58:16GTX580結構がんばるんだな…
年末のセールでELSA一枚買ってみようかな
↓480での結果(過去スレから引用)
[単精度]
15360 bodies, total time for 10 iterations: 73.826 ms
= 31.957 billion interactions per second
= 639.149 single-precision GFLOP/s at 20 flops per interaction
0123デフォルトの名無しさん
2010/12/12(日) 22:49:40自分でコード書いたほうが速いし確実だと思うけどね
0124デフォルトの名無しさん
2010/12/12(日) 22:52:47GeForceをごにょごにょしたらdoubleが4倍速になったりしないかな。
0125デフォルトの名無しさん
2010/12/12(日) 22:54:13あきらめろん
0126デフォルトの名無しさん
2010/12/12(日) 22:55:44行列演算で、CUBLASより速いコードかける?
0127デフォルトの名無しさん
2010/12/12(日) 23:09:540128デフォルトの名無しさん
2010/12/13(月) 22:52:09__constant__ int i;
int j=1;とした場合(コンスタントメモリに配列以外を置きたい場合)になぜ
cudaMemcpyToSymbol("i",&j,sizeof(int),0,cudaMemcpyHostToDevice);
と文字列のようにして渡すのでしょうか
また、外部ファイルにextern __const__ int i;
と書くと定義が二重ですと言われます
複数ファイルにかかれたカーネルからコンスタントメモリを読み込みたいのですがどのようにすればよいのでしょうか
お願いいたします
0129デフォルトの名無しさん
2010/12/13(月) 23:24:44うーん、簡単に言えば、cudaMalloc()でVRAM上のアドレスを得ることができればcudaMemcpy()できる。
一方__constant__の場合はそういうわけにはいかないので、
nvcc独自のちょっとえげつないcudaMemcopyToSymbol()を使わないといけない。
コンパイル単位の分かれたカーネル同士で共有するのは難しいかもねぇ。
0130デフォルトの名無しさん
2010/12/13(月) 23:42:11> R+CUDAの方に感けて,まだやってませんでした。すみません。
R+CUDAってどういう構成なのですが?
やはり、
R+gputools+CUDA+CULA
なのですか?
それとも、Rのビルド時にCUBLASをリンクするのですか?
当方の環境は、
Q9450 + Memory 6GB + GeForce9600GT
Ubuntsu10.04(64bit)
です。gotoBLASで行列演算時にコア4つ使用するところまではやったのですが、
CUDA使うともっと凄くなるのですか?
0131デフォルトの名無しさん
2010/12/13(月) 23:45:440132130
2010/12/14(火) 23:05:18>>118さんがどういう構成でやっているのか知りたいです。
未だ、コンパイルも通らない実行も出来ていないレベルの者です。
RにgotoBLAS組み込んでも、乱数の生成速度に差が出ませんでした。
BLASは乱数を生成するライブラリではないから当然ですが。
CUDAのページのサンプルに
・MersenneTwister
・Monte Carlo Option Pricing with multi-GPU support
などとあるのをみると期待に胸がふくらみます。
0133デフォルトの名無しさん
2010/12/15(水) 21:41:190134デフォルトの名無しさん
2010/12/15(水) 21:44:27BLASがいくら早くても関係ないじゃん、相関には
0135デフォルトの名無しさん
2010/12/15(水) 21:54:23O(n^2)かかるところを32*8=256ビットの乱数を生成させるのに32個の
コアで16回の計算で得られるので効率的になる。
0136デフォルトの名無しさん
2010/12/16(木) 00:26:210137デフォルトの名無しさん
2010/12/16(木) 00:29:550138デフォルトの名無しさん
2010/12/16(木) 07:15:11釣られた?
0139デフォルトの名無しさん
2010/12/16(木) 11:51:04Compute Shaderスレないですか?
0140デフォルトの名無しさん
2010/12/16(木) 15:56:270141デフォルトの名無しさん
2010/12/16(木) 17:50:04ものすごい遅いレスですみませんが、driverAPIなら可能です。
0142デフォルトの名無しさん
2010/12/16(木) 19:24:21BLASが乱数を生成するライブラリでないと言ったら、
後藤さんの名前を汚していることになるのですか?
0143デフォルトの名無しさん
2010/12/16(木) 23:05:19みんななにに使ってるの?
0144デフォルトの名無しさん
2010/12/17(金) 09:23:300145デフォルトの名無しさん
2010/12/17(金) 22:12:472月くらいにはなんとかならないだろか。
0146デフォルトの名無しさん
2010/12/17(金) 22:51:22前者はバージョン2.3
後者はバージョン3.1
が入っているのですが、toolkitが3.xになってからプロファイラの名称が変わったというだけでしょうか?
0147デフォルトの名無しさん
2010/12/17(金) 23:04:450148デフォルトの名無しさん
2010/12/17(金) 23:18:01ウッカリ長時間のCUDAジョブを走らせてしまったとき
中断はどすればよいですか?
0149146
2010/12/17(金) 23:34:33ありがとうございます
ものの本やサイトには/usr/local/cuda/cudaprof/のcudaprofiler使えって書いてあったのですが
変なWarning出て困っていたところでした
0150146
2010/12/17(金) 23:36:32ありがとうございます
ものの本やサイトには/usr/local/cuda/cudaprof/のcudaprofiler使えって書いてあったのですが
変なWarning出て困っていたところでした
0151デフォルトの名無しさん
2010/12/18(土) 00:03:29GPUはずしちゃえよ
0152デフォルトの名無しさん
2010/12/18(土) 00:34:360153デフォルトの名無しさん
2010/12/18(土) 00:44:26九十九ネットストアにあったけど。あと価格コムとかで在庫もってる店あたってみたらいいんじゃない?
0154デフォルトの名無しさん
2010/12/18(土) 00:47:360155デフォルトの名無しさん
2010/12/18(土) 15:04:40最初の実行に時間がかかるのと
同様に最初の実行時にグリッド数、スレッド数が大きめだたとmemcpyでエラーが出る
小さめのグリッド数、スレッド数で最初に実行した後だと
大きめのグリッド数、スレッド数でも正常するんだが…
0156デフォルトの名無しさん
2010/12/18(土) 15:16:14ドライバのバージョンとSDK、TOOLキット、OSの情報も載せないと駄目じゃね?
あとgccか。
0157デフォルトの名無しさん
2010/12/18(土) 15:28:07ドライバ、SDK、ToolKitは最新
gccは4.1.2
今のところ原因は全く不明
ドライバのバージョンを一個前にすると初回実行時の遅延がちょっと早くなる(最新が約2.5s、一個前が約1.8s)
gtx470と465で試してどちらも同じ症状
0158130
2010/12/18(土) 16:44:30速いのは分かったけど、
行列の規模が、本体のメモリ容量と関係なく、ビデオカードのメモリ容量に制約される。
キャッシュイーター向けじゃなかったですね。
0159デフォルトの名無しさん
2010/12/18(土) 17:43:36行列の計算ってのはね、分割できるんですよ。
0160デフォルトの名無しさん
2010/12/18(土) 18:13:31確かに。
分割できなければコアが複数あっても意味がない。
0161デフォルトの名無しさん
2010/12/18(土) 18:55:39結論C2070を買え。
0162130
2010/12/19(日) 18:30:10一般論として言えるのかどうか知らないけど、
行列が(計算結果を含めて)ビデオカードのメモリにスッポリ入るなら劇的な高速化が見込めるけど、
そうでなければ、マルチコアCPU+本体メモリのほうが速度でもメモリ消費量でも有利なようです。
0163128
2010/12/19(日) 20:32:38__const__宣言時とは考えにくいので最初にcudaMemcpyToSymbolしたときでしょうか
デバイスにあるということはコンスタントメモリも終了時に解放したほうがいいと思うのですが、
cudaFree()では解放できないようです
コンスタントメモリはどのように解放すればいいのでしょうか。
0164デフォルトの名無しさん
2010/12/19(日) 21:03:47リンパックベンチでGPU搭載マシーンが上位に来てるんだし
0165デフォルトの名無しさん
2010/12/19(日) 23:49:30SDKのサンプルみればいいんじゃないか?
0166デフォルトの名無しさん
2010/12/20(月) 00:43:41kernelコードだって似たようなものなんじゃ。
0167デフォルトの名無しさん
2010/12/20(月) 11:19:24定数メモリはcudaMalloc()のように「確保する」ということなしに使うようになっています。
その為、cudaMemcpyToSymbol()のような特殊な関数でしか転送できません。
また、「確保する」ことなしに使える以上、「解放する」手続きは用意されていません。
0168デフォルトの名無しさん
2010/12/21(火) 15:33:13CUDA向きの安定ソートってありますか?
16000要素くらいです
0169デフォルトの名無しさん
2010/12/21(火) 15:52:440170デフォルトの名無しさん
2010/12/21(火) 15:54:210171168
2010/12/21(火) 16:18:22マージソートのサンプルがSDKにあったのですね。
チェックしてみます
0172デフォルトの名無しさん
2010/12/21(火) 20:58:05やれやれ.
0173デフォルトの名無しさん
2010/12/21(火) 21:12:11何Wで駄目だった?
0174デフォルトの名無しさん
2010/12/21(火) 21:46:211400W。ディスク系を外においてもなぁ。
0175デフォルトの名無しさん
2010/12/21(火) 21:49:200176デフォルトの名無しさん
2010/12/21(火) 21:59:470177デフォルトの名無しさん
2010/12/21(火) 22:17:39いったことをしないと無理そうだな。
0178デフォルトの名無しさん
2010/12/21(火) 23:19:450179デフォルトの名無しさん
2010/12/22(水) 01:04:00>>172が日本在住でないなら話は別だが。
0180デフォルトの名無しさん
2010/12/22(水) 07:14:57あれは110V/1400W。80GOLDのリタンダント。
日本じゃ実効1100W。
0181デフォルトの名無しさん
2010/12/22(水) 22:34:500182デフォルトの名無しさん
2010/12/22(水) 23:25:26なら、「1100W電源」として売るべきだろ。
排気量2000ccの車を3000ccだと言って売る自動車屋などありえない。
そのあり得ないことを長年やってきたのが、PC電源代理業者だ。
0183デフォルトの名無しさん
2010/12/23(木) 00:59:59supermicroは日本向けのサイトを持っていない。仕方あるまい。
ちなみに法人で入れたときはこの説明はあった。
www.supermicro.com/products/system/4U/7046/SYS-7046GT-TRF.cfm
0184デフォルトの名無しさん
2010/12/24(金) 20:02:02例えばGTX480は
Fermi-ComputeCapability2.0-GF100-GTX480
でよいのでしょうか
それともアーキテクチャとCCは1対1対応しているのでしょうか
また、Fermiの分類位置はここでよいのでしょうか
0185デフォルトの名無しさん
2010/12/24(金) 22:08:13は正しいけど、FermiはC.C.2.X全部を含む
GFの"F"はFermi、GTXの"T"はなぜかTeslaだったはず
0186デフォルトの名無しさん
2010/12/24(金) 22:20:210187デフォルトの名無しさん
2010/12/24(金) 23:41:29CUDA目的ならComputeCapabilityだけ意識すればよいかと。
あとはパフォーマンスが違うだけだから。
0188デフォルトの名無しさん
2010/12/24(金) 23:43:40GeForce 7800GTXさんはCUDA対象外じゃ・・・・・・
0189デフォルトの名無しさん
2010/12/25(土) 00:54:10toolktのインストールは必須ですか?
LD_LIBRARY_PATH /usr/local/cuda/lib64
を通さなくてはならないということはtoolkit必須なのかな
0190デフォルトの名無しさん
2010/12/25(土) 01:00:31そのdllをexeと同じフォルダにおけばいいんじゃないか?
きっとlib64内のなんかのライブラリつかってんでしょ。cutilとかかな
0191デフォルトの名無しさん
2010/12/25(土) 02:39:520192デフォルトの名無しさん
2010/12/25(土) 09:47:13もちろんツールキットのライブラリも必要。
0193デフォルトの名無しさん
2010/12/27(月) 19:25:09CentOS 5.5(64bit)でCUDAを導入したくて
・Developer Drivers for Linux (260.19.26)
・CUDA Toolkit for RedHat Enterprise Linux 5.5 (64bit版)
をインストールし、とりあえずSDK code samplesをコンパイルして動くか確かめたのですが
コンパイル後のファイルを実行しようとすると
"cannot execute binary file"のメッセージが出てきてプログラムが実行できません。
この問題を解決したいのですがどこを確認してどう改善したらよいのでしょうか?
ちなみにカーネルのバージョンは2.6.9-55.0.2.ELsmpです。
0194デフォルトの名無しさん
2010/12/27(月) 19:40:290195デフォルトの名無しさん
2010/12/27(月) 20:29:330196デフォルトの名無しさん
2010/12/27(月) 22:11:09例えば、ディストリビューションのnvidiaドライバなんかが残っていない?
その上に、最新ドライバをインストールしていて、
コンパイル時と実行時に別の共有ライブラリを参照していたりとか。
ていうか、自分も、ubuntuでx-updatesのドライバと最新ドライバが両方入っていて挙動不審。
ubuntuでもOS再インストールかw
0197193
2010/12/28(火) 00:14:59返信ありがとうございます。
ちょっとその辺を調べてみます。
後日、問題が解決したらフィードバックします。
0198デフォルトの名無しさん
2010/12/28(火) 07:48:230199デフォルトの名無しさん
2010/12/28(火) 16:15:180200デフォルトの名無しさん
2010/12/28(火) 22:47:32スレッド数が192を超えたあたりで処理が行われていません。
スレッド数が512より小さくても処理が行われないことがあるのでしょうか?
グラボは9800GTで11レジスタ、共有メモリを36byte使用してます。
0201デフォルトの名無しさん
2010/12/28(火) 22:48:39「処理が行なわれない」ってどういう意味さ。
0202デフォルトの名無しさん
2010/12/28(火) 23:01:14説明不足でした。
処理を1000回行い、平均処理時間を計測しているのですが、
スレッド数が192を超えると0(ms)になってしまいます。
0203デフォルトの名無しさん
2010/12/28(火) 23:19:140204デフォルトの名無しさん
2010/12/29(水) 00:34:350205デフォルトの名無しさん
2010/12/29(水) 00:38:37エラーが発生するような状況になっていないかどうかを確認・・・・・かな。
これが起きているとそれ以降の(タイマーを含めた)全ての処理がエラーになる。
0206デフォルトの名無しさん
2010/12/29(水) 01:31:39cudaThreadSynchronize();
cudaGetLastError();
0207デフォルトの名無しさん
2010/12/29(水) 14:09:27助言ありがとうございます。
まずはエラーを確認してみたいと思います。
0208デフォルトの名無しさん
2010/12/29(水) 19:33:12カーネルの引数を構造体のポインタにまとめると解決する。
0209デフォルトの名無しさん
2010/12/30(木) 20:27:57lspci | grep -i nVidia
でPCIバスデバイス情報確認したら
VGA compatible controller: nVidia Corporation: Unknown device …
ってコンピューターさんがおっしゃられたんだが…
0210デフォルトの名無しさん
2010/12/31(金) 00:21:26/sbin/update-pciids
してからコンピュータさんにもう一度お伺いたててご覧なさい
0211デフォルトの名無しさん
2010/12/31(金) 05:27:050212デフォルトの名無しさん
2011/01/01(土) 11:59:190213デフォルトの名無しさん
2011/01/04(火) 09:25:33GeforceをCUDA専用、RADEONをグラフィック専用とすることは出来ますか?
0214デフォルトの名無しさん
2011/01/04(火) 09:34:04GeforceとRADEONの2枚を挿して
GeforceをCUDA専用、RADEONをグラフィック専用とすることは出来ますか?
0215デフォルトの名無しさん
2011/01/04(火) 09:41:440216デフォルトの名無しさん
2011/01/04(火) 10:36:530217デフォルトの名無しさん
2011/01/04(火) 10:52:31cudaSetDeviceFlags(cudaDeviceBlockingSync);
0218デフォルトの名無しさん
2011/01/04(火) 10:53:37cudaSetDeviceFlags(cudaDeviceBlockingSync);
が有効になるプログラムとならないプログラムがあるのですがどこを
調べれば良いでしょうか。
0219デフォルトの名無しさん
2011/01/04(火) 12:31:38そして「有効」かどうかをCPU負荷で判定しているのかどうかも。
cudaSetDevice()を除くすべてのCUDA関数より先にcudaSetDeviceFlags()を
呼び出しているかどうか以外には特に調べるべきところはありません。
0220デフォルトの名無しさん
2011/01/05(水) 18:57:39CUDA_C_Best_Practices_Guide.pdf
CUDA_C_Programming_Guide.pdf
Introductory CUDA Technical Training Courses VolumeI.pdf
0221デフォルトの名無しさん
2011/01/05(水) 20:31:24ただこれ内容が古かったと思うからどんな仕組みで動いてるかを把握するくらいでいいと思う
必要になったらまた読み返せばいい。
次にCUDA_C_Programming_Guide.pdfで基本を学ぶ。
これは読み返すことが多くなると思う。
その次にCUDA_C_Best_Practices_Guide.pdfで良い実装の仕方を学ぶ。
これは実際にプログラム組みながらでいいと思う。
あと余談だがfermi tuning guideも読むといい。
英語が苦手なら「はじめてのCUDAプログラミング」や「CUDAプログラミング実践講座 - 超並列プロセッサにおけるプログラミング手法」
を読んでから必要なところをPDFで追加で読むと理解が早いと思う。
0222デフォルトの名無しさん
2011/01/05(水) 21:02:090223デフォルトの名無しさん
2011/01/05(水) 21:15:200224デフォルトの名無しさん
2011/01/05(水) 21:16:37残念ながらありません
0225デフォルトの名無しさん
2011/01/06(木) 00:39:2510くらい増えるのですが、なぜですか?
0226デフォルトの名無しさん
2011/01/06(木) 22:50:400227デフォルトの名無しさん
2011/01/08(土) 18:48:55ローカルメモリを使うのが最良なのでしょうか?
シェアードメモリを使うと、同時に立ち上げられるスレッド数が減って
逆に遅くなりました。
0228デフォルトの名無しさん
2011/01/09(日) 21:38:49そういうのはプログラムによって変わるだろ
とりあえずptxファイル嫁
>>227
>各スレッドで数百バイトのワーキングエリアが必要な場合
これが今以上減らせないならそもそもCUDAに向いてない
0229デフォルトの名無しさん
2011/01/10(月) 17:57:08L2キャッシュは大体何clockかかりますか?
0230デフォルトの名無しさん
2011/01/10(月) 19:32:12L2は大体150
0231デフォルトの名無しさん
2011/01/11(火) 09:01:340232デフォルトの名無しさん
2011/01/13(木) 11:27:07nvccはcuでないと、cudaコードとしてコンパイルしてくれないのですが、
拡張子がcppのファイルをcudaソースファイルとしてコンパイルするオプションはありませんか?
0233デフォルトの名無しさん
2011/01/13(木) 19:39:510234デフォルトの名無しさん
2011/01/13(木) 20:42:06亀レスですが,Driver APIを用いて出来ました.
ありがとうございました.
0235デフォルトの名無しさん
2011/01/14(金) 20:47:45概略をのせてくれまいか?
cudaAAA()で○○を設定
cudaBBB()でXを取得して
pthread_AAA()を使ってXXXする
くらいの事でもいいから
0236234
2011/01/15(土) 13:34:59簡単にで良いなら.
// 前処理
#pragma omp parallel
{
cuInit(); // 初期化
cuDeviceGet(); // GPUへのハンドルを取得
cuCtxCreate(); // コンテキストを作成し,ホストスレッドにバインド
mallocとか色々
cuCtxPopCurrent(); // コンテキストをアンバインド
}
glutMainLoop();
// カーネル実行
#pragma omp parallel
{
cuCtxPushCurrent(); // バインド
カーネル実行
cuCtxPopCurrent(); // アンバインド
}
関数の引数は省略しました.
0237デフォルトの名無しさん
2011/01/15(土) 17:57:20ありがとう。参考にさせてもらいます
0238デフォルトの名無しさん
2011/01/15(土) 21:30:19キャッシュはおまけ程度って考えでおk?
ていうか各スレッドが違うグローバルメモリにアクセスするなら同じカーネル内ではキャッシュ効かないよね
0239デフォルトの名無しさん
2011/01/15(土) 22:14:42キャッシュ汚染はひどいからな
マルチスレッドとキャッシュは相性が非常に悪い
0240デフォルトの名無しさん
2011/01/15(土) 23:23:54速度が10倍は違う。本当に使いにくいよ。
メモリアクセスがネックの(単純な)アプリで、shared memoryの使いまわしのできないアプリでは、
マルチスレッドを上回るのは難しい。
0241デフォルトの名無しさん
2011/01/15(土) 23:49:58FermiならCPUの2倍くらい出るってことは結構ある。
非Fermiならローカルメモリ使ったら即使い物にならない速度だったけど、
Fermiならキャッシュがなんとかしてくれたり。
0242デフォルトの名無しさん
2011/01/16(日) 18:58:37betaを(double)0に指定しても、
内部では真面目に係数をかけているのでしょうか?
参考(PDF注意)
http://developer.download.nvidia.com/compute/cuda/3_2_prod/toolkit/docs/CUBLAS_Library.pdf
0243デフォルトの名無しさん
2011/01/16(日) 19:01:420244デフォルトの名無しさん
2011/01/16(日) 19:13:170245デフォルトの名無しさん
2011/01/16(日) 21:42:18ありがとうございます。
FORTRANのコードを見る限り、ちゃんと係数の値によって
処理変えているようでした。
>>244
ありがとうございます。
実際に測定したのですが、係数の値を変えても
計算時間に明確な差は現れませんでした。
GTX480 1枚, CUDA 3.2, Windows 7 64bit
A=7500x8000の行列,B=8000x7500,C=7500x7500の行列で
C=alpha * A*B + beta * Cを計算
alpha = 1, beta = 0を15回で85.079秒
alpha = 10, beta = 10を15回で85.072秒
FORTRANのコードのような実装になっていない気がするのですが。
0246デフォルトの名無しさん
2011/01/16(日) 21:45:25をひたすらマイナーチェンジしていくのが一番早いと思う。
スクラッチはマイナーチェンジの経験を積んでからで十分。
0247デフォルトの名無しさん
2011/01/17(月) 05:28:45GPUがPTX読むみたいな事が書いてあるんだけど
0248デフォルトの名無しさん
2011/01/17(月) 06:02:220249デフォルトの名無しさん
2011/01/17(月) 07:27:520250デフォルトの名無しさん
2011/01/17(月) 10:47:310251デフォルトの名無しさん
2011/01/17(月) 20:15:090252デフォルトの名無しさん
2011/01/19(水) 09:33:50どう頑張ってもnvccでコンパイルしたサンプルファイルが実行できませーん
って開発以前の問題で精神折れそうだわ…
0253デフォルトの名無しさん
2011/01/19(水) 13:17:05#include <stdio.h>
#include <stdlib.h>
#include <string.h>
struct record { /* 1人分のデータ */
char name[32]; /* 氏名 */
int score; /* 得点 */
struct record *next; /* 次のデータへのポインタ */
};
void printlist( struct record *list ) {
if ( list != NULL ) { /* リストが空でなければ */
printf( "%16s:%4d\n", list->name, list->score );
printlist( list->next ); /* 先頭の次からリストの表示 */
}
}
0254デフォルトの名無しさん
2011/01/19(水) 13:59:28実行できないならできないなりにエラーメッセージが帰ってくるのではないでしょうか
CUDAコアとはスカラプロセッサのことでいいのでしょうか
FermiのみCUDAコアという呼び方をするのですか?
0255デフォルトの名無しさん
2011/01/19(水) 14:09:09実行できないならできないなりにエラーメッセージが帰ってくるのではないでしょうか
CUDAコアとはスカラプロセッサのことでいいのでしょうか
FermiのみCUDAコアという呼び方をするのですか?
0256デフォルトの名無しさん
2011/01/19(水) 21:03:35SFUを使えないのでしょうか
倍精度浮動小数点の計算は単精度浮動小数点演算機が2クロックで計算するのですよね
同じようにSFUが単精度計算分の2週でやってくれるということはないのでしょうか
0257デフォルトの名無しさん
2011/01/19(水) 21:10:310258デフォルトの名無しさん
2011/01/19(水) 21:57:030259デフォルトの名無しさん
2011/01/19(水) 22:20:50fast mathを使えば?
0260デフォルトの名無しさん
2011/01/20(木) 00:39:31そもそも倍精度使ってるくせに精度低いSFU使いたいって時点で
本当に倍精度が必要か考えるべき。
0261256
2011/01/20(木) 02:09:04そんなに精度が低いのかー
0262デフォルトの名無しさん
2011/01/20(木) 02:55:130263デフォルトの名無しさん
2011/01/21(金) 02:52:36元々あるソースにCUDAを取り入れてるんですが構造体を丸々GPU側にコピーはできないんでしょうか?
できるのであれば誰かご教授お願いします。
0264デフォルトの名無しさん
2011/01/21(金) 03:12:270265デフォルトの名無しさん
2011/01/21(金) 03:30:01一応いろいろ試したんですが
Warning : Cannot tell what pointer points to, assuming global memory space
ていう警告が出ていておそらくデータの転送ができていないんです・・・。
普通のC言語を勉強すればできるんですかね・・・?
0266デフォルトの名無しさん
2011/01/21(金) 03:40:270267デフォルトの名無しさん
2011/01/21(金) 04:09:34ポインタのほかにもポインタのポインタも入っているのでそれもGPU側に送らないといけないいうことであってますか?
そうすると結構量があるので必要な部分のみを送るほうが効率がいいんですかね・・・?
0268デフォルトの名無しさん
2011/01/21(金) 05:42:58構造体を送る前に,構造体の中で指してるアドレスをGPU側のアドレスに
書き変えないといけない.
0269デフォルトの名無しさん
2011/01/21(金) 06:51:09なんとなくですが皆さんのおかげで理解できました。
がんばってやってみます。
0270デフォルトの名無しさん
2011/01/21(金) 07:06:460271デフォルトの名無しさん
2011/01/24(月) 01:53:291>ptxas C:/Temp/tmpxft_00000fac_00000000-4_cuda.ptx, line 146; error : State space incorrect for instruction 'st'
1>ptxas C:/Temp/tmpxft_00000fac_00000000-4_cuda.ptx, line 147; error : State space incorrect for instruction 'st'
1>ptxas =8・ Ptx assembly aborted due to errors
いきなりこんなエラーが出たんですがどういう意味ですか?
解決法を知っている方教えてください!!
0272デフォルトの名無しさん
2011/01/24(月) 08:07:36http://en.wikipedia.org/wiki/Comparison_of_Nvidia_graphics_processing_units
誰かこのページでGPUの浮動小数点計算速度の比較に使われてるプログラム知ってる人いませんか?
0273デフォルトの名無しさん
2011/01/24(月) 10:18:04ざっと見ただけではこのページのどこに浮動小数点計算速度の比較が書いてあるのか分からなかったのだが。
いずれにしてもそこに載っているスペックは全て理論ピーク性能かと。
0274デフォルトの名無しさん
2011/01/24(月) 10:27:17いきなりすぎて難しいのでとりあえず中間ファイルの.ptxファイルを保存して146行目を調べて。
あとはSDKに含まれるプログラムとかは正常にコンパイルできるのかどうかも。
0275デフォルトの名無しさん
2011/01/24(月) 13:59:18invalide device function とは何が原因で起こるエラーなんでしょうか?
初歩的な質問ですがどなたか教えてください。
0276デフォルトの名無しさん
2011/01/24(月) 14:38:37device で実行しているfunctionがinvalideなんだろ
でもinvalideじゃなくてinvalidじゃね?
0277275
2011/01/24(月) 15:13:26invalidは「無効な」という意味で捉えています。
デバイス(__global__)の関数名とホストで実行する関数名は同じにしているのですが、なぜ無効と言われるのかがわからないのです。当然ですが、何かの関数名と被らないようには名前をつけています。
どなたかお願いします。
0278デフォルトの名無しさん
2011/01/24(月) 15:38:250279デフォルトの名無しさん
2011/01/24(月) 15:56:170280275
2011/01/24(月) 16:12:51画像処理をCUDAでやっており、今は、メディアンフィルタを作っています。
OpenCVを使いたかったので、ここの方法(ttp://blogs.yahoo.co.jp/simasaki2001/18576477.html)を参考に簡単なサンプルを先生が作成し、それを基に学生の自分がプログラムを作っています。
現在、__global__ void cuda_median() というデバイスの関数をcuda_median<<<ndrange_global, ndrange_local>>>(入力画像, 入力画像のポインタ, 出力画像, 出力画像のポインタ)と呼び出しています。
ndrange_globalとndrange_localはDim3型で
ndrange_local.x= atoi( argv[3] );
ndrange_local.y= 1;
ndrange_local.z= 1;
ndrange_global.x= 1;
ndrange_global.y= (int)dev.multiProcessorCount; <-これは14となっています
ndrange_global.z= 1;
となっています。この部分についてよく理解できなかったので先生に質問しましたが、「人の作ったのだからよく知らない。」と言われました。現在もしっかり理解しているとは言えません。
__device__関数も一つ作成して使っていますが、__global__からしか使用していません。
0281デフォルトの名無しさん
2011/01/24(月) 16:21:56device内のfunctionでのエラーではないのか?
0282275
2011/01/24(月) 16:37:21device関数を呼ぶ前の処理におけるエラーチェックではエラーは発生していないので、device関数でエラーが起こっているのではないかと考えています。
0283デフォルトの名無しさん
2011/01/24(月) 16:57:03ttp://4.bp.blogspot.com/_nq8CjirJSCA/TDWRC5XR8vI/AAAAAAAACZY/pyHDbs2AA0E/s1600/gpgpu.png
上にあるようなCPUとGPUの浮動小数点演算性能の比較図を見るとGPUの方が圧倒的に高いように見えるのですが
性能はGPUの方が上という事でしょうか
それともただ単に1コアのクロック周波数から理論的な演算性能を算出してコア数分かけた値で、
GPUはコアが多いからそう見えているのでしょうか
さらにメモリの転送速度もGPUの方が速いのですよね。
他の点で総合的に見たらCPUのほうがいいということなのでしょうか
0284デフォルトの名無しさん
2011/01/24(月) 17:01:06>GPUはコアが多いからそう見えているのでしょうか
>さらにメモリの転送速度もGPUの方が速いのですよね。
>他の点で総合的に見たらCPUのほうがいいということなのでしょうか
それで合ってる
0285デフォルトの名無しさん
2011/01/24(月) 17:21:26まあ、簡単に言うと計算の仕方とメモリアクセスの仕方について、GPUの方が制約がきついってこと
もっとも、CPUでも高い性能を出すためには厳しい制約が付くけど
0286283
2011/01/24(月) 18:18:42ありがとうございます。
つまり全てのコアを常に使い切ればこれだけの性能が出るので、
メモリアクセスやのレイテンシや転送時の待ち時間に
準備の出来た他の演算(warp)を次々にやらせていって
いかに常にコアをフル回転させるか、がCUDAプログラミング。
warp内では実時間でも完全に同期して実行されるため
warp内分岐は、分岐を通らないwarpは完全に遊んでしまう(コアは空いているのに他のwarpのグループを実行できない)
ので遅くなってしまう(というかピーク性能が出せない)。
メモリアクセス時には他のwarpを計算すればいいが、次のwarpもメモリアクスになり、最初のwarpもまだアクセス中だったときには
3つめのwarpを計算する。このようにどんどん他のwarpを処理していくが、一度に起動できるwarpは上限がある(カーネル関数で使用するレジスタや共有メモリで決まる)
ため演算に対してメモリアクセス時間が長いと、起動上限数使い切ってしまい、
それでも最初のwarpがメモリアクセス中だとコアが遊んでしまうので、
メモリアクセス時間を減らしたり(メモリアクセス回数を減らす、共有メモリを有効利用する、
バンクコンフリクトを回避して共有メモリに早くアクセスできるようにする、コアレスアクセスでグローバルメモリに早くアクセスする)、
同時起動可能なwarp数を増やす(Occupancyうんぬん)
ことで見た目全コアが降る回転するようにする。
ということでしょうか
0287デフォルトの名無しさん
2011/01/24(月) 19:19:41ついこの前出たOpenCV 2.2でNPP経由でCUDA対応したから、それ使ってみたら?
exeではなくCMakeから自前ビルドしないといけないけど
0288デフォルトの名無しさん
2011/01/25(火) 10:09:24SDK内のプログラムはきちんと動作しています。
.ptxファイルは保存方法がぐぐっても良くわからなくて・・・。
どうやったら保存できますか?
0289アク禁長いから久々の書き込み
2011/01/25(火) 12:05:07状況がよく判らんが、__device__な関数は__device__か__global__な関数からしか呼んじゃダメよ。
それと、<<<>>>で呼んだ関数で起きたエラーは、次にAPIを呼んだときに戻されるのでご用心。
つまり、以下の場合cudaThreadSynchronize()の戻り値がエラーになる。
--
func<<<grid, block>>>();
cudaThreadSynchronize();
>>271
取り敢えず、-ptxつけてコンパイルして味噌。
0290デフォルトの名無しさん
2011/01/25(火) 12:51:53コマンドプロンプトでnvcc -ptx cuda.cuであってますか?
普段VSでコンパイルしてるんでやり方がいまいちなんですが・・・。
0291283
2011/01/25(火) 13:02:29ttp://www.4gamer.net/games/099/G009929/20100309044/images/pic2.gif
のUniform Cacheでしょうか
容量と読み込み速度からしてSM内かと思ったのですが
全てのスレッドから読めるという点ではSM外のL2キャッシュなのでしょうか
SM内だとしたら全SMがそれぞれ同じデータをコピーして保持するというのは少々効率が悪い気もします
そもそも大容量なはずのVRAMに転送するコンスタントメモリが
64kBと小容量に決められているという事はキャッシュも同サイズ用意されていて
ヒット率100%が保証されているということなのでしょうか
0292デフォルトの名無しさん
2011/01/28(金) 08:08:16staticの代わりになるようなものはCUDAに存在するんでしょうか?
やっぱりグローバルメモリに領域を確保して使うしかないんでしょうか?
0293デフォルトの名無しさん
2011/01/29(土) 02:47:23すまん質問の意味がわからん
0294デフォルトの名無しさん
2011/01/29(土) 21:17:57すいません、ちょっと日本語がおかしいですね・・・
要するに、GPU上のdevice関数内で、
関数が終わっても変数の値が残るようにしたいんですけど、
それをするにはcudaMemcpyでグローバルメモリに領域確保して、
そこに値を保存するようにするしかないのでしょうか?
0295デフォルトの名無しさん
2011/01/29(土) 21:38:16そもそもインライン展開されるんじゃなかったっけ
0296デフォルトの名無しさん
2011/01/29(土) 23:19:040297デフォルトの名無しさん
2011/02/01(火) 19:48:05a,b,c|a,b,c|a,b,c
こう置かれるからハーフワープの各スレッドがaにアクセスすると
コアレスアクセスできないということですか?
配列の構造体.a[n],.b[n],.c[n]を作った方が良いという事でしょうか
0298デフォルトの名無しさん
2011/02/01(火) 19:56:410299デフォルトの名無しさん
2011/02/02(水) 19:08:00と聞くのですが、以下の二つの状況だとどちらがいいでしょうか?GTX480を用いています。
1.
Enter your resource usage:
Threads Per Block 256
Registers Per Thread 32
Shared Memory Per Block (bytes) 3072
GPU Occupancy Data is displayed here and in the graphs:
Active Threads per Multiprocessor 1024
Active Warps per Multiprocessor 32
Active Thread Blocks per Multiprocessor 4
Occupancy of each Multiprocessor 67%
2.
Enter your resource usage:
Threads Per Block 1024
Registers Per Thread 32
Shared Memory Per Block (bytes) 12288
GPU Occupancy Data is displayed here and in the graphs:
Active Threads per Multiprocessor 1024
Active Warps per Multiprocessor 32
Active Thread Blocks per Multiprocessor 1
Occupancy of each Multiprocessor 67%
1と2の違いはThread Per Blockを四倍にしたかどうかです。
つまりSM内のThread Block数が多いほうが単純にいいのか、BLOCK内のThread数が多いほうがいいのか、
あるいはどちらも同じなのか教えてください。
0300デフォルトの名無しさん
2011/02/02(水) 19:38:03計測してパフォーマンスのいい方にすればいいとおもうが、
1の方があるスレッドブロックが同期待ちになってるあいだに他のスレッドブロックが
実行可能な余地があるぶん性能低下はしにくいと思う
0301デフォルトの名無しさん
2011/02/02(水) 20:23:06回答ありです。
言葉足らずでした。。。
計測して実験したのですが、1のほうがパフォーマンスが上でした。
いい言葉が見つからないのですが、
1も2もブロック数は違うけど同じスレッド数ですので
メモリ隠ぺいのためにスレッド数が多ければそれでいいと思っていたので
性能は同じになると思っていました。
スレッドブロックが同期待ちというのは、syncした場合におこる現象でしょうか?
調べてみましたがやはりSMあたりのスレッド数が同じなら
SMあたりのスレッドブロック数が多いほうがいいというのが
いまいちわかりません。。。
よろしくお願いします
0302デフォルトの名無しさん
2011/02/02(水) 20:35:42__syncthreads();
はブロック内のスレッドを同期する。
全部のスレッドが1024として
1ブロック1024スレっどの場合、__syncthreads();すると全部のスレッドが同期される。
その間どのスレッドも実行できない。
4ブロックで1ブロック当たり256スレッドの場合、__syncthreads();しても256スレッドだけ同期されて
他の3ブロック分のスレッドは実行可能。
>SMあたりのスレッドブロック数が多いほうがいい
スレッドブロックはSM単位に割り当てられるから「SMあたりのスレッドブロック数が多いほうがいい」
というのはActive Thread Blocks per Multiprocessorが高いほうがいいというのと同義
0303デフォルトの名無しさん
2011/02/02(水) 20:53:25--ptxas-options=-vでlmemが出ていなかったので
ローカルメモリは使っていないと思っていたのですが
local〜が0でないカーネル(3種類)はすべてレジスタ/スレッドが63でした。
偶然同じ数になっていると思ってましたが
レジスタは63が上限でそれ以降はローカルメモリが使われるってこと?
0304デフォルトの名無しさん
2011/02/02(水) 23:02:21レジスタの容量を調べてみろ
例えばレジスタの本数が32768で1ブロックあたり512スレッド使うなら
1スレッドが使用できるレジスタは64以下になる、256スレッドなら128以下
ローカルメモリが使われるなんて状況は最悪だからそのプログラムは糞だ
シェアードメモリ余ってんならそっちも有効に使え
0305デフォルトの名無しさん
2011/02/02(水) 23:30:24CPU実行の方が遅いこともあるから必ずしも悪ではない.
0306デフォルトの名無しさん
2011/02/03(木) 17:37:00sharedメモリ使う都合上、最後の段階は1ブロックでやるじゃん
でも1ブロックのみってことは他のsm使わないから無駄してるよね
最後の段階はcpuに転送したほうがはやいってこと?
0307デフォルトの名無しさん
2011/02/03(木) 18:32:560308デフォルトの名無しさん
2011/02/04(金) 01:25:470309303
2011/02/04(金) 03:03:46レジスタ32768ですがスレッドは192/ブロックです。
そもそもActive Thread Blocks per Multiprocessorを2以上にするように
1スレッドで使用しているレジスタ数とシェアードメモリを見て
スレッド数を決めているので足りなくなるなんて事はないはず
--ptxas-options=-vでlmemの項が出ないんだから使ってないと思っているのですが
レジスタ/スレッドに上限なんてないですよね?
0310デフォルトの名無しさん
2011/02/04(金) 03:34:56いっそのことスレッド数を8とか32とかにして実行してどうなるか見てみればいいんじゃないか
なにか他の原因かもよ?
0311デフォルトの名無しさん
2011/02/04(金) 03:35:390312303
2011/02/04(金) 04:33:53GF100,Fermiです
8スレッド/ブロックにして(ブロック数を増やしたので総スレッド数は変わらず)
みましたがlocal load/storeは増えました。
カーネル内の変数を(なにかに使用するように)無駄に増やしてみましたが
レジスタ数が63以上にならなかったので上限があるのかも、と思った次第です。
皆さんはレジスタ数が64を超える事はありますでしょうか
0313デフォルトの名無しさん
2011/02/04(金) 05:01:16あとフェルミならローカルメモリはキャッシュが効くんではなかったかと思ったがどうでしたっけ?
あとコンパイルオプションでレジスタの使用個数指定してたりしない?
0314303
2011/02/04(金) 05:18:06もちろん--maxrregcountは指定していませんでした
ところで
--maxrregcount 64にすると
ptxas warning : Too big maxrregcount value specified 64, will be ignored
とでました。
--maxrregcount 63にすると何も出ずコンパイルされます。
0315デフォルトの名無しさん
2011/02/04(金) 05:23:420316デフォルトの名無しさん
2011/02/04(金) 05:42:170317デフォルトの名無しさん
2011/02/04(金) 12:19:260318デフォルトの名無しさん
2011/02/04(金) 18:43:37int idx = blockIdx.x * blockDim.x + threadIdx.x;
a[idx];
ってアクセスしたらきちんと32bitずつglobalから読みだしますか?
それとも128bitずつですか?
0320デフォルトの名無しさん
2011/02/06(日) 13:23:49経験値からじゃなくて何か根拠のある数字はないのか?
0321303
2011/02/06(日) 13:51:37同時に使ってるであろう変数を1個減らす→63
もういっこへらす→63
繰り返す
へらす→62
やっと62になったということは減らした変数の数*2本分(doubleなので)=26本のレジスタが
63から溢れ出てローカルに保存されているんじゃないか
→本来なら90くらい使われてるんだろう
まあ一概に変数1個=2本とは言えないんでしょうけど
他に表示されないレジスタ数を数える方法思いつきませんでした
0322デフォルトの名無しさん
2011/02/06(日) 15:52:57192スレッド/blockで動作してるなら
32768/192=170.66
で1スレッド170は理論的につかえるはずなのに他になにか制限みたいのがあるのかね
もしかして待機しているactive threadの分が消費されるのか?つまり
32768/192/ActiveThreadPerBlockが1スレッドの使えるレジスタ数になるのだろうか
0323303
2011/02/06(日) 16:46:28逆にactive threadの数がレジスタ数で決まるんじゃないでしょうか
192tread/block(= 6warp/block) → 48maxWarp/(6w/b) =8block
0shared/block → 49152maxShared/(0s/b) = ∞block
63register/tread(=12096register/block) → 32768maxRegister/(12096r/b) =2block
Min(8,∞,2,8) = 2blockActive = 12warpActive → 12warp/48maxWarpActive =25% occupancy
それとは関係なくFermiではregister/threadが上限63で
それ以上はローカルにまわされる上に--ptxas-options=-vでlmem表示されない
のだと解釈しました。
0324322
2011/02/06(日) 16:57:10GT is the thread allocation granularity, equal to 256 for devices of compute
capability 1.0 and 1.1, and 512 for devices of compute capability 1.2 and 1.3,
and 64 for devices of compute capability 2.x.
ってあるがこれか?
0325303
2011/02/06(日) 17:40:251ブロックで使用されるレジスタRbは2.xでは以下のように計算する
Rb=ceil((Register/thread)*(thread/warp),GT)*(warp/block)
すなわち
Rb=ceil(r*32,64)*(warp/block)
ceil(x,y)はxをyの整数倍に切り上げたもの
r=18,w/b=6のとき
Rb=(18*32→9*64)*6=3456
r=17,w/b=6のとき
Rb=(17*32→9*64)*6=3456
つまり1warpが使用するレジスタは量子的で、その単位GT=64
→rは2の倍数で使わないと損
ということかと。r/スレッドの上限とは関係なさそうですかね。
0326デフォルトの名無しさん
2011/02/07(月) 02:36:58省略せずに計算しているようだ。
http://www.netlib.org/blas/dgemm.fに反する実装って許されるの?
0327デフォルトの名無しさん
2011/02/07(月) 08:00:090328デフォルトの名無しさん
2011/02/07(月) 15:59:16公式のフォーラムで本当にそうなのか、
もしそうならどのような理由からか聞いたほうがいいんじゃないか?
もしかしたら自分が勘違いしている部分があるかもしれないし
Nvidiaの間違いなら修正があるだろうし。
0329303
2011/02/07(月) 16:11:41up to 63 registers
ttp://www.beyond3d.com/content/reviews/55/7
Maximum theoretical register allocation per thread is 63 32-bit registers.
このへんに63までと書いてありますが公式にはないのかな...
0330デフォルトの名無しさん
2011/02/07(月) 23:23:08デバッグのためのprintfが使えなくて困っています。
プロジェクト>プロパティ>Cuda Build Rule v3.0.0>General>GPU Architecture
の値をsm_20としたいのですが、無効な値となり設定できません。
(CUDA VS Wizard最新版がFermiより前のリリースのため)
GPU Architectureの箇所を<親またはプロジェクトの規定値を継承>とし
Extra Optionsに-arch sm_20と指定しても-arch sm_20 -arch sm_10となり
ビルドできません。
この状況を経験&打破された方がいらっしゃいましたらアドバイスをお願いします。。。
【環境】
Windows7 x64
Visual Studio 2008 Express
CUDA Toolkit&SDK 3.1(32bit)
CUDA VS Wizard 2.2(32bit)
Geforce GTX 580
0331デフォルトの名無しさん
2011/02/07(月) 23:48:13あくまでCUDA VS Wizardが不具合ならという前提だけど
sdkのtemplateプロジェクトを元に必要な箇所を修正してプロジェクトを
作りなおしてみるという手もある。
0332330
2011/02/08(火) 00:19:27ありがとうございます。頂いたアドバイスを頼りに解決しました。
templateプロジェクトを開く>プロパティ>あれ?CUDA Bulid Rule v3.0.14&sm_20できる>google先生
>先生「Cuda.rulesコピれ」>!SDK導入後コピーするの忘れてた!>�@から�AにCuda.rulesをコピー>ヽ(´ー`)ノ
�@C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK\C\common
�AC:\Program Files (x86)\Microsoft Visual Studio 9.0\VC\VCProjectDefaults
ちゃんとCUDA VS Wizardのプロジェクトでもsm_20を指定できました。
CUDA VS Wizardの問題ではなく、Cuda.rulesの格納し忘れでした。失礼いたしました。
0333デフォルトの名無しさん
2011/02/10(木) 23:09:14吐かれてしまうのですが、どういった理由が考えられますか?
最初に最大サイズの行列サイズでallocした後、
そのdeviceメモリ領域に行列データを転送して演算、
結果をhostに書き戻して、さきほどの領域にまた別のサイズ(最大よりは小さい)の行列を
hostから持ってきて演算・・・のような処理を繰り返しています。
0334デフォルトの名無しさん
2011/02/15(火) 02:14:46このスレのみなさんのおかげです
質問にいっぱい答えてくれてありがとう
0335デフォルトの名無しさん
2011/02/16(水) 01:01:240336デフォルトの名無しさん
2011/02/16(水) 05:21:54CUDA toolkitのサンプルのVecAddで、floatを全てdoubleに置き換えた上で
1.ベクトル要素数Nをどこまで大きくできるか
2.CPUによるVecAddとの速度差はどの程度か
を調べました。
※GPU=GTS450 1GB、CPU=AthlonII-X4 640@定格。CPUのVecAddは単純な3行ループで非マルチスレッド。
1.16*1000*1000までセーフ、17*1000*1000は実行時エラー(debug)。
つまりcudaMalloc()量にして384MBはセーフ、408MBはアウトでした(GTS450 1GB)。
VRAMいっぱいまで使えるわけではないのですね。
2.clock()による計測ですが、
GPU→トータル 437 clock(内訳:cudaMalloc(ABC)=250, cudaMemcpy(AとBをhost2dev)=125, VecAdd=0, Memcpy(Cをdev2host)=62)
CPU→172 clock
でした(debug。releaseでも比はだいたい同じ)。
最初は内訳を出さず437と172をみてガッカリしたのですが、内訳を見てマジスンマセンwでした。
メモリ確保・転送のオーバーヘッドがそれなりにあるということは、
「host関数から高速な関数としてdevice関数を(頻繁に)呼び出す」のではなく、
極力device関数内だけで処理をさせるようにしないとうまみが少ないのですね。
0337デフォルトの名無しさん
2011/02/17(木) 02:47:04入門書?応用書?
0338デフォルトの名無しさん
2011/02/17(木) 03:13:250339デフォルトの名無しさん
2011/02/17(木) 04:43:560340デフォルトの名無しさん
2011/02/17(木) 14:56:50そうそれ。2/14って書いてあったけど翻訳かな
0341デフォルトの名無しさん
2011/02/17(木) 17:23:30翻訳だよ。
内容は入門書
英文マニュアル読めるならいらない。
ストリーム、アトミック演算、OpenGLやD3Dとの連携のヒント
複数GPUを使う方法、ゼロコピーメモリ、ページロックメモリ
テクスチャメモリとコンスタントメモリの使用法などが書いてある。
最適化については記述されてない。つまり
ブランチダイバージェンス、GRAMのコアレスアクセス、シェアードメモリのバンクコンフリクト
の説明はない
0342デフォルトの名無しさん
2011/02/17(木) 22:51:23ttp://www.impressjapan.jp/books/2978
NVIDIA社シニアエンジニアが執筆! CUDAでハードウェア性能を引き出し、計算時間を劇的に減らそう
本書では、汎用GPU(GPGPU)上でのCUDA Cプログラミングの基礎からはじめ、スレッド間の連携、
コンスタントメモリ/テクスチャメモリの使用、アトミック/ストリームの活用、マルチGPUの利用、
DirectX/OpenGLとの相互運用などを解説。CUDA Cプログラミングのさまざまな重要ポイントや
注意点を取り上げています。サンプルを交えて解説しており、
CUDA Cの手法が具体的かつ明確にわかります。初級から中級への橋渡しともなる
CUDA Cプログラマ必携の翻訳書です。
\2,940 (本体 \2,800+税)
発売日:2011/02/14発売
0343336
2011/02/19(土) 01:05:39clock()によるGPUの実行時間は取得方法が誤っておりました。
スーパーコンピューティングニュース[東京大学情報基盤センタースーパーコンピューティング部門]
http://www.cc.u-tokyo.ac.jp/publication/news/
にリンクが張られてある大島聡史先生のPDF:
「これからの並列計算のためのGPGPU連載講座 (�V) GPGPUプログラミング環境CUDA最適化編」
(201005_gpgpu2.pdf)に
>funcname<<<xyz, xyz>>>(args1, args2);
の形の呼び出しは非同期であるためこの前後で時刻を取得しても
GPU上での計算時間は得られないと書かれてありました。
だから0clockになっていたんですね…。
(少なくとも処理後のcudaMalloc()の完了時点では処理は終わっているから
62clock未満ではあるのでしょうけど)
同pdfに、GPUの実行時間を調べられるCUDAプロファイラなるものが紹介されて
いるのでこの使い方も覚えていこうと思います。
このpdfのシリーズはすごくわかりやすいと思いました。
こういうものが公開されているのはすごくありがたいです。
0344デフォルトの名無しさん
2011/02/19(土) 08:58:530345デフォルトの名無しさん
2011/02/19(土) 11:05:26ttp://accc.riken.jp/HPC/training.html
の下の方のCUDA。
0346デフォルトの名無しさん
2011/02/19(土) 14:08:01そういうあなたにCuda by exampleをすすめる
0347デフォルトの名無しさん
2011/02/20(日) 01:43:34.230348デフォルトの名無しさん
2011/02/20(日) 01:54:08.170349デフォルトの名無しさん
2011/02/20(日) 02:46:20.30アレに必要なことは大体書いてある。
あとは実行してプロファイルとったり
nvccでcubin吐いてenvytoolsで確認したり。
0350デフォルトの名無しさん
2011/02/20(日) 05:37:37.970351デフォルトの名無しさん
2011/02/20(日) 08:59:35.860352デフォルトの名無しさん
2011/02/20(日) 18:02:58.211つの大きなGPUとして見え1つのホストスレッドで扱えるのでしょうか
それともSLIしていない場合と同じように複数のGPUとしてみえ、枚数分のホストスレッドで管理してやる
必要があるんでしょうか。
0353デフォルトの名無しさん
2011/02/20(日) 18:24:47.060354デフォルトの名無しさん
2011/02/20(日) 19:00:08.93ゲーム等ではSLIできてCUDAではできないんですね。
手軽さより自分で制御できる事を優先してるのかな
0355デフォルトの名無しさん
2011/02/20(日) 19:07:14.32むしろSLIがうまく分担してるんじゃない
0356デフォルトの名無しさん
2011/02/20(日) 19:39:29.660357デフォルトの名無しさん
2011/02/20(日) 19:44:13.09とくにCUDAスレの住民にはメリット大きいんじゃない?
0358343
2011/02/20(日) 22:40:30.94すみません、最初から読んだはずだったんですが、読み取れなかったか見逃していたみたいです。
>>345
情報ありがとうございます、これちょくちょく更新されているみたいですね。チェックしておきます。
>>346
まだ「はじめてのCUDAプログラミング」を買おうかどうか迷っているレベルなんですが、
行き詰まったらそれの購入も含めて考えます!
>>349
確かにProgramming Guideは英語の割にはわかりやすいと思います。
cuda cubinで検索したら「CUDA テクニカル トレーニング」という
日本語のpdf(内容はプレゼン用スライド・2008)がみつかりました。
www.nvidia.co.jp/docs/IO/59373/VolumeI.pdf
nVidia公式にも日本語のドキュメントが増えてくれるとありがたいです。
一方この状況の中で翻訳したものをwikiで公開してくださっている方もいて
ただただ感謝です(Fermi Compatibility GuideとFermi Tuning Guide)。
0359デフォルトの名無しさん
2011/02/20(日) 23:40:26.16まだなら「はじめてのCUDAプログラミング」を先に勧める。
英文が苦手ならさっさと本買ったほうが時間の節約になるし
実際のプログラムに時間を使った方が有益だと俺は思うよ。
0360358
2011/02/21(月) 00:02:16.28そうですね…約2500円で省略できる時間のほうが有意義のように思います。
よし…買おう…bk1の300円ギフト券が手に入ったら…。
0361デフォルトの名無しさん
2011/02/21(月) 10:12:48.81>前に聞いたときはSLIはなんらCUDAに関与しないとのことだったけど。
使えるっぽいぞ
http://jp.slizone.com/object/sli_cuda_learn_jp.html
0362デフォルトの名無しさん
2011/02/22(火) 00:44:44.56最初にDirectXとGPUのチェックがあるだけで
特にSLI独自のコーディングは見当たらないな
かといって複数のGPUを制御するようなスレッド関係のコードもない
それらからSLI組んでれば1つの大きなGPUとしてプログラムからは扱えるんじゃないかと
思ったんだがどうなんだろうか。
0363デフォルトの名無しさん
2011/02/22(火) 23:08:18.71なんだか滅茶苦茶な翻訳だな。
で、そこを読む限りPhysxとの連係には触れられているけど
それ以上の情報はないし、英語サイトにはその記述さえ見当たらないね。
うーん、どうなんだろう
0364デフォルトの名無しさん
2011/02/22(火) 23:39:43.060365364
2011/02/22(火) 23:40:49.41ごめんなさい
0366デフォルトの名無しさん
2011/02/24(木) 22:29:20.970367デフォルトの名無しさん
2011/02/24(木) 22:52:09.13たしかに載ってるが概要の紹介程度で、SLIを使う方法や
SLIからどうGPUが見え、どのように制御するか(orできるか)は書かれてない。
索引から逆引きしても10行くらいしか書かれてないじゃん。
0368デフォルトの名無しさん
2011/02/25(金) 05:40:01.780369デフォルトの名無しさん
2011/02/25(金) 06:17:44.62あとはGPUが対応してるか調べろよ
0370デフォルトの名無しさん
2011/02/25(金) 14:27:17.760371デフォルトの名無しさん
2011/02/28(月) 19:10:18.87アプリを動かすと”cudart32_32_16.dll”が無いと怒られます。
cudart32_32_16.dll で使われている関数等々をスタティックリンクする方法があれば教えていただけないでしょうか?
CUDA 3.2を使っています。
0372デフォルトの名無しさん
2011/02/28(月) 19:49:36.32-L /usr/local/cuda/lib64/ -lcudart
じゃだめ?
パスは適当
0373デフォルトの名無しさん
2011/03/01(火) 00:11:54.00リンカのオプションの追加の「依存ファイル」で「cudart.lib」は指定してあります。
ツールキットのcudart.lib のファイルサイズは約41KBのものが1つのみ。
cudart32_32_16.dll のファイルサイズは約375KBです。
スタティックリンク用のcudart.libが無いということでしょうか???
開発環境を詳しく書きます
[開発環境]
OS: Win7 64bit
IDE: Visual Studio C++ 2008 Express
CUDA Toolkit: ver 3.2 32bitバージョン
0374デフォルトの名無しさん
2011/03/01(火) 00:47:03.69CUDA RuntimeのDLLは再配布が許可されているみたいだし
単に他の環境で動かしたいだけなら、DLL同梱でいい気がする。
ttp://developer.download.nvidia.com/compute/cuda/3_2_prod/toolkit/docs/EULA.txt
0375デフォルトの名無しさん
2011/03/01(火) 02:29:40.71ほむ!
ありがとうございます。
インストールしたときについてくるライセンスではそこら辺がよくわからなかったので助かります。
とりあえず、DLLも同梱しようと思います。(スタティックリンクしたいけど・・・)
0376デフォルトの名無しさん
2011/03/01(火) 03:33:24.52http://www.4gamer.net/games/076/G007660/20110228018/
Fermi持ってたら恩恵受けられるのかな。なんかよくわからんけど楽しみだ。
>ちなみにGPUDirect 2.0とThrustをフルサポートするGPUは,Fermiアーキテクチャ限定となる。
0377デフォルトの名無しさん
2011/03/01(火) 03:56:33.42・GPUDirect 2.0 …スパコンには良さそう?
・Unified Virtual Addressing…どれが何のメモリなのかごっちゃになりそう?
・Thrust…STL。。。どうでもよさそう?
0378デフォルトの名無しさん
2011/03/01(火) 23:32:35.21ちょっと行き詰まりましたのでできれば2点ご教授いただきたく…
1.
CUDA 3.2 SDKについてくるtemplateプロジェクトのひな形を使ってアプリを書いているのですが、
プロジェクトの[プロパティ]→[構成プロパティ]→[CUDA Runtime API]→[GPU]
→[NVCC Compilation Type]で"Generate 32 bit .cubin file (-m32 -cubin)"を
選択してコンパイルしたところ
"1>nvcc fatal : redefinition of argument 'machine'"
と怒られました。おそらく"--machine 32" と、cubinオプションを選択したときに付く、"-m 32" の重複が怒られていると思うのですが
"-m 32"はcubin形式をプルダウン選択で選ぶと自動的に付いて、
"--machine 32"はどこで付くのかわからない状況です(Win32アプリには自動で付く???)。
どうすれば競合を解消できるでしょうか?
2.
仕方がないので、とりあえずコマンドラインから
-Xptxas="-v" -cubin --ptxas-options -v
を指定してコンパイルしたところ
無事にcubinファイルが出来たのですが、このcubinファイルはバイナリー形式で私の目では読むことが出来ませんでした。
このcubinをテキスト形式で読むにはどうしたら良いでしょうか?
(CUDA 3.0からcubinは一旦バイナリ形式で出力されるようになったようです)
[環境]
CUDA Toolkit, SDK 3.2
0379デフォルトの名無しさん
2011/03/02(水) 00:46:11.360380デフォルトの名無しさん
2011/03/02(水) 13:54:56.94今からくだ勉強したい場合はCUDA4.0がオススメですか?
0381デフォルトの名無しさん
2011/03/02(水) 15:15:15.23むしろcuda2.2くらいから。
基礎ができてから発展にすすめ
0382デフォルトの名無しさん
2011/03/02(水) 15:23:23.730383デフォルトの名無しさん
2011/03/02(水) 15:39:44.51CUDA 4.0でも基本的にこれまでの機能を全て備えている。
また古い世代のGPUは余裕がある場合だけ追うべき。
0384デフォルトの名無しさん
2011/03/02(水) 15:51:19.20同意
つか、
>cuda2.2くらい
"くらい"なんて中途半端な答え方してるのかがワカラン
0385デフォルトの名無しさん
2011/03/03(木) 23:20:08.40readelf でcubinみましたが、いまいちよくわからなかったのでptxみます。
まあ、ptxもよくわからないのですが…
0386デフォルトの名無しさん
2011/03/04(金) 11:17:18.220387デフォルトの名無しさん
2011/03/04(金) 20:05:59.060388デフォルトの名無しさん
2011/03/06(日) 02:29:43.510389デフォルトの名無しさん
2011/03/08(火) 02:57:27.52どなたかAMBER(Assisted Model Building with Energy Refinement)というソフトをCUDAで動かしてる人いませんか?
このソフトを
Cent OS5.5 x86_64(カーネル 2.6.18-194.el5)
インテルコンパイラーとの併用で動かそうと何度も試みているのですがなかなか上手くいきません。
不具合の原因はGCCのバージョンとインテルコンパイラーのバージョンの組み合わせにあると睨んでいるのですが…
もし、動かせている方がいらしたらGCCのバージョンとインテルコンパイラーのバージョンの組み合わせををヒントとして教えてほしいのですが
0390デフォルトの名無しさん
2011/03/08(火) 22:01:42.39コンパイルでエラーなのかリンカでエラーなのかわからん。
取り合えずsdkのサンプルでiccで動かしてみたら?それで問題がわかるでしょ。
それが出来ないなら、
そもそもnvcc自体がgccのラッパだし、
素直にgcc使えば?
0391デフォルトの名無しさん
2011/03/09(水) 22:09:17.79適切な関数があれば教えていただけないでしょうか。
0392デフォルトの名無しさん
2011/03/09(水) 22:12:55.970393デフォルトの名無しさん
2011/03/09(水) 22:49:28.90フロントエンドというかデスクトップというか、とにかくグラフィックスの描画が遅くなる(当たり前ですが…)のと、
GPUの温度がかなり上がるので、
デバイス関数のループの中に、ちびっとだけSleep(0) 的なものを入れたいのですが、
関数的なものは用意されていないっぽいですかね。
例えば、どれか1つだけスレッドの処理ループを長引かせて、__SyncThreads() でもしたらいいのかしらん…???
0394デフォルトの名無しさん
2011/03/09(水) 22:53:11.30熱いから休憩するって発想は俺含めてないんじゃね?
0395デフォルトの名無しさん
2011/03/09(水) 23:41:01.500396デフォルトの名無しさん
2011/03/09(水) 23:58:40.99例えば、
__global__ void test()
{
...
for (i = 0; i < LOOP; i++)
{
__syncThreads();
...
// Sleep(0) 的なもの???
if (blockIdx.x == 0 && blockIdx.y == 0)
{
// 何か冗長な処理
}
}
}
って感じじゃSleep(0)の代わりにならないですかね?
0397デフォルトの名無しさん
2011/03/10(木) 00:06:55.92結局熱くなってしまうんだが
0398デフォルトの名無しさん
2011/03/10(木) 00:16:02.21そうなんですか…
ビジーウェイト調べてみます。
0399デフォルトの名無しさん
2011/03/10(木) 04:14:12.92カーネルを分割すればいいんじゃね?
わざわざ遅くする理由を思いつかないから試す気にもなれないけど。
0400デフォルトの名無しさん
2011/03/10(木) 09:03:17.86これが正解だろう。
でもそもそも、ちょっと休ませたくらいじゃ熱なんて変わらんだろうから、ドライバの設定でクロックを下げたらいい。
もしくは、放熱を強化するかだろう。
0401デフォルトの名無しさん
2011/03/17(木) 13:38:21.43GTX460を2枚刺ししたPC2台と、4枚刺ししたPC1台のどちらを組む方が良いでしょうか?
4枚刺しですと、電源のとりまわしと冷却に不安を感じております。
0402デフォルトの名無しさん
2011/03/17(木) 14:33:27.940403402
2011/03/17(木) 17:10:47.410404デフォルトの名無しさん
2011/03/17(木) 17:47:16.68単純に、書かれた内容だけで判断する限り前者の方が計算能力が高い。
但し、後者により強力なCPUとメモリを搭載する場合は一概に言えない。
0405デフォルトの名無しさん
2011/03/21(月) 23:08:34.74GPU間の通信は無いです。
>>403-404
GPU間の通信が無く、CPU側の処理も軽いので
4枚刺しによる恩恵が特になさそうなので
2台で組んだ方がよさそうですね。
ありがとうございました。
0406デフォルトの名無しさん
2011/03/22(火) 00:55:46.70長い目で見ると、DSPとかで食いつないでた信号処理屋的に失業の危機じゃん…
0407デフォルトの名無しさん
2011/03/22(火) 05:48:58.160408デフォルトの名無しさん
2011/03/22(火) 10:14:13.39カスタムチップによる高速演算ボードを出していたところは尽く撤退しました。
共有メモリが16バンクに分割されているってのは�@�Aのどっちだろ。
�@マルチプロセッサのもつ16384Byteが16分割されてる
�Aブロックに割り当てた共有メモリが16分割される
ハードウェアが固定されてそうだから�@の気がするんだけど…。
0410福盛俊明
2011/03/23(水) 02:45:25.670411デフォルトの名無しさん
2011/03/23(水) 13:39:51.91GeForce GTX 580を二枚搭載したPC(CUDA 4.0 RC インストール済み)を使って,
GPUDirect 2.0を用いたプログラムを作成しようと思っています.
cudaDeviceCanAccessPeerという関数を用いて,
GPUDirect 2.0が使用可能かどうか調べたのですが,
結果が0(=出来ない)ということでした.
FermiでもGeForceではダメだということでしょうか?
それとも,GPUDirect 2.0は Peer-to-Peer Memory Access(←私が勝手にGPUDirect 2.0だと思っている)
とは別のものなのでしょうか?
試された方はここに居らっしゃいますでしょうか?
0412デフォルトの名無しさん
2011/03/23(水) 17:48:37.13確かにTeslaシリーズでは1が、GeForceシリーズでは0が返る。
だけど実際にはGeForceでも通信はできる。
UVAだかなんだかに対応するGPUなら問題ないのかも。
0413デフォルトの名無しさん
2011/03/23(水) 23:13:07.00例えばSMの上限ワープが32のところに16個のワープにしても、
メモリアクセスのようなレイテンシがないなら問題ない気がするんだが。
0414デフォルトの名無しさん
2011/03/24(木) 11:46:09.16CUDAを使って、ゲームエンジンのUnrealとかcryを使ってる方いらっしゃいますか。
参考になるサイトを教えてほしいのですが
0415デフォルトの名無しさん
2011/03/24(木) 20:35:29.92論理演算部分のアセンブラ化による高速化は望めるでしょうか?
それとも、レジスタの使い方がコンパイラの方がうまくて、
より高速なアセンブラを書くのは難しいでしょうか?
__global__ void test(int *a, int *b, int *c)
{
// 初期化処理
// メインのループ
for (i = 0; i < LOOP; i++)
{
a[i] = (b[i]<<10 | (b[i] >> 22)) ^ c[i];
}
}
0416デフォルトの名無しさん
2011/03/24(木) 21:10:44.15そもそもメモリアクセスがボトルネックにならないほど多くの演算があるのかどうか。
0417デフォルトの名無しさん
2011/03/24(木) 21:50:24.370418デフォルトの名無しさん
2011/03/25(金) 12:10:15.17また、コンパイル時にptxファイルに明示的にコメントを挿入するコマンドがあれば教えていただけないでしょうか。
例えば、test.cu に、「__PTX_COMMENT("foo");」という行を記述すると、ptxファイルに
// foo
といった行が挿入されるという具合のものを探しています。
0419デフォルトの名無しさん
2011/03/25(金) 13:36:10.03ptxにブロックごとに行番号情報は含まれているのでソースのどこに相当するかは分かるのでは。
特定の場所が知りたければ例えば__syncthreads()を10個入れとくとか。
0420デフォルトの名無しさん
2011/03/25(金) 15:16:09.980421デフォルトの名無しさん
2011/03/25(金) 17:03:42.70なるほど。例えば、ダミーでfoo0()とか、foo1()とか関数を作ってコメント代わりに呼び出せばよいと。
>>420
コマンドとしては無いということですね。
0422デフォルトの名無しさん
2011/03/28(月) 20:51:27.79CUT_DEVICE_INIT(argc, argv);
cufftHandle plan;
cufftPlan1d(&plan, 4096, CUFFT_C2C, 16);
cufftDestroy(plan);
↓エラー
First-chance exception at 0x000007fefd96aa7d in CudaTest.exe: Microsoft C++ exception: cudaError_enum at memory location 0x001ef830..
なずぇ・・・。cufftExecC2C走らせるとFFTは実行して変換結果も正しいのだが謎す。
0423デフォルトの名無しさん
2011/03/29(火) 11:17:10.20VRAMは未だ余っている筈なのにエラーになったり、作ったプランを使わずに破棄するとこけたり……
って、>422はそのパターンか?
0424デフォルトの名無しさん
2011/03/30(水) 14:28:25.63cufftExecC2C実行してもやっぱでる。
海外フォーラムみてたらver 3.2が怪しくて ver3.0にしたら消えたって書いてあったんで
ver3.0にしてみたがやっぱでる。うーん・・・。
0425デフォルトの名無しさん
2011/03/30(水) 16:09:16.17うーん、SDKとcufftの組み合わせで相性あるしなぁ。
あと、サイズが2の冪かどうかでも変わるかもしれない。
つーか、3.1よりは3.2の方がましだった気がするんだが……
0426デフォルトの名無しさん
2011/04/02(土) 17:59:42.31VC++でダイアログから.cuの関数使うようなサンプルってなんかない?
0427デフォルトの名無しさん
2011/04/02(土) 18:01:41.56おまいさんは関数を分割することもできんのかね。
0428デフォルトの名無しさん
2011/04/02(土) 19:08:36.61それはライブラリ作んなくてもできるんかい?
0429427
2011/04/02(土) 19:39:35.97extern "C" つけたらできたや。ごめんごめん。
0430デフォルトの名無しさん
2011/04/02(土) 20:08:34.70お前はだれ?
そんで誰にレス支店の?
0431425
2011/04/02(土) 21:07:05.730432デフォルトの名無しさん
2011/04/04(月) 17:39:36.21CU_SAFE_CALL(cuInit(0));
CU_SAFE_CALL(cuDeviceGet(&device, 0));
CU_SAFE_CALL(cuGLCtxCreate(&context, 0, device));
と初期化すると,cuGLCtxCreateで
Cuda driver error 3e7
というエラーが発生してしまいます.
詳しい方居らっしゃいますでしょうか?
当方の環境は
Windows xp 32bit
CUDA 4.0 RC
ドライバは270.51を使用しています.
0433デフォルトの名無しさん
2011/04/04(月) 22:25:10.79CU_SAFE_CALL(cuGLInit(0));
CU_SAFE_CALL(cuDeviceGet(&device, 0));
CU_SAFE_CALL(cuGLCtxCreate(&context, 0, device));
0434デフォルトの名無しさん
2011/04/04(月) 22:26:11.471点だけ。
uGLInit(0)→uGLInit()
0435432
2011/04/05(火) 11:10:29.98この通りにしましたところ,今度はcuGLInitで
Cuda driver error c9
が発生しました...
0436デフォルトの名無しさん
2011/04/05(火) 17:48:55.53Intelコンパイラ:"icl.exe"に行わせたいのですが、どのようにコマンドを記述したら良いでしょうか?
環境:
Windows 7
CUDA 3.2
0437デフォルトの名無しさん
2011/04/05(火) 17:59:21.74もっといい方法あるんかいな?
0438デフォルトの名無しさん
2011/04/05(火) 20:34:55.574/7 予測 http://up3.viploader.net/ippan/src/vlippan198234.jpg
発表はドイツ気象庁 http://www.dwd.de/
0439デフォルトの名無しさん
2011/04/05(火) 21:13:26.380440デフォルトの名無しさん
2011/04/05(火) 23:50:34.550441デフォルトの名無しさん
2011/04/06(水) 02:41:41.960442デフォルトの名無しさん
2011/04/06(水) 07:54:08.270443デフォルトの名無しさん
2011/04/06(水) 10:06:08.15さすがにそれはちょっと違うと思うけど
気象系シミュレーションでGPUつかってるところって結構あるよね?
0444デフォルトの名無しさん
2011/04/06(水) 17:09:45.350445デフォルトの名無しさん
2011/04/06(水) 17:21:47.830446デフォルトの名無しさん
2011/04/06(水) 18:02:23.22cuInit
cuDeviceGet
cuCtxCreate
cuGLInit
cuGLCtxCreate
じゃね?
0448デフォルトの名無しさん
2011/04/07(木) 18:05:26.230450デフォルトの名無しさん
2011/04/07(木) 19:40:03.310451432
2011/04/08(金) 08:52:38.77大変御迷惑おかけしました.
自分の無知さに絶望しました.
OpenGLの初期化後に
CU_SAFE_CALL(cuInit(0));
CU_SAFE_CALL(cuDeviceGet(&device, 0));
CU_SAFE_CALL(cuGLCtxCreate(&context, 0, device));
で動作させることが出来ました.
御手数おかけしました.
0452デフォルトの名無しさん
2011/04/08(金) 20:02:33.080453デフォルトの名無しさん
2011/04/08(金) 21:02:09.880454デフォルトの名無しさん
2011/04/09(土) 16:43:22.41cutil.hにもcuda.hにも定義されてねーのだが。
0455デフォルトの名無しさん
2011/04/09(土) 17:05:25.290456デフォルトの名無しさん
2011/04/09(土) 17:10:51.07,;f ヽ
i: i
| | ///;ト,
| ^ ^ ) ////゙l゙l;
(. >ノ(、_, )ヽ、} l .i .! |
,,∧ヽ !-=ニ=- | │ | .|
/\..\\`ニニ´ !, { .ノ.ノ
/ \ \ ̄ ̄ ̄../ / .|
0457デフォルトの名無しさん
2011/04/10(日) 01:00:12.430458デフォルトの名無しさん
2011/04/10(日) 11:48:25.160459デフォルトの名無しさん
2011/04/10(日) 12:13:51.20メイン ドット クー?
0460デフォルトの名無しさん
2011/04/10(日) 12:41:02.43フォーラム見ても遅い遅いと書いてあるし、どうしようもないのかこれ?
0461デフォルトの名無しさん
2011/04/10(日) 18:44:20.120462デフォルトの名無しさん
2011/04/10(日) 18:48:15.640463デフォルトの名無しさん
2011/04/10(日) 19:12:02.260464デフォルトの名無しさん
2011/04/10(日) 19:56:18.220465デフォルトの名無しさん
2011/04/11(月) 20:09:52.50walking
working
star wars
それぞれなんて読む?
0466デフォルトの名無しさん
2011/04/11(月) 20:23:49.350467デフォルトの名無しさん
2011/04/11(月) 20:38:03.05ワーニング
ワーキング
ウォーキング
スターウォーズ
0468デフォルトの名無しさん
2011/04/11(月) 20:41:31.720469デフォルトの名無しさん
2011/04/11(月) 20:42:58.660470デフォルトの名無しさん
2011/04/11(月) 21:04:45.770471デフォルトの名無しさん
2011/04/11(月) 22:09:14.040472デフォルトの名無しさん
2011/04/11(月) 22:31:58.48わかるわ
油断するとIdxとIdyがあると思っちゃう
threadIdx.y
xのy???
スレッドはフルスペルなんたからthreadIndexでいいじゃねえか
0473デフォルトの名無しさん
2011/04/11(月) 23:56:43.330474デフォルトの名無しさん
2011/04/12(火) 22:35:25.63スレッドに分けて暗黙にコアレスアクセスさせると、どうもコードが汚くなる。
0475デフォルトの名無しさん
2011/04/13(水) 01:04:12.90計算結果を最終的にOpenGLで表示するような場合だと、無駄なメモリコピーは仕方ないものとして
計算結果をホスト側に逐一戻したほうがよさげ?
0476デフォルトの名無しさん
2011/04/19(火) 21:32:44.770477デフォルトの名無しさん
2011/04/20(水) 01:10:42.960478デフォルトの名無しさん
2011/04/20(水) 01:11:01.88(CUDAのカーネル関数の中で、MMXのPADDDのような32bit整数の4並列計算のようなことができなかと思っています)
0479デフォルトの名無しさん
2011/04/20(水) 01:16:07.37CAL+IL
0482デフォルトの名無しさん
2011/04/20(水) 11:51:01.49イミフ
0483デフォルトの名無しさん
2011/04/20(水) 20:42:07.79えと、CUDA用には無いということでしょうか?
add.u32
を2並列(64bitレジスタ)、あるいは4並列(128bitレジスタ)で実行するような命令セットはないでしょうか?
0484デフォルトの名無しさん
2011/04/20(水) 21:24:43.20そのためのSIMT
0485デフォルトの名無しさん
2011/04/20(水) 22:27:22.88何かを誤解しているようだが、CUDAには必要ないから無い。
SSEを使わないと100%の性能が出せないとかいうCPUのような制約は無い。
add.u32という命令ならほっといてもWarp単位の32スレッドが(ほぼ同時に)実行する。
全命令が32並列のSIMD命令になっていると考えればよい。(32ビットデータの場合)
0486デフォルトの名無しさん
2011/04/20(水) 22:50:57.29CUDAがSIMTなのか、SIMDなのか、SPMDなのかはさておいて…
昔の記述だとCUDAアーキテクチャのレジスタは32bitとかだったと思いますが、
今時のNVIDAのGPUには64bitレジスタ、あるいは128bitレジスタがあって、それを使って
padd.u32 みたいな命令により、「1スレッド内の演算において」32bit整数の4並列処理が出来るとかないでしょうか?
ということなのですが、いかがでしょうか?
Streamでは128bitレジスタで32bit4並列の演算ができるみたいのですが・・・
0487デフォルトの名無しさん
2011/04/21(木) 04:36:54.190488デフォルトの名無しさん
2011/04/21(木) 08:16:13.34そんなものはないから誤解だと言ってるんだが。
AMDは1スレッドが4 or 5演算のVLIWで動作するから4並列のSIMD命令化の意味があるが
CUDAは1スレッドが各種演算ユニットを1ずつしかもたない(or複数スレッドが共有)CUDA Coreで
実行されるので128bit型を用いたとしても4命令になるだけ。
0489デフォルトの名無しさん
2011/04/21(木) 09:31:19.38誤解というか、質問なのですが・・・
今どきのCUDAでもレジスタレベルのSIMD命令は無く、
ハードウエア的に、レジスタも演算器も(今のところは)そのような用途には向いていないということですね。
ありがとうございましたゞ
0490デフォルトの名無しさん
2011/04/21(木) 09:35:39.70まだ試験実装だからコア数が少ないけどメインメモリ直結だからけっこう早そうなんだが
0491デフォルトの名無しさん
2011/04/21(木) 09:59:44.72イミフな反応が多々あったけどね
0492デフォルトの名無しさん
2011/04/21(木) 10:04:26.920493デフォルトの名無しさん
2011/04/21(木) 21:59:22.93論理レジスタとか、物理レジスタとかいう話しですか???
>>486への根本的な回答はどうなりますか?
0494デフォルトの名無しさん
2011/04/21(木) 22:29:01.490495デフォルトの名無しさん
2011/04/21(木) 22:39:09.84streamsdkには専用の命令があると?
cudaはそんなものは勝手にやると言ってる人が居るけど、どうやってやるんだ?
cudaにはそういう命令は無いよ
内部で処理することが出来るとは思えないけどな
少なくともユーザープログラムからやる方法はないよ
内部的にはそういう命令も持ってるかもしれないけど
というかそういうアセンブラレベルの最適化をしなくていいようにしたのがCUDA言語なわけで
複雑化する要因はパフォーマンスを犠牲にしてでも排除されるだろう
どうしてもやりたかったら64bit変数を使って自前でやるしかないだろう
0496デフォルトの名無しさん
2011/04/21(木) 23:07:04.51>>486 への回答は”No”ということですね。
ありがとうございます。
将来的には"Yes"になると思っているのですが、そこら辺が「根本的に解ってない」と言われる所以ですかね。
「根本的に解ってない」らしいので、何がなんだかさっぱりですが・・・
# 今、スクラッチのCUDA用アプリで、SSEで書かれたCPU用アプリの約5倍(CPU側はSandyBridgeで4スレッドで動作)の速度が出ているのですが、
CPU側がきちんとAVXに対応してきたらこの差はかなり縮まりそうなんですよね・・・
0497デフォルトの名無しさん
2011/04/21(木) 23:11:08.25まずそっからじゃないのかね?
0498デフォルトの名無しさん
2011/04/21(木) 23:14:37.36さっきからすげー鼻につくんだけど
教わる態度じゃないよお前
教えてんの俺じゃないけどさ
0499デフォルトの名無しさん
2011/04/21(木) 23:22:41.50腹癒せしているか言い訳探しているかってところだろ。
今日たまたま計測していたんだが、cufftでもfftwfの3倍程度だった。
汎用ルーチンじゃこんなもんだよ。
0500デフォルトの名無しさん
2011/04/21(木) 23:51:10.29CUDAの将来的な見通しについて知りたかった部分もあったので
失礼に見えた部分があるかもしれません。
ごめんなさい。
0501デフォルトの名無しさん
2011/04/22(金) 00:46:21.02単純なコアの数を増やす方向に進むだろうなあ。
0502デフォルトの名無しさん
2011/04/22(金) 00:51:40.280503デフォルトの名無しさん
2011/04/22(金) 01:02:26.93おまけにGF104は倍精度をサポートしているコア数は1/3だし。
更にHPC向けではなくグラフィック向けの物は制限かかっているらしいし。
倍精度が必要になるなら仕方ないだろうけど、32bit変数で足りるのに64bit変数を使うのはデメリットの方が圧倒的に多そう。
0504デフォルトの名無しさん
2011/04/22(金) 01:08:10.32労力掛かるわりには環境依存甚だしいわ
4倍くらいだったら別に我慢出来ないようなものなんてないし
マルチコアCPUで普通に組のが一番いいという結論に達した
0505デフォルトの名無しさん
2011/04/22(金) 01:11:09.84計算したい内容を簡潔に書いてみ
計算内容によってGPUに向き不向きがあるから
0506デフォルトの名無しさん
2011/04/22(金) 01:12:19.97膨大なスレッド数と、その切り替えの軽量さを活かして
メモリアクセスのレイテンシの隠蔽とかだっけ。
>>502
スレッドあたりの変数が多いと苦労するね。
そういう場合はシェアードメモリをうまく使って切り抜けるしかないのかな。
0507デフォルトの名無しさん
2011/04/22(金) 01:15:58.96メモリアクセス考えると結局16要素単位の処理の塊になってしまう。
0508デフォルトの名無しさん
2011/04/22(金) 01:19:14.16共有使ってもたらん
Sharedも増やせ!キャッシュも増やせ!
0509デフォルトの名無しさん
2011/04/22(金) 01:31:49.640510デフォルトの名無しさん
2011/04/22(金) 07:57:21.85早くintel、amd,nvidiaで統一してくれよ。
機械語レベルでも、ライブラリレベルでもいいからさ。
0511デフォルトの名無しさん
2011/04/22(金) 10:14:49.72CPUレベルでスイッチングしてるだけみたいだから
ほとんどオーバーヘッドがないみたいだぞ
十分標準として使えそうなんだが
0512デフォルトの名無しさん
2011/04/22(金) 10:24:09.66だから何度もスレッド生成するような処理はかえって遅くなる
0513デフォルトの名無しさん
2011/04/22(金) 10:26:56.960514デフォルトの名無しさん
2011/04/22(金) 10:42:12.62おまえ>>504だろ
試行レベルが同一
0515デフォルトの名無しさん
2011/04/22(金) 11:01:13.45お前、いろいろと残念な人だな
0516デフォルトの名無しさん
2011/04/22(金) 11:09:33.71つか下らんこと書くならよそでやれヴぉけ
0517デフォルトの名無しさん
2011/04/22(金) 11:16:10.78な。
煽るなら突っ込まれる隙作っちゃダメだ
0518デフォルトの名無しさん
2011/04/22(金) 11:20:23.650519デフォルトの名無しさん
2011/04/22(金) 11:21:44.53カーネル側でスレッドの生成回数を出来るだけ減らす工夫が必要だったり。
0520デフォルトの名無しさん
2011/04/22(金) 11:55:51.71スレッドを切り替えるときにスタックポインタ変えたりするコンテクストスイッチの負荷が問題
0521デフォルトの名無しさん
2011/04/22(金) 13:09:56.12cpuの話をここに書いてる?
0522デフォルトの名無しさん
2011/04/22(金) 13:20:03.480523デフォルトの名無しさん
2011/04/22(金) 13:23:34.120524デフォルトの名無しさん
2011/04/22(金) 13:55:34.91だからスレッドの性質については知っておく必要がある
CUDAを使いこなすのにアップアップのアホどもには関係ないけど
0525デフォルトの名無しさん
2011/04/22(金) 14:04:51.88GPGPUでやるもんだと思ってたんだが違うのか?
0526デフォルトの名無しさん
2011/04/22(金) 14:11:42.11単純に並列度が高ければGPUでOKってもんじゃない
Opteronの48コアマシンがお手頃価格で手に入る時代だぜ?
0527デフォルトの名無しさん
2011/04/22(金) 14:13:49.170528デフォルトの名無しさん
2011/04/22(金) 14:15:39.880529デフォルトの名無しさん
2011/04/22(金) 17:16:29.77> Opteronの48コアマシンがお手頃価格
今、おいくら万円くらい?
0530デフォルトの名無しさん
2011/04/22(金) 17:24:19.240531デフォルトの名無しさん
2011/04/22(金) 17:48:16.700532デフォルトの名無しさん
2011/04/22(金) 17:49:59.020533デフォルトの名無しさん
2011/04/22(金) 17:50:02.420534デフォルトの名無しさん
2011/04/22(金) 17:56:21.72個人とか研究室とかだと絶対的金額から手が出せないかもしれないけど。
0535デフォルトの名無しさん
2011/04/22(金) 17:58:31.180536デフォルトの名無しさん
2011/04/22(金) 18:06:31.47えーと、0.5テラFlopsくらい?
0537デフォルトの名無しさん
2011/04/22(金) 18:59:41.35欲しいけど家庭用のコンセントで大丈夫だろうか
似たようなスペックでGPGPU構築したらいくらぐらいかかるのかな
0.5テラFlopはしょぼくないか
0538デフォルトの名無しさん
2011/04/22(金) 20:43:24.14家庭用コンセントじゃ無理
グラボ刺さるからGPGPUできるよ
0539デフォルトの名無しさん
2011/04/22(金) 21:43:26.64このマザーで自作すると・・・
http://akiba.kakaku.com/pc/1103/31/131500.php
130000+100000*4+12800*16+100000+50000+α=884800+α
くらい?
だとしたら100万円ほどで完成品が買えるなら良心的な価格な気がするね。
下のマザーボードだと気合いで8本GPU載せられるんだろか?
http://akiba.kakaku.com/pc/1104/20/213000.php
0540デフォルトの名無しさん
2011/04/22(金) 21:59:24.97> 下のマザーボードだと気合いで8本GPU載せられるんだろか?
どこかに引っかかりそう。ラック筐体次第なのかな
0541デフォルトの名無しさん
2011/04/22(金) 22:17:54.30まあ、高周波シュミレソフトで1000万いくから
性能気にしないでいくなら自作できそうだけど・・・
0542デフォルトの名無しさん
2011/04/22(金) 23:17:49.25そのマザーのPCIeスロットはx8+x1+x8+x1+x8+x1+x8+x1だから
ホスト・デバイス間の帯域不足に悩みそう。
演算性能さえあればいいプログラムなら何とかなるのかもしれないけど。
というかハイエンドグラボは複数スロット占有が一般的だし、水冷のグラボが必須になりそうw
0543デフォルトの名無しさん
2011/04/22(金) 23:27:34.87P6T7 WS SuperComputer
http://www.unitycorp.co.jp/asus/motherboard/intel/lga1366/p6t7_ws_super/index.html
と
EVGA GeForce GTX 580 FTW Hydro Copper 2
http://jp.evga.com/products/moreInfo.asp?pn=015-P3-1589-AR
になるのかな。
電源の確保がものすごく大変そうだけどw
0544デフォルトの名無しさん
2011/04/22(金) 23:28:08.92"Error in profiler data file'(動かすディレクトリ)/temp_compute_profiler_0_0.csv' at line number 1. No column found"
と出て計測結果がうまく出ません
何が原因なのでしょうか?
0545デフォルトの名無しさん
2011/04/23(土) 00:17:34.23キーを押さなきゃ終わらないプログラムをプロファイラに渡してはまったことがあるZE
0546デフォルトの名無しさん
2011/04/23(土) 00:24:57.17Visual Profiler使ってないからよく判らんが、その作業ディレクトリはレガシーな名前かい?
空白やら何かの所為で巧くアクセスできないってことはないよね?
0547デフォルトの名無しさん
2011/04/23(土) 00:39:41.60外部から画像を読み込んでいますがキーは押さなくても終了するプログラムです
>>546
作業ディレクトリと同じディレクトリに入っている計測したい実行ファイルは
プロファイルを開始すると実行されるようなのでおそらくディレクトリにアクセスは出来ています
プロファイラはプログラムを実行してから計測をしているようなのですがこの計測経過が100%になったら先ほどのエラーがでてしまいます
temp_compute_profiler_0_0.csvは自動生成のファイルの用です
0548デフォルトの名無しさん
2011/04/23(土) 00:56:29.31エラーがなんで出てるかはさっぱりわからんので、自分がこけたときの経験を列挙するぐらいしかできんのだけど、
Visual Profilerのデフォルト設定では30秒以内に処理が終わらない実行ファイルは実行を途中でうち切って
解析結果出力せず、えんえん解析を繰り返すってのはあったぜ。
0549デフォルトの名無しさん
2011/04/23(土) 00:58:45.66そこ毎回適当に300に変えてたわ
しかしなんで一回でぜんぶできないかな
0550デフォルトの名無しさん
2011/04/23(土) 01:02:49.32画面表示による影響を平滑化するために複数回実行時間を計測してるんじゃないかと。
0551デフォルトの名無しさん
2011/04/23(土) 14:34:55.7916×16の二次元のスレッド作ったとして
threadIdx.x = 0〜15 が半ワープに入るのか
threadIdx.y = 0〜15 が半ワープに入るのか、どっちだべ。
0552デフォルトの名無しさん
2011/04/23(土) 14:56:49.140553デフォルトの名無しさん
2011/04/23(土) 15:03:38.310554デフォルトの名無しさん
2011/04/23(土) 15:42:37.530555デフォルトの名無しさん
2011/04/23(土) 22:47:46.900556デフォルトの名無しさん
2011/04/24(日) 00:29:55.820557デフォルトの名無しさん
2011/04/24(日) 01:23:47.09AVXがまさにLarrabeeのインターフェースらしい
0558デフォルトの名無しさん
2011/04/24(日) 02:12:08.17デバイス側関数もブレークポイント仕掛けられるって聞いたのにできねえ。
リンクするライブラリ変えるとかしなきゃいかんの?
0559デフォルトの名無しさん
2011/04/24(日) 02:25:06.200560デフォルトの名無しさん
2011/04/24(日) 02:31:36.91Nsightを使わない理由は、NsightインストールしたらCUDA3.3がインストールされて、
でCUDA3.3のcuFFTの動作が怪しかったんで、NsightやめてCUDA3.0にしてる。
NsightでもCUDA3.0使えるのかもしれないけど、使い方調べるのがめんどくさかったので
環境構築に慣れたNsight使わないやり方でやってる。
0561デフォルトの名無しさん
2011/04/24(日) 02:44:49.490562デフォルトの名無しさん
2011/04/24(日) 03:35:32.04使えるとしたら、Sandy BridgeのGPU用のコーディングには何を使えば良いですか?
0563デフォルトの名無しさん
2011/04/24(日) 03:40:42.99SBってGPGPUに対応してるの?
SDKとか出てるって聞いたことないんだけど。
0564デフォルトの名無しさん
2011/04/24(日) 03:53:11.32CUDAに対応してるって話はみたことないから、CUDAとは同時に使えないと思われる。
OpenCLスレで聞いたほうがいいかも。
0565デフォルトの名無しさん
2011/04/24(日) 03:57:39.89>CUDAに対応してるって話はみたことないから、CUDAとは同時に使えないと思われる。
対応していない=同時に使えないという考えは改めた方位が良い。
0566デフォルトの名無しさん
2011/04/24(日) 06:20:12.37グラフィックチップとしては完全に機能停止されるだろうけど
演算装置として稼動するんだろうか
0567558
2011/04/24(日) 11:04:44.800568デフォルトの名無しさん
2011/04/24(日) 12:25:30.290569デフォルトの名無しさん
2011/04/24(日) 16:01:57.36Nsight使うとこういうのの原因すぐに特定できるようになんの?
0570デフォルトの名無しさん
2011/04/24(日) 23:11:39.8464bit用のバイナリを作るには、CUDAも64-bit版をインストールする必要があるでしょうか?
それとも、CUDAの32-bit版でも nvcc のオプションの切り替えで32bit用と64bit用の両方のバイナリを生成することが出来るでしょうか?
出来れば、WindowsXP(32bit) の環境で64bit用バイナリも作りたいと思っています。
0571デフォルトの名無しさん
2011/04/24(日) 23:26:06.00>64bit用のバイナリを作るには、CUDAも64-bit版をインストールする必要があるでしょうか
たしかそうだったはず
>CUDAの32-bit版でも nvcc のオプションの切り替えで32bit用と64bit用の両方のバイナリを生成することが出来るでしょうか?
生成したい処理系のSDKをインストールしないといけなかったはず
切り替えはできないと思う
2つ(32と64)いれればOKなのかは不明
0572デフォルトの名無しさん
2011/04/24(日) 23:42:26.64レスありがとうございます。
ややこしそうなので、32bit 環境、64bit 環境それぞれを用意し、それぞれ用のバイナリを作ろうと思います。
0573デフォルトの名無しさん
2011/04/25(月) 10:06:00.860574デフォルトの名無しさん
2011/04/27(水) 08:28:47.19コンパイルできません。
因みにC++は初心者ですがよろしくお願いします。juria_gpu.cu
http://codepad.org/UtZeyAkV
julia_gpu.cu(44): error: calling a host function("cuComplex::cuComplex") from a
__device__/__global__ function("julia") is not allowed
0575デフォルトの名無しさん
2011/04/27(水) 08:35:16.69基本がわかってないと思うから基本からよく勉強だな
0576デフォルトの名無しさん
2011/04/27(水) 08:43:59.77書いて無いし、サンプルコードなので動くはずなのですが、なぜか動きません。
http://codepad.org/oDh0YFM5
cuComplexをjulia関数から呼ぶのは許されていないという意味なのでしょうか?
0577デフォルトの名無しさん
2011/04/27(水) 08:50:02.31全部晒せよ
0578デフォルトの名無しさん
2011/04/27(水) 09:03:00.78ここです。よろしくお願いします。
0579デフォルトの名無しさん
2011/04/27(水) 09:07:22.66まずそういう情報を先に出せよ
0580デフォルトの名無しさん
2011/04/27(水) 09:11:49.07知っている方とお見受けしますので安心です。
0581デフォルトの名無しさん
2011/04/27(水) 09:22:12.28nvccのオプションは?
0582デフォルトの名無しさん
2011/04/27(水) 09:27:35.630583デフォルトの名無しさん
2011/04/27(水) 09:31:28.050584デフォルトの名無しさん
2011/04/27(水) 09:33:25.86上のソース名は間違いだよな?それでやってどのエラーメッセージが出るんだ?
julia_gpu.cu(44): error: calling a host function("cuComplex::cuComplex") from a
__device__/__global__ function("julia") is not allowed
がでるのか?
0585デフォルトの名無しさん
2011/04/27(水) 09:38:07.66です。間違いです。julia_cup.cuの場合はこちらでも通ります。
ただGPUを使おうとするとコンパイル時にエラーが出ます。
cuComplex c(-0.8, 0.156);
cuComplex a(jx, jy);
と
a = a * a + c;
の部分でエラーが出ます。原因がわかりません。
0586デフォルトの名無しさん
2011/04/27(水) 09:40:32.600587デフォルトの名無しさん
2011/04/27(水) 09:51:50.54ここです。
0588デフォルトの名無しさん
2011/04/27(水) 09:57:53.82これじゃ -lcutil32とかないじゃん なけりゃダメだろ
0589デフォルトの名無しさん
2011/04/27(水) 10:01:49.76そちらでも同じようなエラーが出るか確かめてほしいです。
もしエラーが出なかったら私の環境がおかしいのかも知れません。
Linuxでの開発環境とWindowsでは少し違うと思うのでWindowsでも
確認してもらいたいです。
0590デフォルトの名無しさん
2011/04/27(水) 10:05:27.84無いのですが・・・
0591デフォルトの名無しさん
2011/04/27(水) 10:13:19.95あ、そういうこと
じゃあソースが壊れてるんでなければそっちの環境がおかしそうだね
環境の再構築したほうがよさげ
0592デフォルトの名無しさん
2011/04/27(水) 10:33:39.34nvcc -lglut julia_gpu.cu → OK
nvcc julia_gpu.cu → コンパイルOK、リンクNG
0593デフォルトの名無しさん
2011/04/27(水) 11:22:50.87nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2011 NVIDIA Corporation
Built on Thu_Mar_24_14:53:10_PDT_2011
Cuda compilation tools, release 4.0, V0.2.1221
こちらの開発環境ですが、nvccのバージョンを教えてください。
環境を再インストールしましたがやっぱりだめです。
構文エラーとして解釈されます。
0594デフォルトの名無しさん
2011/04/27(水) 11:28:41.584.0RC2?はまだ早いんでね?
0595デフォルトの名無しさん
2011/04/27(水) 11:40:43.34(最新は1.3)に対応しているみたいなので古いほうをインストール
してみます。
0596デフォルトの名無しさん
2011/04/27(水) 11:45:26.50ありがとうございました。
0597デフォルトの名無しさん
2011/04/27(水) 13:22:17.13何をどうしたら解決したか書いてほしいな
0598デフォルトの名無しさん
2011/04/27(水) 13:30:13.834.0RC→3.2で解決したようだ
0599デフォルトの名無しさん
2011/04/27(水) 13:44:28.802台とかあって3.2と4.0の環境があったりするなら切り分けしやすいから良いかもしれないけど。
0600デフォルトの名無しさん
2011/04/27(水) 15:49:38.3516*32の行列に長さ16の配列をXORの積をしたいのですが可能でしょうか?
イメージ的には32個のブロックに16個のスレッドという感じで考えています。
行列計算なので実装はそんなに難しくないと思うのですが、具体例があまり
見当たらず苦心しています。はじめてからまだ三日なのですが、解る方が
いらしたらソースを見せてほしいです。ヒントだけでも良いです。
よろしくお願いします。
0601デフォルトの名無しさん
2011/04/27(水) 16:03:05.24これどういう意味?もう少し計算を具体的に書いてくれ
0602デフォルトの名無しさん
2011/04/27(水) 16:34:04.46[abcd]*[{efgh}{ijkl}{mnop}[qrst}{uvwx}{yz12}{3456}{7890}]=
a^e&b^f&c^g&d^h=e
a^i&b^j&c^k&d^l=i
...
このように計算したいです。
並列計算なら全体をばらして8個のブロックに4個のスレッドを当てれば
4サイクルで実行できると思うのですが何か間違えていますでしょうか?
よろしくお願いします。
0603デフォルトの名無しさん
2011/04/27(水) 16:40:44.57o=FG[a[j]]^GF[u1.m[j]];
p=FG[b[j]]^GF[u.m[j]];
for(i=0;i<16;i++){
d1[j]^=t[o][h1[p][i]];
d2[j]^=t[o][h2[p][i]];
}
buf[j]^=d1[j];
buf[j+16]^=d2[j];
}
この処理を並列化したいです。
0604デフォルトの名無しさん
2011/04/27(水) 17:08:33.14ideoneかcodepadにC版の関数をアップロードしてくれ。
そうしてくれたらcudaでどこをどう並列化するか考えてみるよ。
0605デフォルトの名無しさん
2011/04/27(水) 17:11:57.15http://codepad.org/15hfT9HS
0606デフォルトの名無しさん
2011/04/27(水) 18:53:48.98見ても解らないだとか汚いと罵るだけ。
どうせやる気無いんだろ?
0607デフォルトの名無しさん
2011/04/27(水) 18:56:58.77うん。
忙しいからね(´・∀・` )
0608デフォルトの名無しさん
2011/04/27(水) 19:19:58.01またお前か。
いい加減自分の手を動かしたらどうだ。
時間がないなら対価を払って時間を買え。
0609デフォルトの名無しさん
2011/04/27(水) 19:50:55.960610デフォルトの名無しさん
2011/04/27(水) 19:53:32.57何怒ってんだ?
その程度のコードなら普通にできるだろ
バカなの?
0611デフォルトの名無しさん
2011/04/27(水) 19:57:29.620612デフォルトの名無しさん
2011/04/27(水) 20:06:58.92自分でやる気が無いっていう態度が滲み出てるんだよ。
0613デフォルトの名無しさん
2011/04/27(水) 20:14:43.33ここならもっと詳しい人が居ると思うから聞いてみただけ。
何で最初からやる気が無いなんて決め付けるの?
0614デフォルトの名無しさん
2011/04/27(水) 20:18:05.38もう少し考えた過程が見えないと回答もしづらい罠
0615デフォルトの名無しさん
2011/04/27(水) 20:19:27.54分からない、分からない、一個教えてもらうとすぐ次が分からないで
自分で調べる気なんてなさそうに見えるんだが。
0616デフォルトの名無しさん
2011/04/27(水) 21:13:16.83Intel Parallel Studio 使って並列化。
↓
逆アセンブルして、コード解析してCUDA化。
↓
(゜д゜)ウマァ
0617デフォルトの名無しさん
2011/04/27(水) 21:18:58.59CUDA 3.2 SDKにあるVC++ 2008用のtemplateプロジェクトを、VC++ 2010用に変換しようとしましたがうまくいきません。。。
0618604
2011/04/27(水) 21:32:18.42>並列化したいのはループ使っている部分だけなのでそこだけです。
>http://codepad.org/15hfT9HS
悪い悪い、デートしてて遅くなった。
で、このコードのどこをcuda化したいんだ?
未だ帰り道だから、んな長いの見てられないんだ。
0619デフォルトの名無しさん
2011/04/27(水) 22:18:04.420620デフォルトの名無しさん
2011/04/28(木) 00:21:09.33Windows版のnvccはVisual Studio 2005か2008のコンパイラが前提になっているらしいから
VC++ 2010用のテンプレートだけではうまくいかないと思う。
それらのコンパイラを用意するとなったら、VS2008で作ったほうが楽な気もする・・・
0621デフォルトの名無しさん
2011/04/28(木) 01:56:26.69きたないコードだなあ・・・。
0622デフォルトの名無しさん
2011/04/28(木) 02:06:27.26俺ならこうだな。
for (i = 0 ; i < 16; i++) {
for (j = 0; j < 16; j++) {
}
}
主なコーディング作法が3つくらいあるから、そのどれかにしろ。
0623デフォルトの名無しさん
2011/04/28(木) 02:08:04.55スペースきれい!
ふしぎ!
0624デフォルトの名無しさん
2011/04/28(木) 02:10:39.48あまりにもイミフなんだけど、とりあえず並列化したいなら
関数呼び出し側
func<<<<16, 16>>>(d1, d2, t, h1, h2, buf);
処理内容
__global__ void func(hoge* d1, hogehoge* d2, hogehogehoge* t, hogehogehogehoge* h1, hogehogehogehogehoge* h2, hogehogehogehogehoge* buf){
int o = FG[a[blockIdx.x]^GF[u1.m[blockIdx.x.]];
int p = FG[b[blockIdx.x]^GR[u.m[blockIdx.x]];
d1[blockIdx.x] ^= t[o][h1[p][threadIdx.x]];
d2[blockIdx.x] ^= t[o][h2[p][threadIdx.x]];
buf[blockIdx.x] ^= d1[blockIdx.x];
buf[blockIdx.x + 16] ^= d2[blockIdx.x];
}
酔っ払いながらコード書いたから間違いある気がするけど、
CUDAの根本的な文法も理解しようとせず使おうとしてるから、バレないし問題ないよね。
0625デフォルトの名無しさん
2011/04/28(木) 20:52:23.73VC2010はCUDA 4.0 RC2から対応で、SDKにテンプレートが付属してる。
うちの環境だとMSBuild関連が正しくインストールされなかったのでひと手間必要だったけど。
0626デフォルトの名無しさん
2011/04/29(金) 23:49:38.06VSで開発してるんだが、どうもレジスタが足りないとFirst chance exceptionエラーが出る気がしてきた。
0627デフォルトの名無しさん
2011/04/29(金) 23:58:25.76祈るってアプローチが間違いとは言わないけど、
同じ計算を2並列で行うとか、間違いなく構成を変えられる方法があれば
演算器を入れ替えて2回行うとかって手もあるよ。
2倍時間がかかってもまだCPUには勝てるでしょ。
0628デフォルトの名無しさん
2011/04/30(土) 00:29:11.48ググると出来ないというソースが目立ちますが...
X58マザー使っています。
0629デフォルトの名無しさん
2011/04/30(土) 00:35:04.100630デフォルトの名無しさん
2011/04/30(土) 01:12:10.48cuda4はどうか知らないが、それ以前のcudaではSLIだからより速くできると言うことはない。
勿論GPU2枚挿しに対応したアプリなら、2枚刺した方が速くはなるがSLIには関係ないようだ。
0631デフォルトの名無しさん
2011/04/30(土) 03:38:04.89引数はsharedメモリに確保されるってどっかで見たことあるけど…
sharedメモリつかってないプログラムだしなあ。
使ってるのはGeForce9800 - G92コアなんだけど、
コアに上限の24ワープ割り当てるVSがエラー吐いて動かなくなるんだけど、
控え見えに16ワープ放りこむと問題なく動く。
cuFFTとかの標準ライブラリでも動作が怪しかったりするのは
レジスタ周りの取り扱いのせいなんじゃって気がしてきてるのですよ。
0632デフォルトの名無しさん
2011/04/30(土) 08:15:07.06それ以前はシェアードだったかな
関数内変数はレジスタで不足したら勝手にグローバルにいくんでしょ
ワープ数が上限ぴったりならエラーは吐かないと思うんだけど
0633デフォルトの名無しさん
2011/04/30(土) 08:23:40.360634デフォルトの名無しさん
2011/04/30(土) 15:46:08.21そんな曖昧な知識で推測する前にPDFちゃんと読むんだ
0635デフォルトの名無しさん
2011/04/30(土) 16:50:47.84了解したぜ
0636デフォルトの名無しさん
2011/04/30(土) 23:22:46.13ありがとう。
SLI環境はもう少し様子見します。。
0637デフォルトの名無しさん
2011/05/02(月) 12:42:12.87Fermiに変えたが早すぎワロタ。
0638デフォルトの名無しさん
2011/05/02(月) 13:14:07.57キャッシュはでかい
0639デフォルトの名無しさん
2011/05/08(日) 01:40:31.320640デフォルトの名無しさん
2011/05/08(日) 05:13:16.28nVIDIAのページにもELSAのページにもなくなってしまったのですが。
0641CUDA初心者
2011/05/09(月) 11:25:03.07FFT.exe の 0x7c812afb で初回の例外が発生しました: Microsoft C++ の例外: cudaError_enum (メモリの場所 0x0012ae3c)。
というエラーが出てしまい困っています。ホストからデバイスへのメモリのコピー等
の簡単なプログラムは動きます。ちなみにGeforce210です。
0642デフォルトの名無しさん
2011/05/09(月) 21:02:35.890643デフォルトの名無しさん
2011/05/10(火) 09:55:41.040644デフォルトの名無しさん
2011/05/10(火) 22:29:43.44多分それ、cuFFTのバグ。そのエラーコードはメモリアクセス例外のときに出る。
俺の場合9800GTから550Tiに買い換えたらエラーが消えた。
cuFFTがデバイスメモリ使い過ぎて、デバイスが積んでるメモリ量によっては
エラーを吐くとかが原因の気がしてるんだが、本当の理由はわからん。
NVidia側のバグフィックスを待つぐらいしか現状で対策はないと思う。
0645デフォルトの名無しさん
2011/05/11(水) 00:23:29.41cuFFTの内部で結構デバイスメモリを使うようだから、ボードの半分以下のFFTしかかけられない。
0646デフォルトの名無しさん
2011/05/11(水) 19:42:26.40FFTなんて本来ライブラリの側で対処すべき次元のもので、
ユーザー側がデバッグに煩わされるべきものではないがな。
0647デフォルトの名無しさん
2011/05/11(水) 19:57:25.35もしバグだとおもうならフォーラムで問い合わせたり報告したりしたほうが良いんじゃないか?
バグなら修正してもらえる可能性があるし、そうじゃないとしたら自分の間違いだとわかる
0648デフォルトの名無しさん
2011/05/11(水) 23:13:45.45deviceQueryを実行しようとすると
Error: API mismatch: the NVIDIA kernel module has version 270.29,
but this NVIDIA driver component has version 270.40. Please make
sure that the kernel module and all NVIDIA driver components
have the same version.
と出てしまい実行出来ません。どうしたら良いでしょうか?
0649デフォルトの名無しさん
2011/05/11(水) 23:20:38.40http://www.alc.co.jp/
0650デフォルトの名無しさん
2011/05/11(水) 23:33:32.11ありがとうございます
0651デフォルトの名無しさん
2011/05/11(水) 23:35:55.99ありがとうございます
0652デフォルトの名無しさん
2011/05/12(木) 00:05:11.00フォーラムには既にあったと思う。
0653CUDA初心者
2011/05/12(木) 15:13:05.97ありがとうございます。
550Ti持ってるのでそっちでやってみようと思います。
0654デフォルトの名無しさん
2011/05/16(月) 10:53:35.28Xwindowの解像度が640*480固定になるのはどうしてなのかしら?
0655デフォルトの名無しさん
2011/05/16(月) 10:57:58.450656デフォルトの名無しさん
2011/05/16(月) 11:00:18.750657デフォルトの名無しさん
2011/05/18(水) 12:36:20.91The Nouveau kernel driver is currently in use by your system.
というエラーが出て、先に進めません。
どうしたらNouveau kernel driverをとりのぞけますか?
0658デフォルトの名無しさん
2011/05/18(水) 12:42:47.54https://help.ubuntu.com/community/BinaryDriverHowto/Nvidia
0659デフォルトの名無しさん
2011/05/18(水) 13:18:46.93早速の返信ありがとうございます。
sudo apt-get --purge remove xserver-xorg-video-nouveau
と打って、再起動もしてみましたが、やはり同じメッセージが帰ってきます。
ほかに、何か手立てはないものでしょうか?
0660デフォルトの名無しさん
2011/05/18(水) 21:39:19.83俺がLive CDで試したときは
kernelの引数にnouveau.modeset=0として起動後に
modprobe -r nouveauした。
そもそもnouveauがモジュールではなくて組み込みになってるなんてことはないよね?
0661デフォルトの名無しさん
2011/05/19(木) 04:06:30.550662デフォルトの名無しさん
2011/05/19(木) 04:12:46.190663デフォルトの名無しさん
2011/05/19(木) 08:14:50.0816人17脚で走ること
ただし、一番端っこの人の番号に制限あり
0664デフォルトの名無しさん
2011/05/19(木) 08:28:38.02なんだかすごく遅くなりそうな例えだなw
0665デフォルトの名無しさん
2011/05/19(木) 09:19:43.1516冊まで借りられるから自分の欲しいタイトル1冊と隣に並んでるの15冊借りて帰って来るの
そん中にフラットメイト15人が欲しい本が入ってたらラッキー
何人か足りなかったらまたまた誰か派遣して16冊借りてこなくちゃなんない
一回で16人全員満足したり16冊全部が有効に使われたらうれしいな
俺ら棚のブックエンドの右からまとめてとってく癖があるから
司書さんはいいかんじに本並べといてね
家にあればこっちのもんだから欲しい本被っててもおkおk
てのはどうだろう
0666デフォルトの名無しさん
2011/05/19(木) 09:33:17.23食堂のカウンターに並んでいる先頭の16人が皆「カレーライスください」という。
食堂のおばちゃんが「皆カレーかい?鍋ごとあげるから皆で分けてね。」
もちろん正確ではないが速くなりそうな例ではある。
0667デフォルトの名無しさん
2011/05/19(木) 15:22:10.76皆さんCUDA版プログラムとCPU版プログラムの速度比較は
どのようにされてますでしょうか?
CUDA3.0以降はエミュレーションモードがなくなったようなので、
別にCPU版プログラムも作って比較するのがよいでしょうか?
NVIDIAの回し者でない限り、iccでカリカリにチューンしたプログラムと比較するべし。
iccでチューニングした上で、更にその先の為にcudaを使うのだから。
まして、エミュレーションだなんてなんの悪夢だか。
0670デフォルトの名無しさん
2011/05/19(木) 21:39:40.22メモリ読み書きするなら32byte, 64byte, 128byte単位にしてください。
これだけ。
0671デフォルトの名無しさん
2011/05/20(金) 14:19:17.96サンプルを実行すると
cudart32_31_9.dllが見つからなかったため、このアプリケーションを開始できませんでした
。cudartを間違えて消してしまったと思われるのですが、どうやったら解消されるのでしょうか?
ぐぐってみたのですが、わかりませんでした。
0672デフォルトの名無しさん
2011/05/20(金) 14:31:37.090673デフォルトの名無しさん
2011/05/20(金) 15:41:59.990674デフォルトの名無しさん
2011/05/20(金) 17:07:36.680675デフォルトの名無しさん
2011/05/20(金) 17:37:59.88よかったな、cudartを間違って消してしまったわけじゃないことが判ったじゃないか。
0676デフォルトの名無しさん
2011/05/27(金) 03:06:11.390677デフォルトの名無しさん
2011/05/27(金) 11:09:49.930678デフォルトの名無しさん
2011/05/28(土) 08:21:44.260679デフォルトの名無しさん
2011/05/28(土) 08:56:58.210680デフォルトの名無しさん
2011/05/28(土) 21:17:04.74device query含め、CUDA Cのサンプルが全部動かない。
OpenCLとDirectComputeは動く。
なんでだろう。
0682デフォルトの名無しさん
2011/05/29(日) 11:48:30.13$(CUDA_PATH)\extras\visual_studio_integration\rules
にあるファイルをコピーしたのですが、これを設定してプロジェクトを読み込みなおすと
「要素 <UsingTask> 内の属性 "AssemblyFile" の値 "$(CudaBuildTasksPath)" を評価した結果 "" は無効です」
といったエラーが出てプロジェクトが読み込めないようです…。
0683デフォルトの名無しさん
2011/05/29(日) 21:12:48.05CUDAで高速化してちょ、といわれたんだけど殴っていい?
0684デフォルトの名無しさん
2011/05/29(日) 21:18:05.50いや無理にでもやれよ
それが仕事ってもんだ
0685デフォルトの名無しさん
2011/05/29(日) 21:27:29.77んで、ゼロコピーを使うと遅くなるケースが出てきたんで、
仕方なく使わないことにしたんだけど、
やっぱりcudaMallocやcudaMemcpyの時間がもったいないんで
どうにかしたいんだけど、なんか知恵ある?
0686デフォルトの名無しさん
2011/05/29(日) 21:53:26.07どう頑張っても並列化できないを証明すれば良いじゃん.
なかなか難しいよ.
0687デフォルトの名無しさん
2011/05/29(日) 22:08:28.81ゼロコピーって何よ。
何をしたいのか判らんけど、cudaMemset()も使えない内容なの?
0688682
2011/05/29(日) 22:53:30.43先ほどあった投稿で解決することができました
http://forums.nvidia.com/index.php?showtopic=201433
0689デフォルトの名無しさん
2011/05/29(日) 23:36:05.38マジレスするとゼロコピーとは、cudaHostGetDevicePointer()を使うやつのことだす。
cudaMallocHost()でホスト側のメモリを確保しとかないとダメだけど。
0690デフォルトの名無しさん
2011/05/30(月) 01:01:06.08処で、勿体無いと言うけどcudaMalloc()って時間掛かる?
0691デフォルトの名無しさん
2011/06/01(水) 01:01:48.66プログラミングガイドによるとfloatにおけるmulとmadの速度は同一だとあるのですが、
for(i=0;i<N;i++)
a = b*c;
と
for(i=0;i<N;i++)
a = d*(b+c);
では明らかに前者のほうが短い時間で処理を行っています。
dをコンスタンとメモリや定数にしても同じでした。
いったいどういうことなんでしょうか?
助けてください。
0692デフォルトの名無しさん
2011/06/01(水) 01:07:21.88どこのプログラミングガイドだよ教えろよ教えてくださいお願いします
0693デフォルトの名無しさん
2011/06/01(水) 01:12:04.190694デフォルトの名無しさん
2011/06/01(水) 01:14:43.17GPU : GeForce GTX 285M
CUDA Toolkit : 2.2
です。
>>692
NVIDIA CUDA Programming Guide 2.2です。
第5章Instruction Performance の 5.1.1.1 Arithmetic Instructions より
Throughput of single-precision floating-point add, multiply, and multiply-add is 8
operations per clock cycle.
だそうです。
0695デフォルトの名無しさん
2011/06/01(水) 01:27:04.68ごめんなさい書き間違えました。
for(i=0;i<N;i++)
a = d*b+c;
です。
0697デフォルトの名無しさん
2011/06/02(木) 01:22:02.76a = b*c+d*e
とかじゃないのん?畳込みとか行列積とかで頻出する計算です。
0698デフォルトの名無しさん
2011/06/02(木) 01:30:56.84a = a + (b*c)
だったわ。
a+= (b*c)
でもええかもしれんが。
0699デフォルトの名無しさん
2011/06/02(木) 03:02:35.87270.81でないとCUDAランタイムAPIが変な動作をするみたい
つまり現時点ではPC用のドライバしかない
0700デフォルトの名無しさん
2011/06/02(木) 11:07:00.65「Unified Virtual Addressing (UVA) やGPUDirect? v2.0を特徴とするCUDA 4.0を使ったアプリケーション用サポートの追加」
と書いておきながら動かないなんてふざけた話だわ
0701デフォルトの名無しさん
2011/06/02(木) 22:17:32.06でも世の中に出てる日本語のCUDAの本って4冊ともFermi以前のアーキテクチャでの説明だからsharedメモリが16kbyteだったり、1ブロックあたりのスレッド数が512だったりするんで困ります。
なんとかして。
0702691
2011/06/02(木) 22:33:23.38たぶんそのせいで遅くなっているんだと思います。
ご迷惑かけてすみません。
0703デフォルトの名無しさん
2011/06/02(木) 22:59:17.85ちょうど新しいドライバが出たみたいです
0704デフォルトの名無しさん
2011/06/02(木) 23:33:53.87英語も勉強汁
0705デフォルトの名無しさん
2011/06/02(木) 23:44:11.150706デフォルトの名無しさん
2011/06/03(金) 00:27:22.61Fermi用の正しい値が分かる人なら何も困らないような気が・・・・。
0707デフォルトの名無しさん
2011/06/03(金) 22:26:32.06あれ、結局、どう活用していかわからん。
0708デフォルトの名無しさん
2011/06/04(土) 10:50:31.43Fermiの方が簡単になってるからいんじゃね
>>707
ローカルメモリ使われてたりバンクコンフリトあるかをチェックすんじゃね
0709デフォルトの名無しさん
2011/06/04(土) 11:02:21.920710デフォルトの名無しさん
2011/06/04(土) 11:18:35.390711デフォルトの名無しさん
2011/06/04(土) 20:06:16.650712デフォルトの名無しさん
2011/06/04(土) 23:19:09.390713デフォルトの名無しさん
2011/06/05(日) 00:30:29.730714デフォルトの名無しさん
2011/06/05(日) 00:56:21.390715デフォルトの名無しさん
2011/06/05(日) 04:04:47.930716デフォルトの名無しさん
2011/06/05(日) 10:23:06.12少なくとも小中で算数(数学)や社会、英語などを並行して勉強してきたはず
0717デフォルトの名無しさん
2011/06/05(日) 10:42:05.41CUDAのプログラム記法以外に知るべき事が多い。と言うかむしろそちらが大半じゃないか。並行してできる、がんばれ
0718デフォルトの名無しさん
2011/06/05(日) 18:19:38.330719デフォルトの名無しさん
2011/06/05(日) 18:43:41.04CUDA BY EXAMPLEのサンプルをいくつか試したがエラーをはく。
まだ3.2使うわ
0720デフォルトの名無しさん
2011/06/05(日) 19:14:22.50インストール後に検索しても見つからないから、変だな〜と思った。
Direct Compute の Example が動かないのもよくわからん。
0721デフォルトの名無しさん
2011/06/06(月) 09:41:55.580722デフォルトの名無しさん
2011/06/06(月) 13:37:59.320723デフォルトの名無しさん
2011/06/06(月) 13:54:50.090724デフォルトの名無しさん
2011/06/06(月) 21:16:41.99それを知らなかったので、PATHの設定やmake実行さえ一苦労でした。
人によっては、ディレクトリの概念が難しい場合もあるかもしれません。
0725デフォルトの名無しさん
2011/06/06(月) 21:26:48.74初心者に手にを得る代物じゃね〜べよ
0726デフォルトの名無しさん
2011/06/06(月) 21:32:48.380727デフォルトの名無しさん
2011/06/08(水) 14:58:17.390728デフォルトの名無しさん
2011/06/08(水) 16:43:07.02OpenCLはOpenGLと同様、自助努力が基本になる
0729デフォルトの名無しさん
2011/06/09(木) 05:38:12.860730デフォルトの名無しさん
2011/06/09(木) 06:45:26.860731デフォルトの名無しさん
2011/06/09(木) 10:03:14.93亀だけど、cufftPlanは対象とするデータのサイズと同じサイズをバッファとして確保するみたい。だから、二倍以上メモリが空いてないとこける。
0732デフォルトの名無しさん
2011/06/09(木) 17:23:54.07そういうのを自分で考えたり調べたりする姿勢や行動
0733デフォルトの名無しさん
2011/06/09(木) 18:41:25.08CUDAプログラミング入門とかフリーであるもの
見たほうがそこらの書籍よりかなり役立つよ
OpenCLは少し大変だけどCPUならdouble型
オプションあるからいいんだよね
CUDAにもあるのかな?
0734デフォルトの名無しさん
2011/06/09(木) 18:58:20.290735デフォルトの名無しさん
2011/06/11(土) 18:37:21.68使っての 並列FPU高速化ってどのくらい大変ですか?
一般のプログラムと英語ドキュメント読解に問題はないです
0736デフォルトの名無しさん
2011/06/11(土) 18:40:16.570737デフォルトの名無しさん
2011/06/11(土) 19:04:07.230738デフォルトの名無しさん
2011/06/12(日) 06:56:49.050739デフォルトの名無しさん
2011/06/12(日) 19:39:36.69CUDAを勉強してみようかと少し取り組んでみましたが、
C言語の基本も知らぬプログラミング初心者なので
C言語から勉強しようと思ってます。
0740デフォルトの名無しさん
2011/06/12(日) 19:49:59.75故にCUDAに触れることも出来なくなる。
0741デフォルトの名無しさん
2011/06/12(日) 19:55:24.780742デフォルトの名無しさん
2011/06/12(日) 20:12:36.74C言語の文法を知ることは必須としても、まずプログラムがどういうもので何が出来るものなのか、って視点からも眺めた方がいいよ。
文法を覚えることに溺れないように
0743デフォルトの名無しさん
2011/06/12(日) 20:32:53.39統計解析ソフトだから乱数の生成は得意だと思うけど、
ループが苦手だから本当にモンテカルロが必要な場面でが役立たずかも。
Rユーザから見ると、curandライブラリがRから呼び出せたら嬉しい。
0744デフォルトの名無しさん
2011/06/13(月) 16:13:05.080745デフォルトの名無しさん
2011/06/13(月) 17:09:26.420746デフォルトの名無しさん
2011/06/13(月) 20:52:40.31/ \
【資料見た?】 NO → なら、ねぇよ
\
NO → 死ね
0747デフォルトの名無しさん
2011/06/13(月) 22:35:02.62頭痛いです。
0748デフォルトの名無しさん
2011/06/13(月) 22:46:41.93/ \
【探した?】 NO → それより僕と踊りませんか?
\
NO → 死ね
0749デフォルトの名無しさん
2011/06/13(月) 22:56:54.15くだすれなんだし並列に書いてみてはどうか
0750デフォルトの名無しさん
2011/06/13(月) 23:03:14.51ifが多すぎてパフォーマンスが落ちると思う
0751デフォルトの名無しさん
2011/06/13(月) 23:07:55.03嫁>>1
0752デフォルトの名無しさん
2011/06/14(火) 09:21:34.640753デフォルトの名無しさん
2011/06/14(火) 09:45:27.25日本語の書籍購入するのが早いとおもうよ。
・ASCII.technologies (アスキードットテクノロジーズ) 2010年 08月号
・はじめてのCUDAプログラミング
・CUDA by Example 汎用GPUプログラミング入門
この辺がお勧め。自分はこの順で読んだ。
でもプログラミングガイドとか読むのは避けられないはず。
・CUDAプログラミング実践講座 - 超並列プロセッサにおけるプログラミング手法
これは買ってないけど良さそう。
0754デフォルトの名無しさん
2011/06/14(火) 09:47:57.13上から三つ目は一つ目と二つ目が参考資料だっからな
アスキーは触りとしてはわかりやすかった
どれにしろFermiについて書いてないけどね
0755デフォルトの名無しさん
2011/06/14(火) 09:55:59.010756デフォルトの名無しさん
2011/06/14(火) 20:59:49.06AMDが発表したメインストリームAPU「Llano」のアーキテクチャ
ttp://pc.watch.impress.co.jp/docs/column/kaigai/20110614_452948.html
>もう1つのバスは「Fusion Compute Link (Onion)」で、こちらはCPUとのコヒーレントバスになっている。
>このバスを使うと、GPUコアが、これまでできなかったCPUキャッシュへのスヌープをできるようになる。
>このOnionバスを使うことで、GPUでの汎用コンピューティング時に、
>CPUとGPUの間での無駄なメモリコピーを排除するゼロコピーが実現できる。
0757デフォルトの名無しさん
2011/06/14(火) 22:14:25.100758デフォルトの名無しさん
2011/06/14(火) 22:50:31.080759デフォルトの名無しさん
2011/06/14(火) 22:59:19.26計算の方が時間かかってるんだけど
もっと最適化できそうとか
0760デフォルトの名無しさん
2011/06/15(水) 01:56:49.380761デフォルトの名無しさん
2011/06/15(水) 14:58:40.10助けてください。
0762デフォルトの名無しさん
2011/06/15(水) 16:01:33.65720
0763デフォルトの名無しさん
2011/06/15(水) 17:10:58.62レスありがとうございます。
cutilとshrUtilsは既にビルドはしているんですが、開けないとでるんですよ。。
0764デフォルトの名無しさん
2011/06/15(水) 21:41:20.65メモリコピーの時間がかかってでも並列演算したいって問題がほとんどだから
別におかしくないんじゃない。
並列計算が早すぎてメモリコピーがネックのように見えるってことはあるけど。
0765デフォルトの名無しさん
2011/06/15(水) 22:05:59.38同時にはアクセスできないよね、きっと。
でも大した問題でもないか。
0766デフォルトの名無しさん
2011/06/16(木) 03:21:57.35http://blogs.nvidia.com/2011/06/microsoft-going-all-in-on-gpu-computing/?sf1642229=1
0767デフォルトの名無しさん
2011/06/16(木) 06:03:44.90アスキーは2冊あって2010/08はFermiについても掲載されてるよ
そのほかにOpenCL、Direct compute, ATI streamについても載ってる
0768デフォルトの名無しさん
2011/06/16(木) 15:34:19.10今大学でCUDAの勉強を始めた者です。
現在貸し出されたPCに入っていたのがGeForce7600GSで、
CUDAの公式を見る限りだとCUDAに対応していないようですが
どうにかして動かすことは出来ないでしょうか?
一応一番初期のVer1.0〜最新の4.0までひと通り試してみたのですが
サンプルのdeviceQueryを動かしてもCUDA Driver Versionは0.0のままでした。
もし可能でしたら少しでも助言を・・・
無理なら無理と言ってくだされば大学側に申請して
可能なGPUの購入を考えてます。
0769デフォルトの名無しさん
2011/06/16(木) 15:37:06.49対応してない物が仮に動いたとしても、それが正確な動作かわからんから勉強にならんぞ
0770デフォルトの名無しさん
2011/06/16(木) 15:54:45.30Amazon EC2
GPU買わなくてもおk
0771デフォルトの名無しさん
2011/06/16(木) 16:13:17.23安いのなら1万切るから自分でグラボ買ってもいいんじゃない?
0772デフォルトの名無しさん
2011/06/16(木) 16:48:57.04大学研究室ならキッチリ金出してくれて580くらいはポンと買ってくれるだろうから自分で買う意味ないだろう
0773デフォルトの名無しさん
2011/06/16(木) 18:37:25.670774デフォルトの名無しさん
2011/06/16(木) 22:59:07.51授業?独学?
研究のわけないよね
GPGPUの研究でDX9のGPUとか遠回しなアカハラだw
0775デフォルトの名無しさん
2011/06/17(金) 04:29:44.26それかここのpdfがお勧め
ttp://accc.riken.jp/HPC/training.html
GPUと大量の配列使うならCUDAが楽だよねえ
自作のテキストだけど、インストール部分だけなら役に立つかもしれない
ttps://docs.google.com/viewer?a=v&pid=explorer&chrome=true&srcid=0B3RsNc5-fK5OZWY4MDVjYTItOTM2MS00NDc3LWE3NjEtZTc5YmQyMjg4Y2Q4&hl=ja&authkey=CJHji7AH
0776デフォルトの名無しさん
2011/06/17(金) 07:42:07.25横からだが、これは面白そう 読んでみるわ
0777デフォルトの名無しさん
2011/06/17(金) 10:15:21.27全部印刷しておいても良いかもしれん
0778デフォルトの名無しさん
2011/06/19(日) 02:32:29.65参考になりそうだ。DLできないの?
0779デフォルトの名無しさん
2011/06/19(日) 10:39:21.500780デフォルトの名無しさん
2011/06/19(日) 22:33:54.54�@felmiにはキャッシュが搭載されているみたいですがプログラム中で明示しなくてもキャッシュにより高速化するのですか?
�AGPUでcuda化して高速化が見込めるアプリケーションが載ってるサイトとか知りませんか?ソースコードつきで。できればC言語だとうれしいです
0781デフォルトの名無しさん
2011/06/19(日) 22:41:52.94だれかこたえてあげれや
0782デフォルトの名無しさん
2011/06/19(日) 22:54:32.14使いこなすのは無理だよ。
0783デフォルトの名無しさん
2011/06/19(日) 23:07:34.93一応は調べました
せめて�@だけでも答えてもらえると助かります
�Aに関しては自分もいくつか見つけましたが他の方がどういったサイトを使ってるか気になったので聞いてみました
0784デフォルトの名無しさん
2011/06/19(日) 23:08:08.80綴り間違い-10
sage+3
投稿時間+10
内容に言及しなくても-7点だ
もうちょいがんばれ
0785デフォルトの名無しさん
2011/06/19(日) 23:31:02.350786デフォルトの名無しさん
2011/06/19(日) 23:52:40.95プッ
0787デフォルトの名無しさん
2011/06/19(日) 23:52:54.57キャッシュによって高速化(何と比べて?)するかはアプリによる。
ttp://journal.mycom.co.jp/articles/2010/07/21/fermi_cache/index.html
こんな話もあるから、キャッシュが如何使われるか位は
意識しないと駄目だし、究極を目指せば結局コアレスアクセスに
なってキャッシュの意味は無くなるかも知れない。
で、そこまで頑張れば微妙にteslaの方が速くなる。
0788デフォルトの名無しさん
2011/06/20(月) 02:53:19.67びっくらこいた
高速化できる例はCUDA SDKのサンプルみればよろし
0789デフォルトの名無しさん
2011/06/20(月) 02:55:44.71G80/G92のように総レジスタ数の少ないGPUはもっと遅くなる。大体1/10
0790デフォルトの名無しさん
2011/06/20(月) 07:37:03.230791デフォルトの名無しさん
2011/06/22(水) 16:10:37.96どういう原理なのかご存知の方がいらっしゃいましたら教えていただけませんか。
また、ご存知でしたらその記載がどこにあるか教えていただけませんか。
お願いします。
0792デフォルトの名無しさん
2011/06/22(水) 18:40:23.890793デフォルトの名無しさん
2011/06/23(木) 15:54:07.78ありがとうございます。
0794やむちゃ
2011/06/24(金) 22:13:50.670795デフォルトの名無しさん
2011/06/30(木) 16:21:00.38HD6990単体で362Mtrip/sくらいだよな。
単純に倍なら722Mtrip/s?
0796デフォルトの名無しさん
2011/06/30(木) 19:12:34.59GPUを使う場合ってターゲットのnVidia、AMD(ATI)GPU用専用バイナリコードを作成するの?
(両方で動くようにするには両方用のバイナリを用意しないと駄目?)
それともJAVAのbytecodeみたいなのを作成して、それで両方のnVidia、AMD(ATI)GPUに対応
出来るようになっているの?
0797デフォルトの名無しさん
2011/07/01(金) 03:02:06.82→ (ネイティブコードB) → [チップ]GTX560
→ (ネイティブコードC) → [チップ]……
→ [ATIドライバー] → (ネイティブコードD) → [チップ]知らね
→ (ネイティブコードE) → [チップ]
→ (ネイティブコードF) → [チップ]
0798デフォルトの名無しさん
2011/07/01(金) 03:10:03.120799デフォルトの名無しさん
2011/07/01(金) 09:31:07.350800デフォルトの名無しさん
2011/07/02(土) 00:44:21.66Javaみたいな抽象度の高いバイトコードも解析が容易だしで、
小手先のノウハウで大幅に速度や精度が変わっちゃう商売には使いづらい世の中になってきた。
0801デフォルトの名無しさん
2011/07/03(日) 00:25:09.93CUDAプログラミングでVRAMの大きさって気にする必要がある?
まぁ、作るプログラムによるんだろうけど、1GB以上食うのはたぶん確実。
GTX 560とかで2GBのカードもあるので、差が出るならそっちを選んでおきたいかな…と。
0802デフォルトの名無しさん
2011/07/03(日) 01:03:13.540803801
2011/07/03(日) 07:38:06.67やっぱし計算に必要な分はメモリがいるのね。
ちと調べたら
CUDA4.0からUnified Virtual Addressing(UVA)なるものが登場していて、
コード的にはVRAM容量は気にしなくていいみたいだけど、
性能的には差が出る、と理解しておきます。
0804デフォルトの名無しさん
2011/07/03(日) 18:11:44.57研究で有限要素法を使っていますが、あまりにも計算に時間がかかるので、GPUを使えればと思っています。
とりあえず、使っているプログラム(fortran)がGPUで処理できるような内容か否かを知りたいのですが、
CUDA入門書などを読んでも、プログラミングについて書いてあるだけで、
具体的にどのような処理が不可能かは見つけられませんでした。
プログラム自体は、ソースを書き換えれば並列化が可能です。
cpuに出来てGPUに出来ない処理とは具体的にどのような処理なのでしょうか?
gpuは単精度しか扱えないとの記述を見ましたが、単純にreal*8で定義した変数をreal*4に書き換えることで
gpuでも可能になるということでしょうか?
長文失礼しました、ご回答よろしくお願いします。
0805デフォルトの名無しさん
2011/07/03(日) 18:32:01.01http://developer.nvidia.com/cuda-toolkit-40
http://gpu.fixstars.com/index.php/CUDA_%E6%96%87%E6%B3%95_1
http://www.kumikomi.net/archives/2008/10/22gpu2.php
http://tech.ckme.co.jp/cuda.shtml
http://www.codelogy.org/quillpen/?entry=62
0806デフォルトの名無しさん
2011/07/03(日) 22:17:41.750807デフォルトの名無しさん
2011/07/04(月) 18:01:43.46ど素人だけど
有限要素法がCUDAに向くかグルグルしてみる。
Fortran用いてCUDAで有限要素法している大学の研究室あるカモ
0808デフォルトの名無しさん
2011/07/04(月) 18:14:51.81逐次処理なんだからさ
0809デフォルトの名無しさん
2011/07/04(月) 21:19:29.48SOR法のような逐次処理であっても
並列化することは可能。
共役勾配法ならふつうに並列化できそう。
0810デフォルトの名無しさん
2011/07/04(月) 22:15:17.340811804
2011/07/04(月) 23:13:48.80>>805
ちょっとまだ読めていませんが、頑張って読ませていただきます。
>>807-810
なるほど、ググってみます。
ただ、私が使っているのは、構成式の計算が非常に複雑なのと、
多相混合体を計算しているので、その部分が逐次処理でも早くなるのでは??と考えています。
あと、12スレッドでCPUを100%使用してくれているので、
単純にマルチスレッドで行けるのかなと甘い考えを持っていました。
0812デフォルトの名無しさん
2011/07/05(火) 19:50:07.25はじめてのcudaプログラミングに書いてある方法では奇数はおろか2の乗数の配列しか扱えないんで
0813デフォルトの名無しさん
2011/07/05(火) 21:06:02.22もっと良い方法ありそうだけど,ホスト側に転送して計算させるよりは圧倒的に高速になるから,
深く考えずに実装してる…
0814デフォルトの名無しさん
2011/07/05(火) 23:14:25.80GPUComputingSDKのreductionのソースを参考に改良。
0815デフォルトの名無しさん
2011/07/06(水) 20:49:49.53バイトニックソートの時は邪魔にならんようにめちゃくちゃでかい数セットしてるwwwwww
0816デフォルトの名無しさん
2011/07/07(木) 02:20:00.74Optimizing Parallel Reduction in CUDAの#6
(Kernel 6)までしか訳してない。
#7(Kernel 7)でさらに速くなり、データ数の上限
もアップする。
0817デフォルトの名無しさん
2011/07/07(木) 12:23:25.6310000000
1000000
100000
100
1
づつreductionさせればどんな要素数でもできそうだけどめんどくさそう
0818デフォルトの名無しさん
2011/07/07(木) 13:12:12.120819デフォルトの名無しさん
2011/07/07(木) 14:28:48.64例えば約百万回の実行を行うのに1024*1024と512*512を4回では環境によって
前者が速かったり後者が速かったりするんだけど
実測。
0821デフォルトの名無しさん
2011/07/07(木) 18:54:52.10返答サンクス
でも意味不明w
0822デフォルトの名無しさん
2011/07/08(金) 07:01:46.60自分で答え書いてると思うけど,環境によって最適値が変わるので自分の環境で試行錯誤して最適値を決定するのが一番いいと思います
0823デフォルトの名無しさん
2011/07/08(金) 09:47:06.220824デフォルトの名無しさん
2011/07/11(月) 21:28:58.35まぁまぁ速い?
http://www.youtube.com/watch?v=3dEN9JQ3R0U
0825デフォルトの名無しさん
2011/07/11(月) 22:20:08.16各ピクセルにデータいれたものをテクスチャとして渡して、ピクセルシェーダになんか演算させて、画像として帰ってきたものを1ピクセルずつ読み取るってかんじ?
0826デフォルトの名無しさん
2011/07/11(月) 22:23:45.830827デフォルトの名無しさん
2011/07/11(月) 22:47:01.34早いのか遅いのかわからんwww
そもそもは何年ごろに確立されたの?
0828デフォルトの名無しさん
2011/07/11(月) 22:56:54.62(主流はOpenCLとかになりそうだけど)
0829デフォルトの名無しさん
2011/07/11(月) 23:22:50.92提唱時期は知らんが、>>825 はGPGPUの「古典」だ
0830デフォルトの名無しさん
2011/07/11(月) 23:37:20.93thx
その古典についてもうちょい詳しく知りたい
0831デフォルトの名無しさん
2011/07/12(火) 00:05:46.80http://www.4gamer.net/games/076/G007660/20090206031/
大まかな話ならこのあたりとかじゃないかな。
0832デフォルトの名無しさん
2011/07/12(火) 00:27:14.22ありがとう!
スレ違いすまん
しかしCUDAはそういうの意識せずにできるんだな
プログラマブルシェーダってなに?ってレベルだったけどちゃんと高速化できたわ
GPUの構造は多少知らなきゃいけないけど
0833デフォルトの名無しさん
2011/07/12(火) 09:38:30.960834デフォルトの名無しさん
2011/07/12(火) 14:31:44.47MINPACKがそれらのBLASで動けば、非線型最適化が瞬時に出来るかも
しれないのに。
0835デフォルトの名無しさん
2011/07/13(水) 11:28:00.27atomic関数を使った.cuソースのコンパイルがどうしても通りません。
"identifier atomicAdd is undefined"が出てしまいます。
環境はvisual studio 2008でGTX480(cc = 2.0)です。
cudaフォーラムにあった、
nvcc -arch sm_11オプションコンパイルをVSコマンドプロンプトからコンパイルしようとしても
"Visual Studio configuration file '(null)'could not be found for instration at ~~~
"というエラーメッセージが出てしまいます。
解決方法ご存知の方いませんでしょうか。
0836デフォルトの名無しさん
2011/07/13(水) 11:28:19.09atomic関数を使った.cuソースのコンパイルがどうしても通りません。
"identifier atomicAdd is undefined"が出てしまいます。
環境はvisual studio 2008でGTX480(cc = 2.0)です。
cudaフォーラムにあった、
nvcc -arch sm_11オプションコンパイルをVSコマンドプロンプトからコンパイルしようとしても
"Visual Studio configuration file '(null)'could not be found for instration at ~~~
"というエラーメッセージが出てしまいます。
解決方法ご存知の方いませんでしょうか。
0838デフォルトの名無しさん
2011/07/13(水) 11:48:11.690839デフォルトの名無しさん
2011/07/13(水) 14:05:01.64atomicAddとか使うとき-arch sm_11じゃなく、-arch=sm_11 としているけど
0840デフォルトの名無しさん
2011/07/13(水) 16:52:49.53sm_11_atomic_functions.h をインクルードしてみるとか。
0841デフォルトの名無しさん
2011/07/13(水) 21:38:27.190842デフォルトの名無しさん
2011/07/14(木) 21:10:56.94著者はこのスレの住人なんですね。32人33脚の例の話とか過去ログ読んで経緯理解したよ…
しかしこの本、意味不明な挿絵でいちいち吹くんだが…暗闇からのニョッキリ手ってなんだよw
0843デフォルトの名無しさん
2011/07/14(木) 22:25:02.91日本語で勧められるのはcuda by Exampleが出版される前は
「はじめてのCUDAプログラミング」以外なかった。
今でも内容がfermiにあっていないところは時代遅れだが
初学者にとってリダクションのチューニングの演習過程は参考になる。
0844デフォルトの名無しさん
2011/07/14(木) 23:32:04.91別に本人が勧めてるわけじゃないだろ
アーキテクチャとかはPCゲームの記事が詳しかったな確か
0845デフォルトの名無しさん
2011/07/15(金) 19:32:35.44私もCUDA by Example読み始めて同じところで引っかかったのですが、
今検索するとforumでいくつかヒットしますね。
http://forums.nvidia.com/index.php?showtopic=202480
ただ私の環境では、コンパイル通っても実行すると
> unknown error in CudaByExample/chapter04/julia_gpu.cu at line 89
なんてエラーが出てさっぱりですが…
0846デフォルトの名無しさん
2011/07/18(月) 13:34:23.28もしくは質問者自身何が何だが分からない質問を勇気を持って書き込むスレッドです。
CUDA使いが優しくコメントを返しますが、
お礼はCUDAの布教と初心者の救済をお願いします。
CUDA・HomePage
http://www.nvidia.com/cuda
CUDA Zone
http://www.nvidia.co.jp/object/cuda_home_jp.html
NVIDIA GPU Computing Documentation
http://developer.nvidia.com/nvidia-gpu-computing-documentation
CUDA GPU Computing Forums
http://forums.nvidia.com/index.php?showforum=62
関連スレ
GPGPU#5
http://hibari.2ch.net/test/read.cgi/tech/1281876470/
前スレ
【GPGPU】くだすれCUDAスレ【NVIDIA】
http://pc12.2ch.net/test/read.cgi/tech/1206152032/
【GPGPU】くだすれCUDAスレ pert2【NVIDIA】
http://pc12.2ch.net/test/read.cgi/tech/1254997777/
【GPGPU】くだすれCUDAスレ pert3【NVIDIA】
http://hibari.2ch.net/test/read.cgi/tech/1271587710/
【GPGPU】くだすれCUDAスレ pert4【NVIDIA】
http://hibari.2ch.net/test/read.cgi/tech/1291467433/
0847デフォルトの名無しさん
2011/07/18(月) 13:35:15.94CUDAテクニカルトレーニングI/II
http://www.nvidia.co.jp/docs/IO/59373/VolumeI.pdf
http://www.nvidia.co.jp/docs/IO/59373/VolumeII.pdf
GPUコンピューティング研究会(各回の資料リンク)
http://gpu-computing.gsic.titech.ac.jp/Japanese/Lecture/index.html
ソフテック・PGIテクニカル情報・コラム(CUDAに関する話題複数あり)
http://www.softek.co.jp/SPG/Pgi/TIPS/para_guide.html
CUDAベンチ
http://wataco.air-nifty.com/syacho/2008/02/cuda_2044.html
KNOPPIX for CUDA
http://www.yasuoka.mech.keio.ac.jp/cuda/
0848デフォルトの名無しさん
2011/07/18(月) 13:36:26.200849デフォルトの名無しさん
2011/07/18(月) 14:40:05.02CUDAフォーラム(日本語): http://forum.nvidia.co.jp
ってなくなっちゃったの?
http://www.nvidia.co.jp/object/io_1223650778428.html
http://www.okwave.co.jp/news/press_log/2008/1010.html
0850デフォルトの名無しさん
2011/07/18(月) 14:46:05.65基本、学術計算的なことじゃない限り利用価値があまりないし
英語できる奴は本家のForumで事が足りるし、そうでない奴はここに聞きに来ている悪寒。
0852デフォルトの名無しさん
2011/07/18(月) 23:24:52.800853デフォルトの名無しさん
2011/07/19(火) 00:25:03.510854デフォルトの名無しさん
2011/07/19(火) 00:30:48.550855デフォルトの名無しさん
2011/07/19(火) 00:32:52.240856デフォルトの名無しさん
2011/07/19(火) 03:18:04.57どういった理由ででしょうか?
0857デフォルトの名無しさん
2011/07/19(火) 19:19:37.60DSP使ったり、FPGAで並列回路組んでたのをこっちに乗り換えてヤッホーイな所ってあるのかな
0859デフォルトの名無しさん
2011/07/20(水) 01:59:01.950860デフォルトの名無しさん
2011/07/20(水) 02:21:10.61VS使わないフリー環境でのコンパイル方法を紹介してる
0861デフォルトの名無しさん
2011/07/20(水) 07:36:42.98早過ぎる最適化は怖くないんですか?
カリカリの最適化した後に、やっぱり違うやり方の方がいいって思っても
大きな変更は億劫になってしまいそうです。
0862デフォルトの名無しさん
2011/07/20(水) 07:54:30.98納期に追われると特に
でも、遊びでやるなら楽しいんじゃないでしょうか
パズルみたいで
0863デフォルトの名無しさん
2011/07/20(水) 11:40:51.26VC Express 使えば良いじゃない。
0864デフォルトの名無しさん
2011/07/20(水) 13:29:38.030865デフォルトの名無しさん
2011/07/20(水) 14:14:21.74compute capability 2.xを満たし、念のためにconcurrentKernelsも確認してConcurrent Kernel Executionが使えることは確認しました。
しかし、いまいち使い方がわかりません。
カーネルを2つ連続で書けば自動的に行われるのか、それとも何かフラグが必要なのか。
また、何か特殊な関数が必要なのか。
どなたか、ご存知の方がいらっしゃいましたらご指導をお願いします。
0866デフォルトの名無しさん
2011/07/20(水) 15:46:27.09Stream
0867デフォルトの名無しさん
2011/07/20(水) 17:12:02.58ありがとうございます。
プログラミングガイドのStreamとその周辺を読んでみたいと思います。
0868デフォルトの名無しさん
2011/07/21(木) 14:37:05.43CLIは必須じゃない
0869デフォルトの名無しさん
2011/07/21(木) 17:22:50.43ブロック、スレッドというのは、カーネル関数を実行するときの<<<M,N>>>で指定すると思うのですが、
グリッドというのは何なんでしょうか?
他の入門用説明を読んでも、1つしかないことを前提に書かれているようですが…
0870デフォルトの名無しさん
2011/07/21(木) 17:29:34.96とにかく用語がおぼえにくいからね
0871デフォルトの名無しさん
2011/07/21(木) 18:02:37.66ブロックの集まりがグリッド。
どちらも論理概念でハードウェア資源と完全に同一ではないので注意。
0872デフォルトの名無しさん
2011/07/22(金) 21:00:42.880873デフォルトの名無しさん
2011/07/22(金) 21:06:35.84http://www.nv-jp-event.jp/gtc-workshop2011/
参加された方、感想をぜひ。
0874デフォルトの名無しさん
2011/07/23(土) 00:52:21.610875デフォルトの名無しさん
2011/07/23(土) 12:51:02.38基調講演だけじゃ…
0876デフォルトの名無しさん
2011/07/23(土) 15:36:44.71boolが1にとは限らなかったりする?
int selector = value1 < value2
int out = value1 * (1 - selector) + value2 * selector
0877デフォルトの名無しさん
2011/07/23(土) 15:49:01.44これって、if-else文より速くなるの?
0878デフォルトの名無しさん
2011/07/23(土) 15:51:24.83GPUでやった事はないけど、
x86,x64じゃ比較演算子は分岐命令が交じる。
0879デフォルトの名無しさん
2011/07/23(土) 16:54:15.000880デフォルトの名無しさん
2011/07/23(土) 17:17:15.22勉強がてらPTXを見てみた。誤ってたら指摘よろしく。
http://developer.download.nvidia.com/compute/DevZone/docs/html/C/doc/ptx_isa_2.3.pdf
int selector = value1 < value2; の部分:
setp.gt.s32 %p1, %r1, %r2;
selp.s32 %r3, 1, 0, %p1;
%r1,%2がそれぞれvalue1,value2,%r3がselector。
selectorが0か1になるのはC++言語仕様通り保証されてるね。
"setp.gt.s32"の部分だけで
int out = value1 < value2 ? value2 : value1;
と変わらないので、>>876のそれ以外は全部無駄だね。
0881880
2011/07/23(土) 17:48:23.97%r2がvalue1, %r1がvalue2だった。
0882デフォルトの名無しさん
2011/07/23(土) 17:51:07.240883デフォルトの名無しさん
2011/07/23(土) 18:15:22.010884デフォルトの名無しさん
2011/07/23(土) 21:28:51.28100MTrips/s以上速くなったw GTX 580で640MTrips/s出てる。
こりゃすごいわ。
0885デフォルトの名無しさん
2011/07/23(土) 21:35:26.38というか、普通に書けばMAXを取る命令が使われるのでそもそも分岐命令なぞ入らないが。
0886デフォルトの名無しさん
2011/07/23(土) 21:46:23.580887デフォルトの名無しさん
2011/07/23(土) 21:48:09.55他のアプリで速度向上したりしてます?
0888デフォルトの名無しさん
2011/07/23(土) 22:59:15.99他のアプリは試していないのでちょっとわかりません。
あまりにも嬉しかったのでスレを立ててしまったw
【トリップ検索】CUDA SHA-1 Tripper【GeForce】
http://hibari.2ch.net/test/read.cgi/software/1311428038/
0889デフォルトの名無しさん
2011/07/24(日) 05:14:37.53そんな後出しでドヤ顔されても…
0890869
2011/07/24(日) 15:59:56.54> ブロックの集合体であるグリッド(Grid)は、デバイス(Device)に対応しています。
とあるので、簡単に言うとGPUカード2枚挿すとグリッド数は2、ということでしょうか。
0891デフォルトの名無しさん
2011/07/24(日) 17:08:31.45グリッドとは一回のカーネル呼び出しで実行される単位みたいなもの。
0892デフォルトの名無しさん
2011/07/30(土) 17:13:04.480893デフォルトの名無しさん
2011/07/31(日) 12:47:17.83注意するところとかはまるポイントある?
0894デフォルトの名無しさん
2011/08/01(月) 18:07:05.670895892
2011/08/01(月) 20:24:42.14894の言うようにググれば出てくる。
VC2010 Express CUDA4 の3つのキーワードで最初に出てくるサイトを参考にした。
変えたのは以下の点
C:\Program Files (x86)\MSBuild\Microsoft.Cpp\v4.0\BuildCustomizations\CUDA 4.0.props
を修正するところ
上記のサイトでは2行コメントアウト1行追加だが、俺は46行目の後にこれを追加した
<CudaClVersion Condition="'$(PlatformToolset)' == 'Windows7.1SDK'">2010</CudaClVersion>
あと、サンプル実行は適当にググってヒットした、CUDA用のHelloWorld。
上記のサイトからリンク辿れるサンプルはそのままでは動かなかった。
(俺の探し方が悪いのかもしれんが)
0896デフォルトの名無しさん
2011/08/03(水) 18:26:44.868月4日に秋葉原で行われるサードウェーブ主催の「GPGPUセミナー2011 夏」に協賛し、「NVIDIA GPU コンピューティング: CUDA 4.0による並列プログラミング」と題して講演します。http://bit.ly/l1ozuB
7月29日 TweetDeckから
0897デフォルトの名無しさん
2011/08/05(金) 13:31:24.47一応できたので書いておきます。
compute capabilityの認識のさせ方
[visual studio 2008を利用する場合]
プロジェクトのプロパティ
→構成プロパティ
→CUDARUNTIMEAPI(カスタムビルド規則で設定したルールの名前)
→GPU
→GPU Architecture の欄に「sm_21」などと追加(←はcompute capabilityが2.1の場合)
[コマンドプロンプトで実行する場合]
�@nvcc -arch=sm_21 hello.cu -o hello などと入力。
�Ahello ←これを実行すると出力がでる
hello.cuは実行したいファイルの名前。もちろん人によって変わる。
-oはコンパイル・リンク後に生成する実行形式ファイルの設定。デフォルトではa.out だったような。
0898デフォルトの名無しさん
2011/08/08(月) 20:42:38.80460か560Tiくらいあったほうがいい?
0899デフォルトの名無しさん
2011/08/08(月) 20:55:39.640900デフォルトの名無しさん
2011/08/08(月) 21:06:15.160901デフォルトの名無しさん
2011/08/08(月) 22:14:03.770902デフォルトの名無しさん
2011/08/08(月) 22:38:23.360903デフォルトの名無しさん
2011/08/09(火) 06:53:52.22普通、プログラミングの基礎を固めた方がいい初心者とは真逆じゃない?
最初から最適化を過度に意識し過ぎると、貧乏性の汚いコードが身についちゃいそう。
0904デフォルトの名無しさん
2011/08/09(火) 16:57:38.45GTX520で貧弱とか何贅沢なことを…と思ったけどGT520のことかな?
勉強なら十分、て言葉を否定する訳じゃないけど、
http://developer.nvidia.com/cuda-gpus
にあるCompute Capabilityも少し気にした方が良い、かも。
0905デフォルトの名無しさん
2011/08/09(火) 17:54:32.00その一覧あってないような気がする
うちのGT430はSDKのサンプルでちゃんとCC2.1と表示されるよ
同一世代ならCCは一緒で性能が違うものだったと思う。
0906デフォルトの名無しさん
2011/08/09(火) 19:57:22.62おお、それは初めて聞く情報!
グラボ詳しくないので同一世代がどれとか知りませんでした…ありがとうございます。
0907デフォルトの名無しさん
2011/08/09(火) 21:29:29.440908デフォルトの名無しさん
2011/08/09(火) 23:29:58.63これって460のほうが580よりCudaに関しては性能上ってこと?
0909デフォルトの名無しさん
2011/08/10(水) 00:58:59.110910デフォルトの名無しさん
2011/08/10(水) 01:07:30.442.0と2.1でアーキテクチャ違うのかな
0911デフォルトの名無しさん
2011/08/10(水) 01:35:07.79で、460と480のフルスペック版が560Tiと580
0912デフォルトの名無しさん
2011/08/10(水) 01:49:07.11SMあたりのコア数やスケジューラの変更ではないでしょうかね。
ttp://pc.watch.impress.co.jp/docs/column/tawada/20100712_380174.html
> CC2.1についてNVIDIAに問い合わせたところ、
> The compute capability is indeed called 2.1 (versus 2.0 in GF100), but it’s the same
> program functionality as Compute 2.0 spec.
> The newer version number simply tells the complier to optimize for GF104 SM design.
> という回答をもらった。要するに「GF100のCC2.0と同じ機能を持つが、
> SMのCUDA Core数などが変化したことによるコンパイラの最適化に適応させたもの」ということで
という話もありますし。
0913デフォルトの名無しさん
2011/08/10(水) 03:06:28.120914デフォルトの名無しさん
2011/08/10(水) 08:02:15.620915デフォルトの名無しさん
2011/08/10(水) 16:40:00.150916デフォルトの名無しさん
2011/08/11(木) 09:35:32.07Tool Kitは動作してサンプルも動くのに・・・
悔しいけど、3.2にロールバックした ><
0917デフォルトの名無しさん
2011/08/12(金) 09:34:32.250919デフォルトの名無しさん
2011/08/12(金) 21:04:48.830920デフォルトの名無しさん
2011/08/16(火) 01:37:33.400921デフォルトの名無しさん
2011/08/16(火) 02:25:21.160922デフォルトの名無しさん
2011/08/16(火) 03:24:06.31それぐらい自分で確認しろよ
ttp://developer.download.nvidia.com/compute/cuda/3_2_prod/toolkit/docs/EULA.txt
と思ったけど、なんか4.0のは見当たらないね・・・
インストールすればライセンスのファイルも入るのかな・・・
0923デフォルトの名無しさん
2011/08/16(火) 10:10:43.62nVidiaもATIも足並みを揃える気無いだろ
技術を相手に公開しなきゃならんし、自社でライブラリをガンガン開発してったほうが効率良いし
0924デフォルトの名無しさん
2011/08/16(火) 11:24:21.15意思の疎通は必ずしも密ではないのではないかと思う。まして、OpenCL対応は……
0925デフォルトの名無しさん
2011/08/17(水) 16:01:16.04今のところ470か480あたりを買おうと思っていますが,CCを重視して560にしたほうが良いという場合は有りうるでしょうか
また,メーカーはELSA安定だとは思うのですが,他の安いメーカー製にした場合問題が発生しやすい等有るのでしょうか
0926デフォルトの名無しさん
2011/08/17(水) 18:55:46.310927デフォルトの名無しさん
2011/08/17(水) 22:25:00.02それにCompute Capabilityは2.0と2.1だと違いがないんじゃないの?
0928デフォルトの名無しさん
2011/08/17(水) 22:27:04.780929デフォルトの名無しさん
2011/08/17(水) 23:19:17.69今さら480を選ぶ理由は価格以外にないと思うが、
それでも570とかにしたほうが幸せな気がする。
メーカーはCUDAに関しては保証してないし交換もきかないからどこでも
一緒なような気がする。
0930デフォルトの名無しさん
2011/08/18(木) 12:12:59.41ELSA: うちのGeForceはNVIDIAに供給しているQuadroと同じ部品を使っているから安くて高品質(キリィッ
他社: うちのGeForceはリファレンスボードの無駄や非効率を改善しているから高性能(キリリィッ
0931デフォルトの名無しさん
2011/08/18(木) 12:22:51.68でもこれ、実際のところ正しいと思うよ。
リファの580をオーバークロックしてCUDAのプログラムを走らせると、
結構な確率でデータが化けるよ。化け方を見ると、明らかに中の演算装置が誤動作してるんだよね。
趣味ならともかく、仕事で使うんなら高いの買わないと駄目でしょ。
0932デフォルトの名無しさん
2011/08/18(木) 12:45:11.47バルタザール メルキオール カスパー
0933デフォルトの名無しさん
2011/08/18(木) 12:59:22.40問題はここじゃね〜のか?普通の頭はそう考えると思うぞ
0934デフォルトの名無しさん
2011/08/18(木) 13:05:39.65> >> リファの580をオーバークロックしてCUDAのプログラムを走らせると、
> 問題はここじゃね〜のか?普通の頭はそう考えると思うぞ
もちろんそうなんだけど、意外にOCの許容範囲が低いなあという印象だったもんで…
0935デフォルトの名無しさん
2011/08/18(木) 13:28:00.76これ間違いじゃねーの?
QuadroじゃなくてTeslaのみだったと思うが。
0936デフォルトの名無しさん
2011/08/18(木) 13:51:17.780937デフォルトの名無しさん
2011/08/18(木) 14:45:44.000938デフォルトの名無しさん
2011/08/19(金) 00:43:52.72町のPCショップとかで買わない方がいいですよとか言っていた
おめーてきとーいってんじゃねえよとぶち切れそうになった
0939デフォルトの名無しさん
2011/08/19(金) 00:59:54.030940デフォルトの名無しさん
2011/08/19(金) 01:00:44.31ttp://www.4gamer.net/games/137/G013737/20110725062/
0941デフォルトの名無しさん
2011/08/20(土) 15:47:24.08動作保証温度とかのテスト結果見れるんなら欲しいのだが。
0942デフォルトの名無しさん
2011/08/20(土) 16:12:43.09geforceよりは厳しいテストをしてますってだけで
そういうデータは公表されてない。
0943デフォルトの名無しさん
2011/08/20(土) 16:15:20.28エラーが出たら無償交換すると言ってるだけでないの?
何がまずいのかさっぱりわからん。
0944デフォルトの名無しさん
2011/08/20(土) 17:44:19.42「CUDA保証します」っていったらCUDAで計算ミスがでて損害が出ても
保証しますって風にも取れるからでしょ。
0945デフォルトの名無しさん
2011/08/20(土) 18:23:35.34(1) カーネルルーチンのコードって実行中はどのメモリに置かれているのですか。
グローバルメモリでないとしたら、サイズ制限が知りたい。
グローバルメモリなら多分命令キャッシュがあると思うので、そのサイズを知りたい。
(2) コンスタントメモリは(一旦コンスタントキャッシュに読み込んでしまえば)
シェアードメモリみたいなバンクコンフリクト無しでかつシェアードメモリ並に低レイテンシなのでしょうか。
(3) あるWarpでシェアードメモリに読み込んだデータは、同一SMの次のWarpで参照できる?
(3) 自スレッドがどのSMで実行されているのか知る方法は?
0946デフォルトの名無しさん
2011/08/20(土) 19:07:38.140947デフォルトの名無しさん
2011/08/20(土) 22:20:49.99(1)と(2)はわからんが(3)はどっちも無理だろうな
次のWarpで参照しようにも次のWarpがどれなのかはスケジューラが毎回適当に振るんだから意味がないでしょ
0948デフォルトの名無しさん
2011/08/20(土) 22:53:29.20(1)何がしたいのか分かりませんが、昔は2M命令という制限があった。よほどのことがない限りは
これを越えないと思います。
(2)低レイテンシですが、バンクが1個しかないので各スレッドが異なるアドレスにアクセスすると
バンクコンフリクト状態になります。
(3.1)質問の意味が分かりません。シェアードメモリを介したデータの受け渡しで反映されるまでの
タイミングに関することでしょうか?そのタイミングを気にする必要がある=__syncthreads()を
入れていない状態と思いますが、あるWarpがシェアードメモリへの書き込みを行い、次にある他の
Warpがシェアードメモリからデータを読み込んだ場合、更新後のデータを受け取れる可能性は
ほぼ0でしょう。
(3.2)%smidという特殊レジスタによって自分がどのSMで実行されているかが分かります。
このレジスタを参照する方法はSDKにあるinlinePTXを参照してください。laneidをsmidに変更するだけです。
0949945
2011/08/21(日) 10:35:01.97レス�dクスだいたい了解
項番ミス失礼。(3.1)は、次の2つの可能性を問うものだったのです
1. コンスタントキャッシュが遅い場合の代替手段。定数はSM別 -- これは、スレッド別定数でなければコンスタントキャッシュでおk理解
2. Warpダイバージェンス対策としてのWarp中断 -- これは誤解だった感じ。できたとしても中断しない場合と同速と思われる。
0950945
2011/08/21(日) 10:42:18.54SM別に処理を分けるはアーキテクチャーを活かした良い方法とは言えませんけども
(TB内のスレッドがどのSMに割り当てられるかは>947の通り予測し難い)
0951デフォルトの名無しさん
2011/08/21(日) 12:42:06.49>>950の方のTB間のシェアードメモリデータの再利用は一応可能です。(Fermiより前は知りませんが)
しかし保証されていませんし、大して速くならないので使わない方が良いと思います。
0952デフォルトの名無しさん
2011/08/21(日) 12:53:46.00PC関係の製品で計算結果の保証なんてどこもやってないじゃん。
HPやDELLやNECの数百万円の鯖を買っても「保証」は無償修理もしくは交換のこと。
つーか、GeforceでCUDA稼働保証があるというその大手ショップの名前を是非教えてほしい。
会社のマシンに試験用に一つ入れたいが、Teslaじゃ稟議が通らない。
0953デフォルトの名無しさん
2011/08/21(日) 12:58:46.13営業の口八丁、もしくは個別契約の餌だろうから聞いても無駄じゃね?
宮崎大みたいに数百とか導入予定なら交換対応してもペイするかもしれんが
1個だけとかじゃそんな話でないでしょ。
0954デフォルトの名無しさん
2011/08/21(日) 13:35:28.11>>938が、そんなでかい金の決定権がある人間のレスに見えるかね?
0955デフォルトの名無しさん
2011/08/21(日) 13:45:08.72というか数百だろうがgeforceでCUDAの演算結果保証してくれるところなんてないと思う。
あるならホント知りたいもんだわ
0956デフォルトの名無しさん
2011/08/21(日) 14:01:06.87ミ川ミミ川三ミ川ミミ
ミ ,.--"゛~~~~''´゛゛゛ミ
ミ川 _ _ l.
ミ川 -´==(`r__i ェl `r
ミ川 ̄ ̄'、_,r' 丶Lノ
ミ川 /( __)) |
ミ川 r―''"入 |
|i 廷廾ニツ j
|ー、____,.ノ
▂▃▃▃▂ ヽ∧∧∧ /
. ██████< 角刈り!> /
███ け ,fj / ∨∨V ヽ /
`█、"ヮノ | /
/⌒ヽヾ'リ、 | /
. { j`ー' ハ // ヽ∧∧∧∧∧∧∨/
k〜'l レヘ. ,r'ス <ソイヤ!ソイヤ!>
| ヽ \ ト、 ヽ-kヾソ < 角刈りソイヤ!>
. l \ `ー‐ゝ-〈/´ / ∨∨∨∨∨∨ヽ
0957デフォルトの名無しさん
2011/08/21(日) 16:51:38.56不良の検証を客に押しつけられるなら、理論上は2割増しの値段で売れば動作保証は可能だ。
0958デフォルトの名無しさん
2011/08/21(日) 16:58:50.99プログラムの実行結果は間違ってるけれど、エラーにはならずに正常終了した風に終わるってこと?
0959デフォルトの名無しさん
2011/08/21(日) 17:03:46.290960デフォルトの名無しさん
2011/08/21(日) 17:45:34.55なんか、計算機の根幹を揺るがすような話だなw
0961デフォルトの名無しさん
2011/08/21(日) 18:03:08.290962デフォルトの名無しさん
2011/08/21(日) 18:13:58.58多少結果が変でも問題ないからとかなんとか。
0963デフォルトの名無しさん
2011/08/21(日) 18:19:25.02長崎大のセンセのプレゼン資料はこれね。
http://www.ssken.gr.jp/MAINSITE/activity/sectionmeeting/sci/2009-2/lecture-1/ppt.pdf
・Geforce GTX295を421枚買ったら、うち72枚が壊れていた。
・壊れていたうち31枚を買い足して、380枚使ってスパコンを作った。
0964デフォルトの名無しさん
2011/08/21(日) 18:31:20.49CPU例外が起きておかしいなと思ったら、CUDAからのデータが化けてたとかあったわ。
0965デフォルトの名無しさん
2011/08/21(日) 18:36:48.43今頃「やっぱりな」とか寝ぼけたことを言っているやつの場合、
そいつの書いたプログラムがバグっている可能性の方がはるかに高いね。
0966デフォルトの名無しさん
2011/08/21(日) 18:37:41.87コカコーラそんなに飲んで大丈夫?(48/54)
0967デフォルトの名無しさん
2011/08/21(日) 18:42:26.76Geforce?それともQuadroやTesla?
GeforceのOCモデルはエラーになるの有名だけど
もしかしてTeslaがクロック落としてるのはそういう側面があるのかな?
0968デフォルトの名無しさん
2011/08/21(日) 18:46:11.00しかも、421枚の72枚ってかなりの割合じゃん
0969デフォルトの名無しさん
2011/08/21(日) 18:49:28.08馬鹿だろお前w GTX 580をOCしたらエラーが出たんだよ。
>>967
そうそう、GeForceだよ。紛らわしいから開発中はOCを切るようにしてるよ。
0970デフォルトの名無しさん
2011/08/21(日) 18:50:31.57> 負荷かけるとってことは、熱的な事情か。一般人向けの使用上はエンコード・デコードを想定して問題が事実上は無いとは言え、酷い話だなww
> しかも、421枚の72枚ってかなりの割合じゃん
グローバルメモリへのアクセスを増やしたらエラーが出やすくなったりとか、
結構いい加減だよ。
0971デフォルトの名無しさん
2011/08/21(日) 18:51:02.90しかしまあ、こんなことに$414kも税金投じたのか
0972デフォルトの名無しさん
2011/08/21(日) 18:53:20.62なるほど、そんな感覚なのか。
プログラム実行で誤ったら暴走だけど、ひたすら多量のデータを算出するだけだから、
間違いの程度問題で、手放し運転も一つの割り切りって考え方なのかね。
0973デフォルトの名無しさん
2011/08/21(日) 18:56:13.90コンスタントメモリでキャッシュされているからと安心してしまい、
バンクコンフリクト状態というのは意外とはまりそうですね。
>>963
39ページ目、そんな挿し方で冷却大丈夫か、と心配になります。
0974デフォルトの名無しさん
2011/08/21(日) 18:58:22.240975デフォルトの名無しさん
2011/08/21(日) 19:11:57.51各スレッドが異なる〜
なんてことあるの?
0976デフォルトの名無しさん
2011/08/21(日) 19:32:57.25記述することは一応可能。
__constant__ int table[] = {0, 16, 8, 24, ...... };
.....
int k;
.....
k += table[threadIdx.x & 0x1f];
0977デフォルトの名無しさん
2011/08/22(月) 01:06:53.21いやたったそれだけですんだってゴードンベル賞もらってたんだが
0978デフォルトの名無しさん
2011/08/22(月) 02:32:29.330979デフォルトの名無しさん
2011/08/22(月) 05:47:31.89無策でのほほんとしてると戦闘機や放射能汚染水処理施設みたいに足元見られてぼったくられることになる
アラブみたいに札束が地下から湧いてくる国ならいいけど、日本はそうじゃないから
0980デフォルトの名無しさん
2011/08/22(月) 15:29:33.90かといって自分で検証プログラムを組んでもどこまでのチェックでいいなのかよくわからないんだよなぁ
(そもそも自分の要求がどのくらいのオーダーであるかを自覚していないというのもあるんだけど…)
長崎大の不良検出プログラムってどういうことをして調べてるんだろう
演算命令を網羅するようになっているんだろうか。
0981デフォルトの名無しさん
2011/08/22(月) 17:02:01.89正答と比較するだけじゃないの?パラメータを変更して何度かやれば大体引っかかると思うが。
0982デフォルトの名無しさん
2011/08/22(月) 17:52:06.650983デフォルトの名無しさん
2011/08/22(月) 19:49:05.45すまんが、自作やったことある人なら誰でも知っているところ、としか。
具体的な連絡先まで晒してしまってもかまわんのだが、連絡するだけ時間の無駄だと思う。
契約取りたいがために適当言っているだけだろうから、本気にするような話じゃ…。
0984デフォルトの名無しさん
2011/08/22(月) 19:53:55.490985デフォルトの名無しさん
2011/08/22(月) 20:23:58.483者多数決はそもそもエラー検知をする為のものじゃないから仕方ないよね
0986デフォルトの名無しさん
2011/08/23(火) 02:43:59.970987デフォルトの名無しさん
2011/08/23(火) 02:54:01.89じゃあたぶん、うちの会社にも営業に来てるとこだね。ちょっとは取引もある。
上でも書いたが、多少値段割増しで交換保証してくれればそれで良いんだよ。
>>980
そういや、不良検出プログラムを公開するとあのセンセ言ってたような。
どうなってるんだろ?
0988デフォルトの名無しさん
2011/08/23(火) 21:11:12.540989デフォルトの名無しさん
2011/08/23(火) 21:17:51.790990デフォルトの名無しさん
2011/08/23(火) 21:44:37.800991デフォルトの名無しさん
2011/08/23(火) 22:02:15.220992デフォルトの名無しさん
2011/08/23(火) 22:04:56.420993デフォルトの名無しさん
2011/08/23(火) 22:12:58.26ttp://hibari.2ch.net/test/read.cgi/tech/1314104886/
0994デフォルトの名無しさん
2011/08/24(水) 00:27:14.54論文には出してるんじゃあるまいか?チェックしてないからほんとに出てるかは分からんが。
不良で一括りにされてるけど、どんな不良内容なのかわからない。このあたりも論文出てるんだろうか?
0995やんやん ◆yanyan72E.
2011/08/24(水) 04:58:14.98サポートするつもりがないから公開しないってだけかも。
0996デフォルトの名無しさん
2011/08/24(水) 11:24:13.010997デフォルトの名無しさん
2011/08/25(木) 00:23:18.600998デフォルトの名無しさん
2011/08/25(木) 00:33:53.150999デフォルトの名無しさん
2011/08/25(木) 00:35:35.691000デフォルトの名無しさん
2011/08/25(木) 00:36:32.8710011001
Over 1000Threadもう書けないので、新しいスレッドを立ててくださいです。。。
レス数が1000を超えています。これ以上書き込みはできません。