トップページ⇒tech
1001コメント314KB

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

■ このスレッドは過去ログ倉庫に格納されています
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:52
関連サイト
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/
0003デフォルトの名無しさん2010/12/04(土) 22:27:14
なるほど
低レベル質問にも対応するスレだったんだな
道理でCUDAは意味が無いという意見が出て来るわけだ
0004デフォルトの名無しさん2010/12/04(土) 23:04:15
いいかげんpertはなおしたほうがいいと思うの…
0005デフォルトの名無しさん2010/12/04(土) 23:09:17
pertって？
ORのPERT？
だとしてもCUDAとの関係がわからん
0006デフォルトの名無しさん2010/12/04(土) 23:32:44
…単純にスペルミスだろ
0007デフォルトの名無しさん2010/12/04(土) 23:43:31
あ〜スレタイのか、納得しますた
0008デフォルトの名無しさん2010/12/05(日) 01:01:35
並列計算の実験でGT240のグラフィックボードを買おうとしてるんですが
買ったほうがいいですか？ちなみに行列の演算を高速化したいというのが
目的です。
0009デフォルトの名無しさん2010/12/05(日) 01:28:17
金をどぶに捨てる気か？
GT430にしろよ・・・
0010デフォルトの名無しさん2010/12/05(日) 01:30:18
このスレってゆとりも対象にしてるの？
0011デフォルトの名無しさん2010/12/05(日) 11:15:34
SQLiteをVRAMに構築して並列で検索してはどうだろう
0012デフォルトの名無しさん2010/12/06(月) 02:44:34
どうしてもvisual profilerのcounter valueの項目のうち、cc2.0以降のactive cycle とactive warps
の意味がわかりません。意味を詳しく教えていただけませんか？
0013デフォルトの名無しさん2010/12/06(月) 17:16:54
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が分からないとかなら青木先生の本でも読めばいい
0014デフォルトの名無しさん2010/12/06(月) 19:37:01
レスありがとうございます。
warpについては理解しています、がいまだcycleがわかりません・・・

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

またこれらの値は大きいほうがいいのか、小さいほうがいいのかもいまいちわからない状況です。
たとえばSummary Tableにでているactive warps/active cycle やretire ipcなども
意味と値が大きいほうがいいのか、がいくら調べてもわかっていない状況です。
よろしければご教授お願いします。
0015デフォルトの名無しさん2010/12/06(月) 21:51:21
サイクルはクロックやHzとほぼ同じ
0016デフォルトの名無しさん2010/12/06(月) 21:58:05
一応聞いとこ
>>14高校生？
0017デフォルトの名無しさん2010/12/06(月) 22:08:20
>>14
プロファイラのマニュアルはちゃんとあるし、
プログラミングガイド、ベストパフォーマンスガイドに基本は全て載ってる
全部公式のガイドだぞ
0018デフォルトの名無しさん2010/12/06(月) 22:26:33
>>16
初心者とゆとりが融合した素晴らしい素材ってことは分かる
0019デフォルトの名無しさん2010/12/06(月) 23:18:43
cycleがクロックだとは分かっていましたが、1cycleにつき1SM内では1warpのある処理を実行できると思っていましたので（正確には2warpを2cycleで）
1cycleで複数のwarpを処理できるという点がいまいちどういう仕組みなのかわかっていませんでした。
Type SMとありましたので、この値はあるSMで計測された値だということがわかりますが、
上記のような認識だとつじつまが合わない（1SMで1warpずつ処理すると思っていたのに、1cycleで最大48warpがアクティブになるという点）ため質問させていただきました。
ややこしく、説明がいまいちになってしまい申し訳ないです。

プロファイラのマニュアルも読みましたが、読んだ上でゆとり質問となり言い返す言葉もありません。
これ以上スレを汚すのもどうかと思いますので質問を〆させていただきます。
レスくださった方ありがとうございました。
0020デフォルトの名無しさん2010/12/06(月) 23:50:52
active warpは実行可能な待ち行列に待機中のwarpのことだろ。
待ち行列が無いと、演算器に空きが出来てもwarpを割り当てる事が
出来ないから、ある程度の長さの待ち行列が常にある事が
(演算器の利用効率を最大化するという意味では)重要になる。

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

0021デフォルトの名無しさん2010/12/07(火) 01:36:55
>>19
質問とはちょっとずれるかもしれないけど、
http://www.softek.co.jp/SPG/Pgi/TIPS/public/accel/gpu-accel2.html
この記事が長いけど結構分かりやすいかもね
ゆとりとかは挨拶だから気にするなよ
0022デフォルトの名無しさん2010/12/07(火) 01:52:08
グラフィックボードを買うのは初めてなんですが、
GT430で探したんですが1スロットで収まる商品が2つしかなく、
評判も安定性も良くないらしいのでELSAのGT240にしようかと思っています。
何かアドバイスがあればよろしくお願いします。
0023デフォルトの名無しさん2010/12/07(火) 02:23:56
このスレの住人にはやはり愛がありますね。
21さんのレスなんか涙がレスが読めなくなるほど。

ゆとりくんは心に刻むんだぞ。
玄人のおぢさんたちに生暖かく見守られているということを。
0024デフォルトの名無しさん2010/12/07(火) 02:39:12
>>22
430使ってるが不具合ないよ。
あと今更fermi以前でPG組む気になれん。
CUDA目的ならCCのバージョンを意識しないと駄目だと思うぞ。
それをわかってて240にするなら何も言わない。
0025デフォルトの名無しさん2010/12/07(火) 02:40:26
キャッシュの有無はでかいね。2スロットは兎も角、長さにも要注意。
0026デフォルトの名無しさん2010/12/07(火) 02:53:55
>>23
目的、予算、制限を明確にしてほしいな
高速処理が至上命題で、予算も制限もないなら
言うまでもなく、Teslaボードをお勧めします
やっぱりfermi世代の方が無難なチョイスなので、
fermiチップのボードだとしてアドバイス


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

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

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

それからWindowsの場合は画面出力用のボードとCUDA用のボードを分けると
デバッグしやすい気がするので、ボード2枚構成というのも考慮の余地があるかも
Linuxはしらない
0027デフォルトの名無しさん2010/12/07(火) 02:58:31
おっと、GT240はfermiじゃないね
GT440だと脳内保管されてた・・・
ワシもfermiでないGT240は推奨しません

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

fermiはキャッシュのおかげでタコなコードでも
そこそこ動いてくれるから
0028262010/12/07(火) 03:09:57
しまったアンカも間違えてた
>>23は間違いで>>22
もう寝ます
0029デフォルトの名無しさん2010/12/07(火) 09:53:37
fermiってなんですか？自分は行列の演算を高速化したいだけなので
詳しいことはわかりませんが、同じコア数でもGT240とGT439ではなぜ
これだけ値段の差があるのか解りません。できれば安いほうを買いた
いのですが予算は10000円程度で考えています。
何かよきアドバイスがあればよろしくお願いします。
0030デフォルトの名無しさん2010/12/07(火) 09:59:41
fermiってやつを買えば簡単なプログラミングで高速演算できる
fermiじゃなきゃ高速演算の手間がかかる
と思えばおk
0031デフォルトの名無しさん2010/12/07(火) 10:11:53
GF-GT430-LE1GHD/1ST
これなんかどうでしょう？騒音がすごいというレビューを見ましたが。
0032デフォルトの名無しさん2010/12/07(火) 12:48:45
>>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しかないな。
0033デフォルトの名無しさん2010/12/07(火) 13:16:57
単精度だけでいいのかな。
最初はIONみたいなので消費電力を気にせずつかえるのがいいと思う。
速いカードの方が当然チューニングの効果がでるからおもしろい。
プロファイラで取れる情報もカードによって違うみたい。
0034デフォルトの名無しさん2010/12/07(火) 13:24:26
テクスチャメモリってどのような利点があるのですか？
添字が小数点の時の線形補完を使ってみたいのですが、
メモリはチップ外にあるということはグローバルメモリと同等の速度なのですよね
キャッシュはされるみたいですが同じデータを何度も利用しなくては意味がないということでしょうか
0035デフォルトの名無しさん2010/12/07(火) 15:47:54
ELSA GLADIAC GT 430 LPと
EVGA GeForce GT 430で迷っています。
どちらがいいかアドバイスをよろしくお願いします。
0036デフォルトの名無しさん2010/12/07(火) 15:54:56
どっちでも変わらないとおもうよ。
印象でいえばELSAは手堅い印象。
EVGAはゲーム向けのチューニングとかしてて精度の点が不安。
予算的に許すならquadro600がいいかもしれない。
ただしELSAでもgeforce,quadroはcudaでの演算結果の保証はしてないから
演算を間違うコアが含まれていても交換対象ではない。
0037デフォルトの名無しさん2010/12/07(火) 16:01:43
因みに整数演算しか行いません。
有限体の行列なので浮動小数点は使いません。
予算は１万円です。コアの精度ってなんですか？
コアが計算を間違えることがあるのですか？
精度を求めるならどちらがいいでしょうか？
質問ばかりで申し訳ありませんが、ご回答よろしくお願いします。
0038デフォルトの名無しさん2010/12/07(火) 16:06:17
予算1万ならELSAでぎりぎり。
0039デフォルトの名無しさん2010/12/07(火) 16:08:52
それ、浮動小数点演算いるんじゃね？
浮動小数点の意味わかってる？
0040デフォルトの名無しさん2010/12/07(火) 16:13:50
floatとかdoubleとかですよね。
自分が計算したいのはunsigned shortなのです。
以前雑誌の特集にCUDAの記事がありNVIDIAのコアには整数演算も
可能であるということが書かれてありました。なので整数の行列も
高速化できるかと思っていました。何か間違いがありましたら
ご指摘をお願いします。
0041デフォルトの名無しさん2010/12/07(火) 16:14:46
quadro600がいいかもしれないといったのは、
演算結果は保証されないけど、geforceより厳しい検査がされているから
より演算結果の値にエラーがでにくいだろうという推測にもとづくもの。

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

精度を求めるというのが演算結果を保証するならという意味ならTESLAしかない。
そこまでの予算を確保できない場合、上記のようなコアが含まれていないかを確認するプログラムを自分で
組んで確認しないといけない。
単純に足し算だけでなく、四則演算、浮動小数点演算、指数関数、超越関数、L1、L2キャッシュ、
VRAMに至るまで正当性を確認する必要がでると思ってる
0042デフォルトの名無しさん2010/12/07(火) 16:21:20
あ、CPUでの演算結果があるなら、同じ計算をしてみて結果が同じかどうかで
簡略な確認としてもいいかもしれない。それで許されるのであれば。
0043デフォルトの名無しさん2010/12/07(火) 16:30:24
CPUの演算結果とGT430の９６個のコアの演算結果を比較して
間違いがないことを実験する必要があるのですね。
同じ計算機なのにどうしてそのような確率的な動作をするのか
理解しかねます。CPUだけが特別に演算が正確なのはなぜでしょう?
GPUも同じ32ビットレジスタなので構造が同じだと思ったのですが。
0044デフォルトの名無しさん2010/12/07(火) 16:41:06
その比較はあくまで簡略な方法だよ。全部のコアが使われる保証だってないんだから。

CPUは元々計算用に作ってあるでしょ？CPUで１＋１＝２を保証しなくてだれが買うよ？
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:10
体だと素数で剰余をとるのですか、どういうコードか教えてくれませんか。
0050デフォルトの名無しさん2010/12/07(火) 16:54:15
長崎大の検査プログラムがほしい。。。。
が、無理なんだろうなあ
0051デフォルトの名無しさん2010/12/07(火) 17:07:36
ちょっと恥ずかしいけど見せますねｗ
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]);

}
0052デフォルトの名無しさん2010/12/07(火) 17:08:51
__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;
}
0053デフォルトの名無しさん2010/12/07(火) 17:13:42
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;
}
0054デフォルトの名無しさん2010/12/07(火) 17:15:17
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;
}
0055デフォルトの名無しさん2010/12/07(火) 17:15:28
メモリ階層ってこの理解でいいのかな。自信がない

レジスタ

L1

共有メモリ

L2


0056デフォルトの名無しさん2010/12/07(火) 17:16:12
>>55
誤爆
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:08
CUDAでINTって糞性能だろ・・・
int(4)なら加算だけでもfloatの4倍かかるぞ
FMAとかも使えないし
0060デフォルトの名無しさん2010/12/07(火) 18:42:12
こっちのほうが早いですね。荒らしてしまい申し訳ないです。

http://sky.geocities.jp/tcshacina/grs1.c
0061デフォルトの名無しさん2010/12/07(火) 18:45:53
>>59
じゃあグラフィックボードで並列計算するのはこの場合逆効果になる
のでしょうか？買わないほうがいいですか？
0062デフォルトの名無しさん2010/12/07(火) 18:50:41
floatで計算すりゃいいんじゃね？
intにこだわる必要あるのか？
0063デフォルトの名無しさん2010/12/07(火) 18:58:39
floatのXORってできるんですか？疑似乱数生成器なのでfloatで一旦
計算しておいて後から一斉にunsigned charに変換することが可能で
しょうか？暗号に使う疑似乱数なので整数で処理しないといけない
という条件があります。
0064デフォルトの名無しさん2010/12/07(火) 19:03:47
楕円暗号入門を手に取ってみたけどやっぱりわからなかった。
あれだけコードがかけるならいいおもちゃだから買えばいいんじゃないですか。
Cellも整数の論理演算にむいていた気がします。６個のタスクを生成できる
ので使いやすい場合もありました。
0065デフォルトの名無しさん2010/12/07(火) 19:05:17
>>63
xorshiftのこと？
であればCURANDってのがあるからそれ使えば乱数生成可？
メルセンヌもSDKについてる
0066デフォルトの名無しさん2010/12/07(火) 19:14:48
普通のCPUでやっているような排他的論理和（XOR）をCUDAもできるのか
という疑問です。できなければfloatのままデータを生成して、後で
CPUにまとめてfloat->intの処理をさせながらXORを計算するという
方法ではうまくいかないでしょうか？マニュアルには論理演算も
可能と書いてありますけど、整数演算の速度が浮動小数点より遅い
というのが気がかりです。

Rubyで作った暗号化モジュールいろいろ
http://sky.geocities.jp/tcshacina/crypt.html
0067デフォルトの名無しさん2010/12/07(火) 19:24:50
だからさXORを乱数生成に使うんじゃねの？
0068デフォルトの名無しさん2010/12/07(火) 19:29:45
使います、そのためには整数でなければならないと思っています。
0069デフォルトの名無しさん2010/12/07(火) 19:32:36
だからCURAND使えばいいべよって書いたんだべ
CURANDのマニュアル読んで目的にかなうかどうか確認し
0070デフォルトの名無しさん2010/12/07(火) 19:34:29
xorshiftってなんですか？自分のオリジナルのアルゴリズムに
どうやってそれを使えばいいのですか？ソースは晒してあるので
自由に見て感想を聞かせてください。
0071デフォルトの名無しさん2010/12/07(火) 19:37:09
悪いが人のソース追っかけるほど暇ないよ
0072デフォルトの名無しさん2010/12/07(火) 19:40:05
私のソースを見たいと言った人はどこへ行ったんだ！？
0073デフォルトの名無しさん2010/12/07(火) 20:49:54
胡散臭い文体見た瞬間に例の暗号の人だってすぐ分かった。
0074デフォルトの名無しさん2010/12/07(火) 21:00:10
＞自分のオリジナルのアルゴリズムにどうやってそれを使えばいいのですか？
それを考えるのがプログラマだろうが！１から１０まで聞いてないで自分でやってみろ
0075デフォルトの名無しさん2010/12/07(火) 22:12:31
使えないと言いたい
0076デフォルトの名無しさん2010/12/07(火) 23:07:11
一日に書き込む疑問文は3つまでにしてくれ。
0077デフォルトの名無しさん2010/12/07(火) 23:28:12
珍しく伸びてるなと思ったら…

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

答えるは簡単だがこの程度のことを人に頼る時点で
そいつにまともなプログラムが作れると思えん
0078デフォルトの名無しさん2010/12/07(火) 23:48:44
>>77
同感。
そういうやつでも数学的素養が優れていたりして特定のモデルの構築とかには
優れた才能を発揮する奴がいるから、そういう所だけは認めざるを得ないこともある
のが悔しい。
0079デフォルトの名無しさん2010/12/07(火) 23:52:59
シミュレーション系で使う人、シミュレート板に新スレ立てたんでよろしく
プログラミングの方法より使い方を話題にしたいなと

GPGPUでシミュレーションを高速化してみる
http://kamome.2ch.net/test/read.cgi/sim/1291687353/
0080デフォルトの名無しさん2010/12/08(水) 01:41:45
>>66
Micro-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:08
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


と、あった
0082デフォルトの名無しさん2010/12/08(水) 07:54:02
だからFermiはfloatとintが同じ速度だとあれほど(ry
0083デフォルトの名無しさん2010/12/09(木) 17:13:20
環境をインストールしてはじめようと思ったんだが
SDKのサンプルtempleteが2008で開けない。

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

って出たからrulesファイル探してるんだけどないんだよな…
SDK再インストールしたけどやっぱりない。たすけて
0084デフォルトの名無しさん2010/12/10(金) 11:21:46
ドライバAPIを使えば，一つのホストスレッドで
複数のGPUに並列でカーネルを実行させることが出来ますか？
0085デフォルトの名無しさん2010/12/10(金) 12:16:29
二つのホストスレッドが必要じゃね?
0086842010/12/10(金) 13:49:23
>>85
無理ですか…

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

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

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

マニュアルのドライバAPIのコンテキストの部分を読んでいたら，
色々命令があるので出来るかもしれないと思い質問させていただきました．
ホストスレッドの数は複数でも問題ないので，
コンテキストの操作に詳しい方いらっしゃいますか？
0087デフォルトの名無しさん2010/12/10(金) 15:36:25
複数GPUの扱いがおかしいんじゃね？
良く見直してみ
0088デフォルトの名無しさん2010/12/10(金) 17:16:41
kernelの実行命令が繰り返されるのはどういう状況のときそうなりますか？

Globalからの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
visualvasicでプログラムを組んでいるものですが
ビルドルールの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:25
エミュレーションって3.1はないんじゃないの
0091デフォルトの名無しさん2010/12/10(金) 23:55:04
では、printf();は設定なしにつかえるということなのでしょうか？
calling a host function from a __device__/__global__ function is not allowed
するとやはりエラーが消えないのですが・・・
0092デフォルトの名無しさん2010/12/11(土) 00:13:47
>>91
おまえ言ってること矛盾してるし、何をやってそういうエラーがでたかの説明してなくて
アドバイスを求めても無理。
vbにprintfとかないよね？どういうことなの？
0093デフォルトの名無しさん2010/12/11(土) 00:14:07
3.1からデバッガが標準でつくようになったのでエミュレーションはもうなくなったはず
Fermi系列なら確かそのまんまprintf使えるけど

windowsならVisual Nsightとかいうのでデバッグできたはず
あとはどうしてもエミュレーション使いたいならcuda2.3使うとか
0094デフォルトの名無しさん2010/12/11(土) 00:16:07
>>93
printfは確かfermi系だけのサポートのはず。
そして2.3はfermiサポートしてない。3.0からだったはず。
0095デフォルトの名無しさん2010/12/11(土) 00:19:12
Prallel Nsightな
0096デフォルトの名無しさん2010/12/11(土) 00:24:18
>>88
同様の現象が起きていて困ってる
GTX480ではinstruction issueが取れなかったんだが、Tesla使い？
0097882010/12/11(土) 01:12:07
>>96
いえ、GTX480です。ASUS製としか覚えてませんが・・・
0098デフォルトの名無しさん2010/12/11(土) 18:03:20
C1060x 4をFermiのGTX580x 4に入れ替えた場合の性能比較を
探しています。

フロートの性能は欲しいがC2050は高すぎるということで。
0099デフォルトの名無しさん2010/12/11(土) 21:19:19
C1060って260相当だっけ？
3DVantageとかの260と580の性能差を4倍したものを求めて見積もればいいんじゃない？
0100デフォルトの名無しさん2010/12/11(土) 22:44:36
これでCPUがシングルコアだったら大笑い。
0101デフォルトの名無しさん2010/12/11(土) 23:25:14
floatじゃなくてdouble茶羽化？
■ このスレッドは過去ログ倉庫に格納されています