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

このスレッドは、他のスレッドでは書き込めない超低レベル、
もしくは質問者自身何が何だが分からない質問を勇気を持って書き込むスレッドです。
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/

関連サイト
CUDA 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/

なるほど
低レベル質問にも対応するスレだったんだな
道理でCUDAは意味が無いという意見が出て来るわけだ

いいかげんpertはなおしたほうがいいと思うの…

pertって？
ORのPERT？
だとしてもCUDAとの関係がわからん

…単純にスペルミスだろ

あ〜スレタイのか、納得しますた

並列計算の実験でGT240のグラフィックボードを買おうとしてるんですが
買ったほうがいいですか？ちなみに行列の演算を高速化したいというのが
目的です。

金をどぶに捨てる気か？
GT430にしろよ・・・

このスレってゆとりも対象にしてるの？

SQLiteをVRAMに構築して並列で検索してはどうだろう

どうしてもvisual profilerのcounter valueの項目のうち、cc2.0以降のactive cycle とactive warps
の意味がわかりません。意味を詳しく教えていただけませんか？

active warp は毎サイクルでのactiveなwarpの数の総和
activeなwarpってのは要はメモリアクセスとかで待ってる状態じゃなく今すぐ実行できるwarpの事
fermiだと0から48まで

active cycleはactiveなwarpが一つでもあるcycleの数
例えば総cycleが10cycleだとして、
毎cycleにactiveなwarpが有ればactive cycleの値は10
10cycleの内、activeなwarpが一つもないcycleが一回あったとしたら9

warpが分からないとかなら青木先生の本でも読めばいい

レスありがとうございます。
warpについては理解しています、がいまだcycleがわかりません・・・

cycleとはある命令を実行するのにかかる時間であり、ある命令はwarp単位で発行されるという認識です。
ですので1cycleで48warpがアクティブになるというのがよくわかりません。

またこれらの値は大きいほうがいいのか、小さいほうがいいのかもいまいちわからない状況です。
たとえばSummary Tableにでているactive warps/active cycle やretire ipcなども
意味と値が大きいほうがいいのか、がいくら調べてもわかっていない状況です。
よろしければご教授お願いします。

サイクルはクロックやHzとほぼ同じ

一応聞いとこ
>>14高校生？

>>14
プロファイラのマニュアルはちゃんとあるし、
プログラミングガイド、ベストパフォーマンスガイドに基本は全て載ってる
全部公式のガイドだぞ

>>16
初心者とゆとりが融合した素晴らしい素材ってことは分かる

cycleがクロックだとは分かっていましたが、1cycleにつき1SM内では1warpのある処理を実行できると思っていましたので（正確には2warpを2cycleで）
1cycleで複数のwarpを処理できるという点がいまいちどういう仕組みなのかわかっていませんでした。
Type SMとありましたので、この値はあるSMで計測された値だということがわかりますが、
上記のような認識だとつじつまが合わない（1SMで1warpずつ処理すると思っていたのに、1cycleで最大48warpがアクティブになるという点）ため質問させていただきました。
ややこしく、説明がいまいちになってしまい申し訳ないです。

プロファイラのマニュアルも読みましたが、読んだ上でゆとり質問となり言い返す言葉もありません。
これ以上スレを汚すのもどうかと思いますので質問を〆させていただきます。
レスくださった方ありがとうございました。

active warpは実行可能な待ち行列に待機中のwarpのことだろ。
待ち行列が無いと、演算器に空きが出来てもwarpを割り当てる事が
出来ないから、ある程度の長さの待ち行列が常にある事が
(演算器の利用効率を最大化するという意味では)重要になる。

fermiの細かいアーキテクチャまでは分からないが、
アレだけレジスタが多いとregisterがマルチバンクになっている可能性が高いから
例えば奇数cycleは奇数IDのwarp、偶数cycleは偶数IDのwarpしか発行できない
とかそういう制限があるかもしれない。
その場合はactive warpだからといって
常に任意のcycleに実行可能とならない場合もあるから
待ち行列は長い方が良い。

>>19
質問とはちょっとずれるかもしれないけど、
http://www.softek.co.jp/SPG/Pgi/TIPS/public/accel/gpu-accel2.html
この記事が長いけど結構分かりやすいかもね
ゆとりとかは挨拶だから気にするなよ

グラフィックボードを買うのは初めてなんですが、
GT430で探したんですが1スロットで収まる商品が2つしかなく、
評判も安定性も良くないらしいのでELSAのGT240にしようかと思っています。
何かアドバイスがあればよろしくお願いします。

このスレの住人にはやはり愛がありますね。
21さんのレスなんか涙がレスが読めなくなるほど。

ゆとりくんは心に刻むんだぞ。
玄人のおぢさんたちに生暖かく見守られているということを。

>>22
430使ってるが不具合ないよ。
あと今更fermi以前でPG組む気になれん。
CUDA目的ならCCのバージョンを意識しないと駄目だと思うぞ。
それをわかってて240にするなら何も言わない。

キャッシュの有無はでかいね。2スロットは兎も角、長さにも要注意。

>>23
目的、予算、制限を明確にしてほしいな
高速処理が至上命題で、予算も制限もないなら
言うまでもなく、Teslaボードをお勧めします
やっぱりfermi世代の方が無難なチョイスなので、
fermiチップのボードだとしてアドバイス


ELSAのGT240は使ってないからよくわからないけれど、
メーカーとしては鉄板の印象。安心していいのでは？
CUDAの場合、仮想メモリは使えないので、やりたい計算で
512MBのメモリだけで、溢れないか見積もることをお勧めします

１スロット幅を譲れない場合、少々値が張るが
Quadroならば１スロット幅のモデルもあるよ．
Quadro4000までは1スロット幅のようだ
市場価格は2万程度から10万弱くらい

2スロットでもよくて，そこそこ安く高性能というなら
やっぱりGTX580あたりがいいんじゃないかな
doubleやECC機能を利用しない限り、Teslaモデルより高速
コスパ最強です

それからWindowsの場合は画面出力用のボードとCUDA用のボードを分けると
デバッグしやすい気がするので、ボード2枚構成というのも考慮の余地があるかも
Linuxはしらない

おっと、GT240はfermiじゃないね
GT440だと脳内保管されてた・・・
ワシもfermiでないGT240は推奨しません

CUDAの勉強が目的ならGT240の方が
勉強になるかもしれない

fermiはキャッシュのおかげでタコなコードでも
そこそこ動いてくれるから

しまったアンカも間違えてた
>>23は間違いで>>22
もう寝ます

fermiってなんですか？自分は行列の演算を高速化したいだけなので
詳しいことはわかりませんが、同じコア数でもGT240とGT439ではなぜ
これだけ値段の差があるのか解りません。できれば安いほうを買いた
いのですが予算は10000円程度で考えています。
何かよきアドバイスがあればよろしくお願いします。

fermiってやつを買えば簡単なプログラミングで高速演算できる
fermiじゃなきゃ高速演算の手間がかかる
と思えばおk

GF-GT430-LE1GHD/1ST
これなんかどうでしょう？騒音がすごいというレビューを見ましたが。

>>29
詳しくはは自分でも調べてほしいのだが、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しかないな。

単精度だけでいいのかな。
最初はIONみたいなので消費電力を気にせずつかえるのがいいと思う。
速いカードの方が当然チューニングの効果がでるからおもしろい。
プロファイラで取れる情報もカードによって違うみたい。

テクスチャメモリってどのような利点があるのですか？
添字が小数点の時の線形補完を使ってみたいのですが、
メモリはチップ外にあるということはグローバルメモリと同等の速度なのですよね
キャッシュはされるみたいですが同じデータを何度も利用しなくては意味がないということでしょうか

ELSA GLADIAC GT 430 LPと
EVGA GeForce GT 430で迷っています。
どちらがいいかアドバイスをよろしくお願いします。

どっちでも変わらないとおもうよ。
印象でいえばELSAは手堅い印象。
EVGAはゲーム向けのチューニングとかしてて精度の点が不安。
予算的に許すならquadro600がいいかもしれない。
ただしELSAでもgeforce,quadroはcudaでの演算結果の保証はしてないから
演算を間違うコアが含まれていても交換対象ではない。

因みに整数演算しか行いません。
有限体の行列なので浮動小数点は使いません。
予算は１万円です。コアの精度ってなんですか？
コアが計算を間違えることがあるのですか？
精度を求めるならどちらがいいでしょうか？
質問ばかりで申し訳ありませんが、ご回答よろしくお願いします。

予算1万ならELSAでぎりぎり。

それ、浮動小数点演算いるんじゃね？
浮動小数点の意味わかってる？

floatとかdoubleとかですよね。
自分が計算したいのはunsigned shortなのです。
以前雑誌の特集にCUDAの記事がありNVIDIAのコアには整数演算も
可能であるということが書かれてありました。なので整数の行列も
高速化できるかと思っていました。何か間違いがありましたら
ご指摘をお願いします。

quadro600がいいかもしれないといったのは、
演算結果は保証されないけど、geforceより厳しい検査がされているから
より演算結果の値にエラーがでにくいだろうという推測にもとづくもの。

演算を間違うコアがふくまれる場合があるというのは簡単に言えば１＋１が２以外の結果になる
コアが含まれる場合があるって意味。

精度を求めるというのが演算結果を保証するならという意味ならTESLAしかない。
そこまでの予算を確保できない場合、上記のようなコアが含まれていないかを確認するプログラムを自分で
組んで確認しないといけない。
単純に足し算だけでなく、四則演算、浮動小数点演算、指数関数、超越関数、L1、L2キャッシュ、
VRAMに至るまで正当性を確認する必要がでると思ってる

あ、CPUでの演算結果があるなら、同じ計算をしてみて結果が同じかどうかで
簡略な確認としてもいいかもしれない。それで許されるのであれば。

CPUの演算結果とGT430の９６個のコアの演算結果を比較して
間違いがないことを実験する必要があるのですね。
同じ計算機なのにどうしてそのような確率的な動作をするのか
理解しかねます。CPUだけが特別に演算が正確なのはなぜでしょう?
GPUも同じ32ビットレジスタなので構造が同じだと思ったのですが。

その比較はあくまで簡略な方法だよ。全部のコアが使われる保証だってないんだから。

CPUは元々計算用に作ってあるでしょ？CPUで１＋１＝２を保証しなくてだれが買うよ？
GPUは元々描画用に作られてて座標計算とかfloatで計算することが多い。
そのときに1.000002が1.000002になってもあんまり影響ないでしょ？
もちろん正確に計算できるほうがいいが、描画だったら多少不正確でも問題はないだろうっていう判断だと思う。
それにCPUに比べてALUが膨大なこともあってすべてのコアの演算精度を保証してたらキリがないでしょ。コスト的に。
そういう用途の人のために厳しいテストをしたTESLAが用意されてる訳。でもその分料金上乗せさせてもらいますよって事だよ。

検算しないといけないとなると計算結果が信用できないということで
並列化のメリットがなくなってしまう気がします。

あ、ごめん。
×　そのときに1.000002が1.000002になっても
○　そのときに1.000001が1.000002になっても

いや、毎回検算しろといってるわけじゃなく、

一回CPUと同等の計算をして、演算結果に謝りがないことの確認をもって
GPUのコアにエラーが発生していないことの簡略確認にすれば？
ってことだよ

浮動小数点はどうでもいいけど、せめて整数演算だけは正確に
計算してもらいたいと祈るばかりですね。

体だと素数で剰余をとるのですか、どういうコードか教えてくれませんか。

長崎大の検査プログラムがほしい。。。。
が、無理なんだろうなあ

ちょっと恥ずかしいけど見せますねｗ
typedef 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]);

}

__m128i s5(__m128i cc){
int 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;
}

int main(int argc,char *argv[]){
int 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;
}

while((read_size=fread(buff,1,SIZE,fq))){
for(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;
}

メモリ階層ってこの理解でいいのかな。自信がない

レジスタ

L1

共有メモリ

L2

>>55
誤爆

メモリ階層ってこの理解でいいのかな。自信がない

レジスタ
｜　　　　　↑
L1　　　　　｜コーヒレント
｜　　　　　↓
共有メモリ
｜　　　　　↑
L2　　　　　｜コーヒレントじゃない
｜　　　　　↓
グローバルメモリ

この資料に書いてあるとかいう情報があるなら
それも教えてもらえると助かります

一部省略してありますが大体こんな感じです。素体は楕円曲線暗号に
使っていますがRubyの実装でC言語にはしてないです。
このプログラムは疑似乱数生成器で基本的には最大32*32の行列演算を
主に行っています。これが並列化できれば32回の計算で256ビットの
乱数を生成できるので効率的になります。
というかなればいいなあと思ってます。

CUDAでINTって糞性能だろ・・・
int(4)なら加算だけでもfloatの4倍かかるぞ
FMAとかも使えないし

こっちのほうが早いですね。荒らしてしまい申し訳ないです。

http://sky.geocities.jp/tcshacina/grs1.c

>>59
じゃあグラフィックボードで並列計算するのはこの場合逆効果になる
のでしょうか？買わないほうがいいですか？

floatで計算すりゃいいんじゃね？
intにこだわる必要あるのか？

floatのXORってできるんですか？疑似乱数生成器なのでfloatで一旦
計算しておいて後から一斉にunsigned charに変換することが可能で
しょうか？暗号に使う疑似乱数なので整数で処理しないといけない
という条件があります。

楕円暗号入門を手に取ってみたけどやっぱりわからなかった。
あれだけコードがかけるならいいおもちゃだから買えばいいんじゃないですか。
Cellも整数の論理演算にむいていた気がします。６個のタスクを生成できる
ので使いやすい場合もありました。

>>63
xorshiftのこと？
であればCURANDってのがあるからそれ使えば乱数生成可？
メルセンヌもSDKについてる

普通のCPUでやっているような排他的論理和（XOR）をCUDAもできるのか
という疑問です。できなければfloatのままデータを生成して、後で
CPUにまとめてfloat->intの処理をさせながらXORを計算するという
方法ではうまくいかないでしょうか？マニュアルには論理演算も
可能と書いてありますけど、整数演算の速度が浮動小数点より遅い
というのが気がかりです。

Rubyで作った暗号化モジュールいろいろ
http://sky.geocities.jp/tcshacina/crypt.html

だからさXORを乱数生成に使うんじゃねの？

使います、そのためには整数でなければならないと思っています。

だからCURAND使えばいいべよって書いたんだべ
CURANDのマニュアル読んで目的にかなうかどうか確認し

xorshiftってなんですか？自分のオリジナルのアルゴリズムに
どうやってそれを使えばいいのですか？ソースは晒してあるので
自由に見て感想を聞かせてください。

悪いが人のソース追っかけるほど暇ないよ

私のソースを見たいと言った人はどこへ行ったんだ！？

胡散臭い文体見た瞬間に例の暗号の人だってすぐ分かった。

＞自分のオリジナルのアルゴリズムにどうやってそれを使えばいいのですか？
それを考えるのがプログラマだろうが！１から１０まで聞いてないで自分でやってみろ

使えないと言いたい

一日に書き込む疑問文は3つまでにしてくれ。

珍しく伸びてるなと思ったら…

いくら超低レベルスレといっても
"ちょっと調べて分かること"を次々質問するのはどうなんだ？

答えるは簡単だがこの程度のことを人に頼る時点で
そいつにまともなプログラムが作れると思えん

>>77
同感。
そういうやつでも数学的素養が優れていたりして特定のモデルの構築とかには
優れた才能を発揮する奴がいるから、そういう所だけは認めざるを得ないこともある
のが悔しい。

シミュレーション系で使う人、シミュレート板に新スレ立てたんでよろしく
プログラミングの方法より使い方を話題にしたいなと

GPGPUでシミュレーションを高速化してみる
http://kamome.2ch.net/test/read.cgi/sim/1291687353/

>>66
Micro-benchmarking the GT200 GPU
http://www.eecg.toronto.edu/~moshovos/CUDA08/arx/microbenchmark_report.pdf

GT200ではint加算、減算、論理演算、シフトのレイテンシ、スループットはfloat積和算と同じ。

CPUと違ってレイテンシが24もあるが、並列度が十分にあれば普通に使える。

Fermiでの資料を探したが見つからなかった。誰か頼む。

Arithmetic Instruction
Int 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


と、あった

だからFermiはfloatとintが同じ速度だとあれほど(ry

環境をインストールしてはじめようと思ったんだが
SDKのサンプルtempleteが2008で開けない。

カスタム ビルド規則 ファイル 'C:\Program Files\Microsoft Visual Studio 9.0\VC\VCProjectDefaults\NvCudaRuntimeApi.rules' が見つからないか、または読み込めませんでした。

って出たからrulesファイル探してるんだけどないんだよな…
SDK再インストールしたけどやっぱりない。たすけて

ドライバAPIを使えば，一つのホストスレッドで
複数のGPUに並列でカーネルを実行させることが出来ますか？

二つのホストスレッドが必要じゃね?

>>85
無理ですか…

シミュレーションを複数のGPUで実行して，
結果をOpenGLで可視化したいと思っています(出力は一つから)．
ランタイムAPIとOpenMPを組み合わせて作ろうと思っていました．

しかし，その場合malloc等の前処理の段階で
複数のホストスレッドでcudaSetDeviceを使ってしまうと，
いざカーネルを実行しようとしたときに，マスターホストスレッド以外のホストスレッドの
cudaSetDeviceで"もう既にセット済"というエラーが出てしまいます．

OpenGLのイベント待ちのために，
複数スレッド→マスタスレッドのみ→複数スレッド→マスタスレッドのみ→…
という状態を繰り返すのですが，1スレッドに戻っても
マスタスレッド以外で設定したコンテキストが消えないのが原因のようです．

マニュアルのドライバAPIのコンテキストの部分を読んでいたら，
色々命令があるので出来るかもしれないと思い質問させていただきました．
ホストスレッドの数は複数でも問題ないので，
コンテキストの操作に詳しい方いらっしゃいますか？

複数GPUの扱いがおかしいんじゃね？
良く見直してみ

kernelの実行命令が繰り返されるのはどういう状況のときそうなりますか？

GlobalからのReadキャッシュをL1キャッシュにおいて有効にしたときと無効にしたときにL1キャッシュが無効のほうが早い場合があり、
Visual Profilerで調べてみるとinstruction　issueが93152（L1有効）と67666(L1無効)
で差があり、replayを含まないinstruction executed が54592(有効無効どちらも同じ値)
ですのでreplayに問題があると思ったのですが、どういう状況下でreplayが起こるのかいまいちわかりません。

replayはコアレッシングできていないreadのときにおこると思っていたのですが、間違っているなら教えてください。

visualvasicでプログラムを組んでいるものですが
ビルドルールのemulation modeにて、はい (-deviceemu -D_DEVICEEMU)を選択しても、
下記エラー分が消えません。
calling a host function from a __device__/__global__ function is not allowed

どうすればよいのでしょう？
CUDAのバージョンは3.1です。

エミュレーションって3.1はないんじゃないの

では、printf();は設定なしにつかえるということなのでしょうか？
calling a host function from a __device__/__global__ function is not allowed
するとやはりエラーが消えないのですが・・・

>>91
おまえ言ってること矛盾してるし、何をやってそういうエラーがでたかの説明してなくて
アドバイスを求めても無理。
vbにprintfとかないよね？どういうことなの？

3.1からデバッガが標準でつくようになったのでエミュレーションはもうなくなったはず
Fermi系列なら確かそのまんまprintf使えるけど

windowsならVisual Nsightとかいうのでデバッグできたはず
あとはどうしてもエミュレーション使いたいならcuda2.3使うとか

>>93
printfは確かfermi系だけのサポートのはず。
そして2.3はfermiサポートしてない。3.0からだったはず。

Prallel Nsightな

>>88
同様の現象が起きていて困ってる
GTX480ではinstruction issueが取れなかったんだが、Tesla使い？

>>96
いえ、GTX480です。ASUS製としか覚えてませんが・・・

C1060x 4をFermiのGTX580x 4に入れ替えた場合の性能比較を
探しています。

フロートの性能は欲しいがC2050は高すぎるということで。

C1060って260相当だっけ？
3DVantageとかの260と580の性能差を4倍したものを求めて見積もればいいんじゃない？

これでCPUがシングルコアだったら大笑い。

floatじゃなくてdouble茶羽化？

性能比較とかどんなアプリ想定しているか書いていない時点でなんだかなあ

とりあえずsoftekでググると何か出てくるかもしれんよ

CPUはGulftwon/3.33GHz*2発で，メモリが残念な16GBです。
アプリ？はR+CUDAかMATLABで多体シミュレーションです。
けどRって言語だよな。

スレ違いな気がするけど、MATLAB 2010bのGPUArrayって使ってみましたか？

3.0以上でfermi系ならカーネル内でprintf対応してるのか
誰かfermi系で128台ぐらいのGPUクラスター作ってくんねーかな

128台でprintfするの？

>>105
ttp://tsubame.gsic.titech.ac.jp/ja/hardware-architecture
東工大にでも逝けや

>>103
その条件でfloatの性能求めるならたいして変わらないんじゃないかな
doubleなら別だけど

>>108
メモリの転送が非同期で双方向通信できるとか
コンカレントカーネル実行ができるとか、自動キャッシュがあるとか
のおかげで最適化（上記の機能を使うように修正）すればかなり早くなると思う。
修正なしでそのままならちょっと〜けっこう早くなると思う。

でもR+CUDAを知らないので無責任な推測です。

最適化しなければ、物凄い速くなるけど
最適化してしまえばそんなに性能変わらないんじゃね。

>>110
C1060に最適化したものをそのままFermiで動かしても差はそんなにないだろうって事だと思うけど
それでもコアレスの制限が緩和された影響と自動キャッシュのおかげで結構速くなると思う。

さらにFermiのストリームとかを使用してデータ転送の最適化を行えば段違いになると思う。
特に4枚挿してるということからx8*4の転送速度になって転送がボトルネックになりそうだから
そこの最適化は効いてくるんじゃないかと思ってる

ところで103のMBはガルフ対応だからPCI-e2.0対応だよね？

ごめんストリームはFermi以前から対応してたみたいね。

ストリームは1.3だっけ? C1060って2.0だっけ?

>>113
ストリームはCC依存でtoolkit依存のようだが？

>>113
ストリームはCC依存ではなくtoolkit依存のようだが？

メモリバウンドならその話もわかるが
多体問題って演算バウンドだったんじゃないの確か
GTX580に交換するうまみはそれほど無い気がする

ということはtesla->fermiの性能向上ってメモリ関係（キャッシュ含む）とdoubleだけってこと？
floatなら計算パワーだけで考えると大して性能向上してなかったの？

floatだけならメモリ周りのボトルネックを解消して全体的にFlopsが上がりやすくなりましたって
だけだったのか・・・・

考えるより動けということでGTX580x4枚の見積もり取ってます。
20万＋なら自分の財布でやってしまえ。
ちなみにMathematica8もCUDA対応で年始位にデリバリーとか何とか。

>>104
>MATLAB 2010bのGPUArray
R+CUDAの方に感けて，まだやってませんでした。すみません。

C1060と580のnbodyのベンチをさがしてみた

C1060　443.4GFLOPS
580 447.154 single-precision GFLOP/s at 20 flops per interaction

580は過去スレからc1060はググった結果
なにこれ！？

ごめん↑のC2050だった

> Compute 2.0 CUDA device: [GeForce GTX 580]
16384 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

>>121
GTX580結構がんばるんだな…
年末のセールで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

>>118
自分でコード書いたほうが速いし確実だと思うけどね

昔、GeForceをごにょごにょしたらQuadroになったように、
GeForceをごにょごにょしたらdoubleが4倍速になったりしないかな。

最近のはならないようになっている
あきらめろん

>>123
行列演算で、CUBLASより速いコードかける？

何でそんな不毛な煽りするかね

コンスタントメモリについて教えてください
__constant__ int i;
int j=1;とした場合(コンスタントメモリに配列以外を置きたい場合)になぜ
cudaMemcpyToSymbol("i",&j,sizeof(int),0,cudaMemcpyHostToDevice);
と文字列のようにして渡すのでしょうか

また、外部ファイルにextern __const__ int i;
と書くと定義が二重ですと言われます
複数ファイルにかかれたカーネルからコンスタントメモリを読み込みたいのですがどのようにすればよいのでしょうか
お願いいたします

>>128
うーん、簡単に言えば、cudaMalloc()でVRAM上のアドレスを得ることができればcudaMemcpy()できる。
一方__constant__の場合はそういうわけにはいかないので、
nvcc独自のちょっとえげつないcudaMemcopyToSymbol()を使わないといけない。
コンパイル単位の分かれたカーネル同士で共有するのは難しいかもねぇ。

>>118
> R+CUDAの方に感けて，まだやってませんでした。すみません。

R+CUDAってどういう構成なのですが？
やはり、

R+gputools+CUDA+CULA

なのですか？

それとも、Rのビルド時にCUBLASをリンクするのですか？

当方の環境は、

Q9450 + Memory 6GB + GeForce9600GT
Ubuntsu10.04(64bit)

です。gotoBLASで行列演算時にコア4つ使用するところまではやったのですが、
CUDA使うともっと凄くなるのですか？

速くなるかどうかはケースバイケースでアプリの内容によるだろ

>>131
>>118さんがどういう構成でやっているのか知りたいです。
未だ、コンパイルも通らない実行も出来ていないレベルの者です。

RにgotoBLAS組み込んでも、乱数の生成速度に差が出ませんでした。
BLASは乱数を生成するライブラリではないから当然ですが。

CUDAのページのサンプルに
・MersenneTwister
・Monte Carlo Option Pricing with multi-GPU support
などとあるのをみると期待に胸がふくらみます。

胸は膨らみブラ破け

だいたい乱数で時間かかるって相関乱数の相関部分じゃね〜の？
BLASがいくら早くても関係ないじゃん、相関には

乱数の作り方にもよるけど、誤り訂正符号を使った方式では符号化に
O(n^2)かかるところを32*8=256ビットの乱数を生成させるのに32個の
コアで16回の計算で得られるので効率的になる。

３流プログラマがgotoさんの名前を出して汚すなよ。プライドが少しでもあるならな。

・・・

後藤blasだろ
釣られた？

次買うのＡＴＩになりそ・・・
Compute Shaderスレないですか？

ATIやる人はOpenCLかCAL+ILじゃないの？

>>86

ものすごい遅いレスですみませんが、driverAPIなら可能です。

>>136
BLASが乱数を生成するライブラリでないと言ったら、
後藤さんの名前を汚していることになるのですか？

Textureメモリってふつうのより早いの？
みんななにに使ってるの？

ググレ

GTX580，4枚まとめ買い目指して奔走してますが，ない。
2月くらいにはなんとかならないだろか。

CUDA Visual Profiler と Compute Visual Profilerの違いは何でしょうか
前者はバージョン2.3
後者はバージョン3.1
が入っているのですが、toolkitが3.xになってからプロファイラの名称が変わったというだけでしょうか？

でしょ。

質問です
ウッカリ長時間のCUDAジョブを走らせてしまったとき
中断はどすればよいですか？

>>147
ありがとうございます
ものの本やサイトには/usr/local/cuda/cudaprof/のcudaprofiler使えって書いてあったのですが
変なWarning出て困っていたところでした

>>147
ありがとうございます
ものの本やサイトには/usr/local/cuda/cudaprof/のcudaprofiler使えって書いてあったのですが
変なWarning出て困っていたところでした

>>148
GPUはずしちゃえよ

kill すればいいんじゃね？

>>145
九十九ネットストアにあったけど。あと価格コムとかで在庫もってる店あたってみたらいいんじゃない？

補足、九十九にあったのはエルザのね。

Linuxドライバにバグないか？
最初の実行に時間がかかるのと
同様に最初の実行時にグリッド数、スレッド数が大きめだたとmemcpyでエラーが出る
小さめのグリッド数、スレッド数で最初に実行した後だと
大きめのグリッド数、スレッド数でも正常するんだが…

>>155
ドライバのバージョンとSDK、TOOLキット、OSの情報も載せないと駄目じゃね？
あとgccか。

OS：CentOS5.5
ドライバ、SDK、ToolKitは最新
ｇｃｃは4.1.2
今のところ原因は全く不明
ドライバのバージョンを一個前にすると初回実行時の遅延がちょっと早くなる（最新が約2.5s、一個前が約1.8ｓ）
gtx470と465で試してどちらも同じ症状

ようやくgputoolsが動いた。

速いのは分かったけど、
行列の規模が、本体のメモリ容量と関係なく、ビデオカードのメモリ容量に制約される。
キャッシュイーター向けじゃなかったですね。

>>158
行列の計算ってのはね、分割できるんですよ。

>>159
確かに。
分割できなければコアが複数あっても意味がない。

でもね、いったんメインメモリに移して計算するのはもったいない。

結論C2070を買え。

結局、部分行列作成と計算結果の結合で、処理にかかる時間が逆転してしまいました。

一般論として言えるのかどうか知らないけど、
行列が（計算結果を含めて）ビデオカードのメモリにスッポリ入るなら劇的な高速化が見込めるけど、
そうでなければ、マルチコアCPU＋本体メモリのほうが速度でもメモリ消費量でも有利なようです。

コンスタントメモリはいつデバイスに確保されるのでしょうか
__const__宣言時とは考えにくいので最初にcudaMemcpyToSymbolしたときでしょうか

デバイスにあるということはコンスタントメモリも終了時に解放したほうがいいと思うのですが、
cudaFree()では解放できないようです

コンスタントメモリはどのように解放すればいいのでしょうか。

行列計算なら転送を隠せるはず
リンパックベンチでGPU搭載マシーンが上位に来てるんだし

>>163
SDKのサンプルみればいいんじゃないか？

__constant__を気にするなら
kernelコードだって似たようなものなんじゃ。

>>163
定数メモリはcudaMalloc()のように「確保する」ということなしに使うようになっています。
その為、cudaMemcpyToSymbol()のような特殊な関数でしか転送できません。
また、「確保する」ことなしに使える以上、「解放する」手続きは用意されていません。

安定ソートが欲しいのですが
CUDA向きの安定ソートってありますか？
16000要素くらいです

NVIDIAにあるバイトニック使ってみ

SDKにmergesortってあるがそれが同じ奴だ

バイトニックは安定じゃないようなのですが、
マージソートのサンプルがSDKにあったのですね。
チェックしてみます

gtx580を4基積んだ．電源が音を上げた．
やれやれ．

>>172
何Wで駄目だった？

>>173
1400W。ディスク系を外においてもなぁ。

SC748あたりなら・・・

7046GT-TRFね。

1400Wで駄目だとなると電源2つ積めるケースにして連動させるケーブルで繋いでと
いったことをしないと無理そうだな。

無茶しやがって…

1400Wとか言って売っている時点で粗悪電源確定。
>>172が日本在住でないなら話は別だが。

>>179
あれは110V/1400W。80GOLDのリタンダント。
日本じゃ実効1100W。

大容量PC電源はエアコン引っこ抜いて200Vから取ろう。

>>180
なら、「1100W電源」として売るべきだろ。

排気量2000ccの車を3000ccだと言って売る自動車屋などありえない。
そのあり得ないことを長年やってきたのが、PC電源代理業者だ。

>>182
supermicroは日本向けのサイトを持っていない。仕方あるまい。
ちなみに法人で入れたときはこの説明はあった。
www.supermicro.com/products/system/4U/7046/SYS-7046GT-TRF.cfm

GPUの分類について教えてください
例えばGTX480は
Fermi-ComputeCapability2.0-GF100-GTX480
でよいのでしょうか
それともアーキテクチャとCCは1対1対応しているのでしょうか
また、Fermiの分類位置はここでよいのでしょうか

ComputeCapability2.0-GF100-GTX480
は正しいけど、FermiはC.C.2.X全部を含む
GFの"F"はFermi、GTXの"T"はなぜかTeslaだったはず

てめぇGeForce 7800GTXさんに謝れよ

>>184
CUDA目的ならComputeCapabilityだけ意識すればよいかと。
あとはパフォーマンスが違うだけだから。

>>186
GeForce 7800GTXさんはCUDA対象外じゃ・・・・・・

コンパイルされたCUDAプログラムを実行するだけの場合、ドライバのみでいけますか？
toolktのインストールは必須ですか？
LD_LIBRARY_PATH /usr/local/cuda/lib64
を通さなくてはならないということはtoolkit必須なのかな

ライブラリ使っててもstaticならいらないはずだし、dllつかってるなら
そのdllをexeと同じフォルダにおけばいいんじゃないか？
きっとlib64内のなんかのライブラリつかってんでしょ。cutilとかかな

DLL再配布して良かったっけ

Linuxならドライバ（カーネルモジュールを含む）も必要なので再配布だけでは解決しない。
もちろんツールキットのライブラリも必要。

すみません、教えてください。
CentOS 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です。

環境設定は？

それはCentOS 4.5のようだが。

>>193
例えば、ディストリビューションのnvidiaドライバなんかが残っていない？
その上に、最新ドライバをインストールしていて、
コンパイル時と実行時に別の共有ライブラリを参照していたりとか。

ていうか、自分も、ubuntuでx-updatesのドライバと最新ドライバが両方入っていて挙動不審。
ubuntuでもOS再インストールかｗ

>>196

返信ありがとうございます。
ちょっとその辺を調べてみます。
後日、問題が解決したらフィードバックします。

スルーされた。勝手にしてくれ。

いじけるなよｗ

とある処理のスレッド数と時間を関係を調べているのですが、
スレッド数が192を超えたあたりで処理が行われていません。

スレッド数が512より小さくても処理が行われないことがあるのでしょうか？

グラボは9800GTで11レジスタ、共有メモリを36byte使用してます。

>>200
「処理が行なわれない」ってどういう意味さ。

>>201
説明不足でした。

処理を1000回行い、平均処理時間を計測しているのですが、
スレッド数が192を超えると0(ms)になってしまいます。

何かが間違っているんだろうね。

まずはエラー確認してからだろ。

スレッド数を増やしたことによって確保されていない領域にアクセスするような
エラーが発生するような状況になっていないかどうかを確認・・・・・かな。
これが起きているとそれ以降の（タイマーを含めた）全ての処理がエラーになる。

まずは
cudaThreadSynchronize();
cudaGetLastError();

>>203-206
助言ありがとうございます。

まずはエラーを確認してみたいと思います。

たぶんリソース不足でしょ。自分もよくそうなる。
カーネルの引数を構造体のポインタにまとめると解決する。

GTX460ってCUDAは使えるんですかね？
lspci | grep -i nVidia
でPCIバスデバイス情報確認したら
VGA compatible controller: nVidia Corporation: Unknown device …
ってコンピューターさんがおっしゃられたんだが…

>>209
/sbin/update-pciids
してからコンピュータさんにもう一度お伺いたててご覧なさい

ドライバをインストールしてないんじゃないの？

正月セールで狙う1万5000円くらいのグラボでオススメのものをおせーて

GeforceとRADEONを2枚挿して
GeforceをCUDA専用、RADEONをグラフィック専用とすることは出来ますか？

日本語的に複数の意味に取れるので訂正

GeforceとRADEONの2枚を挿して
GeforceをCUDA専用、RADEONをグラフィック専用とすることは出来ますか？

no problem

ありがとうございます

cudaSetDeviceFlags(cudaDeviceBlockingSync);

ほとんど同じソースで
cudaSetDeviceFlags(cudaDeviceBlockingSync);
が有効になるプログラムとならないプログラムがあるのですがどこを
調べれば良いでしょうか。

ほとんど同じソースならどこが違うのかを書いてください。
そして「有効」かどうかをCPU負荷で判定しているのかどうかも。

cudaSetDevice()を除くすべてのCUDA関数より先にcudaSetDeviceFlags()を
呼び出しているかどうか以外には特に調べるべきところはありません。

マニュアル的なpdfが沢山あってどれを読めばいいのかわからないのですがどれから読むのがいいのでしょうか？
CUDA_C_Best_Practices_Guide.pdf
CUDA_C_Programming_Guide.pdf
Introductory CUDA Technical Training Courses VolumeI.pdf

Introductory CUDA Technical Training Courses VolumeI.pdfで軽く概要をつかむ。
ただこれ内容が古かったと思うからどんな仕組みで動いてるかを把握するくらいでいいと思う
必要になったらまた読み返せばいい。

次にCUDA_C_Programming_Guide.pdfで基本を学ぶ。
これは読み返すことが多くなると思う。

その次にCUDA_C_Best_Practices_Guide.pdfで良い実装の仕方を学ぶ。
これは実際にプログラム組みながらでいいと思う。

あと余談だがfermi　tuning guideも読むといい。
英語が苦手なら「はじめてのCUDAプログラミング」や「CUDAプログラミング実践講座 - 超並列プロセッサにおけるプログラミング手法」
を読んでから必要なところをPDFで追加で読むと理解が早いと思う。

フランス語のはないですかね？

ワロタ

>>222
残念ながらありません

cc2.0でコンパイルすると、cc1.3でコンパイルした時よりレジスタ使用量が
10くらい増えるのですが、なぜですか？

なんでだろうねー

Fermiで各スレッドで数百バイトのワーキングエリアが必要な場合、
ローカルメモリを使うのが最良なのでしょうか？
シェアードメモリを使うと、同時に立ち上げられるスレッド数が減って
逆に遅くなりました。

>>225
そういうのはプログラムによって変わるだろ
とりあえずptxファイル嫁

>>227
>各スレッドで数百バイトのワーキングエリアが必要な場合
これが今以上減らせないならそもそもCUDAに向いてない

グローバルメモリのアクセスは400~500clock、レジスタやシェアードメモリ、L1キャッシュはほぼ0らしいですが
L2キャッシュは大体何clockかかりますか？

L1は大体15
L2は大体150

4枚発注ー。

CPU向けコードとGPU向けコードを同一ファイルで作成したいと思っています。
nvccはcuでないと、cudaコードとしてコンパイルしてくれないのですが、
拡張子がcppのファイルをcudaソースファイルとしてコンパイルするオプションはありませんか?

obj同士はnvccでリンクできるから、それぞれコンパイルすればＯＫだよ

>>141
亀レスですが，Driver APIを用いて出来ました．
ありがとうございました．

>>234
概略をのせてくれまいか？

cudaAAA()で○○を設定
cudaBBB()でXを取得して
pthread_AAA()を使ってXXXする

くらいの事でもいいから

>>235
簡単にで良いなら．

// 前処理
#pragma omp parallel
{
　　cuInit(); // 初期化
　　cuDeviceGet(); // GPUへのハンドルを取得
　　cuCtxCreate(); // コンテキストを作成し，ホストスレッドにバインド

　　mallocとか色々

　　cuCtxPopCurrent(); // コンテキストをアンバインド
}

glutMainLoop();

// カーネル実行
#pragma omp parallel
{
　　cuCtxPushCurrent(); // バインド

　　カーネル実行

　　cuCtxPopCurrent(); // アンバインド
}

関数の引数は省略しました．

>>236
ありがとう。参考にさせてもらいます

Fermiもグローバルメモリアクセスはコアレッシングにしたほうがよくって
キャッシュはおまけ程度って考えでおk？
ていうか各スレッドが違うグローバルメモリにアクセスするなら同じカーネル内ではキャッシュ効かないよね

CPUのマルチコア（巨大キャッシュ、低並列）のマルチスレッドですら
キャッシュ汚染はひどいからな

マルチスレッドとキャッシュは相性が非常に悪い

スレッド番号とグロバールメモリーのアドレスが並んでいないと使い物にならないな。
速度が１０倍は違う。本当に使いにくいよ。
メモリアクセスがネックの（単純な）アプリで、shared memoryの使いまわしのできないアプリでは、
マルチスレッドを上回るのは難しい。

しかし、いままでGPUと相性の悪いといわれるプログラムでも
FermiならCPUの2倍くらい出るってことは結構ある。

非Fermiならローカルメモリ使ったら即使い物にならない速度だったけど、
Fermiならキャッシュがなんとかしてくれたり。

cublasDgemmにおいて、係数のalphaを(double)1に、
betaを(double)0に指定しても、
内部では真面目に係数をかけているのでしょうか？

参考(PDF注意)
http://developer.download.nvidia.com/compute/cuda/3_2_prod/toolkit/docs/CUBLAS_Library.pdf

FORTRANのソースを自分で追っかけて確認しろ

速度を測ればすぐに分かりそう

>>243
ありがとうございます。
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のコードのような実装になっていない気がするのですが。

「はじめてのCUDAプログラミング」に出てるコード（もしくは同じ出版社からインターネット上に出てる奴）
をひたすらマイナーチェンジしていくのが一番早いと思う。
スクラッチはマイナーチェンジの経験を積んでからで十分。

PTXって機械語なの？
GPUがPTX読むみたいな事が書いてあるんだけど

ptxはgpuのアセンブラ言語じゃないの？？

PTXはあくまで中間言語。

TMPGEncで使おうと思ったらG80非対応なんだな。。

GTX580を4枚積んでサンプルコード走らせたら電力不足来たー。

CUDAツールキットを導入する際にまずやることって何？
どう頑張ってもnvccでコンパイルしたサンプルファイルが実行できませーん
って開発以前の問題で精神折れそうだわ…

list.cは以下の通りです．
#include <stdio.h>
#include <stdlib.h>
#include <string.h>

struct record { /* １人分のデータ */
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 ); /* 先頭の次からリストの表示 */
}
}

>>252
実行できないならできないなりにエラーメッセージが帰ってくるのではないでしょうか


CUDAコアとはスカラプロセッサのことでいいのでしょうか
FermiのみCUDAコアという呼び方をするのですか？

>>252
実行できないならできないなりにエラーメッセージが帰ってくるのではないでしょうか


CUDAコアとはスカラプロセッサのことでいいのでしょうか
FermiのみCUDAコアという呼び方をするのですか？

Fermiで倍精度浮動小数点の三角関数計算に
SFUを使えないのでしょうか

倍精度浮動小数点の計算は単精度浮動小数点演算機が2クロックで計算するのですよね
同じようにSFUが単精度計算分の2週でやってくれるということはないのでしょうか

geforceは駄目だろ

ドライバに　266.58 WHQL　が来たね

>>256
fast mathを使えば？

fast mathは単精度専用。

そもそも倍精度使ってるくせに精度低いSFU使いたいって時点で
本当に倍精度が必要か考えるべき。

SFU遊ばしておくのはもったいないかな、と思いまして

そんなに精度が低いのかー

応用次第だけど、SFU使うと駄目で使わなければfloatでもOKという例があった。

初歩的なこと聞いてすみません。
元々あるソースにCUDAを取り入れてるんですが構造体を丸々GPU側にコピーはできないんでしょうか？
できるのであれば誰かご教授お願いします。

C言語の勉強をしたらできると思います。

>>264
一応いろいろ試したんですが
Warning : Cannot tell what pointer points to, assuming global memory space
ていう警告が出ていておそらくデータの転送ができていないんです・・・。
普通のC言語を勉強すればできるんですかね・・・？

構造体の中にポインタが入ってて、そのポインタがメインメモリを指してるんじゃないの？

>>266
ポインタのほかにもポインタのポインタも入っているのでそれもGPU側に送らないといけないいうことであってますか？
そうすると結構量があるので必要な部分のみを送るほうが効率がいいんですかね・・・？

もちろんポインタが指してるデータは送らないといけないし，
構造体を送る前に，構造体の中で指してるアドレスをGPU側のアドレスに
書き変えないといけない．

解答ありがとうございます。
なんとなくですが皆さんのおかげで理解できました。
がんばってやってみます。

なんとなくじゃ無理だろうな

少し聞きたいんですが、
1>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
いきなりこんなエラーが出たんですがどういう意味ですか？
解決法を知っている方教えてください！！

http://en.wikipedia.org/wiki/Comparison_of_Nvidia_graphics_processing_units
誰かこのページでGPUの浮動小数点計算速度の比較に使われてるプログラム知ってる人いませんか？

>>272
ざっと見ただけではこのページのどこに浮動小数点計算速度の比較が書いてあるのか分からなかったのだが。

いずれにしてもそこに載っているスペックは全て理論ピーク性能かと。

>>271
いきなりすぎて難しいのでとりあえず中間ファイルの.ptxファイルを保存して146行目を調べて。

あとはSDKに含まれるプログラムとかは正常にコンパイルできるのかどうかも。

お聞きしたいのですが、プログラムを実行したときに発生するエラーで
invalide device function とは何が原因で起こるエラーなんでしょうか？
初歩的な質問ですがどなたか教えてください。

invalide device function
device で実行しているfunctionがinvalideなんだろ
でもinvalideじゃなくてinvalidじゃね？

スペルを間違えていました。invalideではなくinvalidでした。
invalidは「無効な」という意味で捉えています。
デバイス(__global__)の関数名とホストで実行する関数名は同じにしているのですが、なぜ無効と言われるのかがわからないのです。当然ですが、何かの関数名と被らないようには名前をつけています。
どなたかお願いします。

もう少し具体的に書け

_device_な関数を_global_と_device_な関数以外から呼び出してる予感

詳細を書きました。長文です。すみません。
画像処理を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__からしか使用していません。

いつ、どこでのエラーなんだ？
device内のfunctionでのエラーではないのか？

ホストでdevice関数を呼んだ後にcudaGetLastErrorを実行しており、そこでエラーが出ます。
device関数を呼ぶ前の処理におけるエラーチェックではエラーは発生していないので、device関数でエラーが起こっているのではないかと考えています。

ttp://lzhan.wikispaces.com/file/view/gflops.png/30874460/gflops.png
ttp://4.bp.blogspot.com/_nq8CjirJSCA/TDWRC5XR8vI/AAAAAAAACZY/pyHDbs2AA0E/s1600/gpgpu.png

上にあるようなCPUとGPUの浮動小数点演算性能の比較図を見るとGPUの方が圧倒的に高いように見えるのですが
性能はGPUの方が上という事でしょうか

それともただ単に1コアのクロック周波数から理論的な演算性能を算出してコア数分かけた値で、
GPUはコアが多いからそう見えているのでしょうか

さらにメモリの転送速度もGPUの方が速いのですよね。
他の点で総合的に見たらCPUのほうがいいということなのでしょうか

＞それともただ単に1コアのクロック周波数から理論的な演算性能を算出してコア数分かけた値で、
＞GPUはコアが多いからそう見えているのでしょうか

＞さらにメモリの転送速度もGPUの方が速いのですよね。
＞他の点で総合的に見たらCPUのほうがいいということなのでしょうか

それで合ってる

>>283
まあ、簡単に言うと計算の仕方とメモリアクセスの仕方について、GPUの方が制約がきついってこと
もっとも、CPUでも高い性能を出すためには厳しい制約が付くけど

>>284-285
ありがとうございます。

つまり全てのコアを常に使い切ればこれだけの性能が出るので、
メモリアクセスやのレイテンシや転送時の待ち時間に
準備の出来た他の演算(warp)を次々にやらせていって
いかに常にコアをフル回転させるか、がCUDAプログラミング。

warp内では実時間でも完全に同期して実行されるため
warp内分岐は、分岐を通らないwarpは完全に遊んでしまう(コアは空いているのに他のwarpのグループを実行できない)
ので遅くなってしまう(というかピーク性能が出せない)。

メモリアクセス時には他のwarpを計算すればいいが、次のwarpもメモリアクスになり、最初のwarpもまだアクセス中だったときには
3つめのwarpを計算する。このようにどんどん他のwarpを処理していくが、一度に起動できるwarpは上限がある(カーネル関数で使用するレジスタや共有メモリで決まる)
ため演算に対してメモリアクセス時間が長いと、起動上限数使い切ってしまい、
それでも最初のwarpがメモリアクセス中だとコアが遊んでしまうので、

メモリアクセス時間を減らしたり(メモリアクセス回数を減らす、共有メモリを有効利用する、
バンクコンフリクトを回避して共有メモリに早くアクセスできるようにする、コアレスアクセスでグローバルメモリに早くアクセスする)、
同時起動可能なwarp数を増やす(Occupancyうんぬん)
ことで見た目全コアが降る回転するようにする。

ということでしょうか

>>280
ついこの前出たOpenCV 2.2でNPP経由でCUDA対応したから、それ使ってみたら？
exeではなくCMakeから自前ビルドしないといけないけど

>>274
SDK内のプログラムはきちんと動作しています。
.ptxファイルは保存方法がぐぐっても良くわからなくて・・・。
どうやったら保存できますか？

>>282
状況がよく判らんが、__device__な関数は__device__か__global__な関数からしか呼んじゃダメよ。
それと、<<<>>>で呼んだ関数で起きたエラーは、次にAPIを呼んだときに戻されるのでご用心。
つまり、以下の場合cudaThreadSynchronize()の戻り値がエラーになる。
--
func<<<grid, block>>>();
cudaThreadSynchronize();

>>271
取り敢えず、-ptxつけてコンパイルして味噌。

>>289
コマンドプロンプトでnvcc -ptx cuda.cuであってますか？
普段VSでコンパイルしてるんでやり方がいまいちなんですが・・・。

Fermiでコンスタントキャッシュはハードウェア上ではどこにあるのでしょうか

ttp://www.4gamer.net/games/099/G009929/20100309044/images/pic2.gif
のUniform Cacheでしょうか

容量と読み込み速度からしてSM内かと思ったのですが
全てのスレッドから読めるという点ではSM外のL2キャッシュなのでしょうか
SM内だとしたら全SMがそれぞれ同じデータをコピーして保持するというのは少々効率が悪い気もします

そもそも大容量なはずのVRAMに転送するコンスタントメモリが
64kBと小容量に決められているという事はキャッシュも同サイズ用意されていて
ヒット率100％が保証されているということなのでしょうか

ちょっと質問なんですけど、GPU上ではstaticは使えないみたいですけど、
staticの代わりになるようなものはCUDAに存在するんでしょうか？
やっぱりグローバルメモリに領域を確保して使うしかないんでしょうか？

>>292
すまん質問の意味がわからん

>>293
すいません、ちょっと日本語がおかしいですね・・・
要するに、ＧＰＵ上のdevice関数内で、
関数が終わっても変数の値が残るようにしたいんですけど、
それをするにはcudaMemcpyでグローバルメモリに領域確保して、
そこに値を保存するようにするしかないのでしょうか？

それどういうときに使うの？カウンタ？


そもそもインライン展開されるんじゃなかったっけ

根本的に取り組む姿勢が間違っている気がする。

構造体の配列って
a,b,c|a,b,c|a,b,c
こう置かれるからハーフワープの各スレッドがaにアクセスすると
コアレスアクセスできないということですか？
配列の構造体.a[n],.b[n],.c[n]を作った方が良いという事でしょうか

誰だかまるわかりw

CUDA GPU Occupancy Calculator を使って占有率等をみるときにActive Thread Block per Multi Processor が高いほうがいい
と聞くのですが、以下の二つの状況だとどちらがいいでしょうか？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数が多いほうがいいのか、
あるいはどちらも同じなのか教えてください。

>>299
計測してパフォーマンスのいい方にすればいいとおもうが、
１の方があるスレッドブロックが同期待ちになってるあいだに他のスレッドブロックが
実行可能な余地があるぶん性能低下はしにくいと思う

>>300
回答ありです。
言葉足らずでした。。。
計測して実験したのですが、1のほうがパフォーマンスが上でした。

いい言葉が見つからないのですが、
1も2もブロック数は違うけど同じスレッド数ですので
メモリ隠ぺいのためにスレッド数が多ければそれでいいと思っていたので
性能は同じになると思っていました。

スレッドブロックが同期待ちというのは、syncした場合におこる現象でしょうか？
調べてみましたがやはりSMあたりのスレッド数が同じなら
SMあたりのスレッドブロック数が多いほうがいいというのが
いまいちわかりません。。。
よろしくお願いします

>>301
__syncthreads();
はブロック内のスレッドを同期する。

全部のスレッドが1024として
1ブロック1024スレっどの場合、__syncthreads();すると全部のスレッドが同期される。
その間どのスレッドも実行できない。

4ブロックで1ブロック当たり256スレッドの場合、__syncthreads();しても256スレッドだけ同期されて
他の３ブロック分のスレッドは実行可能。

＞SMあたりのスレッドブロック数が多いほうがいい
スレッドブロックはSM単位に割り当てられるから「SMあたりのスレッドブロック数が多いほうがいい」
というのはActive Thread Blocks per Multiprocessorが高いほうがいいというのと同義

Visual Profilerでlocal load localとstoreが数万回起こっているのですが何が原因でしょうか
--ptxas-options=-vでlmemが出ていなかったので
ローカルメモリは使っていないと思っていたのですが
local〜が0でないカーネル(3種類)はすべてレジスタ/スレッドが63でした。
偶然同じ数になっていると思ってましたが
レジスタは63が上限でそれ以降はローカルメモリが使われるってこと？

>>303
レジスタの容量を調べてみろ
例えばレジスタの本数が32768で1ブロックあたり512スレッド使うなら
1スレッドが使用できるレジスタは64以下になる、256スレッドなら128以下

ローカルメモリが使われるなんて状況は最悪だからそのプログラムは糞だ
シェアードメモリ余ってんならそっちも有効に使え

だがローカルメモリ使ったカーネルとCPU実行では，
CPU実行の方が遅いこともあるから必ずしも悪ではない．

総和計算を何段階かにわけたときって
sharedメモリ使う都合上、最後の段階は1ブロックでやるじゃん
でも1ブロックのみってことは他のsm使わないから無駄してるよね
最後の段階はcpuに転送したほうがはやいってこと？

fermi系ならコンカレントカーネルで問題無し

どうせ最後にcpuに持ってくるんだから、充分データ量が減ったところで転送しちゃうのも手ではあるね。

>>304
レジスタ32768ですがスレッドは192/ブロックです。
そもそもActive Thread Blocks per Multiprocessorを2以上にするように
1スレッドで使用しているレジスタ数とシェアードメモリを見て
スレッド数を決めているので足りなくなるなんて事はないはず

--ptxas-options=-vでlmemの項が出ないんだから使ってないと思っているのですが
レジスタ/スレッドに上限なんてないですよね？

そのレジスタ数から考えてfermi系ではないとおもじゅんだが
いっそのことスレッド数を８とか３２とかにして実行してどうなるか見てみればいいんじゃないか
なにか他の原因かもよ？

c/おもじゅんだが/思うんだが

>>310
GF100,Fermiです
8スレッド/ブロックにして(ブロック数を増やしたので総スレッド数は変わらず)
みましたがlocal load/storeは増えました。
カーネル内の変数を(なにかに使用するように)無駄に増やしてみましたが
レジスタ数が63以上にならなかったので上限があるのかも、と思った次第です。
皆さんはレジスタ数が64を超える事はありますでしょうか

今環境ないからできないがnbodyのサンプルとかはプロファイルとったらどうなんだろうね。
あとフェルミならローカルメモリはキャッシュが効くんではなかったかと思ったがどうでしたっけ？
あとコンパイルオプションでレジスタの使用個数指定してたりしない？

>>313
もちろん--maxrregcountは指定していませんでした

ところで
--maxrregcount 64にすると
ptxas warning : Too big maxrregcount value specified 64, will be ignored
とでました。
--maxrregcount 63にすると何も出ずコンパイルされます。

ということは６３に意味はありそうだね。うーんなんだろ。即答できない

Also, the compiler reports total local memory usage per kernel (lmem) when compiling with the --ptxas-options=-v option. Note that some mathematical functions have implementation paths that might access local memory.

Fermiは63までしかないよ。GT200とかは127まで行けたけど。

CC2.0でchar型の配列a[N]に
int idx = blockIdx.x * blockDim.x + threadIdx.x;

a[idx];

ってアクセスしたらきちんと32bitずつglobalから読みだしますか？
それとも128bitずつですか？

>>317
そうだったんですか
ちょっとずつ減らしてったら現在90くらい使ってることが判明
かきなおしかー

64しか使われないのに今90使ってるっておかしくないか

経験値からじゃなくて何か根拠のある数字はないのか？

--ptxas-options=-v→63
同時に使ってるであろう変数を1個減らす→63
もういっこへらす→63
繰り返す
へらす→62
やっと62になったということは減らした変数の数*2本分(doubleなので)=26本のレジスタが
63から溢れ出てローカルに保存されているんじゃないか
→本来なら90くらい使われてるんだろう

まあ一概に変数1個=2本とは言えないんでしょうけど
他に表示されないレジスタ数を数える方法思いつきませんでした

言ってることがようやくわかりました。

192スレッド/blockで動作してるなら
32768/192=170.66
で1スレッド170は理論的につかえるはずなのに他になにか制限みたいのがあるのかね

もしかして待機しているactive threadの分が消費されるのか？つまり
32768/192/ActiveThreadPerBlockが1スレッドの使えるレジスタ数になるのだろうか

>>322
逆に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表示されない
のだと解釈しました。

C programming guide 3.2 p83に
GT 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.

ってあるがこれか？

>>324
1ブロックで使用されるレジスタ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/スレッドの上限とは関係なさそうですかね。

上の方で出てたが、やはりCUBLASは省略すべき係数計算を
省略せずに計算しているようだ。
http://www.netlib.org/blas/dgemm.fに反する実装って許されるの？

許せるに決まっているだろう。

>>326
公式のフォーラムで本当にそうなのか、
もしそうならどのような理由からか聞いたほうがいいんじゃないか？
もしかしたら自分が勘違いしている部分があるかもしれないし
Nvidiaの間違いなら修正があるだろうし。

ttps://people.maths.ox.ac.uk/gilesm/cuda/lec4-2x2.pdf#search='Fermi%20register%2063'
　up to 63 registers
ttp://www.beyond3d.com/content/reviews/55/7
　Maximum theoretical register allocation per thread is 63 32-bit registers.

このへんに63までと書いてありますが公式にはないのかな...

CUDA VS Wizard + Fermiな環境の人いますか？
デバッグのための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

>>330
あくまでCUDA VS Wizardが不具合ならという前提だけど
sdkのtemplateプロジェクトを元に必要な箇所を修正してプロジェクトを
作りなおしてみるという手もある。

>>331さん
ありがとうございます。頂いたアドバイスを頼りに解決しました。

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の格納し忘れでした。失礼いたしました。

CUBLASのcublasSetmatrix()にCUBLAS_STATUS_MAPPING_ERRORを
吐かれてしまうのですが、どういった理由が考えられますか？
最初に最大サイズの行列サイズでallocした後、
そのdeviceメモリ領域に行列データを転送して演算、
結果をhostに書き戻して、さきほどの領域にまた別のサイズ（最大よりは小さい）の行列を
hostから持ってきて演算・・・のような処理を繰り返しています。

卒論書き終わりました
このスレのみなさんのおかげです
質問にいっぱい答えてくれてありがとう

ちゃんと謝辞に書いたか？

他のスレッドでは書き込めない超低レベルなチラ裏で恐縮ですが、
CUDA toolkitのサンプルのVecAddで、floatを全てdoubleに置き換えた上で
１．ベクトル要素数Nをどこまで大きくできるか
２．CPUによるVecAddとの速度差はどの程度か
を調べました。
※GPU=GTS450 1GB、CPU=AthlonII-X4 640＠定格。CPUのVecAddは単純な3行ループで非マルチスレッド。

１．16*1000*1000までセーフ、17*1000*1000は実行時エラー(debug)。
つまりcudaMalloc()量にして384MBはセーフ、408MBはアウトでした（GTS450 1GB)。
VRAMいっぱいまで使えるわけではないのですね。

２．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関数内だけで処理をさせるようにしないとうまみが少ないのですね。

気づいたらなんか本出てたんだけどどうあれ？
入門書？応用書？

最近？どの本のことですか？

CUDA by Example?

>>339
そうそれ。2/14って書いてあったけど翻訳かな

>>340
翻訳だよ。
内容は入門書

英文マニュアル読めるならいらない。
ストリーム、アトミック演算、OpenGLやD3Dとの連携のヒント
複数GPUを使う方法、ゼロコピーメモリ、ページロックメモリ
テクスチャメモリとコンスタントメモリの使用法などが書いてある。

最適化については記述されてない。つまり
ブランチダイバージェンス、GRAMのコアレスアクセス、シェアードメモリのバンクコンフリクト
の説明はない

CUDA by Example 汎用GPUプログラミング入門
ttp://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発売

超低レベルな>>336の訂正で申し訳ないですが、
clock()による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のシリーズはすごくわかりやすいと思いました。
こういうものが公開されているのはすごくありがたいです。

このスレ読めばちゃんと書いてあるのに……

公開されている資料といえばこんなのを見つけたが。
ttp://accc.riken.jp/HPC/training.html
の下の方のCUDA。

>>343
そういうあなたにCuda by exampleをすすめる

著者が必死に宣伝しているように見える

まぁ青木本が1番

SDKのドキュメントが一番だろう。
アレに必要なことは大体書いてある。
あとは実行してプロファイルとったり
nvccでcubin吐いてenvytoolsで確認したり。

日本語でおｋ

nVidia日本法人にまともなエンジニアがいないんだ。日本語だけじゃなんともならん

SLIしてる場合、CUDAからはどういう扱いになるんでしょうか
１つの大きなGPUとして見え１つのホストスレッドで扱えるのでしょうか
それともSLIしていない場合と同じように複数のGPUとしてみえ、枚数分のホストスレッドで管理してやる
必要があるんでしょうか。

前に聞いたときはSLIはなんらCUDAに関与しないとのことだったけど。

ということは複数GPU挿しているときと同様の扱いになるということですか。
ゲーム等ではSLIできてCUDAではできないんですね。
手軽さより自分で制御できる事を優先してるのかな

もともとコアとしては遠くにあるわけだしVRAMも遠いんだから同一扱いは無理でしょ
むしろSLIがうまく分担してるんじゃない

SLIで充分なんだから本来グラフィックスにGTX480みたいなおばけチップはいらないんだよね

でもそのおばけチップでSLIにしたらもっと性能あがるじゃん
とくにCUDAスレの住民にはメリット大きいんじゃない？

>>344
すみません、最初から読んだはずだったんですが、読み取れなかったか見逃していたみたいです。
>>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)。

>>358
まだなら「はじめてのCUDAプログラミング」を先に勧める。
英文が苦手ならさっさと本買ったほうが時間の節約になるし
実際のプログラムに時間を使った方が有益だと俺は思うよ。

>>359
そうですね…約2500円で省略できる時間のほうが有意義のように思います。
よし…買おう…bk1の300円ギフト券が手に入ったら…。

>353
>前に聞いたときはSLIはなんらCUDAに関与しないとのことだったけど。

使えるっぽいぞ
http://jp.slizone.com/object/sli_cuda_learn_jp.html

SDｋにSLID3D10Textureてサンプルがあってソースみたんだが、
最初にDirectXとGPUのチェックがあるだけで
特にSLI独自のコーディングは見当たらないな
かといって複数のGPUを制御するようなスレッド関係のコードもない

それらからSLI組んでれば1つの大きなGPUとしてプログラムからは扱えるんじゃないかと
思ったんだがどうなんだろうか。

>>361
なんだか滅茶苦茶な翻訳だな。
で、そこを読む限りPhysxとの連係には触れられているけど
それ以上の情報はないし、英語サイトにはその記述さえ見当たらないね。
うーん、どうなんだろう

俺の友達も教務情報改造して人数出るようにしたりしてたけど日本人だったぞ

間違えました
ごめんなさい

最近出たインプレスの本にSLIのことが載ってた

>>366
たしかに載ってるが概要の紹介程度で、SLIを使う方法や
SLIからどうGPUが見え、どのように制御するか（orできるか）は書かれてない。

索引から逆引きしても１０行くらいしか書かれてないじゃん。

新型マックじゃ使えないの?

Mac用のドライバもtoolkitもあるだろ
あとはGPUが対応してるか調べろよ

使えない

CUDAのアプリを作り始めたのですが、CUDAのSDKがインストールされていない環境で
アプリを動かすと”cudart32_32_16.dll”が無いと怒られます。

cudart32_32_16.dll で使われている関数等々をスタティックリンクする方法があれば教えていただけないでしょうか？

CUDA 3.2を使っています。

よくわからんけどコンパイル時に
-L /usr/local/cuda/lib64/ -lcudart
じゃだめ？
パスは適当

>>372
リンカのオプションの追加の「依存ファイル」で「cudart.lib」は指定してあります。

ツールキットのcudart.lib のファイルサイズは約41KBのものが１つのみ。
cudart32_32_16.dll のファイルサイズは約375KBです。

スタティックリンク用のcudart.libが無いということでしょうか？？？

開発環境を詳しく書きます

[開発環境]
OS: Win7 64bit
IDE: Visual Studio C++ 2008 Express
CUDA Toolkit: ver 3.2 32bitバージョン

>>371
CUDA RuntimeのDLLは再配布が許可されているみたいだし
単に他の環境で動かしたいだけなら、DLL同梱でいい気がする。

ttp://developer.download.nvidia.com/compute/cuda/3_2_prod/toolkit/docs/EULA.txt

>>374
ほむ！
ありがとうございます。
インストールしたときについてくるライセンスではそこら辺がよくわからなかったので助かります。
とりあえず、DLLも同梱しようと思います。（スタティックリンクしたいけど・・・）

4Gamer.net ― NVIDIA，「CUDA 4.0」の新機能を予告。より広い層の開発者にアピールする新バージョン（CUDA）
http://www.4gamer.net/games/076/G007660/20110228018/

Fermi持ってたら恩恵受けられるのかな。なんかよくわからんけど楽しみだ。
＞ちなみにGPUDirect 2.0とThrustをフルサポートするGPUは，Fermiアーキテクチャ限定となる。

>>376
・GPUDirect 2.0 …スパコンには良さそう？
・Unified Virtual Addressing…どれが何のメモリなのかごっちゃになりそう？
・Thrust…STL。。。どうでもよさそう？

こんにちは。cubinとptx関連の情報が欲しくてごにょごにょいじっていたのですが
ちょっと行き詰まりましたのでできれば２点ご教授いただきたく…

１．
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アプリには自動で付く？？？)。
どうすれば競合を解消できるでしょうか？


２．
仕方がないので、とりあえずコマンドラインから
-Xptxas="-v" -cubin --ptxas-options -v
を指定してコンパイルしたところ
無事にcubinファイルが出来たのですが、このcubinファイルはバイナリー形式で私の目では読むことが出来ませんでした。
このcubinをテキスト形式で読むにはどうしたら良いでしょうか？
(CUDA 3.0からcubinは一旦バイナリ形式で出力されるようになったようです)

[環境]
CUDA Toolkit, SDK 3.2

ptx出力じゃ、用が足らんのけ?

初心者です。
今からくだ勉強したい場合はCUDA4.0がオススメですか？

>>380
むしろcuda2.2くらいから。
基礎ができてから発展にすすめ

381さん有り難うございます。cuda2.2からやってみます。

古いバージョンのCUDAを知る必要は普通はない。
CUDA 4.0でも基本的にこれまでの機能を全て備えている。
また古い世代のGPUは余裕がある場合だけ追うべき。

>>383
同意

つか、
>cuda2.2くらい
"くらい"なんて中途半端な答え方してるのかがワカラン

>>379
readelf でcubinみましたが、いまいちよくわからなかったのでptxみます。
まあ、ptxもよくわからないのですが…

PTXを読んで分からない人にはcubinは無理。

CUDA 4,0 RC が来たよ

Visual Studio 2010 まだ来ないよ

すみません。
どなたかAMBER(Assisted Model Building with Energy Refinement)というソフトをCUDAで動かしてる人いませんか？
このソフトを
Cent OS5.5 x86_64(カーネル 2.6.18-194.el5)
インテルコンパイラーとの併用で動かそうと何度も試みているのですがなかなか上手くいきません。
不具合の原因はGCCのバージョンとインテルコンパイラーのバージョンの組み合わせにあると睨んでいるのですが…
もし、動かせている方がいらしたらGCCのバージョンとインテルコンパイラーのバージョンの組み合わせををヒントとして教えてほしいのですが

>>389
コンパイルでエラーなのかリンカでエラーなのかわからん。
取り合えずsdkのサンプルでiccで動かしてみたら？それで問題がわかるでしょ。
それが出来ないなら、
そもそもnvcc自体がgccのラッパだし、
素直にgcc使えば？

デバイス関数で、VC++でいうところの Sleep(0)、gccでいうところのsleep(0)？？？と同等の動作をする関数を探しているのですが
適切な関数があれば教えていただけないでしょうか。

眠らす必要がなさそうだから無いと思う

>>392
フロントエンドというかデスクトップというか、とにかくグラフィックスの描画が遅くなる（当たり前ですが…）のと、
GPUの温度がかなり上がるので、
デバイス関数のループの中に、ちびっとだけSleep(0) 的なものを入れたいのですが、
関数的なものは用意されていないっぽいですかね。

例えば、どれか１つだけスレッドの処理ループを長引かせて、__SyncThreads() でもしたらいいのかしらん…？？？

利用者の大半はカーネルを一気に計算したい派だとおも
熱いから休憩するって発想は俺含めてないんじゃね？

自分はカーネルを終了させる必要があると思います

>>395

例えば、

__global__ void test()
{
　　...
　　for (i = 0; i < LOOP; i++)
　　{
　　　　__syncThreads();
　　　　...
　　　　// Sleep(0) 的なもの？？？
　　　　if (blockIdx.x == 0 && blockIdx.y == 0)
　　　　{
　　　　　　// 何か冗長な処理
　　　　}
　　}
}

って感じじゃSleep(0)の代わりにならないですかね？

それはビジーウェイトになんね？
結局熱くなってしまうんだが

>>397
そうなんですか…
ビジーウェイト調べてみます。

>>396
カーネルを分割すればいいんじゃね?
わざわざ遅くする理由を思いつかないから試す気にもなれないけど。

>>399
これが正解だろう。
でもそもそも、ちょっと休ませたくらいじゃ熱なんて変わらんだろうから、ドライバの設定でクロックを下げたらいい。
もしくは、放熱を強化するかだろう。

CUDAでGTX460(1.5GHz)、4枚分程の計算能力を求めています。
GTX460を2枚刺ししたPC2台と、4枚刺ししたPC1台のどちらを組む方が良いでしょうか？

4枚刺しですと、電源のとりまわしと冷却に不安を感じております。

GPU間の通信はどれくらいしますか。

そういうのを意識しないなら、なるべくわけた方が安定しそうですね。

>>401
単純に、書かれた内容だけで判断する限り前者の方が計算能力が高い。
但し、後者により強力なCPUとメモリを搭載する場合は一概に言えない。

>>402
GPU間の通信は無いです。

>>403-404
GPU間の通信が無く、CPU側の処理も軽いので
4枚刺しによる恩恵が特になさそうなので
2台で組んだ方がよさそうですね。
ありがとうございました。

GPU計算早すぎ、ワロタ。
長い目で見ると、DSPとかで食いつないでた信号処理屋的に失業の危機じゃん…

GPUメモリのTLBミスが気になるのですが、何か情報ないでしょうか

>>406
カスタムチップによる高速演算ボードを出していたところは尽く撤退しました。

sharedメモリへのバンクコンフリクトは避けたほうが良いって話だけど、
共有メモリが16バンクに分割されているってのは�@�Aのどっちだろ。

�@マルチプロセッサのもつ16384Byteが16分割されてる
�Aブロックに割り当てた共有メモリが16分割される

ハードウェアが固定されてそうだから�@の気がするんだけど…。

アハ〜♪”

GPUDirect 2.0について質問です．

GeForce 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だと思っている)
とは別のものなのでしょうか？

試された方はここに居らっしゃいますでしょうか？

CanAccessPeerは無視してよい。
確かにTeslaシリーズでは1が、GeForceシリーズでは0が返る。
だけど実際にはGeForceでも通信はできる。
UVAだかなんだかに対応するGPUなら問題ないのかも。

ワープってやっぱりSMの上限いっぱいまで詰め込んだほうがいいの？
例えばSMの上限ワープが32のところに16個のワープにしても、
メモリアクセスのようなレイテンシがないなら問題ない気がするんだが。

CUDA初心者です。
CUDAを使って、ゲームエンジンのUnrealとかcryを使ってる方いらっしゃいますか。
参考になるサイトを教えてほしいのですが

例えば、デバイス側の主要な処理が下のような論理演算であった場合、
論理演算部分のアセンブラ化による高速化は望めるでしょうか？

それとも、レジスタの使い方がコンパイラの方がうまくて、
より高速なアセンブラを書くのは難しいでしょうか？

__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];
　　}
}

少なくともその論理演算部分をアセンブラで書く必要はない。
そもそもメモリアクセスがボトルネックにならないほど多くの演算があるのかどうか。

分岐を避けるためにビット演算使いたくなるけど、いまいち貧弱なんだよなあ・・・

英語か日本語でptxファイルの読み方を解説しているドキュメントがあれば教えていただけないでしょうか。

また、コンパイル時にptxファイルに明示的にコメントを挿入するコマンドがあれば教えていただけないでしょうか。
例えば、test.cu に、「__PTX_COMMENT("foo");」という行を記述すると、ptxファイルに
// foo
といった行が挿入されるという具合のものを探しています。

PTXのマニュアルはToolkitに含まれているかと。

ptxにブロックごとに行番号情報は含まれているのでソースのどこに相当するかは分かるのでは。
特定の場所が知りたければ例えば__syncthreads()を１０個入れとくとか。

perl とかでつくれそうだよね

>>419
なるほど。例えば、ダミーでfoo0()とか、foo1()とか関数を作ってコメント代わりに呼び出せばよいと。

>>420
コマンドとしては無いということですね。

CUFFTのハンドル確保するだけでエラーが出る…。↓エラーのでるコード。

CUT_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は実行して変換結果も正しいのだが謎す。

cufftは内部であれこれ怪しげなことしているからなぁ。
VRAMは未だ余っている筈なのにエラーになったり、作ったプランを使わずに破棄するとこけたり……
って、>422はそのパターンか?

>423
cufftExecC2C実行してもやっぱでる。
海外フォーラムみてたらver 3.2が怪しくて ver3.0にしたら消えたって書いてあったんで
ver3.0にしてみたがやっぱでる。うーん・・・。

>>424
うーん、SDKとcufftの組み合わせで相性あるしなぁ。
あと、サイズが2の冪かどうかでも変わるかもしれない。
つーか、3.1よりは3.2の方がましだった気がするんだが……

サンプルはmain関数でCUDAを利用するのばっかだけど、
VC++でダイアログから.cuの関数使うようなサンプルってなんかない？

>>426
おまいさんは関数を分割することもできんのかね。

windowsよく知らんけど
それはライブラリ作んなくてもできるんかい？

>427
extern "C" つけたらできたや。ごめんごめん。

>>429
お前はだれ？
そんで誰にレス支店の？

まちがえちゃった。テヘッ

プログラミングガイドに記載されているように

CU_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を使用しています．

CU_SAFE_CALL(cuInit(0));
CU_SAFE_CALL(cuGLInit(0));
CU_SAFE_CALL(cuDeviceGet(&device, 0));
CU_SAFE_CALL(cuGLCtxCreate(&context, 0, device));

ミスった。
1点だけ。
uGLInit(0)→uGLInit()

>>433-434
この通りにしましたところ，今度はcuGLInitで
Cuda driver error c9
が発生しました...

コマンドラインからnvccでコンパイルする際、普通のC++部分のコンパイルを
Intelコンパイラ："icl.exe"に行わせたいのですが、どのようにコマンドを記述したら良いでしょうか？

環境：
Windows 7
CUDA 3.2

nvccしか使えないところはライブラリ化してやってるけど
もっといい方法あるんかいな？

西日本も福島原発の放射能に曝される。

4/7 予測 http://up3.viploader.net/ippan/src/vlippan198234.jpg
発表はドイツ気象庁 http://www.dwd.de/

CUDAで予測出来ないんすか？

できますん

どっちですか

日本語で予測できるか質問するくらい意味がないことに気付け

>>442
さすがにそれはちょっと違うと思うけど
気象系シミュレーションでGPUつかってるところって結構あるよね？

？

さすがにこれは？？？だな

>>435

cuInit
cuDeviceGet
cuCtxCreate
cuGLInit
cuGLCtxCreate

じゃね?

>>446
やはりcuGLInitで同様のエラーが出ます．

3.2でやってみれ

>>448
3.2でも同様です．

OpenGLは初期化してるよね?

>>450
大変御迷惑おかけしました．
自分の無知さに絶望しました．

OpenGLの初期化後に
CU_SAFE_CALL(cuInit(0));
CU_SAFE_CALL(cuDeviceGet(&device, 0));
CU_SAFE_CALL(cuGLCtxCreate(&context, 0, device));
で動作させることが出来ました．

御手数おかけしました．

シュミレーション

あながち間違ってはいないな

cudaMallocとかcudaFreeとかcudaMemcpyとか、どのヘッダに定義されてんの？
cutil.hにもcuda.hにも定義されてねーのだが。

cuda_runtime.h

　　　　＿＿＿
　　 ,;f　　　　　ヽ 　　　　 　　　
　　i:　　 　　 　　 i 　　
　　|　　　　　　　　|　　///;ﾄ,
　　|　　　　^　　^　)　////ﾞlﾞl; 　　
　　(.　　>ﾉ(､_, )ヽ､} l 　 .i .! | 　　
　 ,,∧ヽ　!-=ﾆ=- | │　 　| .|
／＼..＼＼｀ﾆﾆ´ !,　{　　 .ﾉ.ﾉ
／　 ＼ ＼￣￣￣../　　 / .|

いまだにキューダと読んでしまうのだが、クーダと読む習慣が付いてる人はどれくらいいるんだろう

あなた以外ほとんどがクーダだと思います

main.cu

メイン ドット クー？

OpenGLとの連携が遅いでござる。
フォーラム見ても遅い遅いと書いてあるし、どうしようもないのかこれ？

warpはなんて読む？

ﾜｰﾋﾟｰ

ウォープ

飯野の巣

warning
walking
working
star wars
それぞれなんて読む？

それ、この板で何度も繰り返し振られる話題だよな

>>465
ワーニング
ワーキング
ウォーキング
スターウォーズ

スターワーズだりょ

threadIdxがidのxだといつも勘違いしてしまう

stdio をスタジオだといつも勘違してしまうのとはちょっと違うか

カタカナ表記した時点で正しい読みから外れてる

>>469
わかるわ
油断するとIdxとIdyがあると思っちゃう
threadIdx.y
xのy？？？
スレッドはフルスペルなんたからthreadIndexでいいじゃねえか

寧ろptx出力宜しくtid.x希望w

明示的にグローバルメモリから共有メモリへの読み込みを行う命令ってないのかね。
スレッドに分けて暗黙にコアレスアクセスさせると、どうもコードが汚くなる。

ううーん、デバイス側にメモリ確保しっぱなしだとOpenGLの性能が著しく低下する…。
計算結果を最終的にOpenGLで表示するような場合だと、無駄なメモリコピーは仕方ないものとして
計算結果をホスト側に逐一戻したほうがよさげ？

3.2から4.0 RC2に変更したんだけど，Cuda.rulesってないの？

つかなんでcommon\lib\Win32の中にcutil32.dllがあるんだっていう

CUDAのアセンブラにMMXやSSEのような命令はあるでしょうか？

（CUDAのカーネル関数の中で、MMXのPADDDのような32bit整数の4並列計算のようなことができなかと思っています）

>>478
CAL+IL

2枚差しの場合、2つプログラムを走らせたら、それぞれ別のGPUを使ってくれるのですか？

現状では自分で何番目のGPUを使うか指定する必要があります。

>>477
イミフ

>>479
えと、CUDA用には無いということでしょうか？

add.u32

を2並列(64bitレジスタ)、あるいは4並列(128bitレジスタ)で実行するような命令セットはないでしょうか？

>>483
そのためのSIMT

>>483
何かを誤解しているようだが、CUDAには必要ないから無い。
SSEを使わないと100%の性能が出せないとかいうCPUのような制約は無い。

add.u32という命令ならほっといてもWarp単位の３２スレッドが（ほぼ同時に）実行する。
全命令が３２並列のSIMD命令になっていると考えればよい。（３２ビットデータの場合）

>>484-485
CUDAがSIMTなのか、SIMDなのか、SPMDなのかはさておいて…

昔の記述だとCUDAアーキテクチャのレジスタは32bitとかだったと思いますが、
今時のNVIDAのGPUには64bitレジスタ、あるいは128bitレジスタがあって、それを使って
padd.u32 みたいな命令により、「1スレッド内の演算において」32bit整数の4並列処理が出来るとかないでしょうか？

ということなのですが、いかがでしょうか？

Streamでは128bitレジスタで32bit4並列の演算ができるみたいのですが・・・

そこまでやるならＣＰＵの方が早かったりしてｗ

>>486
そんなものはないから誤解だと言ってるんだが。
AMDは１スレッドが4 or 5演算のVLIWで動作するから4並列のSIMD命令化の意味があるが
CUDAは１スレッドが各種演算ユニットを１ずつしかもたない(or複数スレッドが共有）CUDA Coreで
実行されるので128bit型を用いたとしても４命令になるだけ。

>>488
誤解というか、質問なのですが・・・

今どきのCUDAでもレジスタレベルのSIMD命令は無く、
ハードウエア的に、レジスタも演算器も（今のところは）そのような用途には向いていないということですね。

ありがとうございましたゞ

それよりSandyBridgeのベクタ演算機の性能はどうなんだろうな
まだ試験実装だからコア数が少ないけどメインメモリ直結だからけっこう早そうなんだが

その辺は簡単な報告を以前SSEスレにあげといたよ
イミフな反応が多々あったけどね

>489が根本的に解ってない悪寒

>>492
論理レジスタとか、物理レジスタとかいう話しですか？？？
>>486への根本的な回答はどうなりますか？

とりあえず公式文書を読んでから質問してくれ。

ようするに64bitレジスタに8x8のデータを入れて同時に処理出来るのかってことでしょ
streamsdkには専用の命令があると？
cudaはそんなものは勝手にやると言ってる人が居るけど、どうやってやるんだ？

cudaにはそういう命令は無いよ
内部で処理することが出来るとは思えないけどな
少なくともユーザープログラムからやる方法はないよ
内部的にはそういう命令も持ってるかもしれないけど
というかそういうアセンブラレベルの最適化をしなくていいようにしたのがＣＵＤＡ言語なわけで
複雑化する要因はパフォーマンスを犠牲にしてでも排除されるだろう
どうしてもやりたかったら64bit変数を使って自前でやるしかないだろう

>>495
>>486 への回答は”No”ということですね。
ありがとうございます。

将来的には"Yes"になると思っているのですが、そこら辺が「根本的に解ってない」と言われる所以ですかね。
「根本的に解ってない」らしいので、何がなんだかさっぱりですが・・・

# 今、スクラッチのCUDA用アプリで、SSEで書かれたCPU用アプリの約5倍（CPU側はSandyBridgeで4スレッドで動作）の速度が出ているのですが、
CPU側がきちんとAVXに対応してきたらこの差はかなり縮まりそうなんですよね・・・

ねえ、GPUの速さの要因ってCPUより圧倒的に軽量なスレッド数によるもんじゃないの
まずそっからじゃないのかね？

>>496
さっきからすげー鼻につくんだけど
教わる態度じゃないよお前

教えてんの俺じゃないけどさ

あれだ、自分が書いたcudaのコードが思ったはど速度が出ないから
腹癒せしているか言い訳探しているかってところだろ。
今日たまたま計測していたんだが、cufftでもfftwfの3倍程度だった。
汎用ルーチンじゃこんなもんだよ。

>>498
CUDAの将来的な見通しについて知りたかった部分もあったので
失礼に見えた部分があるかもしれません。

ごめんなさい。

SIMDみたいなものつけて各コアを高機能にするよりは、
単純なコアの数を増やす方向に進むだろうなあ。

レジスタ/SM増やして欲しい…

GF100やGF104の倍精度（64bit）の演算って単精度（32bit）の1/4のスループットとかじゃなかったかな？
おまけにGF104は倍精度をサポートしているコア数は1/3だし。
更にHPC向けではなくグラフィック向けの物は制限かかっているらしいし。

倍精度が必要になるなら仕方ないだろうけど、32bit変数で足りるのに64bit変数を使うのはデメリットの方が圧倒的に多そう。

いろいろ試したがＣＰＵと比較して４倍程度しか上がらないし
労力掛かるわりには環境依存甚だしいわ
４倍くらいだったら別に我慢出来ないようなものなんてないし
マルチコアＣＰＵで普通に組のが一番いいという結論に達した

>>504
計算したい内容を簡潔に書いてみ
計算内容によってＧＰＵに向き不向きがあるから

>>497
膨大なスレッド数と、その切り替えの軽量さを活かして
メモリアクセスのレイテンシの隠蔽とかだっけ。

>>502
スレッドあたりの変数が多いと苦労するね。
そういう場合はシェアードメモリをうまく使って切り抜けるしかないのかな。

CUDAは4nbyteを1要素とした16要素SIMDって感覚だわ。(n = 1,2,3,4)
メモリアクセス考えると結局16要素単位の処理の塊になってしまう。

>>506
共有使ってもたらん
Sharedも増やせ！キャッシュも増やせ!

向いてないよ。

CUDAとかGPGPUの普及の足引っ張ってるだけだよなぁ。
早くintel、amd,nvidiaで統一してくれよ。
機械語レベルでも、ライブラリレベルでもいいからさ。

OpenMPってよく知らないが_beginthreadしてるわけじゃなく
CPUレベルでスイッチングしてるだけみたいだから
ほとんどオーバーヘッドがないみたいだぞ
十分標準として使えそうなんだが

残念ながら生成負荷は大きいよ
だから何度もスレッド生成するような処理はかえって遅くなる

いやいや、ちゃんとスレッドプールしてるって

>>513
おまえ>>504だろ
試行レベルが同一

>>514
お前、いろいろと残念な人だな

と、せ〜いっぱいの切替し
つか下らんこと書くならよそでやれヴぉけ

切り返し
な。
煽るなら突っ込まれる隙作っちゃダメだ

いいからスレ違いだ。

スレッドの生成負荷は軽いと言われているのを鵜呑みするとなかなか性能が出ないこともある気がする。
カーネル側でスレッドの生成回数を出来るだけ減らす工夫が必要だったり。

スレッドの生成負荷はプールすればいいから問題じゃない
スレッドを切り替えるときにスタックポインタ変えたりするコンテクストスイッチの負荷が問題

cudaでスレッドのプールができるのかえ？
cpuの話をここに書いてる？

そうそう、何度指摘してもスレ違いの話を続けているの。

その話、終わりました

どこまでをCPUで計算し、どこまでをGPUで計算するのかはGPGPUで重要
だからスレッドの性質については知っておく必要がある
CUDAを使いこなすのにアップアップのアホどもには関係ないけど

解決を依存性のない単純に並列化可能なアルゴリズムに置き換えられる問題を
GPGPUでやるもんだと思ってたんだが違うのか？

そりゃそうだが、データの転送速度、超越関数を使用するかどうかなど、
単純に並列度が高ければGPUでOKってもんじゃない
Opteronの48コアマシンがお手頃価格で手に入る時代だぜ？

あ、誤解の無いように書いとくが、CUDAやGPGPUを否定してるわけじゃないよ

夏場に熱で吹っ飛んだのれす

>>526
> Opteronの48コアマシンがお手頃価格

今、おいくら万円くらい？

全部こみこみで100マソ切った

高っ

そうか、ごめん……

48コアで何テラFlopsぐらい出んの？

高いけど、業務で使うならテスラを組み込んだサーバとたいして変わらないんじゃない？
個人とか研究室とかだと絶対的金額から手が出せないかもしれないけど。

それは言えてる。

>>533
えーと、0.5テラFlopsくらい？

17万x4(CPU)+12万(ﾏｻﾞﾎﾞ)+HDDﾒﾓﾘｹｰｽか
欲しいけど家庭用のコンセントで大丈夫だろうか
似たようなスペックでGPGPU構築したらいくらぐらいかかるのかな

0.5テラFlopはしょぼくないか

>>537
家庭用コンセントじゃ無理
グラボ刺さるからGPGPUできるよ

>>538
このマザーで自作すると・・・
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

>>539
> 下のマザーボードだと気合いで8本GPU載せられるんだろか？

どこかに引っかかりそう。ラック筐体次第なのかな

マザー自作するツワモノいないかな

まあ、高周波シュミレソフトで１０００万いくから
性能気にしないでいくなら自作できそうだけど・・・

>>539
そのマザーのPCIeスロットはx8+x1+x8+x1+x8+x1+x8+x1だから
ホスト・デバイス間の帯域不足に悩みそう。

演算性能さえあればいいプログラムなら何とかなるのかもしれないけど。

というかハイエンドグラボは複数スロット占有が一般的だし、水冷のグラボが必須になりそうｗ

コスト度外視で1台で最速を目指すなら
P6T7 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
になるのかな。

電源の確保がものすごく大変そうだけどｗ

CUDA Visual Profiler使ってるんですがなんどやっても計測後に
"Error in profiler data file'(動かすディレクトリ)/temp_compute_profiler_0_0.csv' at line number 1. No column found"
と出て計測結果がうまく出ません
何が原因なのでしょうか？

エラーコード見てもわからんが、俺はプログラム中にgetchar()を書いてある
キーを押さなきゃ終わらないプログラムをプロファイラに渡してはまったことがあるZE

>>544
Visual Profiler使ってないからよく判らんが、その作業ディレクトリはレガシーな名前かい?
空白やら何かの所為で巧くアクセスできないってことはないよね?

>>545
外部から画像を読み込んでいますがキーは押さなくても終了するプログラムです

>>546
作業ディレクトリと同じディレクトリに入っている計測したい実行ファイルは
プロファイルを開始すると実行されるようなのでおそらくディレクトリにアクセスは出来ています

プロファイラはプログラムを実行してから計測をしているようなのですがこの計測経過が100%になったら先ほどのエラーがでてしまいます
temp_compute_profiler_0_0.csvは自動生成のファイルの用です

>547
エラーがなんで出てるかはさっぱりわからんので、自分がこけたときの経験を列挙するぐらいしかできんのだけど、
Visual Profilerのデフォルト設定では30秒以内に処理が終わらない実行ファイルは実行を途中でうち切って
解析結果出力せず、えんえん解析を繰り返すってのはあったぜ。

>>548
そこ毎回適当に300に変えてたわ
しかしなんで一回でぜんぶできないかな

グラフィックボードがGPGPUの処理以外にも画面表示の処理をしてる可能性が高いのだから、
画面表示による影響を平滑化するために複数回実行時間を計測してるんじゃないかと。

コアレスアクセス用に半ワープ単位でアクセスしなきゃいかんのだが
16×16の二次元のスレッド作ったとして

threadIdx.x = 0〜15 が半ワープに入るのか
threadIdx.y = 0〜15 が半ワープに入るのか、どっちだべ。

ふつうにかんがえてx

アブノーマルに考えてx

真逆な意見がでてしまったわけだが、間をとってxと考えることにするよ。

Larrabee はいつ出ますか？

Larrabeeはなかった事にされてるような・・・

LarrabeeってSandryのGPU部分のコア数を数百って規模に拡張したもののことだよ
AVXがまさにLarrabeeのインターフェースらしい

VisualStudio使ってEmuDebugの構成にしたら、
デバイス側関数もブレークポイント仕掛けられるって聞いたのにできねえ。
リンクするライブラリ変えるとかしなきゃいかんの？

NsightじゃなくてEmu使う理由は？

Nsightってグラフィックボード１枚指しでも使えんの？
Nsightを使わない理由は、NsightインストールしたらCUDA3.3がインストールされて、
でCUDA3.3のcuFFTの動作が怪しかったんで、NsightやめてCUDA3.0にしてる。
NsightでもCUDA3.0使えるのかもしれないけど、使い方調べるのがめんどくさかったので
環境構築に慣れたNsight使わないやり方でやってる。

3.3?

CUDAとSandy BridgeのGPUは同時に使えますか？

使えるとしたら、Sandy BridgeのGPU用のコーディングには何を使えば良いですか？

>>562
SBってGPGPUに対応してるの？
SDKとか出てるって聞いたことないんだけど。

Sandy BridgeのGPUはOpenCL1.1に対応してるらしいけど
CUDAに対応してるって話はみたことないから、CUDAとは同時に使えないと思われる。
OpenCLスレで聞いたほうがいいかも。

>>564
>CUDAに対応してるって話はみたことないから、CUDAとは同時に使えないと思われる。
対応していない=同時に使えないという考えは改めた方位が良い。

グラボを刺した状態でSandyのGPUコアが有効になるのかって話だよね
グラフィックチップとしては完全に機能停止されるだろうけど
演算装置として稼動するんだろうか

EmuDebugはリンクするライブラリをcudart.lib→cudaartemu.libにかえたらできたや。

自身ぎゃああああああああああああああああああ

VS2008でFirst chance Exceptionがでるが、カーネルコードの何が原因なのかさっぱりわからねえ…
Nsight使うとこういうのの原因すぐに特定できるようになんの？

CUDA 4.0 RC2 をインストールしてみようと思っています。
64bit用のバイナリを作るには、CUDAも64-bit版をインストールする必要があるでしょうか？
それとも、CUDAの32-bit版でも nvcc のオプションの切り替えで32bit用と64bit用の両方のバイナリを生成することが出来るでしょうか？

出来れば、WindowsXP(32bit) の環境で64bit用バイナリも作りたいと思っています。

>>570
＞64bit用のバイナリを作るには、CUDAも64-bit版をインストールする必要があるでしょうか
たしかそうだったはず

＞CUDAの32-bit版でも nvcc のオプションの切り替えで32bit用と64bit用の両方のバイナリを生成することが出来るでしょうか？
生成したい処理系のSDKをインストールしないといけなかったはず
切り替えはできないと思う
2つ（３２と６４）いれればOKなのかは不明

>>571
レスありがとうございます。
ややこしそうなので、32bit 環境、64bit 環境それぞれを用意し、それぞれ用のバイナリを作ろうと思います。

64bitOS環境に64bit Toolkitを入れれば32bit&64bitの両方のバイナリを生成可能。

GT430を使っていますが、CUDAでコンパイルすると次のようなエラーが出て
コンパイルできません。
因みに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

__device__/__global_　この辺良く見直せ！
基本がわかってないと思うから基本からよく勉強だな

エラーの意味がよくわかりません。__device__/__global__なんてどこにも
書いて無いし、サンプルコードなので動くはずなのですが、なぜか動きません。
http://codepad.org/oDh0YFM5
cuComplexをjulia関数から呼ぶのは許されていないという意味なのでしょうか？

globalが無いのかよ？
全部晒せよ

http://ideone.com/s9anZ

ここです。よろしくお願いします。

CUDA_by_exampleのだな
まずそういう情報を先に出せよ

すみません。書くのを忘れていました。
知っている方とお見受けしますので安心です。

特に問題なくコンパイルできたぞ
nvccのオプションは？

nvcc -O julia_global.cu -lcutil32 -cutil32D

インストールの仕方が悪いのでしょうか？

windowsはよく知らんがLinuxではnvcc -lglut julia_cpu.cuで通るんだが
上のソース名は間違いだよな？それでやってどのエラーメッセージが出るんだ？

julia_gpu.cu(44): error: calling a host function("cuComplex::cuComplex") from a
__device__/__global__ function("julia") is not allowed
がでるのか？

julia_gpu.cu

です。間違いです。julia_cup.cuの場合はこちらでも通ります。
ただGPUを使おうとするとコンパイル時にエラーが出ます。

cuComplex c(-0.8, 0.156);
cuComplex a(jx, jy);
と
a = a * a + c;

の部分でエラーが出ます。原因がわかりません。

正確なコンパイルオプションとエラーメッセージを再掲よろ

http://ideone.com/mZviD

ここです。

C:\Users\admin\CUDA\chapter04>nvcc -O julia_gpu.cu
これじゃ -lcutil32とかないじゃん　なけりゃダメだろ

リンクエラーが出るかもしれないですが、今はコンパイルエラーなので
そちらでも同じようなエラーが出るか確かめてほしいです。
もしエラーが出なかったら私の環境がおかしいのかも知れません。
Linuxでの開発環境とWindowsでは少し違うと思うのでWindowsでも
確認してもらいたいです。

ウインドウズアプリを作るのでリナックスでうまく言っても意味が
無いのですが・・・

>>589
あ、そういうこと
じゃあソースが壊れてるんでなければそっちの環境がおかしそうだね
環境の再構築したほうがよさげ

因みにlinux環境で
nvcc -lglut julia_gpu.cu　→　OK
nvcc julia_gpu.cu　→　コンパイルOK、リンクNG

C:\Users\admin\CUDA\cuda_by_example\chapter04>nvcc --version
nvcc: 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のバージョンを教えてください。
環境を再インストールしましたがやっぱりだめです。
構文エラーとして解釈されます。

3.2
4.0RC2？はまだ早いんでね？

そうですね。自分の持っているやつがGT430なのでバージョン1.0
（最新は1.3）に対応しているみたいなので古いほうをインストール
してみます。

出来ました！
ありがとうございました。

>>596
何をどうしたら解決したか書いてほしいな

流れで分かりそうなもんだけど
4.0RC→3.2で解決したようだ

4.0は大きく拡張されてるみたいだし、RCがとれても移行は様子見したほうがよさそうだね
2台とかあって3.2と4.0の環境があったりするなら切り分けしやすいから良いかもしれないけど。

こんにちは。
16*32の行列に長さ16の配列をXORの積をしたいのですが可能でしょうか？
イメージ的には32個のブロックに16個のスレッドという感じで考えています。
行列計算なので実装はそんなに難しくないと思うのですが、具体例があまり
見当たらず苦心しています。はじめてからまだ三日なのですが、解る方が
いらしたらソースを見せてほしいです。ヒントだけでも良いです。
よろしくお願いします。

>>16*32の行列に長さ16の配列をXORの積をしたいのですが
これどういう意味？もう少し計算を具体的に書いてくれ

4*8の例で考えます。
[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サイクルで実行できると思うのですが何か間違えていますでしょうか？
よろしくお願いします。

for(j=0;j<16;j++){
o=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];
}
この処理を並列化したいです。

>>603
ideoneかcodepadにC版の関数をアップロードしてくれ。
そうしてくれたらcudaでどこをどう並列化するか考えてみるよ。

並列化したいのはループ使っている部分だけなのでそこだけです。
http://codepad.org/15hfT9HS

公開したのに反応が無い。
見ても解らないだとか汚いと罵るだけ。
どうせやる気無いんだろ？

>>606
うん。
忙しいからね(´・∀・｀ )

>>603
またお前か。
いい加減自分の手を動かしたらどうだ。
時間がないなら対価を払って時間を買え。

だったらソース見せろとかいうな馬鹿

>>606
何怒ってんだ？
その程度のコードなら普通にできるだろ
バカなの？

見る気も無いくせに見せろと言うほうが馬鹿

見せろって言ったのは俺じゃないし、ソース見る前からお前だって分かった。
自分でやる気が無いっていう態度が滲み出てるんだよ。

やる気があって調べてるんだよ。
ここならもっと詳しい人が居ると思うから聞いてみただけ。
何で最初からやる気が無いなんて決め付けるの？

まぁ質問に具体性がなく丸投げだからな
もう少し考えた過程が見えないと回答もしづらい罠

最初からじゃなくて俺が知っているのはここ8ヶ月くらいだが
分からない、分からない、一個教えてもらうとすぐ次が分からないで
自分で調べる気なんてなさそうに見えるんだが。

>>613
Intel Parallel Studio 使って並列化。
↓
逆アセンブルして、コード解析してCUDA化。
↓
(゜д゜）ｳﾏｧ

VC++ 2010用のtemplateプロジェクト、どこかにないでしょうか？
CUDA 3.2 SDKにあるVC++ 2008用のtemplateプロジェクトを、VC++ 2010用に変換しようとしましたがうまくいきません。。。

>>605
＞並列化したいのはループ使っている部分だけなのでそこだけです。
＞http://codepad.org/15hfT9HS
悪い悪い、デートしてて遅くなった。
で、このコードのどこをcuda化したいんだ?
未だ帰り道だから、んな長いの見てられないんだ。

16人17脚はデートって言うの？

>>617
Windows版のnvccはVisual Studio 2005か2008のコンパイラが前提になっているらしいから
VC++ 2010用のテンプレートだけではうまくいかないと思う。

それらのコンパイラを用意するとなったら、VS2008で作ったほうが楽な気もする・・・

>603
きたないコードだなあ・・・。

>>603
俺ならこうだな。

for (i = 0 ; i < 16; i++) {
　　for (j = 0; j < 16; j++) {
　　}
}


主なコーディング作法が3つくらいあるから、そのどれかにしろ。

>>622
スペースきれい！
ふしぎ！

>603
あまりにもイミフなんだけど、とりあえず並列化したいなら

関数呼び出し側
　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の根本的な文法も理解しようとせず使おうとしてるから、バレないし問題ないよね。

>>617
VC2010はCUDA 4.0 RC2から対応で、SDKにテンプレートが付属してる。
うちの環境だとMSBuild関連が正しくインストールされなかったのでひと手間必要だったけど。

CUDAのコンパイラってレジスタ不足でもコンパイル通っちゃうのか？
VSで開発してるんだが、どうもレジスタが足りないとFirst chance exceptionエラーが出る気がしてきた。

>48
祈るってアプローチが間違いとは言わないけど、
同じ計算を2並列で行うとか、間違いなく構成を変えられる方法があれば
演算器を入れ替えて2回行うとかって手もあるよ。
2倍時間がかかってもまだCPUには勝てるでしょ。

PowerDirectorなどCUDAを使ってエンコードを高速化出来るアプリがありますは、SLI環境ではさらに高速化出来るのでしょうか？
ググると出来ないというソースが目立ちますが...
X58マザー使っています。

レジスタ不足ってローカルメモリ使われるんじゃないの?

>>628
cuda4はどうか知らないが、それ以前のcudaではSLIだからより速くできると言うことはない。
勿論GPU2枚挿しに対応したアプリなら、2枚刺した方が速くはなるがSLIには関係ないようだ。

>629

引数はsharedメモリに確保されるってどっかで見たことあるけど…
sharedメモリつかってないプログラムだしなあ。
使ってるのはGeForce9800 - G92コアなんだけど、
コアに上限の24ワープ割り当てるVSがエラー吐いて動かなくなるんだけど、
控え見えに16ワープ放りこむと問題なく動く。

cuFFTとかの標準ライブラリでも動作が怪しかったりするのは
レジスタ周りの取り扱いのせいなんじゃって気がしてきてるのですよ。

引数はCC2.0からはコンスタントメモリで
それ以前はシェアードだったかな

関数内変数はレジスタで不足したら勝手にグローバルにいくんでしょ
ワープ数が上限ぴったりならエラーは吐かないと思うんだけど

オーバーしても、普通は遅くなるだけだと思うんだけどなぁ。

>>631
そんな曖昧な知識で推測する前にPDFちゃんと読むんだ

>634
了解したぜ

>>630
ありがとう。
SLI環境はもう少し様子見します。。

最近のグラボはATX電源から電源引っ張ってくんだな。
Fermiに変えたが早すぎワロタ。

それ以前からfermiって劇的な変化じゃね
キャッシュはでかい

FlashPlayerもCUDAつかえるの？

Tesla S2070ってなんで抹消されてしまったのでしょうか。
nVIDIAのページにもELSAのページにもなくなってしまったのですが。

XPでFFTをかけるソースがあったのでvisualstudio2008で実行してみたところ
FFT.exe の 0x7c812afb で初回の例外が発生しました: Microsoft C++ の例外: cudaError_enum (メモリの場所 0x0012ae3c)。
というエラーが出てしまい困っています。ホストからデバイスへのメモリのコピー等
の簡単なプログラムは動きます。ちなみにGeforce210です。

debugを人に頼むような人間は一つ解決しても次に同じような所で引っかかるからきりがないわな。

>>641は何も頼んでない。日記を書いてるだけじゃないか？

>641
多分それ、cuFFTのバグ。そのエラーコードはメモリアクセス例外のときに出る。
俺の場合9800GTから550Tiに買い換えたらエラーが消えた。

cuFFTがデバイスメモリ使い過ぎて、デバイスが積んでるメモリ量によっては
エラーを吐くとかが原因の気がしてるんだが、本当の理由はわからん。
NVidia側のバグフィックスを待つぐらいしか現状で対策はないと思う。

cuFFTのバグなら、サイズが2の冪かどうかでも変わるね。
cuFFTの内部で結構デバイスメモリを使うようだから、ボードの半分以下のFFTしかかけられない。

>>642
FFTなんて本来ライブラリの側で対処すべき次元のもので、
ユーザー側がデバッグに煩わされるべきものではないがな。

>>646
もしバグだとおもうならフォーラムで問い合わせたり報告したりしたほうが良いんじゃないか？
バグなら修正してもらえる可能性があるし、そうじゃないとしたら自分の間違いだとわかる

ubuntu 10.04、CUDA4.0でSDKをコンパイルした後に
deviceQueryを実行しようとすると
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.
と出てしまい実行出来ません。どうしたら良いでしょうか？

>>648
http://www.alc.co.jp/

nvidia-current カーネルを削除してSDKを再コンパイルで行けました
ありがとうございます

nvidia-current カーネルを削除してSDKを再コンパイルで行けました
ありがとうございます

>>647
フォーラムには既にあったと思う。

>>644
ありがとうございます。
550Ti持ってるのでそっちでやってみようと思います。

Developer Drivers for Linuxをインストールすると、
Xwindowの解像度が640*480固定になるのはどうしてなのかしら？

おまえがばかだからじゃないか？

xの設定書き換えちゃってるんじゃないの?

Ubuntu 11.04 CUDA 4.0でドライバをインストールしようとすると、
The Nouveau kernel driver is currently in use by your system.
というエラーが出て、先に進めません。
どうしたらNouveau kernel driverをとりのぞけますか?

これ読んでみ
https://help.ubuntu.com/community/BinaryDriverHowto/Nvidia

>>658
早速の返信ありがとうございます。
sudo apt-get --purge remove xserver-xorg-video-nouveau
と打って、再起動もしてみましたが、やはり同じメッセージが帰ってきます。
ほかに、何か手立てはないものでしょうか?

>>657
俺がLive CDで試したときは
kernelの引数にnouveau.modeset=0として起動後に
modprobe -r nouveauした。

そもそもnouveauがモジュールではなくて組み込みになってるなんてことはないよね？

インテルコンパイラには対応していないの?

Coalesced Accessって簡単にいうとどゆうこと

>>662
１６人１７脚で走ること
ただし、一番端っこの人の番号に制限あり

>>663
なんだかすごく遅くなりそうな例えだなｗ

論文書くときにさ、遠くの図書館に借りに行くわけよ
16冊まで借りられるから自分の欲しいタイトル1冊と隣に並んでるの15冊借りて帰って来るの
そん中にフラットメイト15人が欲しい本が入ってたらラッキー
何人か足りなかったらまたまた誰か派遣して16冊借りてこなくちゃなんない
一回で16人全員満足したり16冊全部が有効に使われたらうれしいな

俺ら棚のブックエンドの右からまとめてとってく癖があるから
司書さんはいいかんじに本並べといてね

家にあればこっちのもんだから欲しい本被っててもおkおk


てのはどうだろう

>>663はCoalesingの説明になってないぞ。

食堂のカウンターに並んでいる先頭の１６人が皆「カレーライスください」という。
食堂のおばちゃんが「皆カレーかい？鍋ごとあげるから皆で分けてね。」

もちろん正確ではないが速くなりそうな例ではある。

初歩的な質問で恐縮なのですが、
皆さんCUDA版プログラムとCPU版プログラムの速度比較は
どのようにされてますでしょうか？

CUDA3.0以降はエミュレーションモードがなくなったようなので、
別にCPU版プログラムも作って比較するのがよいでしょうか？

エミュレーションと比較して何の意味があるの？
NVIDIAの回し者でない限り、iccでカリカリにチューンしたプログラムと比較するべし。

CPU版と比較しても仕方が無い。
iccでチューニングした上で、更にその先の為にcudaを使うのだから。
まして、エミュレーションだなんてなんの悪夢だか。

>662
メモリ読み書きするなら32byte, 64byte, 128byte単位にしてください。
これだけ。

SDK,toolkitを再インストールしたところ
サンプルを実行すると
cudart32_31_9.dllが見つからなかったため、このアプリケーションを開始できませんでした
。cudartを間違えて消してしまったと思われるのですが、どうやったら解消されるのでしょうか？
ぐぐってみたのですが、わかりませんでした。

PCを再起動させてみたらどう

最悪でも再インストールで済むだろ。

再インストールしてみたのですが変わりませんでした泣

>>674
よかったな、cudartを間違って消してしまったわけじゃないことが判ったじゃないか。

CUDA 4.0

__CUDACC__ってもう定義されていないの?

みんな業務でCUDAつかってんの？

一部使ってますが。

CUDA4.0、GTX480M + Notebook Developer Drivers for WinVista and Win7 (270.61)では
device query含め、CUDA Cのサンプルが全部動かない。
OpenCLとDirectComputeは動く。
なんでだろう。

坊やだからさ

VS2010とCUDA4.0の環境でrulesファイルを使えた方おりませんか？
$(CUDA_PATH)\extras\visual_studio_integration\rules
にあるファイルをコピーしたのですが、これを設定してプロジェクトを読み込みなおすと
「要素 <UsingTask> 内の属性 "AssemblyFile" の値 "$(CudaBuildTasksPath)" を評価した結果 "" は無効です」
といったエラーが出てプロジェクトが読み込めないようです…。

完全な逐次処理（どう頑張っても並列化できない）のを、
CUDAで高速化してちょ、といわれたんだけど殴っていい？

>>683
いや無理にでもやれよ
それが仕事ってもんだ

ゼロコピーってキャッシュされないでしょ。

んで、ゼロコピーを使うと遅くなるケースが出てきたんで、
仕方なく使わないことにしたんだけど、
やっぱりcudaMallocやcudaMemcpyの時間がもったいないんで
どうにかしたいんだけど、なんか知恵ある？

>>683
どう頑張っても並列化できないを証明すれば良いじゃん．
なかなか難しいよ．

>>685
ゼロコピーって何よ。
何をしたいのか判らんけど、cudaMemset()も使えない内容なの?

解決しました。ちょうど公式フォーラムで同じ問題に遭遇してた人が居たようです。
先ほどあった投稿で解決することができました
http://forums.nvidia.com/index.php?showtopic=201433

もふ。
マジレスするとゼロコピーとは、cudaHostGetDevicePointer()を使うやつのことだす。
cudaMallocHost()でホスト側のメモリを確保しとかないとダメだけど。

ゼロコピーと言うのは知らんかった。
処で、勿体無いと言うけどcudaMalloc()って時間掛かる?

演算速度について質問です。
プログラミングガイドによるとfloatにおけるmulとmadの速度は同一だとあるのですが、
for(i=0;i<N;i++)
a = b*c;

と

for(i=0;i<N;i++)
a = d*(b+c);

では明らかに前者のほうが短い時間で処理を行っています。
dをコンスタンとメモリや定数にしても同じでした。
いったいどういうことなんでしょうか？
助けてください。

>>691
どこのプログラミングガイドだよ教えろよ教えてくださいお願いします

madって(A*B)+Cじゃないの？

すみませんまったく環境とか書いてませんでした。

GPU : 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.

だそうです。

>>693
ごめんなさい書き間違えました。

for(i=0;i<N;i++)
a = d*b+c;

です。

まぁ、ptxを貼ってみろ。

積和演算なら
a = b*c+d*e
とかじゃないのん？畳込みとか行列積とかで頻出する計算です。

DSPの命令見てみたら、積和演算は
a = a + (b*c)

だったわ。
a+= (b*c)

でもええかもしれんが。

>>680
270.81でないとCUDAランタイムAPIが変な動作をするみたい
つまり現時点ではPC用のドライバしかない

270.61に
「Unified Virtual Addressing (UVA) やGPUDirect? v2.0を特徴とするCUDA 4.0を使ったアプリケーション用サポートの追加」
と書いておきながら動かないなんてふざけた話だわ

CUDA勉強はじめました。
でも世の中に出てる日本語のCUDAの本って4冊ともFermi以前のアーキテクチャでの説明だからsharedメモリが16kbyteだったり、1ブロックあたりのスレッド数が512だったりするんで困ります。
なんとかして。

コンスタントメモリや定数の値をレジスタに読み込む時間をまったく無視していました。
たぶんそのせいで遅くなっているんだと思います。
ご迷惑かけてすみません。

>>700
ちょうど新しいドライバが出たみたいです

>>701
英語も勉強汁

oh! i can't read english. :-P

>>701
Fermi用の正しい値が分かる人なら何も困らないような気が・・・・。

Compute Visual Profilerを使いこなしてる人っている？
あれ、結局、どう活用していかわからん。

>>701
Fermiの方が簡単になってるからいんじゃね

>>707
ローカルメモリ使われてたりバンクコンフリトあるかをチェックすんじゃね

北川景子かわゆす

そうか？

あれ？4.0ってcutil64D.libなくなっちゃった？

本当だなくなってる

　

自前でビルドしなさい

今Cを勉強中ですが、CUDAと並行して勉強できますか？

>>715
少なくとも小中で算数（数学）や社会、英語などを並行して勉強してきたはず

>>715
CUDAのプログラム記法以外に知るべき事が多い。と言うかむしろそちらが大半じゃないか。並行してできる、がんばれ

どっちもものにならないと思う。どっちかだけやればものになるものでもないけど。

4.0でコンパイルが通らないの結構多いね
CUDA BY EXAMPLEのサンプルをいくつか試したがエラーをはく。
まだ3.2使うわ

4.0 から cutil32.lib とかって自分でビルドするようになったのね。
インストール後に検索しても見つからないから、変だな〜と思った。
Direct Compute の Example が動かないのもよくわからん。

ある程度C(というかプログラミング)しってなきゃ百パー無理だろ

%を略すな、池沼

。

>>721
それを知らなかったので、PATHの設定やmake実行さえ一苦労でした。
人によっては、ディレクトリの概念が難しい場合もあるかもしれません。

cudaはただでさえ難しいぞ
初心者に手にを得る代物じゃね〜べよ

資料も英語しかないしかなりきつかったわ

OpenCLとどっちが楽ですか？

Cudaの方が楽
OpenCLはOpenGLと同様、自助努力が基本になる

OpenCLとどっちが速いですか？

自助努力って具体的に何？

>>423
亀だけど、cufftPlanは対象とするデータのサイズと同じサイズをバッファとして確保するみたい。だから、二倍以上メモリが空いてないとこける。

>>730
そういうのを自分で考えたり調べたりする姿勢や行動

OpenGL de プログラミングというサイトや
CUDAプログラミング入門とかフリーであるもの
見たほうがそこらの書籍よりかなり役立つよ

OpenCLは少し大変だけどCPUならdouble型
オプションあるからいいんだよね
ＣＵＤＡにもあるのかな?

あった、ごめん

visual studio pro, radeon 6000 台　で ati stream ないし　open cl
使っての　並列FPU高速化ってどのくらい大変ですか？
一般のプログラムと英語ドキュメント読解に問題はないです

スレタイ読む能力はなかったようだな

間違えました。すいません。

漫才かw

>>732
CUDAを勉強してみようかと少し取り組んでみましたが、
C言語の基本も知らぬプログラミング初心者なので
C言語から勉強しようと思ってます。

Rやoctaveに触れると、Cの世界に戻れない。
故にCUDAに触れることも出来なくなる。

Rでモンテカルロ出来るんだっけ？

>>739
C言語の文法を知ることは必須としても、まずプログラムがどういうもので何が出来るものなのか、って視点からも眺めた方がいいよ。
文法を覚えることに溺れないように

>>741
統計解析ソフトだから乱数の生成は得意だと思うけど、
ループが苦手だから本当にモンテカルロが必要な場面でが役立たずかも。

Rユーザから見ると、curandライブラリがRから呼び出せたら嬉しい。

カーネルを起動するのにかかるコストってNVIDIAの資料等に載ってるんでしょうか？

資　料　読　め　よ

　　　　　　　　　　　YES　→　【見つかった？】　─　YES → じゃあ聞くな死ね
　　　　　　　　　／　　　　　　　　　　　　　　　　　 ＼
【資料見た？】　　　　　　　　　　　　　　　　　　　　　　　 NO → なら、ねぇよ
　　　　　　　　　＼
　　　　　　　　　　　 NO → 死ね

日本語の試料ないですか？
頭痛いです。

　　　　　　　　　　　YES　→　【見つかった？】　─　YES → じゃあ聞くな死ね
　　　　　　　　　／　　　　　　　　　　　　　　　　　 ＼
【探した？】　　　　　　　　　　　　　　　　　　　　　　　 NO → それより僕と踊りませんか？
　　　　　　　　　＼
　　　　　　　　　　　 NO → 死ね

>>746
くだすれなんだし並列に書いてみてはどうか

>>749
ifが多すぎてパフォーマンスが落ちると思う

>>750
嫁>>1

>>1の序文が全く無視されているな

>>747
日本語の書籍購入するのが早いとおもうよ。
・ASCII.technologies (アスキードットテクノロジーズ) 2010年 08月号
・はじめてのCUDAプログラミング
・CUDA by Example 汎用GPUプログラミング入門

この辺がお勧め。自分はこの順で読んだ。
でもプログラミングガイドとか読むのは避けられないはず。


・CUDAプログラミング実践講座 - 超並列プロセッサにおけるプログラミング手法
これは買ってないけど良さそう。

>>753
上から三つ目は一つ目と二つ目が参考資料だっからな
アスキーは触りとしてはわかりやすかった
どれにしろFermiについて書いてないけどね

4.0になると、もうどうしましょ

■後藤弘茂のWeekly海外ニュース■
AMDが発表したメインストリーム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の間での無駄なメモリコピーを排除するゼロコピーが実現できる。

あとはメインメモリがGDDR5になるのを待つだけだな．

また妙なことを始めたなAMDw

メモリコピーがネックにならない計算ってむしろおかしい？
計算の方が時間かかってるんだけど
もっと最適化できそうとか

利用しているデータ量が少ないなら別におかしくはないと思うけど。

SDK4.0(x32)のtemplateを別フォルダにコピーしてVS2008でビルドしようとしたらcutil32D.libが開けませんって出ます。
助けてください。

>>761
720

>>762
レスありがとうございます。
cutilとshrUtilsは既にビルドはしているんですが、開けないとでるんですよ。。

>759
メモリコピーの時間がかかってでも並列演算したいって問題がほとんどだから
別におかしくないんじゃない。
並列計算が早すぎてメモリコピーがネックのように見えるってことはあるけど。

メモリのコピー時間は0でも
同時にはアクセスできないよね、きっと。
でも大した問題でもないか。

Microsoft Going All-in on GPU Computing ≪ NVIDIA
http://blogs.nvidia.com/2011/06/microsoft-going-all-in-on-gpu-computing/?sf1642229=1

>>754
アスキーは2冊あって2010/08はFermiについても掲載されてるよ
そのほかにOpenCL、Direct compute, ATI streamについても載ってる

物凄く基礎的な質問ですいません。
今大学でCUDAの勉強を始めた者です。
現在貸し出されたPCに入っていたのがGeForce7600GSで、
CUDAの公式を見る限りだとCUDAに対応していないようですが
どうにかして動かすことは出来ないでしょうか？

一応一番初期のVer1.0〜最新の4.0までひと通り試してみたのですが
サンプルのdeviceQueryを動かしてもCUDA Driver Versionは0.0のままでした。

もし可能でしたら少しでも助言を・・・
無理なら無理と言ってくだされば大学側に申請して
可能なGPUの購入を考えてます。

CUDAの勉強するならCUDAに対応したグラボを用意して貰え
対応してない物が仮に動いたとしても、それが正確な動作かわからんから勉強にならんぞ

cuda x86
Amazon EC2

GPU買わなくてもおｋ

>>768
安いのなら1万切るから自分でグラボ買ってもいいんじゃない？

>>771
大学研究室ならキッチリ金出してくれて580くらいはポンと買ってくれるだろうから自分で買う意味ないだろう

ゲーム用と思われるのもあれだからTesla 20 Cシリーズ買ってもらいなよ

>>768
授業？独学？
研究のわけないよね
GPGPUの研究でDX9のGPUとか遠回しなアカハラだｗ

はじめてのCUDA読んでソースみたほうが早いかもよ
それかここの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

>>775
横からだが、これは面白そう 読んでみるわ

同じく横からだが、ブクマした
全部印刷しておいても良いかもしれん

>>775
参考になりそうだ。DLできないの？

できますん

二つ質問があります
�@felmiにはキャッシュが搭載されているみたいですがプログラム中で明示しなくてもキャッシュにより高速化するのですか？

�AGPUでcuda化して高速化が見込めるアプリケーションが載ってるサイトとか知りませんか？ソースコードつきで。できればC言語だとうれしいです

めんどく
だれかこたえてあげれや

この程度の事を自分で調べられず、聞いてしまう人間には
使いこなすのは無理だよ。

>>782
一応は調べました
せめて�@だけでも答えてもらえると助かります
�Aに関しては自分もいくつか見つけましたが他の方がどういったサイトを使ってるか気になったので聞いてみました

機種依存文字-10
綴り間違い-10
sage+3
投稿時間+10

内容に言及しなくても-7点だ
もうちょいがんばれ

教えるきないなら書き込むなよ、暇人

>>785
ﾌﾟｯ

キャッシュは勝手に使われるが、
キャッシュによって高速化(何と比べて？)するかはアプリによる。

ttp://journal.mycom.co.jp/articles/2010/07/21/fermi_cache/index.html
こんな話もあるから、キャッシュが如何使われるか位は
意識しないと駄目だし、究極を目指せば結局コアレスアクセスに
なってキャッシュの意味は無くなるかも知れない。
で、そこまで頑張れば微妙にteslaの方が速くなる。

ランダムアクセスでテストしてみたが（sm1.1でコンパイル）、コアレスと比較して、260では1/5まで速度が落ちたが、470は2/5だった
びっくらこいた
高速化できる例はCUDA SDKのサンプルみればよろし

補足
G80/G92のように総レジスタ数の少ないGPUはもっと遅くなる。大体1/10

キャッシャが要らなきゃ共有メモリを増やせるし、Fermiは便利だね。

fermiで２ワープを２サイクルで処理していると記載を見かけるのですが、なぜ１サイクル１ワープでなく2サイクル２ワープなのでしょうか？
どういう原理なのかご存知の方がいらっしゃいましたら教えていただけませんか。
また、ご存知でしたらその記載がどこにあるか教えていただけませんか。
お願いします。

「fermi half warp」でググればいくらでも出てくる．

>>792
ありがとうございます。

僕もそれ欲しいです。

ゲフォ最上位とラデ最上位さらにSLIとかCF含めるとトリップ検索じゃどっちが上なんだろう。

HD6990単体で362Mtrip/sくらいだよな。
単純に倍なら722Mtrip/s？

アホな質問で、NvidiaのCUDAには関係ないけど
GPUを使う場合ってターゲットのｎVidia、AMD（ATI)GPU用専用バイナリコードを作成するの？
(両方で動くようにするには両方用のバイナリを用意しないと駄目?)
それともJAVAのbytecodeみたいなのを作成して、それで両方のｎVidia、AMD（ATI)GPUに対応
出来るようになっているの?

(テキストコード) → [DirectX,OpenGL] → [nVidiaドライバー] → （ネイティブコードA） → [チップ]GTX580
　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　→ （ネイティブコードB） → [チップ]GTX560
　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　→ （ネイティブコードC） → [チップ]……
　　　　　　　　　　　　　　　　　　　　　　　→ [ATIドライバー]　　→ （ネイティブコードD） → [チップ]知らね
　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　→ （ネイティブコードE） → [チップ]
　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　→ （ネイティブコードF） → [チップ]

機械語をデコードして制御コードに従ってどうたらこうたら

OpenCL使っとけ

OpenCLなんざ複数の環境で動かそうと思ったらテキストそのものだし、
Javaみたいな抽象度の高いバイトコードも解析が容易だしで、
小手先のノウハウで大幅に速度や精度が変わっちゃう商売には使いづらい世の中になってきた。

ちょっと質問。
CUDAプログラミングでVRAMの大きさって気にする必要がある？
まぁ、作るプログラムによるんだろうけど、1GB以上食うのはたぶん確実。

GTX 560とかで2GBのカードもあるので、差が出るならそっちを選んでおきたいかな…と。

計算に必要なデータがVRAM上に全部乗るならそれ以上はいらんだろ

ありがとう。
やっぱし計算に必要な分はメモリがいるのね。

ちと調べたら
CUDA4.0からUnified Virtual Addressing(UVA)なるものが登場していて、
コード的にはVRAM容量は気にしなくていいみたいだけど、
性能的には差が出る、と理解しておきます。

CUDA初心者なんですが、質問させてください。

研究で有限要素法を使っていますが、あまりにも計算に時間がかかるので、GPUを使えればと思っています。
とりあえず、使っているプログラム(fortran)がGPUで処理できるような内容か否かを知りたいのですが、
CUDA入門書などを読んでも、プログラミングについて書いてあるだけで、
具体的にどのような処理が不可能かは見つけられませんでした。
プログラム自体は、ソースを書き換えれば並列化が可能です。
cpuに出来てGPUに出来ない処理とは具体的にどのような処理なのでしょうか？

gpuは単精度しか扱えないとの記述を見ましたが、単純にreal*8で定義した変数をreal*4に書き換えることで
gpuでも可能になるということでしょうか？

長文失礼しました、ご回答よろしくお願いします。

http://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

fortranからcudaってなんか必要じゃなかったっけ

>>804
ど素人だけど
有限要素法がCUDAに向くかグルグルしてみる。　
Fortran用いてCUDAで有限要素法している大学の研究室あるカモ

有限要素法ってマルチでスレッド走らせても効果あがらんべよ
逐次処理なんだからさ

連立一次方程式を解く部分が反復法で
SOR法のような逐次処理であっても
並列化することは可能。
共役勾配法ならふつうに並列化できそう。

すでにopenMPIかなんかで並列化されてるやつならCUDAでもいける

たくさんのレスありがとうございます。
>>805
ちょっとまだ読めていませんが、頑張って読ませていただきます。
>>807-810
なるほど、ググってみます。
ただ、私が使っているのは、構成式の計算が非常に複雑なのと、
多相混合体を計算しているので、その部分が逐次処理でも早くなるのでは？？と考えています。
あと、12スレッドでCPUを100%使用してくれているので、
単純にマルチスレッドで行けるのかなと甘い考えを持っていました。

reduction計算で任意の配列数ってみなさんどうしてますか？
はじめてのcudaプログラミングに書いてある方法では奇数はおろか2の乗数の配列しか扱えないんで

要素数 2^nで確保して，必要ないところには0をセットして使ってる
もっと良い方法ありそうだけど，ホスト側に転送して計算させるよりは圧倒的に高速になるから，
深く考えずに実装してる…

>>812
GPUComputingSDKのreductionのソースを参考に改良。

同じく0をセットしてる

バイトニックソートの時は邪魔にならんようにめちゃくちゃでかい数セットしてるｗｗｗｗｗｗ

そもそも「はじめてのcudaプログラミング」は
Optimizing Parallel Reduction in CUDAの#6
(Kernel 6)までしか訳してない。
#7(Kernel 7)でさらに速くなり、データ数の上限
もアップする。

配列の要素数が二進数で11100101とかなら先頭から
10000000
1000000
100000
100
1
づつreductionさせればどんな要素数でもできそうだけどめんどくさそう

めんどくさいどころのさわぎじゃないな

モンテカルロなんだけどスレッド数とかブロック数のいい決め方ってある？
例えば約百万回の実行を行うのに1024＊1024と512＊512を4回では環境によって
前者が速かったり後者が速かったりするんだけど

>>819
実測。

>>820
返答サンクス
でも意味不明w

>>819
自分で答え書いてると思うけど，環境によって最適値が変わるので自分の環境で試行錯誤して最適値を決定するのが一番いいと思います

環境ごとに適正なパラメータが違うなら、パラメータ決定の為の評価プログラムを作ればいいじゃない。

Zipのパスワードを1秒間に6億回クラックってどう？
まぁまぁ速い？

http://www.youtube.com/watch?v=3dEN9JQ3R0U

そもそものGPGPUってどんな方法なの？
各ピクセルにデータいれたものをテクスチャとして渡して、ピクセルシェーダになんか演算させて、画像として帰ってきたものを1ピクセルずつ読み取るってかんじ？

それは2005年頃の話

>>826
早いのか遅いのかわからんｗｗｗ
そもそもは何年ごろに確立されたの？

APUが普及すれば，これからもっとGPGPUが盛んになるかな？
(主流はOpenCLとかになりそうだけど)

>>827
提唱時期は知らんが、>>825 はGPGPUの「古典」だ

>>829
thx
その古典についてもうちょい詳しく知りたい

http://journal.mycom.co.jp/articles/2005/09/06/siggraph2/004.html
http://www.4gamer.net/games/076/G007660/20090206031/

大まかな話ならこのあたりとかじゃないかな。

>>831
ありがとう！
スレ違いすまん


しかしCUDAはそういうの意識せずにできるんだな
プログラマブルシェーダってなに？ってレベルだったけどちゃんと高速化できたわ
GPUの構造は多少知らなきゃいけないけど

OpenCLが.NET Framework並みに使い易くなってくれればなぁ・・・

ACML-GPUでもCUBLASでもいいからCygwinかMingwのgfortranで動けばなあ

MINPACKがそれらのBLASで動けば、非線型最適化が瞬時に出来るかも
しれないのに。

質問失礼します。

atomic関数を使った.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 ~~~
"というエラーメッセージが出てしまいます。

解決方法ご存知の方いませんでしょうか。

質問失礼します。

atomic関数を使った.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 ~~~
"というエラーメッセージが出てしまいます。

解決方法ご存知の方いませんでしょうか。

どんなコマンドラインを入力したらどんなメッセージが出たのか詳しく。

instration？

僕は
atomicAddとか使うとき-arch sm_11じゃなく、-arch=sm_11　としているけど

良く分からんが、
sm_11_atomic_functions.h　をインクルードしてみるとか。

コピペもできないのにcudaとか。

何でこのスレでやたら「はじめてのCUDAプログラミング」勧められてるのかと思ったら
著者はこのスレの住人なんですね。32人33脚の例の話とか過去ログ読んで経緯理解したよ…

しかしこの本、意味不明な挿絵でいちいち吹くんだが…暗闇からのニョッキリ手ってなんだよｗ

>>842
日本語で勧められるのはcuda by Exampleが出版される前は
「はじめてのCUDAプログラミング」以外なかった。

今でも内容がfermiにあっていないところは時代遅れだが
初学者にとってリダクションのチューニングの演習過程は参考になる。

住人かどうか知らんが
別に本人が勧めてるわけじゃないだろ

アーキテクチャとかはPCゲームの記事が詳しかったな確か

>>574
私もCUDA by Example読み始めて同じところで引っかかったのですが、
今検索するとforumでいくつかヒットしますね。
http://forums.nvidia.com/index.php?showtopic=202480

ただ私の環境では、コンパイル通っても実行すると
> unknown error in CudaByExample/chapter04/julia_gpu.cu at line 89
なんてエラーが出てさっぱりですが…

このスレッドは、他のスレッドでは書き込めない超低レベル、
もしくは質問者自身何が何だが分からない質問を勇気を持って書き込むスレッドです。
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/

関連サイト
CUDAテクニカルトレーニング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/

他にも参考になるリンクあったら教えろください。

>>2 にもあるけど
CUDAフォーラム（日本語）： 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

こんなマニアックなことする人間はあまりいないんだろ
基本、学術計算的なことじゃない限り利用価値があまりないし

つーか、NVIDIA日本に回答できる奴もまともに運営できる奴も居ないんだから意味無いよ。
英語できる奴は本家のForumで事が足りるし、そうでない奴はここに聞きに来ている悪寒。

産業用途の画像処理分野にGPGPUは劇的な効果だよ

OpenCLのほうがよくね？

同じプログラムを書くとき，OpenCLよりCUDAのほうが早く書き終わる．

CUDAで作ってからOpenCLに移植したほうが楽なのは俺だけ？

>>853
どういった理由ででしょうか？

>>852
DSP使ったり、FPGAで並列回路組んでたのをこっちに乗り換えてﾔｯﾎｰｲな所ってあるのかな

>>857
胡散臭い演算専用ボードを使わずに済んでYahoo!

PhysX を飼ってしまった‥‥‥。

今更だが日経ソフトウェア8月号で
VS使わないフリー環境でのコンパイル方法を紹介してる

GPGPUするぐらいだから、ここにいる人は速度が欲しいだと思いますが
早過ぎる最適化は怖くないんですか？
カリカリの最適化した後に、やっぱり違うやり方の方がいいって思っても
大きな変更は億劫になってしまいそうです。

仕事でやるなら辛いでしょうね
納期に追われると特に

でも、遊びでやるなら楽しいんじゃないでしょうか
パズルみたいで

>>860
VC Express 使えば良いじゃない。

C++/CLIを使えと・・・

Concurrent Kernel Executionについて質問させてください。
compute capability 2.xを満たし、念のためにconcurrentKernelsも確認してConcurrent Kernel Executionが使えることは確認しました。
しかし、いまいち使い方がわかりません。
カーネルを2つ連続で書けば自動的に行われるのか、それとも何かフラグが必要なのか。
また、何か特殊な関数が必要なのか。
どなたか、ご存知の方がいらっしゃいましたらご指導をお願いします。

ヒント
Stream

>>866
ありがとうございます。
プログラミングガイドのStreamとその周辺を読んでみたいと思います。

>>864
CLIは必須じゃない

はじめてのCUDAプログラミング読み始めましたが、ちょっと疑問が。

ブロック、スレッドというのは、カーネル関数を実行するときの<<<M,N>>>で指定すると思うのですが、
グリッドというのは何なんでしょうか？
他の入門用説明を読んでも、1つしかないことを前提に書かれているようですが…

いい質問だ
とにかく用語がおぼえにくいからね

スレッドの集まりがブロック。
ブロックの集まりがグリッド。
どちらも論理概念でハードウェア資源と完全に同一ではないので注意。

みんなCUDAで何やってんの？オナニーか？

今日こんなのあったのね。事前に知ってたら参加してたのに…
http://www.nv-jp-event.jp/gtc-workshop2011/
参加された方、感想をぜひ。

ストリーミングで見れるじゃない。

>>874
基調講演だけじゃ…

分岐さけるために↓みたいなコード書こうかと思ったんだけど、
boolが1にとは限らなかったりする？

int selector = value1 < value2
int out = value1 * (1 - selector) + value2 * selector

>>876
これって、if-else文より速くなるの？

>>876
GPUでやった事はないけど、
x86,x64じゃ比較演算子は分岐命令が交じる。

cudaなんだからptxファイルを出力して見比べればいいじゃん。

>>876
勉強がてら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のそれ以外は全部無駄だね。

ごめん、ちょっと間違えてた。
%r2がvalue1, %r1がvalue2だった。

意味なしなのね。調べてくれてありがとう。

意味無しじゃなくて害悪だろう…

CUDA SHA-1 TripperをCUDA Toolkit 4.0でリコンパイルしただけで
100MTrips/s以上速くなったｗ GTX 580で640MTrips/s出てる。
こりゃすごいわ。

というか、普通に書けばMAXを取る命令が使われるのでそもそも分岐命令なぞ入らないが。

確実に重くなるだろ。

>>884
他のアプリで速度向上したりしてます？

>>887
他のアプリは試していないのでちょっとわかりません。
あまりにも嬉しかったのでスレを立ててしまったｗ

【トリップ検索】CUDA SHA-1 Tripper【GeForce】
http://hibari.2ch.net/test/read.cgi/software/1311428038/

>>885-886
そんな後出しでドヤ顔されても…

http://thinkit.co.jp/story/2010/07/09/1654?page=0,1
> ブロックの集合体であるグリッド（Grid）は、デバイス（Device）に対応しています。
とあるので、簡単に言うとGPUカード2枚挿すとグリッド数は2、ということでしょうか。

グリッド数という概念はない。
グリッドとは一回のカーネル呼び出しで実行される単位みたいなもの。

VC2010Expressで4.0環境やっとできたわ

>>892
注意するところとかはまるポイントある？

ググれば出てくる

>>893
894の言うようにググれば出てくる。
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。
上記のサイトからリンク辿れるサンプルはそのままでは動かなかった。
（俺の探し方が悪いのかもしれんが）

https://twitter.com/#!/NVIDIAJapan/statuses/96772501018906624
8月4日に秋葉原で行われるサードウェーブ主催の「GPGPUセミナー2011 夏」に協賛し、「NVIDIA GPU コンピューティング: CUDA 4.0による並列プログラミング」と題して講演します。http://bit.ly/l1ozuB
7月29日 TweetDeckから

>>835
一応できたので書いておきます。

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　だったような。

cuda勉強しようと思うんだけど、GTX520じゃさすがに貧弱すぎかな
460か560Tiくらいあったほうがいい？

勉強なら十分。

ありがとう明日にでも注文してみます

こころが折れそうです。プログラム初心者なのにCUDA…

それは厳しい。自分はCは長い間使ってたのですんなりと入れたけど。

GPGPUって高速化、最適化のためのものだから、
普通、プログラミングの基礎を固めた方がいい初心者とは真逆じゃない？
最初から最適化を過度に意識し過ぎると、貧乏性の汚いコードが身についちゃいそう。

>>898
GTX520で貧弱とか何贅沢なことを…と思ったけどGT520のことかな？

勉強なら十分、て言葉を否定する訳じゃないけど、
http://developer.nvidia.com/cuda-gpus
にあるCompute Capabilityも少し気にした方が良い、かも。

>>904
その一覧あってないような気がする
うちのGT430はSDKのサンプルでちゃんとCC2.1と表示されるよ
同一世代ならCCは一緒で性能が違うものだったと思う。

>>905
おお、それは初めて聞く情報！
グラボ詳しくないので同一世代がどれとか知りませんでした…ありがとうございます。

GeForce9400GTで練習したけどcore i7 940 1コアと同程度の速度しか出ずにがっかりしたな。

>>904
これって460のほうが580よりCudaに関しては性能上ってこと？

2.0と2.1の優劣は処理の内容によるが460で580に勝つのは無理。

なんていうんだろ、性能じゃなくてバージョン？
2.0と2.1でアーキテクチャ違うのかな

460のアーキテクチャは480の小改良版
で、460と480のフルスペック版が560Tiと580

Compute Capability 2.0と2.1の主な違いは
SMあたりのコア数やスケジューラの変更ではないでしょうかね。

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数などが変化したことによるコンパイラの最適化に適応させたもの」ということで

という話もありますし。

2.0と2.1はほぼ同じようなモノか

1.3→2.0が画期的すぎた

そのためのFermiアーキテクチャだからな

ノートPC（VAIO−F)に、CUDA ドライバ4.0入れたら動画再生出来なくなった。
Tool Kitは動作してサンプルも動くのに・・・
悔しいけど、3.2にロールバックした ><

kepler楽しみだ

OpenCLってどうなっちゃったの？

向こうのスレで聞けば？

SDKについてくる動作必須DLLって自分のバイナリと一緒に配布していいの？

top500

>>920
それぐらい自分で確認しろよ
ttp://developer.download.nvidia.com/compute/cuda/3_2_prod/toolkit/docs/EULA.txt
と思ったけど、なんか4.0のは見当たらないね・・・

インストールすればライセンスのファイルも入るのかな・・・

>>918
nVidiaもATIも足並みを揃える気無いだろ
技術を相手に公開しなきゃならんし、自社でライブラリをガンガン開発してったほうが効率良いし

NVIDIAのCUDAのドライバとSDKとサンプルとCUFFTをメンテナンスしている人間はそれぞれ別で、
意思の疎通は必ずしも密ではないのではないかと思う。まして、OpenCL対応は……

CUDA用のGTXを買おうと思っています
今のところ470か480あたりを買おうと思っていますが，CCを重視して560にしたほうが良いという場合は有りうるでしょうか
また，メーカーはELSA安定だとは思うのですが，他の安いメーカー製にした場合問題が発生しやすい等有るのでしょうか

どちらも無い。

GTX 580をCUDA用に使ってるけど、絶好調だよ。こっちにしなよ。
それにCompute Capabilityは2.0と2.1だと違いがないんじゃないの?

CUDA　4.0に3.2から移行した人の報告よろ

>>925
今さら４８０を選ぶ理由は価格以外にないと思うが、
それでも５７０とかにしたほうが幸せな気がする。
メーカーはCUDAに関しては保証してないし交換もきかないからどこでも
一緒なような気がする。

NVIDIA: CUDAを保証しているのはVGAでは高度な信頼テストをしているうちのQuadroだけ(キリッ
ELSA: うちのGeForceはNVIDIAに供給しているQuadroと同じ部品を使っているから安くて高品質(キリィッ
他社: うちのGeForceはリファレンスボードの無駄や非効率を改善しているから高性能(キリリィッ

> NVIDIA: CUDAを保証しているのはVGAでは高度な信頼テストをしているうちのQuadroだけ(キリッ

でもこれ、実際のところ正しいと思うよ。
リファの580をオーバークロックしてCUDAのプログラムを走らせると、
結構な確率でデータが化けるよ。化け方を見ると、明らかに中の演算装置が誤動作してるんだよね。
趣味ならともかく、仕事で使うんなら高いの買わないと駄目でしょ。

そこで３つの演算器に同じ計算させて多数決ですよ
バルタザール メルキオール カスパー

>> リファの580をオーバークロックしてCUDAのプログラムを走らせると、
問題はここじゃね〜のか？普通の頭はそう考えると思うぞ

>>933
> >> リファの580をオーバークロックしてCUDAのプログラムを走らせると、
> 問題はここじゃね〜のか？普通の頭はそう考えると思うぞ

もちろんそうなんだけど、意外にOCの許容範囲が低いなあという印象だったもんで…

> NVIDIA: CUDAを保証しているのはVGAでは高度な信頼テストをしているうちのQuadroだけ(キリッ
これ間違いじゃねーの？
QuadroじゃなくてTeslaのみだったと思うが。

TeslaはVGAじゃないからねぇ。

それにTeslaも今はVGA端子付いてるじゃん。

某大手PCショップの営業は、うちはGeForceでもCUDA保証します、
町のPCショップとかで買わない方がいいですよとか言っていた
おめーてきとーいってんじゃねえよとぶち切れそうになった

ELSAのQuadroでも保証しないのにgeforceで保証されるはずないよな

ARM，オリジナルGPUコア「Mali」をアピール。統合型シェーダ搭載の「Mali-T604」を搭載した端末は来年後半の登場へ
ttp://www.4gamer.net/games/137/G013737/20110725062/

Quadroの信頼性テストってなにやってんの？
動作保証温度とかのテスト結果見れるんなら欲しいのだが。

>>941
geforceよりは厳しいテストをしてますってだけで
そういうデータは公表されてない。

>>938
エラーが出たら無償交換すると言ってるだけでないの？
何がまずいのかさっぱりわからん。

>>943
「CUDA保証します」っていったらCUDAで計算ミスがでて損害が出ても
保証しますって風にも取れるからでしょ。

質問です。
(1) カーネルルーチンのコードって実行中はどのメモリに置かれているのですか。
　　グローバルメモリでないとしたら、サイズ制限が知りたい。
　　グローバルメモリなら多分命令キャッシュがあると思うので、そのサイズを知りたい。
(2) コンスタントメモリは（一旦コンスタントキャッシュに読み込んでしまえば）
　　シェアードメモリみたいなバンクコンフリクト無しでかつシェアードメモリ並に低レイテンシなのでしょうか。
(3) あるWarpでシェアードメモリに読み込んだデータは、同一SMの次のWarpで参照できる？
(3) 自スレッドがどのSMで実行されているのか知る方法は？

なんでVGAってMTBFや動作温度範囲公表しないんだろ

>>945
(1)と(2)はわからんが(3)はどっちも無理だろうな
次のWarpで参照しようにも次のWarpがどれなのかはスケジューラが毎回適当に振るんだから意味がないでしょ

>>945
(1)何がしたいのか分かりませんが、昔は2M命令という制限があった。よほどのことがない限りは
これを越えないと思います。
(2)低レイテンシですが、バンクが１個しかないので各スレッドが異なるアドレスにアクセスすると
バンクコンフリクト状態になります。
(3.1)質問の意味が分かりません。シェアードメモリを介したデータの受け渡しで反映されるまでの
タイミングに関することでしょうか？そのタイミングを気にする必要がある＝__syncthreads()を
入れていない状態と思いますが、あるWarpがシェアードメモリへの書き込みを行い、次にある他の
Warpがシェアードメモリからデータを読み込んだ場合、更新後のデータを受け取れる可能性は
ほぼ０でしょう。
(3.2)%smidという特殊レジスタによって自分がどのSMで実行されているかが分かります。
このレジスタを参照する方法はSDKにあるinlinePTXを参照してください。laneidをsmidに変更するだけです。

>>947>>948
レス�dクスだいたい了解
項番ミス失礼。(3.1)は、次の2つの可能性を問うものだったのです
1. コンスタントキャッシュが遅い場合の代替手段。定数はSM別 -- これは、スレッド別定数でなければコンスタントキャッシュでおk理解
2. Warpダイバージェンス対策としてのWarp中断 -- これは誤解だった感じ。できたとしても中断しない場合と同速と思われる。

まあ調べたらシェアードメモリ内容の永続期間は同一TB内しか保証されない仕様っぽいので
SM別に処理を分けるはアーキテクチャーを活かした良い方法とは言えませんけども
（TB内のスレッドがどのSMに割り当てられるかは>947の通り予測し難い）

>>949の2.はまだ良く分かりませんが、

>>950の方のTB間のシェアードメモリデータの再利用は一応可能です。（Fermiより前は知りませんが）
しかし保証されていませんし、大して速くならないので使わない方が良いと思います。

>>944
PC関係の製品で計算結果の保証なんてどこもやってないじゃん。
HPやDELLやNECの数百万円の鯖を買っても「保証」は無償修理もしくは交換のこと。

つーか、GeforceでCUDA稼働保証があるというその大手ショップの名前を是非教えてほしい。
会社のマシンに試験用に一つ入れたいが、Teslaじゃ稟議が通らない。

>>952
営業の口八丁、もしくは個別契約の餌だろうから聞いても無駄じゃね？
宮崎大みたいに数百とか導入予定なら交換対応してもペイするかもしれんが
1個だけとかじゃそんな話でないでしょ。

>>953
>>938が、そんなでかい金の決定権がある人間のレスに見えるかね？

どちらにしろ1個とか少数でも保証してくれるところはないと思うけどね。
というか数百だろうがgeforceでCUDAの演算結果保証してくれるところなんてないと思う。
あるならホント知りたいもんだわ

　　　　　　　　　　　,,,,wwwwWwWム,,,
　　　　　　　　　　 ミ川ミミ川三ミ川ミミ
　　　　　　　　　　 ミ ,.--"゛~~~~''´゛゛゛ミ
　　　　　　　　　 　ミ川　　　＿　　　＿ l.
　　　　　　　　　　 ミ川　 -´==(｀r__i ｪl ｀r
　　　　　　　　　　ミ川￣￣'､＿,r'　丶Lﾉ
　　　　　　　　　　ミ川　　　　 /(　 __）） |
　　　　　 　 　　　ミ川　　　　ｒ―''"入　|
　　　　　　 　 　 　｜i　　　　廷廾ﾆツ ｊ
　　　　　　　　　　 　|ー､＿＿＿＿,.ノ
　 　▂▃▃▃▂　　ヽ∧∧∧ 　　　　/
. 　██████＜ 角刈り！＞　／
　 ███ け ,fj　/ ∨∨V ヽ　／
　　 ｀█､"ヮノ　　　　 |　　／
　 /⌒ヽヾ'ﾘ、　　　　 |　/
.　{ 　　j｀ｰ' ﾊ 　 　 　//　ヽ∧∧∧∧∧∧∨/
　 ｋ〜'ｌ 　　ﾚﾍ. 　 ,r'ス　＜ソイヤ！ソイヤ！＞
　 |　ヽ ＼　ト､ ヽ-kヾｿ　＜　 角刈りソイヤ！＞
.　 l　　＼ ｀ｰ‐ゝ-〈/´　　　/ ∨∨∨∨∨∨ヽ

長崎大のセンセによれば、GeforceのCUDA使用時の不良率は2割弱。
不良の検証を客に押しつけられるなら、理論上は2割増しの値段で売れば動作保証は可能だ。

CUDA使用時の不良率ってどう言う事よ？
プログラムの実行結果は間違ってるけれど、エラーにはならずに正常終了した風に終わるってこと？

そんなとこ

本当にか。
なんか、計算機の根幹を揺るがすような話だなw

どんだけずれるのかにもよるな

元々ゲームとかのポリゴン計算するためのものだから
多少結果が変でも問題ないからとかなんとか。

画像処理で1frameのどこかの1dotがまれに色が違っていても誰も気がつかないからね。

長崎大のセンセのプレゼン資料はこれね。
http://www.ssken.gr.jp/MAINSITE/activity/sectionmeeting/sci/2009-2/lecture-1/ppt.pdf

・Geforce GTX295を421枚買ったら、うち72枚が壊れていた。
・壊れていたうち31枚を買い足して、380枚使ってスパコンを作った。

やっぱりなｗ CUDAで負荷かけるととたんに計算結果がおかしくなったりするんだよな。
CPU例外が起きておかしいなと思ったら、CUDAからのデータが化けてたとかあったわ。

>>964
今頃「やっぱりな」とか寝ぼけたことを言っているやつの場合、
そいつの書いたプログラムがバグっている可能性の方がはるかに高いね。

>>963
コカコーラそんなに飲んで大丈夫？（48/54)