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

[表示 : 全て最新50 1-99 101- 201- 301- 401- 501- 601- 701- 2chのread.cgiへ]
Update time : 12/12 04:31 / Filesize : 191 KB / Number-of Response : 793
[このスレッドの書き込みを削除する]
[＋板最近立ったスレ＆熱いスレ一覧 : ＋板最近立ったスレ／記者別一覧] [類似スレッド一覧]


↑キャッシュ検索、類似スレ動作を修正しました、ご迷惑をお掛けしました

1 名前：デフォルトの名無しさん mailto:sage [2010/12/04(土) 21:57:13 ]: このスレッドは、他のスレッドでは書き込めない超低レベル、
もしくは質問者自身何が何だが分からない質問を勇気を持って書き込むスレッドです。
CUDA使いが優しくコメントを返しますが、
お礼はCUDAの布教と初心者の救済をお願いします。

CUDA・HomePage
www.nvidia.com/cuda

関連スレ
GPUで汎用コンピューティングを行うスレ
hibari.2ch.net/test/read.cgi/tech/1167989627/
GPGPU#5
hibari.2ch.net/test/read.cgi/tech/1281876470/

前スレ
【GPGPU】くだすれCUDAスレ【NVIDIA】
pc12.2ch.net/test/read.cgi/tech/1206152032/
【GPGPU】くだすれCUDAスレ pert2【NVIDIA】
pc12.2ch.net/test/read.cgi/tech/1254997777/
【GPGPU】くだすれCUDAスレ pert3【NVIDIA】
hibari.2ch.net/test/read.cgi/tech/1271587710/
2 名前：デフォルトの名無しさん mailto:sage [2010/12/04(土) 21:58:52 ]: 関連サイト
CUDA Zone
www.nvidia.co.jp/object/cuda_home_jp.html
フォーラム
forum.nvidia.co.jp/EokpControl?&event=TE0001

CUDAに触れてみる
chihara.naist.jp/people/STAFF/imura/computer/OpenGL/cuda1/disp_content

cudaさわってみた
gpgpu.jp/article/61432191.html

CUDA のインストール
blog.goo.ne.jp/sdpaninf/e/9533f75438b670a174af345f4a33bd51

NVIDIAの「GeForce 8800 GT(G92)」と次に控える64-bit GPUアーキテクチャ
pc.watch.impress.co.jp/docs/2007/1031/kaigai398.htm

CUDAを使う
tech.ckme.co.jp/cuda.shtml

NVIDIA CUDAを弄ってみたその2
dvd-r.sblo.jp/article/10422960.html

CUDAベンチ
wataco.air-nifty.com/syacho/2008/02/cuda_2044.html

KNOPPIX for CUDA
www.yasuoka.mech.keio.ac.jp/cuda/
3 名前：デフォルトの名無しさん mailto:sage [2010/12/04(土) 22:27:14 ]: なるほど
低レベル質問にも対応するスレだったんだな
道理でCUDAは意味が無いという意見が出て来るわけだ
4 名前：デフォルトの名無しさん mailto:sage [2010/12/04(土) 23:04:15 ]: いいかげんpertはなおしたほうがいいと思うの…
5 名前：デフォルトの名無しさん mailto:sage [2010/12/04(土) 23:09:17 ]: pertって？
ORのPERT？
だとしてもCUDAとの関係がわからん
6 名前：デフォルトの名無しさん mailto:sage [2010/12/04(土) 23:32:44 ]: …単純にスペルミスだろ
7 名前：デフォルトの名無しさん mailto:sage [2010/12/04(土) 23:43:31 ]: あ～スレタイのか、納得しますた
8 名前：デフォルトの名無しさん [2010/12/05(日) 01:01:35 ]: 並列計算の実験でGT240のグラフィックボードを買おうとしてるんですが
買ったほうがいいですか？ちなみに行列の演算を高速化したいというのが
目的です。
9 名前：デフォルトの名無しさん mailto:sage [2010/12/05(日) 01:28:17 ]: 金をどぶに捨てる気か？
GT430にしろよ・・・
10 名前：デフォルトの名無しさん mailto:sage [2010/12/05(日) 01:30:18 ]: このスレってゆとりも対象にしてるの？
11 名前：デフォルトの名無しさん mailto:sage [2010/12/05(日) 11:15:34 ]: SQLiteをVRAMに構築して並列で検索してはどうだろう
12 名前：デフォルトの名無しさん mailto:sage [2010/12/06(月) 02:44:34 ]: どうしてもvisual profilerのcounter valueの項目のうち、cc2.0以降のactive cycle とactive warps
の意味がわかりません。意味を詳しく教えていただけませんか？
13 名前：デフォルトの名無しさん mailto:sage [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が分からないとかなら青木先生の本でも読めばいい
14 名前：デフォルトの名無しさん mailto:sage [2010/12/06(月) 19:37:01 ]: レスありがとうございます。
warpについては理解しています、がいまだcycleがわかりません・・・

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

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

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

fermiの細かいアーキテクチャまでは分からないが、
アレだけレジスタが多いとregisterがマルチバンクになっている可能性が高いから
例えば奇数cycleは奇数IDのwarp、偶数cycleは偶数IDのwarpしか発行できない
とかそういう制限があるかもしれない。
その場合はactive warpだからといって
常に任意のcycleに実行可能とならない場合もあるから
待ち行列は長い方が良い。
21 名前：デフォルトの名無しさん mailto:sage [2010/12/07(火) 01:36:55 ]: >>19
質問とはちょっとずれるかもしれないけど、
www.softek.co.jp/SPG/Pgi/TIPS/public/accel/gpu-accel2.html
この記事が長いけど結構分かりやすいかもね
ゆとりとかは挨拶だから気にするなよ
22 名前：デフォルトの名無しさん [2010/12/07(火) 01:52:08 ]: グラフィックボードを買うのは初めてなんですが、
GT430で探したんですが1スロットで収まる商品が2つしかなく、
評判も安定性も良くないらしいのでELSAのGT240にしようかと思っています。
何かアドバイスがあればよろしくお願いします。
23 名前：デフォルトの名無しさん mailto:sage [2010/12/07(火) 02:23:56 ]: このスレの住人にはやはり愛がありますね。
21さんのレスなんか涙がレスが読めなくなるほど。

ゆとりくんは心に刻むんだぞ。
玄人のおぢさんたちに生暖かく見守られているということを。
24 名前：デフォルトの名無しさん mailto:sage [2010/12/07(火) 02:39:12 ]: >>22
430使ってるが不具合ないよ。
あと今更fermi以前でPG組む気になれん。
CUDA目的ならCCのバージョンを意識しないと駄目だと思うぞ。
それをわかってて240にするなら何も言わない。
25 名前：デフォルトの名無しさん mailto:sage [2010/12/07(火) 02:40:26 ]: キャッシュの有無はでかいね。2スロットは兎も角、長さにも要注意。
26 名前：デフォルトの名無しさん mailto:sage [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はしらない
27 名前：デフォルトの名無しさん mailto:sage [2010/12/07(火) 02:58:31 ]: おっと、GT240はfermiじゃないね
GT440だと脳内保管されてた・・・
ワシもfermiでないGT240は推奨しません

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

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

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

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

CPUは元々計算用に作ってあるでしょ？CPUで１＋１＝２を保証しなくてだれが買うよ？
GPUは元々描画用に作られてて座標計算とかfloatで計算することが多い。
そのときに1.000002が1.000002になってもあんまり影響ないでしょ？
もちろん正確に計算できるほうがいいが、描画だったら多少不正確でも問題はないだろうっていう判断だと思う。
それにCPUに比べてALUが膨大なこともあってすべてのコアの演算精度を保証してたらキリがないでしょ。コスト的に。
そういう用途の人のために厳しいテストをしたTESLAが用意されてる訳。でもその分料金上乗せさせてもらいますよって事だよ。
45 名前：デフォルトの名無しさん [2010/12/07(火) 16:41:22 ]: 検算しないといけないとなると計算結果が信用できないということで
並列化のメリットがなくなってしまう気がします。
46 名前：デフォルトの名無しさん mailto:sage [2010/12/07(火) 16:42:14 ]: あ、ごめん。
×　そのときに1.000002が1.000002になっても
○　そのときに1.000001が1.000002になっても
47 名前：デフォルトの名無しさん mailto:sage [2010/12/07(火) 16:45:03 ]: いや、毎回検算しろといってるわけじゃなく、

一回CPUと同等の計算をして、演算結果に謝りがないことの確認をもって
GPUのコアにエラーが発生していないことの簡略確認にすれば？
ってことだよ
48 名前：デフォルトの名無しさん [2010/12/07(火) 16:48:03 ]: 浮動小数点はどうでもいいけど、せめて整数演算だけは正確に
計算してもらいたいと祈るばかりですね。
49 名前：デフォルトの名無しさん mailto:sage [2010/12/07(火) 16:51:10 ]: 体だと素数で剰余をとるのですか、どういうコードか教えてくれませんか。
50 名前：デフォルトの名無しさん mailto:sage [2010/12/07(火) 16:54:15 ]: 長崎大の検査プログラムがほしい。。。。
が、無理なんだろうなあ
51 名前：デフォルトの名無しさん [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]);

}
52 名前：デフォルトの名無しさん [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;
}
53 名前：デフォルトの名無しさん [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;
}
54 名前：デフォルトの名無しさん [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;
}
55 名前：デフォルトの名無しさん mailto:sage [2010/12/07(火) 17:15:28 ]: メモリ階層ってこの理解でいいのかな。自信がない

レジスタ

L1

共有メモリ

L2
56 名前：デフォルトの名無しさん mailto:sage [2010/12/07(火) 17:16:12 ]: >>55
誤爆
57 名前：デフォルトの名無しさん mailto:sage [2010/12/07(火) 17:18:52 ]: メモリ階層ってこの理解でいいのかな。自信がない

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

この資料に書いてあるとかいう情報があるなら
それも教えてもらえると助かります
58 名前：デフォルトの名無しさん [2010/12/07(火) 17:19:11 ]: 一部省略してありますが大体こんな感じです。素体は楕円曲線暗号に
使っていますがRubyの実装でC言語にはしてないです。
このプログラムは疑似乱数生成器で基本的には最大32*32の行列演算を
主に行っています。これが並列化できれば32回の計算で256ビットの
乱数を生成できるので効率的になります。
というかなればいいなあと思ってます。
59 名前：デフォルトの名無しさん mailto:sage [2010/12/07(火) 18:40:08 ]: CUDAでINTって糞性能だろ・・・
int(4)なら加算だけでもfloatの4倍かかるぞ
FMAとかも使えないし
60 名前：デフォルトの名無しさん [2010/12/07(火) 18:42:12 ]: こっちのほうが早いですね。荒らしてしまい申し訳ないです。

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

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

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

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

GPGPUでシミュレーションを高速化してみる
kamome.2ch.net/test/read.cgi/sim/1291687353/
80 名前：デフォルトの名無しさん mailto:sage [2010/12/08(水) 01:41:45 ]: >>66
Micro-benchmarking the GT200 GPU
www.eecg.toronto.edu/~moshovos/CUDA08/arx/microbenchmark_report.pdf

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

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

Fermiでの資料を探したが見つからなかった。誰か頼む。
81 名前：デフォルトの名無しさん mailto:sage [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

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

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

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

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

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

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

マニュアルのドライバAPIのコンテキストの部分を読んでいたら，
色々命令があるので出来るかもしれないと思い質問させていただきました．
ホストスレッドの数は複数でも問題ないので，
コンテキストの操作に詳しい方いらっしゃいますか？
87 名前：デフォルトの名無しさん mailto:sage [2010/12/10(金) 15:36:25 ]: 複数GPUの扱いがおかしいんじゃね？
良く見直してみ
88 名前：デフォルトの名無しさん mailto:sage [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のときにおこると思っていたのですが、間違っているなら教えてください。
89 名前：デフォルトの名無しさん [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です。
90 名前：デフォルトの名無しさん mailto:sage [2010/12/10(金) 23:51:25 ]: エミュレーションって3.1はないんじゃないの
91 名前：デフォルトの名無しさん [2010/12/10(金) 23:55:04 ]: では、printf();は設定なしにつかえるということなのでしょうか？
calling a host function from a __device__/__global__ function is not allowed
するとやはりエラーが消えないのですが・・・
92 名前：デフォルトの名無しさん mailto:sage [2010/12/11(土) 00:13:47 ]: >>91
おまえ言ってること矛盾してるし、何をやってそういうエラーがでたかの説明してなくて
アドバイスを求めても無理。
vbにprintfとかないよね？どういうことなの？
93 名前：デフォルトの名無しさん mailto:sage [2010/12/11(土) 00:14:07 ]: 3.1からデバッガが標準でつくようになったのでエミュレーションはもうなくなったはず
Fermi系列なら確かそのまんまprintf使えるけど

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

フロートの性能は欲しいがC2050は高すぎるということで。
99 名前：デフォルトの名無しさん mailto:sage [2010/12/11(土) 21:19:19 ]: C1060って260相当だっけ？
3DVantageとかの260と580の性能差を4倍したものを求めて見積もればいいんじゃない？
100 名前：デフォルトの名無しさん mailto:sage [2010/12/11(土) 22:44:36 ]: これでCPUがシングルコアだったら大笑い。
101 名前：デフォルトの名無しさん mailto:sage [2010/12/11(土) 23:25:14 ]: floatじゃなくてdouble茶羽化？
102 名前：デフォルトの名無しさん mailto:sage [2010/12/12(日) 00:01:33 ]: 性能比較とかどんなアプリ想定しているか書いていない時点でなんだかなあ

とりあえずsoftekでググると何か出てくるかもしれんよ
103 名前：デフォルトの名無しさん mailto:sage [2010/12/12(日) 00:58:43 ]: CPUはGulftwon/3.33GHz*2発で，メモリが残念な16GBです。
アプリ？はR+CUDAかMATLABで多体シミュレーションです。
けどRって言語だよな。
104 名前：デフォルトの名無しさん mailto:sage [2010/12/12(日) 02:22:33 ]: スレ違いな気がするけど、MATLAB 2010bのGPUArrayって使ってみましたか？
105 名前：デフォルトの名無しさん mailto:sage [2010/12/12(日) 13:53:06 ]: 3.0以上でfermi系ならカーネル内でprintf対応してるのか
誰かfermi系で128台ぐらいのGPUクラスター作ってくんねーかな
106 名前：デフォルトの名無しさん mailto:sage [2010/12/12(日) 13:57:08 ]: 128台でprintfするの？
107 名前：デフォルトの名無しさん mailto:sage [2010/12/12(日) 14:12:03 ]: >>105
ttp://tsubame.gsic.titech.ac.jp/ja/hardware-architecture
東工大にでも逝けや
108 名前：デフォルトの名無しさん mailto:sage [2010/12/12(日) 14:35:17 ]: >>103
その条件でfloatの性能求めるならたいして変わらないんじゃないかな
doubleなら別だけど
109 名前：デフォルトの名無しさん mailto:sage [2010/12/12(日) 14:47:34 ]: >>108
メモリの転送が非同期で双方向通信できるとか
コンカレントカーネル実行ができるとか、自動キャッシュがあるとか
のおかげで最適化（上記の機能を使うように修正）すればかなり早くなると思う。
修正なしでそのままならちょっと～けっこう早くなると思う。

でもR+CUDAを知らないので無責任な推測です。
110 名前：デフォルトの名無しさん mailto:sage [2010/12/12(日) 16:01:53 ]: 最適化しなければ、物凄い速くなるけど
最適化してしまえばそんなに性能変わらないんじゃね。
111 名前：デフォルトの名無しさん mailto:sage [2010/12/12(日) 16:45:01 ]: >>110
C1060に最適化したものをそのままFermiで動かしても差はそんなにないだろうって事だと思うけど
それでもコアレスの制限が緩和された影響と自動キャッシュのおかげで結構速くなると思う。

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

ところで103のMBはガルフ対応だからPCI-e2.0対応だよね？
112 名前：デフォルトの名無しさん mailto:sage [2010/12/12(日) 16:55:36 ]: ごめんストリームはFermi以前から対応してたみたいね。
113 名前：デフォルトの名無しさん mailto:sage [2010/12/12(日) 17:16:59 ]: ストリームは1.3だっけ? C1060って2.0だっけ?
114 名前：デフォルトの名無しさん mailto:sage [2010/12/12(日) 18:35:23 ]: >>113
ストリームはCC依存でtoolkit依存のようだが？
115 名前：デフォルトの名無しさん mailto:sage [2010/12/12(日) 18:36:43 ]: >>113
ストリームはCC依存ではなくtoolkit依存のようだが？
116 名前：デフォルトの名無しさん mailto:sage [2010/12/12(日) 20:48:29 ]: メモリバウンドならその話もわかるが
多体問題って演算バウンドだったんじゃないの確か
GTX580に交換するうまみはそれほど無い気がする
117 名前：デフォルトの名無しさん mailto:sage [2010/12/12(日) 21:13:38 ]: ということはtesla->fermiの性能向上ってメモリ関係（キャッシュ含む）とdoubleだけってこと？
floatなら計算パワーだけで考えると大して性能向上してなかったの？

floatだけならメモリ周りのボトルネックを解消して全体的にFlopsが上がりやすくなりましたって
だけだったのか・・・・
118 名前：デフォルトの名無しさん [2010/12/12(日) 21:34:38 ]: 考えるより動けということでGTX580x4枚の見積もり取ってます。
20万＋なら自分の財布でやってしまえ。
ちなみにMathematica8もCUDA対応で年始位にデリバリーとか何とか。

>>104
>MATLAB 2010bのGPUArray
R+CUDAの方に感けて，まだやってませんでした。すみません。
119 名前：デフォルトの名無しさん mailto:sage [2010/12/12(日) 21:43:37 ]: C1060と580のnbodyのベンチをさがしてみた

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

580は過去スレからc1060はググった結果
なにこれ！？
120 名前：デフォルトの名無しさん mailto:sage [2010/12/12(日) 21:46:09 ]: ごめん↑のC2050だった
121 名前：デフォルトの名無しさん mailto:sage [2010/12/12(日) 21:46:51 ]: > 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
122 名前：デフォルトの名無しさん mailto:sage [2010/12/12(日) 21:58:16 ]: >>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
123 名前：デフォルトの名無しさん [2010/12/12(日) 22:49:40 ]: >>118
自分でコード書いたほうが速いし確実だと思うけどね
124 名前：デフォルトの名無しさん mailto:sage [2010/12/12(日) 22:52:47 ]: 昔、GeForceをごにょごにょしたらQuadroになったように、
GeForceをごにょごにょしたらdoubleが4倍速になったりしないかな。
125 名前：デフォルトの名無しさん mailto:sage [2010/12/12(日) 22:54:13 ]: 最近のはならないようになっている
あきらめろん
126 名前：デフォルトの名無しさん mailto:sage [2010/12/12(日) 22:55:44 ]: >>123
行列演算で、CUBLASより速いコードかける？
127 名前：デフォルトの名無しさん mailto:sage [2010/12/12(日) 23:09:54 ]: 何でそんな不毛な煽りするかね
128 名前：デフォルトの名無しさん mailto:sage [2010/12/13(月) 22:52:09 ]: コンスタントメモリについて教えてください
__constant__ int i;
int j=1;とした場合(コンスタントメモリに配列以外を置きたい場合)になぜ
cudaMemcpyToSymbol("i",&j,sizeof(int),0,cudaMemcpyHostToDevice);
と文字列のようにして渡すのでしょうか

また、外部ファイルにextern __const__ int i;
と書くと定義が二重ですと言われます
複数ファイルにかかれたカーネルからコンスタントメモリを読み込みたいのですがどのようにすればよいのでしょうか
お願いいたします
129 名前：デフォルトの名無しさん mailto:sage [2010/12/13(月) 23:24:44 ]: >>128
うーん、簡単に言えば、cudaMalloc()でVRAM上のアドレスを得ることができればcudaMemcpy()できる。
一方__constant__の場合はそういうわけにはいかないので、
nvcc独自のちょっとえげつないcudaMemcopyToSymbol()を使わないといけない。
コンパイル単位の分かれたカーネル同士で共有するのは難しいかもねぇ。
130 名前：デフォルトの名無しさん mailto:sage [2010/12/13(月) 23:42:11 ]: >>118
> R+CUDAの方に感けて，まだやってませんでした。すみません。

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

R+gputools+CUDA+CULA

なのですか？

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

当方の環境は、

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

です。gotoBLASで行列演算時にコア4つ使用するところまではやったのですが、
CUDA使うともっと凄くなるのですか？
131 名前：デフォルトの名無しさん mailto:sage [2010/12/13(月) 23:45:44 ]: 速くなるかどうかはケースバイケースでアプリの内容によるだろ
132 名前：130 mailto:sage [2010/12/14(火) 23:05:18 ]: >>131
>>118さんがどういう構成でやっているのか知りたいです。
未だ、コンパイルも通らない実行も出来ていないレベルの者です。

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

CUDAのページのサンプルに
・MersenneTwister
・Monte Carlo Option Pricing with multi-GPU support
などとあるのをみると期待に胸がふくらみます。
133 名前：デフォルトの名無しさん [2010/12/15(水) 21:41:19 ]: 胸は膨らみブラ破け
134 名前：デフォルトの名無しさん mailto:sage [2010/12/15(水) 21:44:27 ]: だいたい乱数で時間かかるって相関乱数の相関部分じゃね～の？
BLASがいくら早くても関係ないじゃん、相関には
135 名前：デフォルトの名無しさん [2010/12/15(水) 21:54:23 ]: 乱数の作り方にもよるけど、誤り訂正符号を使った方式では符号化に
O(n^2)かかるところを32*8=256ビットの乱数を生成させるのに32個の
コアで16回の計算で得られるので効率的になる。
136 名前：デフォルトの名無しさん mailto:sage [2010/12/16(木) 00:26:21 ]: ３流プログラマがgotoさんの名前を出して汚すなよ。プライドが少しでもあるならな。
137 名前：デフォルトの名無しさん mailto:sage [2010/12/16(木) 00:29:55 ]: ・・・
138 名前：デフォルトの名無しさん mailto:sage [2010/12/16(木) 07:15:11 ]: 後藤blasだろ
釣られた？
139 名前：デフォルトの名無しさん mailto:sage [2010/12/16(木) 11:51:04 ]: 次買うのＡＴＩになりそ・・・
Compute Shaderスレないですか？
140 名前：デフォルトの名無しさん mailto:sage [2010/12/16(木) 15:56:27 ]: ATIやる人はOpenCLかCAL+ILじゃないの？
141 名前：デフォルトの名無しさん mailto:sage [2010/12/16(木) 17:50:04 ]: >>86

ものすごい遅いレスですみませんが、driverAPIなら可能です。
142 名前：デフォルトの名無しさん mailto:sage [2010/12/16(木) 19:24:21 ]: >>136
BLASが乱数を生成するライブラリでないと言ったら、
後藤さんの名前を汚していることになるのですか？
143 名前：デフォルトの名無しさん mailto:sage [2010/12/16(木) 23:05:19 ]: Textureメモリってふつうのより早いの？
みんななにに使ってるの？
144 名前：デフォルトの名無しさん mailto:sage [2010/12/17(金) 09:23:30 ]: ググレ
145 名前：デフォルトの名無しさん [2010/12/17(金) 22:12:47 ]: GTX580，4枚まとめ買い目指して奔走してますが，ない。
2月くらいにはなんとかならないだろか。
146 名前：デフォルトの名無しさん mailto:sage [2010/12/17(金) 22:51:22 ]: CUDA Visual Profiler と Compute Visual Profilerの違いは何でしょうか
前者はバージョン2.3
後者はバージョン3.1
が入っているのですが、toolkitが3.xになってからプロファイラの名称が変わったというだけでしょうか？
147 名前：デフォルトの名無しさん mailto:sage [2010/12/17(金) 23:04:45 ]: でしょ。
148 名前：デフォルトの名無しさん mailto:sage [2010/12/17(金) 23:18:01 ]: 質問です
ウッカリ長時間のCUDAジョブを走らせてしまったとき
中断はどすればよいですか？
149 名前：146 mailto:sage [2010/12/17(金) 23:34:33 ]: >>147
ありがとうございます
ものの本やサイトには/usr/local/cuda/cudaprof/のcudaprofiler使えって書いてあったのですが
変なWarning出て困っていたところでした
150 名前：146 mailto:sage [2010/12/17(金) 23:36:32 ]: >>147
ありがとうございます
ものの本やサイトには/usr/local/cuda/cudaprof/のcudaprofiler使えって書いてあったのですが
変なWarning出て困っていたところでした
151 名前：デフォルトの名無しさん mailto:sage [2010/12/18(土) 00:03:29 ]: >>148
GPUはずしちゃえよ
152 名前：デフォルトの名無しさん mailto:sage [2010/12/18(土) 00:34:36 ]: kill すればいいんじゃね？
153 名前：デフォルトの名無しさん mailto:sage [2010/12/18(土) 00:44:26 ]: >>145
九十九ネットストアにあったけど。あと価格コムとかで在庫もってる店あたってみたらいいんじゃない？
154 名前：デフォルトの名無しさん mailto:sage [2010/12/18(土) 00:47:36 ]: 補足、九十九にあったのはエルザのね。
155 名前：デフォルトの名無しさん mailto:sage [2010/12/18(土) 15:04:40 ]: Linuxドライバにバグないか？
最初の実行に時間がかかるのと
同様に最初の実行時にグリッド数、スレッド数が大きめだたとmemcpyでエラーが出る
小さめのグリッド数、スレッド数で最初に実行した後だと
大きめのグリッド数、スレッド数でも正常するんだが…
156 名前：デフォルトの名無しさん mailto:sage [2010/12/18(土) 15:16:14 ]: >>155
ドライバのバージョンとSDK、TOOLキット、OSの情報も載せないと駄目じゃね？
あとgccか。
157 名前：デフォルトの名無しさん mailto:sage [2010/12/18(土) 15:28:07 ]: OS：CentOS5.5
ドライバ、SDK、ToolKitは最新
ｇｃｃは4.1.2
今のところ原因は全く不明
ドライバのバージョンを一個前にすると初回実行時の遅延がちょっと早くなる（最新が約2.5s、一個前が約1.8ｓ）
gtx470と465で試してどちらも同じ症状
158 名前：130 mailto:sage [2010/12/18(土) 16:44:30 ]: ようやくgputoolsが動いた。

速いのは分かったけど、
行列の規模が、本体のメモリ容量と関係なく、ビデオカードのメモリ容量に制約される。
キャッシュイーター向けじゃなかったですね。
159 名前：デフォルトの名無しさん mailto:sage [2010/12/18(土) 17:43:36 ]: >>158
行列の計算ってのはね、分割できるんですよ。
160 名前：デフォルトの名無しさん mailto:sage [2010/12/18(土) 18:13:31 ]: >>159
確かに。
分割できなければコアが複数あっても意味がない。
161 名前：デフォルトの名無しさん mailto:sage [2010/12/18(土) 18:55:39 ]: でもね、いったんメインメモリに移して計算するのはもったいない。

結論C2070を買え。
162 名前：130 mailto:sage [2010/12/19(日) 18:30:10 ]: 結局、部分行列作成と計算結果の結合で、処理にかかる時間が逆転してしまいました。

一般論として言えるのかどうか知らないけど、
行列が（計算結果を含めて）ビデオカードのメモリにスッポリ入るなら劇的な高速化が見込めるけど、
そうでなければ、マルチコアCPU＋本体メモリのほうが速度でもメモリ消費量でも有利なようです。
163 名前：128 mailto:sage [2010/12/19(日) 20:32:38 ]: コンスタントメモリはいつデバイスに確保されるのでしょうか
__const__宣言時とは考えにくいので最初にcudaMemcpyToSymbolしたときでしょうか

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

コンスタントメモリはどのように解放すればいいのでしょうか。
164 名前：デフォルトの名無しさん mailto:sage [2010/12/19(日) 21:03:47 ]: 行列計算なら転送を隠せるはず
リンパックベンチでGPU搭載マシーンが上位に来てるんだし
165 名前：デフォルトの名無しさん mailto:sage [2010/12/19(日) 23:49:30 ]: >>163
SDKのサンプルみればいいんじゃないか？
166 名前：デフォルトの名無しさん mailto:sage [2010/12/20(月) 00:43:41 ]: __constant__を気にするなら
kernelコードだって似たようなものなんじゃ。
167 名前：デフォルトの名無しさん mailto:sage [2010/12/20(月) 11:19:24 ]: >>163
定数メモリはcudaMalloc()のように「確保する」ということなしに使うようになっています。
その為、cudaMemcpyToSymbol()のような特殊な関数でしか転送できません。
また、「確保する」ことなしに使える以上、「解放する」手続きは用意されていません。
168 名前：デフォルトの名無しさん mailto:sage [2010/12/21(火) 15:33:13 ]: 安定ソートが欲しいのですが
CUDA向きの安定ソートってありますか？
16000要素くらいです
169 名前：デフォルトの名無しさん mailto:sage [2010/12/21(火) 15:52:44 ]: NVIDIAにあるバイトニック使ってみ
170 名前：デフォルトの名無しさん mailto:sage [2010/12/21(火) 15:54:21 ]: SDKにmergesortってあるがそれが同じ奴だ
171 名前：168 mailto:sage [2010/12/21(火) 16:18:22 ]: バイトニックは安定じゃないようなのですが、
マージソートのサンプルがSDKにあったのですね。
チェックしてみます
172 名前：デフォルトの名無しさん [2010/12/21(火) 20:58:05 ]: gtx580を4基積んだ．電源が音を上げた．
やれやれ．
173 名前：デフォルトの名無しさん mailto:sage [2010/12/21(火) 21:12:11 ]: >>172
何Wで駄目だった？
174 名前：デフォルトの名無しさん mailto:sage [2010/12/21(火) 21:46:21 ]: >>173
1400W。ディスク系を外においてもなぁ。
175 名前：デフォルトの名無しさん mailto:sage [2010/12/21(火) 21:49:20 ]: SC748あたりなら・・・
176 名前：デフォルトの名無しさん mailto:sage [2010/12/21(火) 21:59:47 ]: 7046GT-TRFね。
177 名前：デフォルトの名無しさん mailto:sage [2010/12/21(火) 22:17:39 ]: 1400Wで駄目だとなると電源2つ積めるケースにして連動させるケーブルで繋いでと
いったことをしないと無理そうだな。
178 名前：デフォルトの名無しさん mailto:sage [2010/12/21(火) 23:19:45 ]: 無茶しやがって…
179 名前：デフォルトの名無しさん mailto:sage [2010/12/22(水) 01:04:00 ]: 1400Wとか言って売っている時点で粗悪電源確定。
>>172が日本在住でないなら話は別だが。
180 名前：デフォルトの名無しさん mailto:sage [2010/12/22(水) 07:14:57 ]: >>179
あれは110V/1400W。80GOLDのリタンダント。
日本じゃ実効1100W。
181 名前：デフォルトの名無しさん mailto:sage [2010/12/22(水) 22:34:50 ]: 大容量PC電源はエアコン引っこ抜いて200Vから取ろう。
182 名前：デフォルトの名無しさん mailto:sage [2010/12/22(水) 23:25:26 ]: >>180
なら、「1100W電源」として売るべきだろ。

排気量2000ccの車を3000ccだと言って売る自動車屋などありえない。
そのあり得ないことを長年やってきたのが、PC電源代理業者だ。
183 名前：デフォルトの名無しさん mailto:sage [2010/12/23(木) 00:59:59 ]: >>182
supermicroは日本向けのサイトを持っていない。仕方あるまい。
ちなみに法人で入れたときはこの説明はあった。
www.supermicro.com/products/system/4U/7046/SYS-7046GT-TRF.cfm
184 名前：デフォルトの名無しさん mailto:sage [2010/12/24(金) 20:02:02 ]: GPUの分類について教えてください
例えばGTX480は
Fermi-ComputeCapability2.0-GF100-GTX480
でよいのでしょうか
それともアーキテクチャとCCは1対1対応しているのでしょうか
また、Fermiの分類位置はここでよいのでしょうか
185 名前：デフォルトの名無しさん mailto:sage [2010/12/24(金) 22:08:13 ]: ComputeCapability2.0-GF100-GTX480
は正しいけど、FermiはC.C.2.X全部を含む
GFの"F"はFermi、GTXの"T"はなぜかTeslaだったはず
186 名前：デフォルトの名無しさん mailto:sage [2010/12/24(金) 22:20:21 ]: てめぇGeForce 7800GTXさんに謝れよ
187 名前：デフォルトの名無しさん mailto:sage [2010/12/24(金) 23:41:29 ]: >>184
CUDA目的ならComputeCapabilityだけ意識すればよいかと。
あとはパフォーマンスが違うだけだから。
188 名前：デフォルトの名無しさん mailto:sage [2010/12/24(金) 23:43:40 ]: >>186
GeForce 7800GTXさんはCUDA対象外じゃ・・・・・・
189 名前：デフォルトの名無しさん mailto:sage [2010/12/25(土) 00:54:10 ]: コンパイルされたCUDAプログラムを実行するだけの場合、ドライバのみでいけますか？
toolktのインストールは必須ですか？
LD_LIBRARY_PATH /usr/local/cuda/lib64
を通さなくてはならないということはtoolkit必須なのかな
190 名前：デフォルトの名無しさん mailto:sage [2010/12/25(土) 01:00:31 ]: ライブラリ使っててもstaticならいらないはずだし、dllつかってるなら
そのdllをexeと同じフォルダにおけばいいんじゃないか？
きっとlib64内のなんかのライブラリつかってんでしょ。cutilとかかな
191 名前：デフォルトの名無しさん mailto:sage [2010/12/25(土) 02:39:52 ]: DLL再配布して良かったっけ
192 名前：デフォルトの名無しさん mailto:sage [2010/12/25(土) 09:47:13 ]: Linuxならドライバ（カーネルモジュールを含む）も必要なので再配布だけでは解決しない。
もちろんツールキットのライブラリも必要。
193 名前：デフォルトの名無しさん [2010/12/27(月) 19:25:09 ]: すみません、教えてください。
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です。
194 名前：デフォルトの名無しさん mailto:sage [2010/12/27(月) 19:40:29 ]: 環境設定は？
195 名前：デフォルトの名無しさん mailto:sage [2010/12/27(月) 20:29:33 ]: それはCentOS 4.5のようだが。
196 名前：デフォルトの名無しさん mailto:sage [2010/12/27(月) 22:11:09 ]: >>193
例えば、ディストリビューションのnvidiaドライバなんかが残っていない？
その上に、最新ドライバをインストールしていて、
コンパイル時と実行時に別の共有ライブラリを参照していたりとか。

ていうか、自分も、ubuntuでx-updatesのドライバと最新ドライバが両方入っていて挙動不審。
ubuntuでもOS再インストールかｗ
197 名前：193 [2010/12/28(火) 00:14:59 ]: >>196

返信ありがとうございます。
ちょっとその辺を調べてみます。
後日、問題が解決したらフィードバックします。
198 名前：デフォルトの名無しさん mailto:sage [2010/12/28(火) 07:48:23 ]: スルーされた。勝手にしてくれ。
199 名前：デフォルトの名無しさん mailto:sage [2010/12/28(火) 16:15:18 ]: いじけるなよｗ
200 名前：デフォルトの名無しさん mailto:sage [2010/12/28(火) 22:47:32 ]: とある処理のスレッド数と時間を関係を調べているのですが、
スレッド数が192を超えたあたりで処理が行われていません。

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

グラボは9800GTで11レジスタ、共有メモリを36byte使用してます。
201 名前：デフォルトの名無しさん mailto:sage [2010/12/28(火) 22:48:39 ]: >>200
「処理が行なわれない」ってどういう意味さ。
202 名前：デフォルトの名無しさん mailto:sage [2010/12/28(火) 23:01:14 ]: >>201
説明不足でした。

処理を1000回行い、平均処理時間を計測しているのですが、
スレッド数が192を超えると0(ms)になってしまいます。
203 名前：デフォルトの名無しさん mailto:sage [2010/12/28(火) 23:19:14 ]: 何かが間違っているんだろうね。
204 名前：デフォルトの名無しさん mailto:sage [2010/12/29(水) 00:34:35 ]: まずはエラー確認してからだろ。
205 名前：デフォルトの名無しさん mailto:sage [2010/12/29(水) 00:38:37 ]: スレッド数を増やしたことによって確保されていない領域にアクセスするような
エラーが発生するような状況になっていないかどうかを確認・・・・・かな。
これが起きているとそれ以降の（タイマーを含めた）全ての処理がエラーになる。
206 名前：デフォルトの名無しさん mailto:sage [2010/12/29(水) 01:31:39 ]: まずは
cudaThreadSynchronize();
cudaGetLastError();
207 名前：デフォルトの名無しさん mailto:sage [2010/12/29(水) 14:09:27 ]: >>203-206
助言ありがとうございます。

まずはエラーを確認してみたいと思います。
208 名前：デフォルトの名無しさん mailto:sage [2010/12/29(水) 19:33:12 ]: たぶんリソース不足でしょ。自分もよくそうなる。
カーネルの引数を構造体のポインタにまとめると解決する。
209 名前：デフォルトの名無しさん [2010/12/30(木) 20:27:57 ]: GTX460ってCUDAは使えるんですかね？
lspci | grep -i nVidia
でPCIバスデバイス情報確認したら
VGA compatible controller: nVidia Corporation: Unknown device …
ってコンピューターさんがおっしゃられたんだが…
210 名前：デフォルトの名無しさん mailto:sage [2010/12/31(金) 00:21:26 ]: >>209
/sbin/update-pciids
してからコンピュータさんにもう一度お伺いたててご覧なさい
211 名前：デフォルトの名無しさん mailto:sage [2010/12/31(金) 05:27:05 ]: ドライバをインストールしてないんじゃないの？
212 名前：デフォルトの名無しさん mailto:sage [2011/01/01(土) 11:59:19 ]: 正月セールで狙う1万5000円くらいのグラボでオススメのものをおせーて
213 名前：デフォルトの名無しさん mailto:sage [2011/01/04(火) 09:25:33 ]: GeforceとRADEONを2枚挿して
GeforceをCUDA専用、RADEONをグラフィック専用とすることは出来ますか？
214 名前：デフォルトの名無しさん [2011/01/04(火) 09:34:04 ]: 日本語的に複数の意味に取れるので訂正

GeforceとRADEONの2枚を挿して
GeforceをCUDA専用、RADEONをグラフィック専用とすることは出来ますか？
215 名前：デフォルトの名無しさん mailto:sage [2011/01/04(火) 09:41:44 ]: no problem
216 名前：デフォルトの名無しさん mailto:sage [2011/01/04(火) 10:36:53 ]: ありがとうございます
217 名前：デフォルトの名無しさん mailto:sage [2011/01/04(火) 10:52:31 ]: cudaSetDeviceFlags(cudaDeviceBlockingSync);
218 名前：デフォルトの名無しさん mailto:sage [2011/01/04(火) 10:53:37 ]: ほとんど同じソースで
cudaSetDeviceFlags(cudaDeviceBlockingSync);
が有効になるプログラムとならないプログラムがあるのですがどこを
調べれば良いでしょうか。
219 名前：デフォルトの名無しさん mailto:sage [2011/01/04(火) 12:31:38 ]: ほとんど同じソースならどこが違うのかを書いてください。
そして「有効」かどうかをCPU負荷で判定しているのかどうかも。

cudaSetDevice()を除くすべてのCUDA関数より先にcudaSetDeviceFlags()を
呼び出しているかどうか以外には特に調べるべきところはありません。
220 名前：デフォルトの名無しさん [2011/01/05(水) 18:57:39 ]: マニュアル的なpdfが沢山あってどれを読めばいいのかわからないのですがどれから読むのがいいのでしょうか？
CUDA_C_Best_Practices_Guide.pdf
CUDA_C_Programming_Guide.pdf
Introductory CUDA Technical Training Courses VolumeI.pdf
221 名前：デフォルトの名無しさん mailto:sage [2011/01/05(水) 20:31:24 ]: 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 名前：デフォルトの名無しさん mailto:sage [2011/01/05(水) 21:02:09 ]: フランス語のはないですかね？
223 名前：デフォルトの名無しさん mailto:sage [2011/01/05(水) 21:15:20 ]: ワロタ
224 名前：デフォルトの名無しさん mailto:sage [2011/01/05(水) 21:16:37 ]: >>222
残念ながらありません
225 名前：デフォルトの名無しさん mailto:sage [2011/01/06(木) 00:39:25 ]: cc2.0でコンパイルすると、cc1.3でコンパイルした時よりレジスタ使用量が
10くらい増えるのですが、なぜですか？
226 名前：デフォルトの名無しさん mailto:sage [2011/01/06(木) 22:50:40 ]: なんでだろうねー
227 名前：デフォルトの名無しさん mailto:sage [2011/01/08(土) 18:48:55 ]: Fermiで各スレッドで数百バイトのワーキングエリアが必要な場合、
ローカルメモリを使うのが最良なのでしょうか？
シェアードメモリを使うと、同時に立ち上げられるスレッド数が減って
逆に遅くなりました。
228 名前：デフォルトの名無しさん mailto:sage [2011/01/09(日) 21:38:49 ]: >>225
そういうのはプログラムによって変わるだろ
とりあえずptxファイル嫁

>>227
>各スレッドで数百バイトのワーキングエリアが必要な場合
これが今以上減らせないならそもそもCUDAに向いてない
229 名前：デフォルトの名無しさん mailto:sage [2011/01/10(月) 17:57:08 ]: グローバルメモリのアクセスは400~500clock、レジスタやシェアードメモリ、L1キャッシュはほぼ0らしいですが
L2キャッシュは大体何clockかかりますか？
230 名前：デフォルトの名無しさん mailto:sage [2011/01/10(月) 19:32:12 ]: L1は大体15
L2は大体150
231 名前：デフォルトの名無しさん mailto:sage [2011/01/11(火) 09:01:34 ]: 4枚発注ー。
232 名前：デフォルトの名無しさん [2011/01/13(木) 11:27:07 ]: CPU向けコードとGPU向けコードを同一ファイルで作成したいと思っています。
nvccはcuでないと、cudaコードとしてコンパイルしてくれないのですが、
拡張子がcppのファイルをcudaソースファイルとしてコンパイルするオプションはありませんか?
233 名前：デフォルトの名無しさん mailto:sage [2011/01/13(木) 19:39:51 ]: obj同士はnvccでリンクできるから、それぞれコンパイルすればＯＫだよ
234 名前：デフォルトの名無しさん mailto:sage [2011/01/13(木) 20:42:06 ]: >>141
亀レスですが，Driver APIを用いて出来ました．
ありがとうございました．
235 名前：デフォルトの名無しさん mailto:sage [2011/01/14(金) 20:47:45 ]: >>234
概略をのせてくれまいか？

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

くらいの事でもいいから
236 名前：234 mailto:sage [2011/01/15(土) 13:34:59 ]: >>235
簡単にで良いなら．

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

　　mallocとか色々

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

glutMainLoop();

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

　　カーネル実行

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

関数の引数は省略しました．
237 名前：デフォルトの名無しさん mailto:sage [2011/01/15(土) 17:57:20 ]: >>236
ありがとう。参考にさせてもらいます
238 名前：デフォルトの名無しさん mailto:sage [2011/01/15(土) 21:30:19 ]: Fermiもグローバルメモリアクセスはコアレッシングにしたほうがよくって
キャッシュはおまけ程度って考えでおk？
ていうか各スレッドが違うグローバルメモリにアクセスするなら同じカーネル内ではキャッシュ効かないよね
239 名前：デフォルトの名無しさん mailto:sage [2011/01/15(土) 22:14:42 ]: CPUのマルチコア（巨大キャッシュ、低並列）のマルチスレッドですら
キャッシュ汚染はひどいからな

マルチスレッドとキャッシュは相性が非常に悪い
240 名前：デフォルトの名無しさん mailto:sage [2011/01/15(土) 23:23:54 ]: スレッド番号とグロバールメモリーのアドレスが並んでいないと使い物にならないな。
速度が１０倍は違う。本当に使いにくいよ。
メモリアクセスがネックの（単純な）アプリで、shared memoryの使いまわしのできないアプリでは、
マルチスレッドを上回るのは難しい。
241 名前：デフォルトの名無しさん mailto:sage [2011/01/15(土) 23:49:58 ]: しかし、いままでGPUと相性の悪いといわれるプログラムでも
FermiならCPUの2倍くらい出るってことは結構ある。

非Fermiならローカルメモリ使ったら即使い物にならない速度だったけど、
Fermiならキャッシュがなんとかしてくれたり。
242 名前：デフォルトの名無しさん mailto:sage [2011/01/16(日) 18:58:37 ]: cublasDgemmにおいて、係数のalphaを(double)1に、
betaを(double)0に指定しても、
内部では真面目に係数をかけているのでしょうか？

参考(PDF注意)
developer.download.nvidia.com/compute/cuda/3_2_prod/toolkit/docs/CUBLAS_Library.pdf
243 名前：デフォルトの名無しさん mailto:sage [2011/01/16(日) 19:01:42 ]: FORTRANのソースを自分で追っかけて確認しろ
244 名前：デフォルトの名無しさん mailto:sage [2011/01/16(日) 19:13:17 ]: 速度を測ればすぐに分かりそう
245 名前：デフォルトの名無しさん mailto:sage [2011/01/16(日) 21:42:18 ]: >>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のコードのような実装になっていない気がするのですが。
246 名前：デフォルトの名無しさん mailto:sage [2011/01/16(日) 21:45:25 ]: 「はじめてのCUDAプログラミング」に出てるコード（もしくは同じ出版社からインターネット上に出てる奴）
をひたすらマイナーチェンジしていくのが一番早いと思う。
スクラッチはマイナーチェンジの経験を積んでからで十分。
247 名前：デフォルトの名無しさん mailto:sage [2011/01/17(月) 05:28:45 ]: PTXって機械語なの？
GPUがPTX読むみたいな事が書いてあるんだけど
248 名前：デフォルトの名無しさん mailto:sage [2011/01/17(月) 06:02:22 ]: ptxはgpuのアセンブラ言語じゃないの？？
249 名前：デフォルトの名無しさん mailto:sage [2011/01/17(月) 07:27:52 ]: PTXはあくまで中間言語。
250 名前：デフォルトの名無しさん mailto:sage [2011/01/17(月) 10:47:31 ]: TMPGEncで使おうと思ったらG80非対応なんだな。。
251 名前：デフォルトの名無しさん mailto:sage [2011/01/17(月) 20:15:09 ]: GTX580を4枚積んでサンプルコード走らせたら電力不足来たー。
252 名前：デフォルトの名無しさん [2011/01/19(水) 09:33:50 ]: CUDAツールキットを導入する際にまずやることって何？
どう頑張ってもnvccでコンパイルしたサンプルファイルが実行できませーん
って開発以前の問題で精神折れそうだわ…
253 名前：デフォルトの名無しさん [2011/01/19(水) 13:17:05 ]: 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 ); /* 先頭の次からリストの表示 */
}
}
254 名前：デフォルトの名無しさん mailto:sage [2011/01/19(水) 13:59:28 ]: >>252
実行できないならできないなりにエラーメッセージが帰ってくるのではないでしょうか

CUDAコアとはスカラプロセッサのことでいいのでしょうか
FermiのみCUDAコアという呼び方をするのですか？
255 名前：デフォルトの名無しさん mailto:sage [2011/01/19(水) 14:09:09 ]: >>252
実行できないならできないなりにエラーメッセージが帰ってくるのではないでしょうか

CUDAコアとはスカラプロセッサのことでいいのでしょうか
FermiのみCUDAコアという呼び方をするのですか？
256 名前：デフォルトの名無しさん mailto:sage [2011/01/19(水) 21:03:35 ]: Fermiで倍精度浮動小数点の三角関数計算に
SFUを使えないのでしょうか

倍精度浮動小数点の計算は単精度浮動小数点演算機が2クロックで計算するのですよね
同じようにSFUが単精度計算分の2週でやってくれるということはないのでしょうか
257 名前：デフォルトの名無しさん mailto:sage [2011/01/19(水) 21:10:31 ]: geforceは駄目だろ
258 名前：デフォルトの名無しさん mailto:sage [2011/01/19(水) 21:57:03 ]: ドライバに　266.58 WHQL　が来たね
259 名前：デフォルトの名無しさん mailto:sage [2011/01/19(水) 22:20:50 ]: >>256
fast mathを使えば？
260 名前：デフォルトの名無しさん mailto:sage [2011/01/20(木) 00:39:31 ]: fast mathは単精度専用。

そもそも倍精度使ってるくせに精度低いSFU使いたいって時点で
本当に倍精度が必要か考えるべき。
261 名前：256 mailto:sage [2011/01/20(木) 02:09:04 ]: SFU遊ばしておくのはもったいないかな、と思いまして

そんなに精度が低いのかー
262 名前：デフォルトの名無しさん mailto:sage [2011/01/20(木) 02:55:13 ]: 応用次第だけど、SFU使うと駄目で使わなければfloatでもOKという例があった。
263 名前：デフォルトの名無しさん [2011/01/21(金) 02:52:36 ]: 初歩的なこと聞いてすみません。
元々あるソースにCUDAを取り入れてるんですが構造体を丸々GPU側にコピーはできないんでしょうか？
できるのであれば誰かご教授お願いします。
264 名前：デフォルトの名無しさん mailto:sage [2011/01/21(金) 03:12:27 ]: C言語の勉強をしたらできると思います。
265 名前：デフォルトの名無しさん [2011/01/21(金) 03:30:01 ]: >>264
一応いろいろ試したんですが
Warning : Cannot tell what pointer points to, assuming global memory space
ていう警告が出ていておそらくデータの転送ができていないんです・・・。
普通のC言語を勉強すればできるんですかね・・・？
266 名前：デフォルトの名無しさん mailto:sage [2011/01/21(金) 03:40:27 ]: 構造体の中にポインタが入ってて、そのポインタがメインメモリを指してるんじゃないの？
267 名前：デフォルトの名無しさん [2011/01/21(金) 04:09:34 ]: >>266
ポインタのほかにもポインタのポインタも入っているのでそれもGPU側に送らないといけないいうことであってますか？
そうすると結構量があるので必要な部分のみを送るほうが効率がいいんですかね・・・？
268 名前：デフォルトの名無しさん mailto:sage [2011/01/21(金) 05:42:58 ]: もちろんポインタが指してるデータは送らないといけないし，
構造体を送る前に，構造体の中で指してるアドレスをGPU側のアドレスに
書き変えないといけない．
269 名前：デフォルトの名無しさん [2011/01/21(金) 06:51:09 ]: 解答ありがとうございます。
なんとなくですが皆さんのおかげで理解できました。
がんばってやってみます。
270 名前：デフォルトの名無しさん mailto:sage [2011/01/21(金) 07:06:46 ]: なんとなくじゃ無理だろうな
271 名前：デフォルトの名無しさん [2011/01/24(月) 01:53:29 ]: 少し聞きたいんですが、
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
いきなりこんなエラーが出たんですがどういう意味ですか？
解決法を知っている方教えてください！！
272 名前：デフォルトの名無しさん [2011/01/24(月) 08:07:36 ]: en.wikipedia.org/wiki/Comparison_of_Nvidia_graphics_processing_units
誰かこのページでGPUの浮動小数点計算速度の比較に使われてるプログラム知ってる人いませんか？
273 名前：デフォルトの名無しさん mailto:sage [2011/01/24(月) 10:18:04 ]: >>272
ざっと見ただけではこのページのどこに浮動小数点計算速度の比較が書いてあるのか分からなかったのだが。

いずれにしてもそこに載っているスペックは全て理論ピーク性能かと。
274 名前：デフォルトの名無しさん mailto:sage [2011/01/24(月) 10:27:17 ]: >>271
いきなりすぎて難しいのでとりあえず中間ファイルの.ptxファイルを保存して146行目を調べて。

あとはSDKに含まれるプログラムとかは正常にコンパイルできるのかどうかも。
275 名前：デフォルトの名無しさん mailto:sage [2011/01/24(月) 13:59:18 ]: お聞きしたいのですが、プログラムを実行したときに発生するエラーで
invalide device function とは何が原因で起こるエラーなんでしょうか？
初歩的な質問ですがどなたか教えてください。
276 名前：デフォルトの名無しさん mailto:sage [2011/01/24(月) 14:38:37 ]: invalide device function
device で実行しているfunctionがinvalideなんだろ
でもinvalideじゃなくてinvalidじゃね？
277 名前：275 mailto:sage [2011/01/24(月) 15:13:26 ]: スペルを間違えていました。invalideではなくinvalidでした。
invalidは「無効な」という意味で捉えています。
デバイス(__global__)の関数名とホストで実行する関数名は同じにしているのですが、なぜ無効と言われるのかがわからないのです。当然ですが、何かの関数名と被らないようには名前をつけています。
どなたかお願いします。
278 名前：デフォルトの名無しさん mailto:sage [2011/01/24(月) 15:38:25 ]: もう少し具体的に書け
279 名前：デフォルトの名無しさん mailto:sage [2011/01/24(月) 15:56:17 ]: _device_な関数を_global_と_device_な関数以外から呼び出してる予感
280 名前：275 mailto:sage [2011/01/24(月) 16:12:51 ]: 詳細を書きました。長文です。すみません。
画像処理をCUDAでやっており、今は、メディアンフィルタを作っています。
OpenCVを使いたかったので、ここの方法（ttp://blogs.yahoo.co.jp/simasaki2001/18576477.html）を参考に簡単なサンプルを先生が作成し、それを基に学生の自分がプログラムを作っています。

現在、__global__ void cuda_median() というデバイスの関数をcuda_median<<<ndrange_global, ndrange_local>>>(入力画像, 入力画像のポインタ, 出力画像, 出力画像のポインタ)と呼び出しています。
ndrange_globalとndrange_localはDim3型で
ndrange_local.x= atoi( argv[3] );
ndrange_local.y= 1;
ndrange_local.z= 1;
ndrange_global.x= 1;
ndrange_global.y= (int)dev.multiProcessorCount; <-これは14となっています
ndrange_global.z= 1;
となっています。この部分についてよく理解できなかったので先生に質問しましたが、「人の作ったのだからよく知らない。」と言われました。現在もしっかり理解しているとは言えません。

__device__関数も一つ作成して使っていますが、__global__からしか使用していません。
281 名前：デフォルトの名無しさん mailto:sage [2011/01/24(月) 16:21:56 ]: いつ、どこでのエラーなんだ？
device内のfunctionでのエラーではないのか？
282 名前：275 mailto:sage [2011/01/24(月) 16:37:21 ]: ホストでdevice関数を呼んだ後にcudaGetLastErrorを実行しており、そこでエラーが出ます。
device関数を呼ぶ前の処理におけるエラーチェックではエラーは発生していないので、device関数でエラーが起こっているのではないかと考えています。
283 名前：デフォルトの名無しさん mailto:sage [2011/01/24(月) 16:57:03 ]: 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のほうがいいということなのでしょうか
284 名前：デフォルトの名無しさん mailto:sage [2011/01/24(月) 17:01:06 ]: ＞それともただ単に1コアのクロック周波数から理論的な演算性能を算出してコア数分かけた値で、
＞GPUはコアが多いからそう見えているのでしょうか

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

それで合ってる
285 名前：デフォルトの名無しさん mailto:sage [2011/01/24(月) 17:21:26 ]: >>283
まあ、簡単に言うと計算の仕方とメモリアクセスの仕方について、GPUの方が制約がきついってこと
もっとも、CPUでも高い性能を出すためには厳しい制約が付くけど
286 名前：283 mailto:sage [2011/01/24(月) 18:18:42 ]: >>284-285
ありがとうございます。

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

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

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

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

ということでしょうか
287 名前：デフォルトの名無しさん mailto:sage [2011/01/24(月) 19:19:41 ]: >>280
ついこの前出たOpenCV 2.2でNPP経由でCUDA対応したから、それ使ってみたら？
exeではなくCMakeから自前ビルドしないといけないけど
288 名前：デフォルトの名無しさん [2011/01/25(火) 10:09:24 ]: >>274
SDK内のプログラムはきちんと動作しています。
.ptxファイルは保存方法がぐぐっても良くわからなくて・・・。
どうやったら保存できますか？
289 名前：アク禁長いから久々の書き込み mailto:sage [2011/01/25(火) 12:05:07 ]: >>282
状況がよく判らんが、__device__な関数は__device__か__global__な関数からしか呼んじゃダメよ。
それと、<<<>>>で呼んだ関数で起きたエラーは、次にAPIを呼んだときに戻されるのでご用心。
つまり、以下の場合cudaThreadSynchronize()の戻り値がエラーになる。
--
func<<<grid, block>>>();
cudaThreadSynchronize();

>>271
取り敢えず、-ptxつけてコンパイルして味噌。
290 名前：デフォルトの名無しさん [2011/01/25(火) 12:51:53 ]: >>289
コマンドプロンプトでnvcc -ptx cuda.cuであってますか？
普段VSでコンパイルしてるんでやり方がいまいちなんですが・・・。
291 名前：283 mailto:sage [2011/01/25(火) 13:02:29 ]: Fermiでコンスタントキャッシュはハードウェア上ではどこにあるのでしょうか

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

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

そもそも大容量なはずのVRAMに転送するコンスタントメモリが
64kBと小容量に決められているという事はキャッシュも同サイズ用意されていて
ヒット率100％が保証されているということなのでしょうか
292 名前：デフォルトの名無しさん mailto:sage [2011/01/28(金) 08:08:16 ]: ちょっと質問なんですけど、GPU上ではstaticは使えないみたいですけど、
staticの代わりになるようなものはCUDAに存在するんでしょうか？
やっぱりグローバルメモリに領域を確保して使うしかないんでしょうか？
293 名前：デフォルトの名無しさん [2011/01/29(土) 02:47:23 ]: >>292
すまん質問の意味がわからん
294 名前：デフォルトの名無しさん mailto:sage [2011/01/29(土) 21:17:57 ]: >>293
すいません、ちょっと日本語がおかしいですね・・・
要するに、ＧＰＵ上のdevice関数内で、
関数が終わっても変数の値が残るようにしたいんですけど、
それをするにはcudaMemcpyでグローバルメモリに領域確保して、
そこに値を保存するようにするしかないのでしょうか？
295 名前：デフォルトの名無しさん mailto:sage [2011/01/29(土) 21:38:16 ]: それどういうときに使うの？カウンタ？

そもそもインライン展開されるんじゃなかったっけ
296 名前：デフォルトの名無しさん mailto:sage [2011/01/29(土) 23:19:04 ]: 根本的に取り組む姿勢が間違っている気がする。
297 名前：デフォルトの名無しさん mailto:sage [2011/02/01(火) 19:48:05 ]: 構造体の配列って
a,b,c|a,b,c|a,b,c
こう置かれるからハーフワープの各スレッドがaにアクセスすると
コアレスアクセスできないということですか？
配列の構造体.a[n],.b[n],.c[n]を作った方が良いという事でしょうか
298 名前：デフォルトの名無しさん mailto:sage [2011/02/01(火) 19:56:41 ]: 誰だかまるわかりw
299 名前：デフォルトの名無しさん mailto:sage [2011/02/02(水) 19:08:00 ]: 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数が多いほうがいいのか、
あるいはどちらも同じなのか教えてください。
300 名前：デフォルトの名無しさん mailto:sage [2011/02/02(水) 19:38:03 ]: >>299
計測してパフォーマンスのいい方にすればいいとおもうが、
１の方があるスレッドブロックが同期待ちになってるあいだに他のスレッドブロックが
実行可能な余地があるぶん性能低下はしにくいと思う
301 名前：デフォルトの名無しさん mailto:sage [2011/02/02(水) 20:23:06 ]: >>300
回答ありです。
言葉足らずでした。。。
計測して実験したのですが、1のほうがパフォーマンスが上でした。

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

スレッドブロックが同期待ちというのは、syncした場合におこる現象でしょうか？
調べてみましたがやはりSMあたりのスレッド数が同じなら
SMあたりのスレッドブロック数が多いほうがいいというのが
いまいちわかりません。。。
よろしくお願いします
302 名前：デフォルトの名無しさん mailto:sage [2011/02/02(水) 20:35:42 ]: >>301
__syncthreads();
はブロック内のスレッドを同期する。

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

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

＞SMあたりのスレッドブロック数が多いほうがいい
スレッドブロックはSM単位に割り当てられるから「SMあたりのスレッドブロック数が多いほうがいい」
というのはActive Thread Blocks per Multiprocessorが高いほうがいいというのと同義
303 名前：デフォルトの名無しさん mailto:sage [2011/02/02(水) 20:53:25 ]: Visual Profilerでlocal load localとstoreが数万回起こっているのですが何が原因でしょうか
--ptxas-options=-vでlmemが出ていなかったので
ローカルメモリは使っていないと思っていたのですが
local～が0でないカーネル(3種類)はすべてレジスタ/スレッドが63でした。
偶然同じ数になっていると思ってましたが
レジスタは63が上限でそれ以降はローカルメモリが使われるってこと？
304 名前：デフォルトの名無しさん mailto:sage [2011/02/02(水) 23:02:21 ]: >>303
レジスタの容量を調べてみろ
例えばレジスタの本数が32768で1ブロックあたり512スレッド使うなら
1スレッドが使用できるレジスタは64以下になる、256スレッドなら128以下

ローカルメモリが使われるなんて状況は最悪だからそのプログラムは糞だ
シェアードメモリ余ってんならそっちも有効に使え
305 名前：デフォルトの名無しさん mailto:sage [2011/02/02(水) 23:30:24 ]: だがローカルメモリ使ったカーネルとCPU実行では，
CPU実行の方が遅いこともあるから必ずしも悪ではない．
306 名前：デフォルトの名無しさん mailto:sage [2011/02/03(木) 17:37:00 ]: 総和計算を何段階かにわけたときって
sharedメモリ使う都合上、最後の段階は1ブロックでやるじゃん
でも1ブロックのみってことは他のsm使わないから無駄してるよね
最後の段階はcpuに転送したほうがはやいってこと？
307 名前：デフォルトの名無しさん mailto:sage [2011/02/03(木) 18:32:56 ]: fermi系ならコンカレントカーネルで問題無し
308 名前：デフォルトの名無しさん mailto:sage [2011/02/04(金) 01:25:47 ]: どうせ最後にcpuに持ってくるんだから、充分データ量が減ったところで転送しちゃうのも手ではあるね。
309 名前：303 mailto:sage [2011/02/04(金) 03:03:46 ]: >>304
レジスタ32768ですがスレッドは192/ブロックです。
そもそもActive Thread Blocks per Multiprocessorを2以上にするように
1スレッドで使用しているレジスタ数とシェアードメモリを見て
スレッド数を決めているので足りなくなるなんて事はないはず

--ptxas-options=-vでlmemの項が出ないんだから使ってないと思っているのですが
レジスタ/スレッドに上限なんてないですよね？
310 名前：デフォルトの名無しさん mailto:sage [2011/02/04(金) 03:34:56 ]: そのレジスタ数から考えてfermi系ではないとおもじゅんだが
いっそのことスレッド数を８とか３２とかにして実行してどうなるか見てみればいいんじゃないか
なにか他の原因かもよ？
311 名前：デフォルトの名無しさん mailto:sage [2011/02/04(金) 03:35:39 ]: c/おもじゅんだが/思うんだが
312 名前：303 mailto:sage [2011/02/04(金) 04:33:53 ]: >>310
GF100,Fermiです
8スレッド/ブロックにして(ブロック数を増やしたので総スレッド数は変わらず)
みましたがlocal load/storeは増えました。
カーネル内の変数を(なにかに使用するように)無駄に増やしてみましたが
レジスタ数が63以上にならなかったので上限があるのかも、と思った次第です。
皆さんはレジスタ数が64を超える事はありますでしょうか
313 名前：デフォルトの名無しさん [2011/02/04(金) 05:01:16 ]: 今環境ないからできないがnbodyのサンプルとかはプロファイルとったらどうなんだろうね。
あとフェルミならローカルメモリはキャッシュが効くんではなかったかと思ったがどうでしたっけ？
あとコンパイルオプションでレジスタの使用個数指定してたりしない？
314 名前：303 mailto:sage [2011/02/04(金) 05:18:06 ]: >>313
もちろん--maxrregcountは指定していませんでした

ところで
--maxrregcount 64にすると
ptxas warning : Too big maxrregcount value specified 64, will be ignored
とでました。
--maxrregcount 63にすると何も出ずコンパイルされます。
315 名前：デフォルトの名無しさん mailto:sage [2011/02/04(金) 05:23:42 ]: ということは６３に意味はありそうだね。うーんなんだろ。即答できない
316 名前：デフォルトの名無しさん mailto:sage [2011/02/04(金) 05:42:17 ]: 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.
317 名前：デフォルトの名無しさん mailto:sage [2011/02/04(金) 12:19:26 ]: Fermiは63までしかないよ。GT200とかは127まで行けたけど。
318 名前：デフォルトの名無しさん mailto:sage [2011/02/04(金) 18:43:37 ]: CC2.0でchar型の配列a[N]に
int idx = blockIdx.x * blockDim.x + threadIdx.x;

a[idx];

ってアクセスしたらきちんと32bitずつglobalから読みだしますか？
それとも128bitずつですか？
319 名前：303 mailto:sage [2011/02/05(土) 08:23:45 ]: >>317
そうだったんですか
ちょっとずつ減らしてったら現在90くらい使ってることが判明
かきなおしかー
320 名前：デフォルトの名無しさん mailto:sage [2011/02/06(日) 13:23:49 ]: 64しか使われないのに今90使ってるっておかしくないか

経験値からじゃなくて何か根拠のある数字はないのか？
321 名前：303 mailto:sage [2011/02/06(日) 13:51:37 ]: --ptxas-options=-v→63
同時に使ってるであろう変数を1個減らす→63
もういっこへらす→63
繰り返す
へらす→62
やっと62になったということは減らした変数の数*2本分(doubleなので)=26本のレジスタが
63から溢れ出てローカルに保存されているんじゃないか
→本来なら90くらい使われてるんだろう

まあ一概に変数1個=2本とは言えないんでしょうけど
他に表示されないレジスタ数を数える方法思いつきませんでした
322 名前：デフォルトの名無しさん mailto:sage [2011/02/06(日) 15:52:57 ]: 言ってることがようやくわかりました。

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

もしかして待機しているactive threadの分が消費されるのか？つまり
32768/192/ActiveThreadPerBlockが1スレッドの使えるレジスタ数になるのだろうか
323 名前：303 mailto:sage [2011/02/06(日) 16:46:28 ]: >>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表示されない
のだと解釈しました。
324 名前：322 mailto:sage [2011/02/06(日) 16:57:10 ]: 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.

ってあるがこれか？
325 名前：303 mailto:sage [2011/02/06(日) 17:40:25 ]: >>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/スレッドの上限とは関係なさそうですかね。
326 名前：デフォルトの名無しさん mailto:sage [2011/02/07(月) 02:36:58 ]: 上の方で出てたが、やはりCUBLASは省略すべき係数計算を
省略せずに計算しているようだ。
www.netlib.org/blas/dgemm.fに反する実装って許されるの？
327 名前：デフォルトの名無しさん mailto:sage [2011/02/07(月) 08:00:09 ]: 許せるに決まっているだろう。
328 名前：デフォルトの名無しさん mailto:sage [2011/02/07(月) 15:59:16 ]: >>326
公式のフォーラムで本当にそうなのか、
もしそうならどのような理由からか聞いたほうがいいんじゃないか？
もしかしたら自分が勘違いしている部分があるかもしれないし
Nvidiaの間違いなら修正があるだろうし。
329 名前：303 mailto:sage [2011/02/07(月) 16:11:41 ]: 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までと書いてありますが公式にはないのかな...
330 名前：デフォルトの名無しさん mailto:sage [2011/02/07(月) 23:23:08 ]: 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
331 名前：デフォルトの名無しさん mailto:sage [2011/02/07(月) 23:48:13 ]: >>330
あくまでCUDA VS Wizardが不具合ならという前提だけど
sdkのtemplateプロジェクトを元に必要な箇所を修正してプロジェクトを
作りなおしてみるという手もある。
332 名前：330 mailto:sage [2011/02/08(火) 00:19:27 ]: >>331さん
ありがとうございます。頂いたアドバイスを頼りに解決しました。

templateプロジェクトを開く＞プロパティ＞あれ？CUDA Bulid Rule v3.0.14＆sm_20できる＞google先生
＞先生「Cuda.rulesコピれ」＞！SDK導入後コピーするの忘れてた！＞①から②にCuda.rulesをコピー＞ヽ(´ー｀)ノ
①C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK\C\common
②C:\Program Files (x86)\Microsoft Visual Studio 9.0\VC\VCProjectDefaults

ちゃんとCUDA VS Wizardのプロジェクトでもsm_20を指定できました。
CUDA VS Wizardの問題ではなく、Cuda.rulesの格納し忘れでした。失礼いたしました。
333 名前：デフォルトの名無しさん mailto:sage [2011/02/10(木) 23:09:14 ]: CUBLASのcublasSetmatrix()にCUBLAS_STATUS_MAPPING_ERRORを
吐かれてしまうのですが、どういった理由が考えられますか？
最初に最大サイズの行列サイズでallocした後、
そのdeviceメモリ領域に行列データを転送して演算、
結果をhostに書き戻して、さきほどの領域にまた別のサイズ（最大よりは小さい）の行列を
hostから持ってきて演算・・・のような処理を繰り返しています。
334 名前：デフォルトの名無しさん mailto:sage [2011/02/15(火) 02:14:46 ]: 卒論書き終わりました
このスレのみなさんのおかげです
質問にいっぱい答えてくれてありがとう
335 名前：デフォルトの名無しさん mailto:sage [2011/02/16(水) 01:01:24 ]: ちゃんと謝辞に書いたか？
336 名前：デフォルトの名無しさん mailto:sage [2011/02/16(水) 05:21:54 ]: 他のスレッドでは書き込めない超低レベルなチラ裏で恐縮ですが、
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関数内だけで処理をさせるようにしないとうまみが少ないのですね。
337 名前：デフォルトの名無しさん mailto:sage [2011/02/17(木) 02:47:04 ]: 気づいたらなんか本出てたんだけどどうあれ？
入門書？応用書？
338 名前：デフォルトの名無しさん mailto:sage [2011/02/17(木) 03:13:25 ]: 最近？どの本のことですか？
339 名前：デフォルトの名無しさん mailto:sage [2011/02/17(木) 04:43:56 ]: CUDA by Example?
340 名前：デフォルトの名無しさん mailto:sage [2011/02/17(木) 14:56:50 ]: >>339
そうそれ。2/14って書いてあったけど翻訳かな
341 名前：デフォルトの名無しさん mailto:sage [2011/02/17(木) 17:23:30 ]: >>340
翻訳だよ。
内容は入門書

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

最適化については記述されてない。つまり
ブランチダイバージェンス、GRAMのコアレスアクセス、シェアードメモリのバンクコンフリクト
の説明はない
342 名前：デフォルトの名無しさん mailto:sage [2011/02/17(木) 22:51:23 ]: 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発売
343 名前：336 mailto:sage [2011/02/19(土) 01:05:39 ]: 超低レベルな>>336の訂正で申し訳ないですが、
clock()によるGPUの実行時間は取得方法が誤っておりました。

スーパーコンピューティングニュース[東京大学情報基盤センタースーパーコンピューティング部門]
www.cc.u-tokyo.ac.jp/publication/news/
にリンクが張られてある大島聡史先生のPDF：
「これからの並列計算のためのGPGPU連載講座 (Ⅲ) GPGPUプログラミング環境CUDA最適化編」
（201005_gpgpu2.pdf）に
＞funcname<<<xyz, xyz>>>(args1, args2);
の形の呼び出しは非同期であるためこの前後で時刻を取得しても
GPU上での計算時間は得られないと書かれてありました。
だから0clockになっていたんですね…。
（少なくとも処理後のcudaMalloc()の完了時点では処理は終わっているから
62clock未満ではあるのでしょうけど）

同pdfに、GPUの実行時間を調べられるCUDAプロファイラなるものが紹介されて
いるのでこの使い方も覚えていこうと思います。
このpdfのシリーズはすごくわかりやすいと思いました。
こういうものが公開されているのはすごくありがたいです。
344 名前：デフォルトの名無しさん mailto:sage [2011/02/19(土) 08:58:53 ]: このスレ読めばちゃんと書いてあるのに……
345 名前：デフォルトの名無しさん [2011/02/19(土) 11:05:26 ]: 公開されている資料といえばこんなのを見つけたが。
ttp://accc.riken.jp/HPC/training.html
の下の方のCUDA。
346 名前：デフォルトの名無しさん mailto:sage [2011/02/19(土) 14:08:01 ]: >>343
そういうあなたにCuda by exampleをすすめる
347 名前：デフォルトの名無しさん mailto:sage [2011/02/20(日) 01:43:34.23 ]: 著者が必死に宣伝しているように見える
348 名前：デフォルトの名無しさん mailto:sage [2011/02/20(日) 01:54:08.17 ]: まぁ青木本が1番
349 名前：デフォルトの名無しさん mailto:sage [2011/02/20(日) 02:46:20.30 ]: SDKのドキュメントが一番だろう。
アレに必要なことは大体書いてある。
あとは実行してプロファイルとったり
nvccでcubin吐いてenvytoolsで確認したり。
350 名前：デフォルトの名無しさん mailto:sage [2011/02/20(日) 05:37:37.97 ]: 日本語でおｋ
351 名前：デフォルトの名無しさん mailto:sage [2011/02/20(日) 08:59:35.86 ]: nVidia日本法人にまともなエンジニアがいないんだ。日本語だけじゃなんともならん
352 名前：デフォルトの名無しさん mailto:sage [2011/02/20(日) 18:02:58.21 ]: SLIしてる場合、CUDAからはどういう扱いになるんでしょうか
１つの大きなGPUとして見え１つのホストスレッドで扱えるのでしょうか
それともSLIしていない場合と同じように複数のGPUとしてみえ、枚数分のホストスレッドで管理してやる
必要があるんでしょうか。
353 名前：デフォルトの名無しさん mailto:sage [2011/02/20(日) 18:24:47.06 ]: 前に聞いたときはSLIはなんらCUDAに関与しないとのことだったけど。
354 名前：デフォルトの名無しさん mailto:sage [2011/02/20(日) 19:00:08.93 ]: ということは複数GPU挿しているときと同様の扱いになるということですか。
ゲーム等ではSLIできてCUDAではできないんですね。
手軽さより自分で制御できる事を優先してるのかな
355 名前：デフォルトの名無しさん mailto:sage [2011/02/20(日) 19:07:14.32 ]: もともとコアとしては遠くにあるわけだしVRAMも遠いんだから同一扱いは無理でしょ
むしろSLIがうまく分担してるんじゃない
356 名前：デフォルトの名無しさん mailto:sage [2011/02/20(日) 19:39:29.66 ]: SLIで充分なんだから本来グラフィックスにGTX480みたいなおばけチップはいらないんだよね
357 名前：デフォルトの名無しさん mailto:sage [2011/02/20(日) 19:44:13.09 ]: でもそのおばけチップでSLIにしたらもっと性能あがるじゃん
とくにCUDAスレの住民にはメリット大きいんじゃない？
358 名前：343 mailto:sage [2011/02/20(日) 22:40:30.94 ]: >>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)。
359 名前：デフォルトの名無しさん mailto:sage [2011/02/20(日) 23:40:26.16 ]: >>358
まだなら「はじめてのCUDAプログラミング」を先に勧める。
英文が苦手ならさっさと本買ったほうが時間の節約になるし
実際のプログラムに時間を使った方が有益だと俺は思うよ。
360 名前：358 mailto:sage [2011/02/21(月) 00:02:16.28 ]: >>359
そうですね…約2500円で省略できる時間のほうが有意義のように思います。
よし…買おう…bk1の300円ギフト券が手に入ったら…。
361 名前：デフォルトの名無しさん mailto:sage [2011/02/21(月) 10:12:48.81 ]: >353
>前に聞いたときはSLIはなんらCUDAに関与しないとのことだったけど。

使えるっぽいぞ
jp.slizone.com/object/sli_cuda_learn_jp.html
362 名前：デフォルトの名無しさん mailto:sage [2011/02/22(火) 00:44:44.56 ]: SDｋにSLID3D10Textureてサンプルがあってソースみたんだが、
最初にDirectXとGPUのチェックがあるだけで
特にSLI独自のコーディングは見当たらないな
かといって複数のGPUを制御するようなスレッド関係のコードもない

それらからSLI組んでれば1つの大きなGPUとしてプログラムからは扱えるんじゃないかと
思ったんだがどうなんだろうか。
363 名前：デフォルトの名無しさん mailto:sage [2011/02/22(火) 23:08:18.71 ]: >>361
なんだか滅茶苦茶な翻訳だな。
で、そこを読む限りPhysxとの連係には触れられているけど
それ以上の情報はないし、英語サイトにはその記述さえ見当たらないね。
うーん、どうなんだろう
364 名前：デフォルトの名無しさん mailto:sage [2011/02/22(火) 23:39:43.06 ]: 俺の友達も教務情報改造して人数出るようにしたりしてたけど日本人だったぞ
365 名前：364 mailto:sage [2011/02/22(火) 23:40:49.41 ]: 間違えました
ごめんなさい
366 名前：デフォルトの名無しさん mailto:sage [2011/02/24(木) 22:29:20.97 ]: 最近出たインプレスの本にSLIのことが載ってた
367 名前：デフォルトの名無しさん mailto:sage [2011/02/24(木) 22:52:09.13 ]: >>366
たしかに載ってるが概要の紹介程度で、SLIを使う方法や
SLIからどうGPUが見え、どのように制御するか（orできるか）は書かれてない。

索引から逆引きしても１０行くらいしか書かれてないじゃん。
368 名前：デフォルトの名無しさん [2011/02/25(金) 05:40:01.78 ]: 新型マックじゃ使えないの?
369 名前：デフォルトの名無しさん mailto:sage [2011/02/25(金) 06:17:44.62 ]: Mac用のドライバもtoolkitもあるだろ
あとはGPUが対応してるか調べろよ
370 名前：デフォルトの名無しさん [2011/02/25(金) 14:27:17.76 ]: 使えない
371 名前：デフォルトの名無しさん [2011/02/28(月) 19:10:18.87 ]: CUDAのアプリを作り始めたのですが、CUDAのSDKがインストールされていない環境で
アプリを動かすと”cudart32_32_16.dll”が無いと怒られます。

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

CUDA 3.2を使っています。
372 名前：デフォルトの名無しさん mailto:sage [2011/02/28(月) 19:49:36.32 ]: よくわからんけどコンパイル時に
-L /usr/local/cuda/lib64/ -lcudart
じゃだめ？
パスは適当
373 名前：デフォルトの名無しさん mailto:sage [2011/03/01(火) 00:11:54.00 ]: >>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バージョン
374 名前：デフォルトの名無しさん mailto:sage [2011/03/01(火) 00:47:03.69 ]: >>371
CUDA RuntimeのDLLは再配布が許可されているみたいだし
単に他の環境で動かしたいだけなら、DLL同梱でいい気がする。

ttp://developer.download.nvidia.com/compute/cuda/3_2_prod/toolkit/docs/EULA.txt
375 名前：デフォルトの名無しさん mailto:sage [2011/03/01(火) 02:29:40.71 ]: >>374
ほむ！
ありがとうございます。
インストールしたときについてくるライセンスではそこら辺がよくわからなかったので助かります。
とりあえず、DLLも同梱しようと思います。（スタティックリンクしたいけど・・・）
376 名前：デフォルトの名無しさん mailto:sage [2011/03/01(火) 03:33:24.52 ]: 4Gamer.net ― NVIDIA，「CUDA 4.0」の新機能を予告。より広い層の開発者にアピールする新バージョン（CUDA）
www.4gamer.net/games/076/G007660/20110228018/

Fermi持ってたら恩恵受けられるのかな。なんかよくわからんけど楽しみだ。
＞ちなみにGPUDirect 2.0とThrustをフルサポートするGPUは，Fermiアーキテクチャ限定となる。
377 名前：デフォルトの名無しさん mailto:sage [2011/03/01(火) 03:56:33.42 ]: >>376
・GPUDirect 2.0 …スパコンには良さそう？
・Unified Virtual Addressing…どれが何のメモリなのかごっちゃになりそう？
・Thrust…STL。。。どうでもよさそう？
378 名前：デフォルトの名無しさん [2011/03/01(火) 23:32:35.21 ]: こんにちは。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
379 名前：デフォルトの名無しさん mailto:sage [2011/03/02(水) 00:46:11.36 ]: ptx出力じゃ、用が足らんのけ?
380 名前：デフォルトの名無しさん [2011/03/02(水) 13:54:56.94 ]: 初心者です。
今からくだ勉強したい場合はCUDA4.0がオススメですか？
381 名前：デフォルトの名無しさん mailto:sage [2011/03/02(水) 15:15:15.23 ]: >>380
むしろcuda2.2くらいから。
基礎ができてから発展にすすめ
382 名前：デフォルトの名無しさん [2011/03/02(水) 15:23:23.73 ]: 381さん有り難うございます。cuda2.2からやってみます。
383 名前：デフォルトの名無しさん mailto:sage [2011/03/02(水) 15:39:44.51 ]: 古いバージョンのCUDAを知る必要は普通はない。
CUDA 4.0でも基本的にこれまでの機能を全て備えている。
また古い世代のGPUは余裕がある場合だけ追うべき。
384 名前：デフォルトの名無しさん mailto:sage [2011/03/02(水) 15:51:19.20 ]: >>383
同意

つか、
>cuda2.2くらい
"くらい"なんて中途半端な答え方してるのかがワカラン
385 名前：デフォルトの名無しさん mailto:sage [2011/03/03(木) 23:20:08.40 ]: >>379
readelf でcubinみましたが、いまいちよくわからなかったのでptxみます。
まあ、ptxもよくわからないのですが…
386 名前：デフォルトの名無しさん mailto:sage [2011/03/04(金) 11:17:18.22 ]: PTXを読んで分からない人にはcubinは無理。
387 名前：デフォルトの名無しさん mailto:sage [2011/03/04(金) 20:05:59.06 ]: CUDA 4,0 RC が来たよ
388 名前：デフォルトの名無しさん mailto:sage [2011/03/06(日) 02:29:43.51 ]: Visual Studio 2010 まだ来ないよ
389 名前：デフォルトの名無しさん [2011/03/08(火) 02:57:27.52 ]: すみません。
どなたかAMBER(Assisted Model Building with Energy Refinement)というソフトをCUDAで動かしてる人いませんか？
このソフトを
Cent OS5.5 x86_64(カーネル 2.6.18-194.el5)
インテルコンパイラーとの併用で動かそうと何度も試みているのですがなかなか上手くいきません。
不具合の原因はGCCのバージョンとインテルコンパイラーのバージョンの組み合わせにあると睨んでいるのですが…
もし、動かせている方がいらしたらGCCのバージョンとインテルコンパイラーのバージョンの組み合わせををヒントとして教えてほしいのですが
390 名前：デフォルトの名無しさん mailto:sage [2011/03/08(火) 22:01:42.39 ]: >>389
コンパイルでエラーなのかリンカでエラーなのかわからん。
取り合えずsdkのサンプルでiccで動かしてみたら？それで問題がわかるでしょ。
それが出来ないなら、
そもそもnvcc自体がgccのラッパだし、
素直にgcc使えば？
391 名前：デフォルトの名無しさん [2011/03/09(水) 22:09:17.79 ]: デバイス関数で、VC++でいうところの Sleep(0)、gccでいうところのsleep(0)？？？と同等の動作をする関数を探しているのですが
適切な関数があれば教えていただけないでしょうか。
392 名前：デフォルトの名無しさん mailto:sage [2011/03/09(水) 22:12:55.97 ]: 眠らす必要がなさそうだから無いと思う
393 名前：デフォルトの名無しさん mailto:sage [2011/03/09(水) 22:49:28.90 ]: >>392
フロントエンドというかデスクトップというか、とにかくグラフィックスの描画が遅くなる（当たり前ですが…）のと、
GPUの温度がかなり上がるので、
デバイス関数のループの中に、ちびっとだけSleep(0) 的なものを入れたいのですが、
関数的なものは用意されていないっぽいですかね。

例えば、どれか１つだけスレッドの処理ループを長引かせて、__SyncThreads() でもしたらいいのかしらん…？？？
394 名前：デフォルトの名無しさん mailto:sage [2011/03/09(水) 22:53:11.30 ]: 利用者の大半はカーネルを一気に計算したい派だとおも
熱いから休憩するって発想は俺含めてないんじゃね？
395 名前：デフォルトの名無しさん mailto:sage [2011/03/09(水) 23:41:01.50 ]: 自分はカーネルを終了させる必要があると思います
396 名前：デフォルトの名無しさん mailto:sage [2011/03/09(水) 23:58:40.99 ]: >>395

例えば、

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

って感じじゃSleep(0)の代わりにならないですかね？
397 名前：デフォルトの名無しさん mailto:sage [2011/03/10(木) 00:06:55.92 ]: それはビジーウェイトになんね？
結局熱くなってしまうんだが
398 名前：デフォルトの名無しさん mailto:sage [2011/03/10(木) 00:16:02.21 ]: >>397
そうなんですか…
ビジーウェイト調べてみます。
399 名前：デフォルトの名無しさん mailto:sage [2011/03/10(木) 04:14:12.92 ]: >>396
カーネルを分割すればいいんじゃね?
わざわざ遅くする理由を思いつかないから試す気にもなれないけど。
400 名前：デフォルトの名無しさん mailto:sage [2011/03/10(木) 09:03:17.86 ]: >>399
これが正解だろう。
でもそもそも、ちょっと休ませたくらいじゃ熱なんて変わらんだろうから、ドライバの設定でクロックを下げたらいい。
もしくは、放熱を強化するかだろう。
401 名前：デフォルトの名無しさん [2011/03/17(木) 13:38:21.43 ]: CUDAでGTX460(1.5GHz)、4枚分程の計算能力を求めています。
GTX460を2枚刺ししたPC2台と、4枚刺ししたPC1台のどちらを組む方が良いでしょうか？

4枚刺しですと、電源のとりまわしと冷却に不安を感じております。
402 名前：デフォルトの名無しさん mailto:sage [2011/03/17(木) 14:33:27.94 ]: GPU間の通信はどれくらいしますか。
403 名前：402 mailto:sage [2011/03/17(木) 17:10:47.41 ]: そういうのを意識しないなら、なるべくわけた方が安定しそうですね。
404 名前：デフォルトの名無しさん mailto:sage [2011/03/17(木) 17:47:16.68 ]: >>401
単純に、書かれた内容だけで判断する限り前者の方が計算能力が高い。
但し、後者により強力なCPUとメモリを搭載する場合は一概に言えない。
405 名前：デフォルトの名無しさん mailto:sage [2011/03/21(月) 23:08:34.74 ]: >>402
GPU間の通信は無いです。

>>403-404
GPU間の通信が無く、CPU側の処理も軽いので
4枚刺しによる恩恵が特になさそうなので
2台で組んだ方がよさそうですね。
ありがとうございました。
406 名前：デフォルトの名無しさん mailto:sage [2011/03/22(火) 00:55:46.70 ]: GPU計算早すぎ、ワロタ。
長い目で見ると、DSPとかで食いつないでた信号処理屋的に失業の危機じゃん…
407 名前：デフォルトの名無しさん mailto:sage [2011/03/22(火) 05:48:58.16 ]: GPUメモリのTLBミスが気になるのですが、何か情報ないでしょうか
408 名前：デフォルトの名無しさん mailto:sage [2011/03/22(火) 10:14:13.39 ]: >>406
カスタムチップによる高速演算ボードを出していたところは尽く撤退しました。
409 名前： [―{}@{}@{}-] デフォルトの名無しさん mailto:sage [2011/03/22(火) 23:46:07.04 ]: sharedメモリへのバンクコンフリクトは避けたほうが良いって話だけど、
共有メモリが16バンクに分割されているってのは①②のどっちだろ。

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

ハードウェアが固定されてそうだから①の気がするんだけど…。
410 名前：福盛俊明 [2011/03/23(水) 02:45:25.67 ]: アハ～♪”
411 名前：デフォルトの名無しさん mailto:sage [2011/03/23(水) 13:39:51.91 ]: 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だと思っている)
とは別のものなのでしょうか？

試された方はここに居らっしゃいますでしょうか？
412 名前：デフォルトの名無しさん mailto:sage [2011/03/23(水) 17:48:37.13 ]: CanAccessPeerは無視してよい。
確かにTeslaシリーズでは1が、GeForceシリーズでは0が返る。
だけど実際にはGeForceでも通信はできる。
UVAだかなんだかに対応するGPUなら問題ないのかも。
413 名前：デフォルトの名無しさん mailto:sage [2011/03/23(水) 23:13:07.00 ]: ワープってやっぱりSMの上限いっぱいまで詰め込んだほうがいいの？
例えばSMの上限ワープが32のところに16個のワープにしても、
メモリアクセスのようなレイテンシがないなら問題ない気がするんだが。
414 名前：デフォルトの名無しさん [2011/03/24(木) 11:46:09.16 ]: CUDA初心者です。
CUDAを使って、ゲームエンジンのUnrealとかcryを使ってる方いらっしゃいますか。
参考になるサイトを教えてほしいのですが
415 名前：デフォルトの名無しさん [2011/03/24(木) 20:35:29.92 ]: 例えば、デバイス側の主要な処理が下のような論理演算であった場合、
論理演算部分のアセンブラ化による高速化は望めるでしょうか？

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

__global__ void test(int *a, int *b, int *c)
{
　　// 初期化処理
　　
　　// メインのループ
　　for (i = 0; i < LOOP; i++)
　　{
　　　　a[i] = (b[i]<<10 | (b[i] >> 22)) ^ c[i];
　　}
}
416 名前：デフォルトの名無しさん mailto:sage [2011/03/24(木) 21:10:44.15 ]: 少なくともその論理演算部分をアセンブラで書く必要はない。
そもそもメモリアクセスがボトルネックにならないほど多くの演算があるのかどうか。
417 名前：デフォルトの名無しさん mailto:sage [2011/03/24(木) 21:50:24.37 ]: 分岐を避けるためにビット演算使いたくなるけど、いまいち貧弱なんだよなあ・・・
418 名前：デフォルトの名無しさん [2011/03/25(金) 12:10:15.17 ]: 英語か日本語でptxファイルの読み方を解説しているドキュメントがあれば教えていただけないでしょうか。

また、コンパイル時にptxファイルに明示的にコメントを挿入するコマンドがあれば教えていただけないでしょうか。
例えば、test.cu に、「__PTX_COMMENT("foo");」という行を記述すると、ptxファイルに
// foo
といった行が挿入されるという具合のものを探しています。
419 名前：デフォルトの名無しさん mailto:sage [2011/03/25(金) 13:36:10.03 ]: PTXのマニュアルはToolkitに含まれているかと。

ptxにブロックごとに行番号情報は含まれているのでソースのどこに相当するかは分かるのでは。
特定の場所が知りたければ例えば__syncthreads()を１０個入れとくとか。
420 名前：デフォルトの名無しさん mailto:sage [2011/03/25(金) 15:16:09.98 ]: perl とかでつくれそうだよね
421 名前：デフォルトの名無しさん mailto:sage [2011/03/25(金) 17:03:42.70 ]: >>419
なるほど。例えば、ダミーでfoo0()とか、foo1()とか関数を作ってコメント代わりに呼び出せばよいと。

>>420
コマンドとしては無いということですね。
422 名前：デフォルトの名無しさん mailto:sage [2011/03/28(月) 20:51:27.79 ]: 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は実行して変換結果も正しいのだが謎す。
423 名前：デフォルトの名無しさん mailto:sage [2011/03/29(火) 11:17:10.20 ]: cufftは内部であれこれ怪しげなことしているからなぁ。
VRAMは未だ余っている筈なのにエラーになったり、作ったプランを使わずに破棄するとこけたり……
って、>422はそのパターンか?
424 名前：デフォルトの名無しさん mailto:sage [2011/03/30(水) 14:28:25.63 ]: >423
cufftExecC2C実行してもやっぱでる。
海外フォーラムみてたらver 3.2が怪しくて ver3.0にしたら消えたって書いてあったんで
ver3.0にしてみたがやっぱでる。うーん・・・。
425 名前：デフォルトの名無しさん mailto:sage [2011/03/30(水) 16:09:16.17 ]: >>424
うーん、SDKとcufftの組み合わせで相性あるしなぁ。
あと、サイズが2の冪かどうかでも変わるかもしれない。
つーか、3.1よりは3.2の方がましだった気がするんだが……
426 名前：デフォルトの名無しさん mailto:sage [2011/04/02(土) 17:59:42.31 ]: サンプルはmain関数でCUDAを利用するのばっかだけど、
VC++でダイアログから.cuの関数使うようなサンプルってなんかない？
427 名前：デフォルトの名無しさん mailto:sage [2011/04/02(土) 18:01:41.56 ]: >>426
おまいさんは関数を分割することもできんのかね。
428 名前：デフォルトの名無しさん mailto:sage [2011/04/02(土) 19:08:36.61 ]: windowsよく知らんけど
それはライブラリ作んなくてもできるんかい？
429 名前：427 mailto:sage [2011/04/02(土) 19:39:35.97 ]: >427
extern "C" つけたらできたや。ごめんごめん。
430 名前：デフォルトの名無しさん mailto:sage [2011/04/02(土) 20:08:34.70 ]: >>429
お前はだれ？
そんで誰にレス支店の？
431 名前：425 mailto:sage [2011/04/02(土) 21:07:05.73 ]: まちがえちゃった。テヘッ
432 名前：デフォルトの名無しさん mailto:sage [2011/04/04(月) 17:39:36.21 ]: プログラミングガイドに記載されているように

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を使用しています．
433 名前：デフォルトの名無しさん mailto:sage [2011/04/04(月) 22:25:10.79 ]: CU_SAFE_CALL(cuInit(0));
CU_SAFE_CALL(cuGLInit(0));
CU_SAFE_CALL(cuDeviceGet(&device, 0));
CU_SAFE_CALL(cuGLCtxCreate(&context, 0, device));
434 名前：デフォルトの名無しさん mailto:sage [2011/04/04(月) 22:26:11.47 ]: ミスった。
1点だけ。
uGLInit(0)→uGLInit()
435 名前：432 mailto:sage [2011/04/05(火) 11:10:29.98 ]: >>433-434
この通りにしましたところ，今度はcuGLInitで
Cuda driver error c9
が発生しました...
436 名前：デフォルトの名無しさん [2011/04/05(火) 17:48:55.53 ]: コマンドラインからnvccでコンパイルする際、普通のC++部分のコンパイルを
Intelコンパイラ："icl.exe"に行わせたいのですが、どのようにコマンドを記述したら良いでしょうか？

環境：
Windows 7
CUDA 3.2
437 名前：デフォルトの名無しさん mailto:sage [2011/04/05(火) 17:59:21.74 ]: nvccしか使えないところはライブラリ化してやってるけど
もっといい方法あるんかいな？
438 名前：デフォルトの名無しさん mailto:sage [2011/04/05(火) 20:34:55.57 ]: 西日本も福島原発の放射能に曝される。

4/7 予測 up3.viploader.net/ippan/src/vlippan198234.jpg
発表はドイツ気象庁 www.dwd.de/
439 名前：デフォルトの名無しさん mailto:sage [2011/04/05(火) 21:13:26.38 ]: CUDAで予測出来ないんすか？
440 名前：デフォルトの名無しさん mailto:sage [2011/04/05(火) 23:50:34.55 ]: できますん
441 名前：デフォルトの名無しさん mailto:sage [2011/04/06(水) 02:41:41.96 ]: どっちですか
442 名前：デフォルトの名無しさん mailto:sage [2011/04/06(水) 07:54:08.27 ]: 日本語で予測できるか質問するくらい意味がないことに気付け
443 名前：デフォルトの名無しさん mailto:sage [2011/04/06(水) 10:06:08.15 ]: >>442
さすがにそれはちょっと違うと思うけど
気象系シミュレーションでGPUつかってるところって結構あるよね？
444 名前：デフォルトの名無しさん mailto:sage [2011/04/06(水) 17:09:45.35 ]: ？
445 名前：デフォルトの名無しさん mailto:sage [2011/04/06(水) 17:21:47.83 ]: さすがにこれは？？？だな
446 名前：デフォルトの名無しさん mailto:sage [2011/04/06(水) 18:02:23.22 ]: >>435

cuInit
cuDeviceGet
cuCtxCreate
cuGLInit
cuGLCtxCreate

じゃね?
447 名前：432 mailto:sage [2011/04/07(木) 17:59:04.18 ]: >>446
やはりcuGLInitで同様のエラーが出ます．
448 名前：デフォルトの名無しさん mailto:sage [2011/04/07(木) 18:05:26.23 ]: 3.2でやってみれ
449 名前：432 mailto:sage [2011/04/07(木) 18:20:08.71 ]: >>448
3.2でも同様です．
450 名前：デフォルトの名無しさん mailto:sage [2011/04/07(木) 19:40:03.31 ]: OpenGLは初期化してるよね?
451 名前：432 mailto:sage [2011/04/08(金) 08:52:38.77 ]: >>450
大変御迷惑おかけしました．
自分の無知さに絶望しました．

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

御手数おかけしました．
452 名前：デフォルトの名無しさん mailto:sage [2011/04/08(金) 20:02:33.08 ]: シュミレーション
453 名前：デフォルトの名無しさん mailto:sage [2011/04/08(金) 21:02:09.88 ]: あながち間違ってはいないな
454 名前：デフォルトの名無しさん mailto:sage [2011/04/09(土) 16:43:22.41 ]: cudaMallocとかcudaFreeとかcudaMemcpyとか、どのヘッダに定義されてんの？
cutil.hにもcuda.hにも定義されてねーのだが。
455 名前：デフォルトの名無しさん mailto:sage [2011/04/09(土) 17:05:25.29 ]: cuda_runtime.h
456 名前：デフォルトの名無しさん mailto:sage [2011/04/09(土) 17:10:51.07 ]: 　　　　＿＿＿
　　 ,;f　　　　　ヽ　　　　　　　
　　i:　　　　　　 i 　　
　　|　　　　　　　　|　　///;ﾄ,
　　|　　　　^　　^　)　////ﾞlﾞl; 　　
　　(.　　>ﾉ(､_, )ヽ､} l 　 .i .! | 　　
　 ,,∧ヽ　!-=ﾆ=- | │　　| .|
／＼..＼＼｀ﾆﾆ´ !,　{　　 .ﾉ.ﾉ
／　＼＼￣￣￣../　　 / .|
457 名前：デフォルトの名無しさん mailto:sage [2011/04/10(日) 01:00:12.43 ]: いまだにキューダと読んでしまうのだが、クーダと読む習慣が付いてる人はどれくらいいるんだろう
458 名前：デフォルトの名無しさん mailto:sage [2011/04/10(日) 11:48:25.16 ]: あなた以外ほとんどがクーダだと思います
459 名前：デフォルトの名無しさん mailto:sage [2011/04/10(日) 12:13:51.20 ]: main.cu

メインドットクー？
460 名前：デフォルトの名無しさん mailto:sage [2011/04/10(日) 12:41:02.43 ]: OpenGLとの連携が遅いでござる。
フォーラム見ても遅い遅いと書いてあるし、どうしようもないのかこれ？
461 名前：デフォルトの名無しさん mailto:sage [2011/04/10(日) 18:44:20.12 ]: warpはなんて読む？
462 名前：デフォルトの名無しさん mailto:sage [2011/04/10(日) 18:48:15.64 ]: ﾜｰﾋﾟｰ
463 名前：デフォルトの名無しさん mailto:sage [2011/04/10(日) 19:12:02.26 ]: ウォープ
464 名前：デフォルトの名無しさん mailto:sage [2011/04/10(日) 19:56:18.22 ]: 飯野の巣
465 名前：デフォルトの名無しさん mailto:sage [2011/04/11(月) 20:09:52.50 ]: warning
walking
working
star wars
それぞれなんて読む？
466 名前：デフォルトの名無しさん mailto:sage [2011/04/11(月) 20:23:49.35 ]: それ、この板で何度も繰り返し振られる話題だよな
467 名前：デフォルトの名無しさん mailto:sage [2011/04/11(月) 20:38:03.05 ]: >>465
ワーニング
ワーキング
ウォーキング
スターウォーズ
468 名前：デフォルトの名無しさん mailto:sage [2011/04/11(月) 20:41:31.72 ]: スターワーズだりょ
469 名前：デフォルトの名無しさん mailto:sage [2011/04/11(月) 20:42:58.66 ]: threadIdxがidのxだといつも勘違いしてしまう
470 名前：デフォルトの名無しさん mailto:sage [2011/04/11(月) 21:04:45.77 ]: stdio をスタジオだといつも勘違してしまうのとはちょっと違うか
471 名前：デフォルトの名無しさん mailto:sage [2011/04/11(月) 22:09:14.04 ]: カタカナ表記した時点で正しい読みから外れてる
472 名前：デフォルトの名無しさん mailto:sage [2011/04/11(月) 22:31:58.48 ]: >>469
わかるわ
油断するとIdxとIdyがあると思っちゃう
threadIdx.y
xのy？？？
スレッドはフルスペルなんたからthreadIndexでいいじゃねえか
473 名前：デフォルトの名無しさん mailto:sage [2011/04/11(月) 23:56:43.33 ]: 寧ろptx出力宜しくtid.x希望w
474 名前：デフォルトの名無しさん mailto:sage [2011/04/12(火) 22:35:25.63 ]: 明示的にグローバルメモリから共有メモリへの読み込みを行う命令ってないのかね。
スレッドに分けて暗黙にコアレスアクセスさせると、どうもコードが汚くなる。
475 名前：デフォルトの名無しさん mailto:sage [2011/04/13(水) 01:04:12.90 ]: ううーん、デバイス側にメモリ確保しっぱなしだとOpenGLの性能が著しく低下する…。
計算結果を最終的にOpenGLで表示するような場合だと、無駄なメモリコピーは仕方ないものとして
計算結果をホスト側に逐一戻したほうがよさげ？
476 名前：デフォルトの名無しさん [2011/04/19(火) 21:32:44.77 ]: 3.2から4.0 RC2に変更したんだけど，Cuda.rulesってないの？
477 名前：デフォルトの名無しさん mailto:sage [2011/04/20(水) 01:10:42.96 ]: つかなんでcommon\lib\Win32の中にcutil32.dllがあるんだっていう
478 名前：デフォルトの名無しさん [2011/04/20(水) 01:11:01.88 ]: CUDAのアセンブラにMMXやSSEのような命令はあるでしょうか？

（CUDAのカーネル関数の中で、MMXのPADDDのような32bit整数の4並列計算のようなことができなかと思っています）
479 名前：デフォルトの名無しさん mailto:sage [2011/04/20(水) 01:16:07.37 ]: >>478
CAL+IL
480 名前：忍法帖【Lv=11,xxxPT】 mailto:sage [2011/04/20(水) 04:19:17.84 ]: 2枚差しの場合、2つプログラムを走らせたら、それぞれ別のGPUを使ってくれるのですか？
481 名前：忍法帖【Lv=15,xxxPT】 mailto:sage [2011/04/20(水) 07:08:07.39 ]: 現状では自分で何番目のGPUを使うか指定する必要があります。
482 名前：デフォルトの名無しさん mailto:sage [2011/04/20(水) 11:51:01.49 ]: >>477
イミフ
483 名前：デフォルトの名無しさん [2011/04/20(水) 20:42:07.79 ]: >>479
えと、CUDA用には無いということでしょうか？

add.u32

を2並列(64bitレジスタ)、あるいは4並列(128bitレジスタ)で実行するような命令セットはないでしょうか？
484 名前：デフォルトの名無しさん mailto:sage [2011/04/20(水) 21:24:43.20 ]: >>483
そのためのSIMT
485 名前：デフォルトの名無しさん [2011/04/20(水) 22:27:22.88 ]: >>483
何かを誤解しているようだが、CUDAには必要ないから無い。
SSEを使わないと100%の性能が出せないとかいうCPUのような制約は無い。

add.u32という命令ならほっといてもWarp単位の３２スレッドが（ほぼ同時に）実行する。
全命令が３２並列のSIMD命令になっていると考えればよい。（３２ビットデータの場合）
486 名前：デフォルトの名無しさん [2011/04/20(水) 22:50:57.29 ]: >>484-485
CUDAがSIMTなのか、SIMDなのか、SPMDなのかはさておいて…

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

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

Streamでは128bitレジスタで32bit4並列の演算ができるみたいのですが・・・
487 名前：デフォルトの名無しさん mailto:sage [2011/04/21(木) 04:36:54.19 ]: そこまでやるならＣＰＵの方が早かったりしてｗ
488 名前：デフォルトの名無しさん [2011/04/21(木) 08:16:13.34 ]: >>486
そんなものはないから誤解だと言ってるんだが。
AMDは１スレッドが4 or 5演算のVLIWで動作するから4並列のSIMD命令化の意味があるが
CUDAは１スレッドが各種演算ユニットを１ずつしかもたない(or複数スレッドが共有）CUDA Coreで
実行されるので128bit型を用いたとしても４命令になるだけ。
489 名前：デフォルトの名無しさん mailto:sage [2011/04/21(木) 09:31:19.38 ]: >>488
誤解というか、質問なのですが・・・

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

ありがとうございましたゞ
490 名前：デフォルトの名無しさん mailto:sage [2011/04/21(木) 09:35:39.70 ]: それよりSandyBridgeのベクタ演算機の性能はどうなんだろうな
まだ試験実装だからコア数が少ないけどメインメモリ直結だからけっこう早そうなんだが
491 名前：デフォルトの名無しさん mailto:sage [2011/04/21(木) 09:59:44.72 ]: その辺は簡単な報告を以前SSEスレにあげといたよ
イミフな反応が多々あったけどね
492 名前：デフォルトの名無しさん mailto:sage [2011/04/21(木) 10:04:26.92 ]: >489が根本的に解ってない悪寒
493 名前：デフォルトの名無しさん mailto:sage [2011/04/21(木) 21:59:22.93 ]: >>492
論理レジスタとか、物理レジスタとかいう話しですか？？？
>>486への根本的な回答はどうなりますか？
494 名前：デフォルトの名無しさん mailto:sage [2011/04/21(木) 22:29:01.49 ]: とりあえず公式文書を読んでから質問してくれ。
495 名前：デフォルトの名無しさん mailto:sage [2011/04/21(木) 22:39:09.84 ]: ようするに64bitレジスタに8x8のデータを入れて同時に処理出来るのかってことでしょ
streamsdkには専用の命令があると？
cudaはそんなものは勝手にやると言ってる人が居るけど、どうやってやるんだ？

cudaにはそういう命令は無いよ
内部で処理することが出来るとは思えないけどな
少なくともユーザープログラムからやる方法はないよ
内部的にはそういう命令も持ってるかもしれないけど
というかそういうアセンブラレベルの最適化をしなくていいようにしたのがＣＵＤＡ言語なわけで
複雑化する要因はパフォーマンスを犠牲にしてでも排除されるだろう
どうしてもやりたかったら64bit変数を使って自前でやるしかないだろう
496 名前：デフォルトの名無しさん mailto:sage [2011/04/21(木) 23:07:04.51 ]: >>495
>>486 への回答は”No”ということですね。
ありがとうございます。

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

# 今、スクラッチのCUDA用アプリで、SSEで書かれたCPU用アプリの約5倍（CPU側はSandyBridgeで4スレッドで動作）の速度が出ているのですが、
CPU側がきちんとAVXに対応してきたらこの差はかなり縮まりそうなんですよね・・・
497 名前：デフォルトの名無しさん mailto:sage [2011/04/21(木) 23:11:08.25 ]: ねえ、GPUの速さの要因ってCPUより圧倒的に軽量なスレッド数によるもんじゃないの
まずそっからじゃないのかね？
498 名前：デフォルトの名無しさん mailto:sage [2011/04/21(木) 23:14:37.36 ]: >>496
さっきからすげー鼻につくんだけど
教わる態度じゃないよお前

教えてんの俺じゃないけどさ
499 名前：デフォルトの名無しさん mailto:sage [2011/04/21(木) 23:22:41.50 ]: あれだ、自分が書いたcudaのコードが思ったはど速度が出ないから
腹癒せしているか言い訳探しているかってところだろ。
今日たまたま計測していたんだが、cufftでもfftwfの3倍程度だった。
汎用ルーチンじゃこんなもんだよ。
500 名前：デフォルトの名無しさん mailto:sage [2011/04/21(木) 23:51:10.29 ]: >>498
CUDAの将来的な見通しについて知りたかった部分もあったので
失礼に見えた部分があるかもしれません。

ごめんなさい。
501 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 00:46:21.02 ]: SIMDみたいなものつけて各コアを高機能にするよりは、
単純なコアの数を増やす方向に進むだろうなあ。
502 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 00:51:40.28 ]: レジスタ/SM増やして欲しい…
503 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 01:02:26.93 ]: GF100やGF104の倍精度（64bit）の演算って単精度（32bit）の1/4のスループットとかじゃなかったかな？
おまけにGF104は倍精度をサポートしているコア数は1/3だし。
更にHPC向けではなくグラフィック向けの物は制限かかっているらしいし。

倍精度が必要になるなら仕方ないだろうけど、32bit変数で足りるのに64bit変数を使うのはデメリットの方が圧倒的に多そう。
504 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 01:08:10.32 ]: いろいろ試したがＣＰＵと比較して４倍程度しか上がらないし
労力掛かるわりには環境依存甚だしいわ
４倍くらいだったら別に我慢出来ないようなものなんてないし
マルチコアＣＰＵで普通に組のが一番いいという結論に達した
505 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 01:11:09.84 ]: >>504
計算したい内容を簡潔に書いてみ
計算内容によってＧＰＵに向き不向きがあるから
506 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 01:12:19.97 ]: >>497
膨大なスレッド数と、その切り替えの軽量さを活かして
メモリアクセスのレイテンシの隠蔽とかだっけ。

>>502
スレッドあたりの変数が多いと苦労するね。
そういう場合はシェアードメモリをうまく使って切り抜けるしかないのかな。
507 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 01:15:58.96 ]: CUDAは4nbyteを1要素とした16要素SIMDって感覚だわ。(n = 1,2,3,4)
メモリアクセス考えると結局16要素単位の処理の塊になってしまう。
508 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 01:19:14.16 ]: >>506
共有使ってもたらん
Sharedも増やせ！キャッシュも増やせ!
509 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 01:31:49.64 ]: 向いてないよ。
510 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 07:57:21.85 ]: CUDAとかGPGPUの普及の足引っ張ってるだけだよなぁ。
早くintel、amd,nvidiaで統一してくれよ。
機械語レベルでも、ライブラリレベルでもいいからさ。
511 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 10:14:49.72 ]: OpenMPってよく知らないが_beginthreadしてるわけじゃなく
CPUレベルでスイッチングしてるだけみたいだから
ほとんどオーバーヘッドがないみたいだぞ
十分標準として使えそうなんだが
512 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 10:24:09.66 ]: 残念ながら生成負荷は大きいよ
だから何度もスレッド生成するような処理はかえって遅くなる
513 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 10:26:56.96 ]: いやいや、ちゃんとスレッドプールしてるって
514 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 10:42:12.62 ]: >>513
おまえ>>504だろ
試行レベルが同一
515 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 11:01:13.45 ]: >>514
お前、いろいろと残念な人だな
516 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 11:09:33.71 ]: と、せ～いっぱいの切替し
つか下らんこと書くならよそでやれヴぉけ
517 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 11:16:10.78 ]: 切り返し
な。
煽るなら突っ込まれる隙作っちゃダメだ
518 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 11:20:23.65 ]: いいからスレ違いだ。
519 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 11:21:44.53 ]: スレッドの生成負荷は軽いと言われているのを鵜呑みするとなかなか性能が出ないこともある気がする。
カーネル側でスレッドの生成回数を出来るだけ減らす工夫が必要だったり。
520 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 11:55:51.71 ]: スレッドの生成負荷はプールすればいいから問題じゃない
スレッドを切り替えるときにスタックポインタ変えたりするコンテクストスイッチの負荷が問題
521 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 13:09:56.12 ]: cudaでスレッドのプールができるのかえ？
cpuの話をここに書いてる？
522 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 13:20:03.48 ]: そうそう、何度指摘してもスレ違いの話を続けているの。
523 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 13:23:34.12 ]: その話、終わりました
524 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 13:55:34.91 ]: どこまでをCPUで計算し、どこまでをGPUで計算するのかはGPGPUで重要
だからスレッドの性質については知っておく必要がある
CUDAを使いこなすのにアップアップのアホどもには関係ないけど
525 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 14:04:51.88 ]: 解決を依存性のない単純に並列化可能なアルゴリズムに置き換えられる問題を
GPGPUでやるもんだと思ってたんだが違うのか？
526 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 14:11:42.11 ]: そりゃそうだが、データの転送速度、超越関数を使用するかどうかなど、
単純に並列度が高ければGPUでOKってもんじゃない
Opteronの48コアマシンがお手頃価格で手に入る時代だぜ？
527 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 14:13:49.17 ]: あ、誤解の無いように書いとくが、CUDAやGPGPUを否定してるわけじゃないよ
528 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 14:15:39.88 ]: 夏場に熱で吹っ飛んだのれす
529 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 17:16:29.77 ]: >>526
> Opteronの48コアマシンがお手頃価格

今、おいくら万円くらい？
530 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 17:24:19.24 ]: 全部こみこみで100マソ切った
531 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 17:48:16.70 ]: 高っ
532 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 17:49:59.02 ]: そうか、ごめん……
533 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 17:50:02.42 ]: 48コアで何テラFlopsぐらい出んの？
534 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 17:56:21.72 ]: 高いけど、業務で使うならテスラを組み込んだサーバとたいして変わらないんじゃない？
個人とか研究室とかだと絶対的金額から手が出せないかもしれないけど。
535 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 17:58:31.18 ]: それは言えてる。
536 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 18:06:31.47 ]: >>533
えーと、0.5テラFlopsくらい？
537 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 18:59:41.35 ]: 17万x4(CPU)+12万(ﾏｻﾞﾎﾞ)+HDDﾒﾓﾘｹｰｽか
欲しいけど家庭用のコンセントで大丈夫だろうか
似たようなスペックでGPGPU構築したらいくらぐらいかかるのかな

0.5テラFlopはしょぼくないか
538 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 20:43:24.14 ]: >>537
家庭用コンセントじゃ無理
グラボ刺さるからGPGPUできるよ
539 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 21:43:26.64 ]: >>538
このマザーで自作すると・・・
akiba.kakaku.com/pc/1103/31/131500.php

130000+100000*4+12800*16+100000+50000+α=884800+α

くらい？

だとしたら100万円ほどで完成品が買えるなら良心的な価格な気がするね。

下のマザーボードだと気合いで8本GPU載せられるんだろか？
akiba.kakaku.com/pc/1104/20/213000.php
540 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 21:59:24.97 ]: >>539
> 下のマザーボードだと気合いで8本GPU載せられるんだろか？

どこかに引っかかりそう。ラック筐体次第なのかな
541 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 22:17:54.30 ]: マザー自作するツワモノいないかな

まあ、高周波シュミレソフトで１０００万いくから
性能気にしないでいくなら自作できそうだけど・・・
542 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 23:17:49.25 ]: >>539
そのマザーのPCIeスロットはx8+x1+x8+x1+x8+x1+x8+x1だから
ホスト・デバイス間の帯域不足に悩みそう。

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

というかハイエンドグラボは複数スロット占有が一般的だし、水冷のグラボが必須になりそうｗ
543 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 23:27:34.87 ]: コスト度外視で1台で最速を目指すなら
P6T7 WS SuperComputer
www.unitycorp.co.jp/asus/motherboard/intel/lga1366/p6t7_ws_super/index.html
と
EVGA GeForce GTX 580 FTW Hydro Copper 2
jp.evga.com/products/moreInfo.asp?pn=015-P3-1589-AR
になるのかな。

電源の確保がものすごく大変そうだけどｗ
544 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 23:28:08.92 ]: CUDA Visual Profiler使ってるんですがなんどやっても計測後に
"Error in profiler data file'(動かすディレクトリ)/temp_compute_profiler_0_0.csv' at line number 1. No column found"
と出て計測結果がうまく出ません
何が原因なのでしょうか？
545 名前：デフォルトの名無しさん mailto:sage [2011/04/23(土) 00:17:34.23 ]: エラーコード見てもわからんが、俺はプログラム中にgetchar()を書いてある
キーを押さなきゃ終わらないプログラムをプロファイラに渡してはまったことがあるZE
546 名前：デフォルトの名無しさん mailto:sage [2011/04/23(土) 00:24:57.17 ]: >>544
Visual Profiler使ってないからよく判らんが、その作業ディレクトリはレガシーな名前かい?
空白やら何かの所為で巧くアクセスできないってことはないよね?
547 名前：デフォルトの名無しさん mailto:sage [2011/04/23(土) 00:39:41.60 ]: >>545
外部から画像を読み込んでいますがキーは押さなくても終了するプログラムです

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

プロファイラはプログラムを実行してから計測をしているようなのですがこの計測経過が100%になったら先ほどのエラーがでてしまいます
temp_compute_profiler_0_0.csvは自動生成のファイルの用です
548 名前：デフォルトの名無しさん mailto:sage [2011/04/23(土) 00:56:29.31 ]: >547
エラーがなんで出てるかはさっぱりわからんので、自分がこけたときの経験を列挙するぐらいしかできんのだけど、
Visual Profilerのデフォルト設定では30秒以内に処理が終わらない実行ファイルは実行を途中でうち切って
解析結果出力せず、えんえん解析を繰り返すってのはあったぜ。
549 名前：デフォルトの名無しさん mailto:sage [2011/04/23(土) 00:58:45.66 ]: >>548
そこ毎回適当に300に変えてたわ
しかしなんで一回でぜんぶできないかな
550 名前：デフォルトの名無しさん mailto:sage [2011/04/23(土) 01:02:49.32 ]: グラフィックボードがGPGPUの処理以外にも画面表示の処理をしてる可能性が高いのだから、
画面表示による影響を平滑化するために複数回実行時間を計測してるんじゃないかと。
551 名前：デフォルトの名無しさん mailto:sage [2011/04/23(土) 14:34:55.79 ]: コアレスアクセス用に半ワープ単位でアクセスしなきゃいかんのだが
16×16の二次元のスレッド作ったとして

threadIdx.x = 0～15 が半ワープに入るのか
threadIdx.y = 0～15 が半ワープに入るのか、どっちだべ。
552 名前：デフォルトの名無しさん mailto:sage [2011/04/23(土) 14:56:49.14 ]: ふつうにかんがえてx
553 名前：デフォルトの名無しさん mailto:sage [2011/04/23(土) 15:03:38.31 ]: アブノーマルに考えてx
554 名前：デフォルトの名無しさん mailto:sage [2011/04/23(土) 15:42:37.53 ]: 真逆な意見がでてしまったわけだが、間をとってxと考えることにするよ。
555 名前：デフォルトの名無しさん mailto:sage [2011/04/23(土) 22:47:46.90 ]: Larrabee はいつ出ますか？
556 名前：デフォルトの名無しさん mailto:sage [2011/04/24(日) 00:29:55.82 ]: Larrabeeはなかった事にされてるような・・・
557 名前：デフォルトの名無しさん mailto:sage [2011/04/24(日) 01:23:47.09 ]: LarrabeeってSandryのGPU部分のコア数を数百って規模に拡張したもののことだよ
AVXがまさにLarrabeeのインターフェースらしい
558 名前：デフォルトの名無しさん mailto:sage [2011/04/24(日) 02:12:08.17 ]: VisualStudio使ってEmuDebugの構成にしたら、
デバイス側関数もブレークポイント仕掛けられるって聞いたのにできねえ。
リンクするライブラリ変えるとかしなきゃいかんの？
559 名前：デフォルトの名無しさん mailto:sage [2011/04/24(日) 02:25:06.20 ]: NsightじゃなくてEmu使う理由は？
560 名前：デフォルトの名無しさん mailto:sage [2011/04/24(日) 02:31:36.91 ]: Nsightってグラフィックボード１枚指しでも使えんの？
Nsightを使わない理由は、NsightインストールしたらCUDA3.3がインストールされて、
でCUDA3.3のcuFFTの動作が怪しかったんで、NsightやめてCUDA3.0にしてる。
NsightでもCUDA3.0使えるのかもしれないけど、使い方調べるのがめんどくさかったので
環境構築に慣れたNsight使わないやり方でやってる。
561 名前：デフォルトの名無しさん mailto:sage [2011/04/24(日) 02:44:49.49 ]: 3.3?
562 名前：デフォルトの名無しさん mailto:sage [2011/04/24(日) 03:35:32.04 ]: CUDAとSandy BridgeのGPUは同時に使えますか？

使えるとしたら、Sandy BridgeのGPU用のコーディングには何を使えば良いですか？
563 名前：デフォルトの名無しさん mailto:sage [2011/04/24(日) 03:40:42.99 ]: >>562
SBってGPGPUに対応してるの？
SDKとか出てるって聞いたことないんだけど。
564 名前：デフォルトの名無しさん mailto:sage [2011/04/24(日) 03:53:11.32 ]: Sandy BridgeのGPUはOpenCL1.1に対応してるらしいけど
CUDAに対応してるって話はみたことないから、CUDAとは同時に使えないと思われる。
OpenCLスレで聞いたほうがいいかも。
565 名前：デフォルトの名無しさん mailto:sage [2011/04/24(日) 03:57:39.89 ]: >>564
>CUDAに対応してるって話はみたことないから、CUDAとは同時に使えないと思われる。
対応していない=同時に使えないという考えは改めた方位が良い。
566 名前：デフォルトの名無しさん mailto:sage [2011/04/24(日) 06:20:12.37 ]: グラボを刺した状態でSandyのGPUコアが有効になるのかって話だよね
グラフィックチップとしては完全に機能停止されるだろうけど
演算装置として稼動するんだろうか
567 名前：558 mailto:sage [2011/04/24(日) 11:04:44.80 ]: EmuDebugはリンクするライブラリをcudart.lib→cudaartemu.libにかえたらできたや。
568 名前：デフォルトの名無しさん mailto:sage [2011/04/24(日) 12:25:30.29 ]: 自身ぎゃああああああああああああああああああ
569 名前：デフォルトの名無しさん mailto:sage [2011/04/24(日) 16:01:57.36 ]: VS2008でFirst chance Exceptionがでるが、カーネルコードの何が原因なのかさっぱりわからねえ…
Nsight使うとこういうのの原因すぐに特定できるようになんの？
570 名前：デフォルトの名無しさん mailto:sage [2011/04/24(日) 23:11:39.84 ]: CUDA 4.0 RC2 をインストールしてみようと思っています。
64bit用のバイナリを作るには、CUDAも64-bit版をインストールする必要があるでしょうか？
それとも、CUDAの32-bit版でも nvcc のオプションの切り替えで32bit用と64bit用の両方のバイナリを生成することが出来るでしょうか？

出来れば、WindowsXP(32bit) の環境で64bit用バイナリも作りたいと思っています。
571 名前：デフォルトの名無しさん mailto:sage [2011/04/24(日) 23:26:06.00 ]: >>570
＞64bit用のバイナリを作るには、CUDAも64-bit版をインストールする必要があるでしょうか
たしかそうだったはず

＞CUDAの32-bit版でも nvcc のオプションの切り替えで32bit用と64bit用の両方のバイナリを生成することが出来るでしょうか？
生成したい処理系のSDKをインストールしないといけなかったはず
切り替えはできないと思う
2つ（３２と６４）いれればOKなのかは不明
572 名前：デフォルトの名無しさん mailto:sage [2011/04/24(日) 23:42:26.64 ]: >>571
レスありがとうございます。
ややこしそうなので、32bit 環境、64bit 環境それぞれを用意し、それぞれ用のバイナリを作ろうと思います。
573 名前：デフォルトの名無しさん [2011/04/25(月) 10:06:00.86 ]: 64bitOS環境に64bit Toolkitを入れれば32bit&64bitの両方のバイナリを生成可能。
574 名前：デフォルトの名無しさん [2011/04/27(水) 08:28:47.19 ]: GT430を使っていますが、CUDAでコンパイルすると次のようなエラーが出て
コンパイルできません。
因みにC++は初心者ですがよろしくお願いします。juria_gpu.cu

codepad.org/UtZeyAkV

julia_gpu.cu(44): error: calling a host function("cuComplex::cuComplex") from a
__device__/__global__ function("julia") is not allowed
575 名前：デフォルトの名無しさん mailto:sage [2011/04/27(水) 08:35:16.69 ]: __device__/__global_　この辺良く見直せ！
基本がわかってないと思うから基本からよく勉強だな
576 名前：デフォルトの名無しさん [2011/04/27(水) 08:43:59.77 ]: エラーの意味がよくわかりません。__device__/__global__なんてどこにも
書いて無いし、サンプルコードなので動くはずなのですが、なぜか動きません。
codepad.org/oDh0YFM5
cuComplexをjulia関数から呼ぶのは許されていないという意味なのでしょうか？
577 名前：デフォルトの名無しさん mailto:sage [2011/04/27(水) 08:50:02.31 ]: globalが無いのかよ？
全部晒せよ
578 名前：デフォルトの名無しさん [2011/04/27(水) 09:03:00.78 ]: ideone.com/s9anZ

ここです。よろしくお願いします。
579 名前：デフォルトの名無しさん mailto:sage [2011/04/27(水) 09:07:22.66 ]: CUDA_by_exampleのだな
まずそういう情報を先に出せよ
580 名前：デフォルトの名無しさん [2011/04/27(水) 09:11:49.07 ]: すみません。書くのを忘れていました。
知っている方とお見受けしますので安心です。
581 名前：デフォルトの名無しさん mailto:sage [2011/04/27(水) 09:22:12.28 ]: 特に問題なくコンパイルできたぞ
nvccのオプションは？
582 名前：デフォルトの名無しさん [2011/04/27(水) 09:27:35.63 ]: nvcc -O julia_global.cu -lcutil32 -cutil32D
583 名前：デフォルトの名無しさん [2011/04/27(水) 09:31:28.05 ]: インストールの仕方が悪いのでしょうか？
584 名前：デフォルトの名無しさん mailto:sage [2011/04/27(水) 09:33:25.86 ]: 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
がでるのか？
585 名前：デフォルトの名無しさん [2011/04/27(水) 09:38:07.66 ]: julia_gpu.cu

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

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

の部分でエラーが出ます。原因がわかりません。
586 名前：デフォルトの名無しさん mailto:sage [2011/04/27(水) 09:40:32.60 ]: 正確なコンパイルオプションとエラーメッセージを再掲よろ
587 名前：デフォルトの名無しさん [2011/04/27(水) 09:51:50.54 ]: ideone.com/mZviD

ここです。
588 名前：デフォルトの名無しさん mailto:sage [2011/04/27(水) 09:57:53.82 ]: C:\Users\admin\CUDA\chapter04>nvcc -O julia_gpu.cu
これじゃ -lcutil32とかないじゃん　なけりゃダメだろ
589 名前：デフォルトの名無しさん [2011/04/27(水) 10:01:49.76 ]: リンクエラーが出るかもしれないですが、今はコンパイルエラーなので
そちらでも同じようなエラーが出るか確かめてほしいです。
もしエラーが出なかったら私の環境がおかしいのかも知れません。
Linuxでの開発環境とWindowsでは少し違うと思うのでWindowsでも
確認してもらいたいです。
590 名前：デフォルトの名無しさん [2011/04/27(水) 10:05:27.84 ]: ウインドウズアプリを作るのでリナックスでうまく言っても意味が
無いのですが・・・
591 名前：デフォルトの名無しさん mailto:sage [2011/04/27(水) 10:13:19.95 ]: >>589
あ、そういうこと
じゃあソースが壊れてるんでなければそっちの環境がおかしそうだね
環境の再構築したほうがよさげ
592 名前：デフォルトの名無しさん mailto:sage [2011/04/27(水) 10:33:39.34 ]: 因みにlinux環境で
nvcc -lglut julia_gpu.cu　→　OK
nvcc julia_gpu.cu　→　コンパイルOK、リンクNG
593 名前：デフォルトの名無しさん [2011/04/27(水) 11:22:50.87 ]: 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のバージョンを教えてください。
環境を再インストールしましたがやっぱりだめです。
構文エラーとして解釈されます。
594 名前：デフォルトの名無しさん mailto:sage [2011/04/27(水) 11:28:41.58 ]: 3.2
4.0RC2？はまだ早いんでね？
595 名前：デフォルトの名無しさん [2011/04/27(水) 11:40:43.34 ]: そうですね。自分の持っているやつがGT430なのでバージョン1.0
（最新は1.3）に対応しているみたいなので古いほうをインストール
してみます。
596 名前：デフォルトの名無しさん [2011/04/27(水) 11:45:26.50 ]: 出来ました！
ありがとうございました。
597 名前：デフォルトの名無しさん mailto:sage [2011/04/27(水) 13:22:17.13 ]: >>596
何をどうしたら解決したか書いてほしいな
598 名前：デフォルトの名無しさん mailto:sage [2011/04/27(水) 13:30:13.83 ]: 流れで分かりそうなもんだけど
4.0RC→3.2で解決したようだ
599 名前：デフォルトの名無しさん mailto:sage [2011/04/27(水) 13:44:28.80 ]: 4.0は大きく拡張されてるみたいだし、RCがとれても移行は様子見したほうがよさそうだね
2台とかあって3.2と4.0の環境があったりするなら切り分けしやすいから良いかもしれないけど。
600 名前：デフォルトの名無しさん [2011/04/27(水) 15:49:38.35 ]: こんにちは。
16*32の行列に長さ16の配列をXORの積をしたいのですが可能でしょうか？
イメージ的には32個のブロックに16個のスレッドという感じで考えています。
行列計算なので実装はそんなに難しくないと思うのですが、具体例があまり
見当たらず苦心しています。はじめてからまだ三日なのですが、解る方が
いらしたらソースを見せてほしいです。ヒントだけでも良いです。
よろしくお願いします。
601 名前：デフォルトの名無しさん mailto:sage [2011/04/27(水) 16:03:05.24 ]: >>16*32の行列に長さ16の配列をXORの積をしたいのですが
これどういう意味？もう少し計算を具体的に書いてくれ
602 名前：デフォルトの名無しさん [2011/04/27(水) 16:34:04.46 ]: 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サイクルで実行できると思うのですが何か間違えていますでしょうか？
よろしくお願いします。
603 名前：デフォルトの名無しさん [2011/04/27(水) 16:40:44.57 ]: 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];
}
この処理を並列化したいです。
604 名前：デフォルトの名無しさん mailto:sage [2011/04/27(水) 17:08:33.14 ]: >>603
ideoneかcodepadにC版の関数をアップロードしてくれ。
そうしてくれたらcudaでどこをどう並列化するか考えてみるよ。
605 名前：デフォルトの名無しさん [2011/04/27(水) 17:11:57.15 ]: 並列化したいのはループ使っている部分だけなのでそこだけです。
codepad.org/15hfT9HS
606 名前：デフォルトの名無しさん [2011/04/27(水) 18:53:48.98 ]: 公開したのに反応が無い。
見ても解らないだとか汚いと罵るだけ。
どうせやる気無いんだろ？
607 名前：デフォルトの名無しさん mailto:sage [2011/04/27(水) 18:56:58.77 ]: >>606
うん。
忙しいからね(´・∀・｀ )
608 名前：デフォルトの名無しさん mailto:sage [2011/04/27(水) 19:19:58.01 ]: >>603
またお前か。
いい加減自分の手を動かしたらどうだ。
時間がないなら対価を払って時間を買え。
609 名前：デフォルトの名無しさん [2011/04/27(水) 19:50:55.96 ]: だったらソース見せろとかいうな馬鹿
610 名前：デフォルトの名無しさん mailto:sage [2011/04/27(水) 19:53:32.57 ]: >>606
何怒ってんだ？
その程度のコードなら普通にできるだろ
バカなの？
611 名前：デフォルトの名無しさん [2011/04/27(水) 19:57:29.62 ]: 見る気も無いくせに見せろと言うほうが馬鹿
612 名前：デフォルトの名無しさん mailto:sage [2011/04/27(水) 20:06:58.92 ]: 見せろって言ったのは俺じゃないし、ソース見る前からお前だって分かった。
自分でやる気が無いっていう態度が滲み出てるんだよ。
613 名前：デフォルトの名無しさん [2011/04/27(水) 20:14:43.33 ]: やる気があって調べてるんだよ。
ここならもっと詳しい人が居ると思うから聞いてみただけ。
何で最初からやる気が無いなんて決め付けるの？
614 名前：デフォルトの名無しさん mailto:sage [2011/04/27(水) 20:18:05.38 ]: まぁ質問に具体性がなく丸投げだからな
もう少し考えた過程が見えないと回答もしづらい罠
615 名前：デフォルトの名無しさん mailto:sage [2011/04/27(水) 20:19:27.54 ]: 最初からじゃなくて俺が知っているのはここ8ヶ月くらいだが
分からない、分からない、一個教えてもらうとすぐ次が分からないで
自分で調べる気なんてなさそうに見えるんだが。
616 名前：デフォルトの名無しさん mailto:sage [2011/04/27(水) 21:13:16.83 ]: >>613
Intel Parallel Studio 使って並列化。
↓
逆アセンブルして、コード解析してCUDA化。
↓
(゜д゜）ｳﾏｧ
617 名前：デフォルトの名無しさん mailto:sage [2011/04/27(水) 21:18:58.59 ]: VC++ 2010用のtemplateプロジェクト、どこかにないでしょうか？
CUDA 3.2 SDKにあるVC++ 2008用のtemplateプロジェクトを、VC++ 2010用に変換しようとしましたがうまくいきません。。。
618 名前：604 mailto:sage [2011/04/27(水) 21:32:18.42 ]: >>605
＞並列化したいのはループ使っている部分だけなのでそこだけです。
＞codepad.org/15hfT9HS
悪い悪い、デートしてて遅くなった。
で、このコードのどこをcuda化したいんだ?
未だ帰り道だから、んな長いの見てられないんだ。
619 名前：デフォルトの名無しさん mailto:sage [2011/04/27(水) 22:18:04.42 ]: 16人17脚はデートって言うの？
620 名前：デフォルトの名無しさん mailto:sage [2011/04/28(木) 00:21:09.33 ]: >>617
Windows版のnvccはVisual Studio 2005か2008のコンパイラが前提になっているらしいから
VC++ 2010用のテンプレートだけではうまくいかないと思う。

それらのコンパイラを用意するとなったら、VS2008で作ったほうが楽な気もする・・・
621 名前：デフォルトの名無しさん mailto:sage [2011/04/28(木) 01:56:26.69 ]: >603
きたないコードだなあ・・・。
622 名前：デフォルトの名無しさん mailto:sage [2011/04/28(木) 02:06:27.26 ]: >>603
俺ならこうだな。

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

主なコーディング作法が3つくらいあるから、そのどれかにしろ。
623 名前：デフォルトの名無しさん mailto:sage [2011/04/28(木) 02:08:04.55 ]: >>622
スペースきれい！
ふしぎ！
624 名前：デフォルトの名無しさん mailto:sage [2011/04/28(木) 02:10:39.48 ]: >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の根本的な文法も理解しようとせず使おうとしてるから、バレないし問題ないよね。
625 名前：デフォルトの名無しさん mailto:sage [2011/04/28(木) 20:52:23.73 ]: >>617
VC2010はCUDA 4.0 RC2から対応で、SDKにテンプレートが付属してる。
うちの環境だとMSBuild関連が正しくインストールされなかったのでひと手間必要だったけど。
626 名前：デフォルトの名無しさん mailto:sage [2011/04/29(金) 23:49:38.06 ]: CUDAのコンパイラってレジスタ不足でもコンパイル通っちゃうのか？
VSで開発してるんだが、どうもレジスタが足りないとFirst chance exceptionエラーが出る気がしてきた。
627 名前：デフォルトの名無しさん mailto:sage [2011/04/29(金) 23:58:25.76 ]: >48
祈るってアプローチが間違いとは言わないけど、
同じ計算を2並列で行うとか、間違いなく構成を変えられる方法があれば
演算器を入れ替えて2回行うとかって手もあるよ。
2倍時間がかかってもまだCPUには勝てるでしょ。
628 名前：デフォルトの名無しさん mailto:sage [2011/04/30(土) 00:29:11.48 ]: PowerDirectorなどCUDAを使ってエンコードを高速化出来るアプリがありますは、SLI環境ではさらに高速化出来るのでしょうか？
ググると出来ないというソースが目立ちますが...
X58マザー使っています。
629 名前：デフォルトの名無しさん mailto:sage [2011/04/30(土) 00:35:04.10 ]: レジスタ不足ってローカルメモリ使われるんじゃないの?
630 名前：デフォルトの名無しさん mailto:sage [2011/04/30(土) 01:12:10.48 ]: >>628
cuda4はどうか知らないが、それ以前のcudaではSLIだからより速くできると言うことはない。
勿論GPU2枚挿しに対応したアプリなら、2枚刺した方が速くはなるがSLIには関係ないようだ。
631 名前：デフォルトの名無しさん mailto:sage [2011/04/30(土) 03:38:04.89 ]: >629

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

cuFFTとかの標準ライブラリでも動作が怪しかったりするのは
レジスタ周りの取り扱いのせいなんじゃって気がしてきてるのですよ。
632 名前：デフォルトの名無しさん mailto:sage [2011/04/30(土) 08:15:07.06 ]: 引数はCC2.0からはコンスタントメモリで
それ以前はシェアードだったかな

関数内変数はレジスタで不足したら勝手にグローバルにいくんでしょ
ワープ数が上限ぴったりならエラーは吐かないと思うんだけど
633 名前：デフォルトの名無しさん mailto:sage [2011/04/30(土) 08:23:40.36 ]: オーバーしても、普通は遅くなるだけだと思うんだけどなぁ。
634 名前：デフォルトの名無しさん mailto:sage [2011/04/30(土) 15:46:08.21 ]: >>631
そんな曖昧な知識で推測する前にPDFちゃんと読むんだ
635 名前：デフォルトの名無しさん mailto:sage [2011/04/30(土) 16:50:47.84 ]: >634
了解したぜ
636 名前：デフォルトの名無しさん mailto:sage [2011/04/30(土) 23:22:46.13 ]: >>630
ありがとう。
SLI環境はもう少し様子見します。。
637 名前：デフォルトの名無しさん mailto:sage [2011/05/02(月) 12:42:12.87 ]: 最近のグラボはATX電源から電源引っ張ってくんだな。
Fermiに変えたが早すぎワロタ。
638 名前：デフォルトの名無しさん mailto:sage [2011/05/02(月) 13:14:07.57 ]: それ以前からfermiって劇的な変化じゃね
キャッシュはでかい
639 名前：デフォルトの名無しさん mailto:sage [2011/05/08(日) 01:40:31.32 ]: FlashPlayerもCUDAつかえるの？
640 名前：デフォルトの名無しさん mailto:sage [2011/05/08(日) 05:13:16.28 ]: Tesla S2070ってなんで抹消されてしまったのでしょうか。
nVIDIAのページにもELSAのページにもなくなってしまったのですが。
641 名前：CUDA初心者 [2011/05/09(月) 11:25:03.07 ]: XPでFFTをかけるソースがあったのでvisualstudio2008で実行してみたところ
FFT.exe の 0x7c812afb で初回の例外が発生しました: Microsoft C++ の例外: cudaError_enum (メモリの場所 0x0012ae3c)。
というエラーが出てしまい困っています。ホストからデバイスへのメモリのコピー等
の簡単なプログラムは動きます。ちなみにGeforce210です。
642 名前：デフォルトの名無しさん mailto:sage [2011/05/09(月) 21:02:35.89 ]: debugを人に頼むような人間は一つ解決しても次に同じような所で引っかかるからきりがないわな。
643 名前：デフォルトの名無しさん mailto:sage [2011/05/10(火) 09:55:41.04 ]: >>641は何も頼んでない。日記を書いてるだけじゃないか？
644 名前：デフォルトの名無しさん mailto:sage [2011/05/10(火) 22:29:43.44 ]: >641
多分それ、cuFFTのバグ。そのエラーコードはメモリアクセス例外のときに出る。
俺の場合9800GTから550Tiに買い換えたらエラーが消えた。

cuFFTがデバイスメモリ使い過ぎて、デバイスが積んでるメモリ量によっては
エラーを吐くとかが原因の気がしてるんだが、本当の理由はわからん。
NVidia側のバグフィックスを待つぐらいしか現状で対策はないと思う。
645 名前：デフォルトの名無しさん mailto:sage [2011/05/11(水) 00:23:29.41 ]: cuFFTのバグなら、サイズが2の冪かどうかでも変わるね。
cuFFTの内部で結構デバイスメモリを使うようだから、ボードの半分以下のFFTしかかけられない。
646 名前：デフォルトの名無しさん mailto:sage [2011/05/11(水) 19:42:26.40 ]: >>642
FFTなんて本来ライブラリの側で対処すべき次元のもので、
ユーザー側がデバッグに煩わされるべきものではないがな。
647 名前：デフォルトの名無しさん mailto:sage [2011/05/11(水) 19:57:25.35 ]: >>646
もしバグだとおもうならフォーラムで問い合わせたり報告したりしたほうが良いんじゃないか？
バグなら修正してもらえる可能性があるし、そうじゃないとしたら自分の間違いだとわかる
648 名前：デフォルトの名無しさん mailto:sage [2011/05/11(水) 23:13:45.45 ]: 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.
と出てしまい実行出来ません。どうしたら良いでしょうか？
649 名前：デフォルトの名無しさん mailto:sage [2011/05/11(水) 23:20:38.40 ]: >>648
www.alc.co.jp/
650 名前：デフォルトの名無しさん mailto:sage [2011/05/11(水) 23:33:32.11 ]: nvidia-current カーネルを削除してSDKを再コンパイルで行けました
ありがとうございます
651 名前：デフォルトの名無しさん mailto:sage [2011/05/11(水) 23:35:55.99 ]: nvidia-current カーネルを削除してSDKを再コンパイルで行けました
ありがとうございます
652 名前：デフォルトの名無しさん mailto:sage [2011/05/12(木) 00:05:11.00 ]: >>647
フォーラムには既にあったと思う。
653 名前：CUDA初心者 [2011/05/12(木) 15:13:05.97 ]: >>644
ありがとうございます。
550Ti持ってるのでそっちでやってみようと思います。
654 名前：デフォルトの名無しさん [2011/05/16(月) 10:53:35.28 ]: Developer Drivers for Linuxをインストールすると、
Xwindowの解像度が640*480固定になるのはどうしてなのかしら？
655 名前：デフォルトの名無しさん mailto:sage [2011/05/16(月) 10:57:58.45 ]: おまえがばかだからじゃないか？
656 名前：デフォルトの名無しさん mailto:sage [2011/05/16(月) 11:00:18.75 ]: xの設定書き換えちゃってるんじゃないの?
657 名前：デフォルトの名無しさん [2011/05/18(水) 12:36:20.91 ]: Ubuntu 11.04 CUDA 4.0でドライバをインストールしようとすると、
The Nouveau kernel driver is currently in use by your system.
というエラーが出て、先に進めません。
どうしたらNouveau kernel driverをとりのぞけますか?
658 名前：デフォルトの名無しさん mailto:sage [2011/05/18(水) 12:42:47.54 ]: これ読んでみ
https://help.ubuntu.com/community/BinaryDriverHowto/Nvidia
659 名前：デフォルトの名無しさん [2011/05/18(水) 13:18:46.93 ]: >>658
早速の返信ありがとうございます。
sudo apt-get --purge remove xserver-xorg-video-nouveau
と打って、再起動もしてみましたが、やはり同じメッセージが帰ってきます。
ほかに、何か手立てはないものでしょうか?
660 名前：デフォルトの名無しさん mailto:sage [2011/05/18(水) 21:39:19.83 ]: >>657
俺がLive CDで試したときは
kernelの引数にnouveau.modeset=0として起動後に
modprobe -r nouveauした。

そもそもnouveauがモジュールではなくて組み込みになってるなんてことはないよね？
661 名前：デフォルトの名無しさん [2011/05/19(木) 04:06:30.55 ]: インテルコンパイラには対応していないの?
662 名前：デフォルトの名無しさん mailto:sage [2011/05/19(木) 04:12:46.19 ]: Coalesced Accessって簡単にいうとどゆうこと
663 名前：デフォルトの名無しさん mailto:sage [2011/05/19(木) 08:14:50.08 ]: >>662
１６人１７脚で走ること
ただし、一番端っこの人の番号に制限あり
664 名前：デフォルトの名無しさん mailto:sage [2011/05/19(木) 08:28:38.02 ]: >>663
なんだかすごく遅くなりそうな例えだなｗ
665 名前：デフォルトの名無しさん mailto:sage [2011/05/19(木) 09:19:43.15 ]: 論文書くときにさ、遠くの図書館に借りに行くわけよ
16冊まで借りられるから自分の欲しいタイトル1冊と隣に並んでるの15冊借りて帰って来るの
そん中にフラットメイト15人が欲しい本が入ってたらラッキー
何人か足りなかったらまたまた誰か派遣して16冊借りてこなくちゃなんない
一回で16人全員満足したり16冊全部が有効に使われたらうれしいな

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

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

てのはどうだろう
666 名前：デフォルトの名無しさん mailto:sage [2011/05/19(木) 09:33:17.23 ]: >>663はCoalesingの説明になってないぞ。

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

もちろん正確ではないが速くなりそうな例ではある。
667 名前：デフォルトの名無しさん [2011/05/19(木) 15:22:10.76 ]: 初歩的な質問で恐縮なのですが、
皆さんCUDA版プログラムとCPU版プログラムの速度比較は
どのようにされてますでしょうか？

CUDA3.0以降はエミュレーションモードがなくなったようなので、
別にCPU版プログラムも作って比較するのがよいでしょうか？
668 名前：忍法帖【Lv=25,xxxPT】 mailto:sage [2011/05/19(木) 16:17:52.26 ]: エミュレーションと比較して何の意味があるの？
NVIDIAの回し者でない限り、iccでカリカリにチューンしたプログラムと比較するべし。
669 名前：忍法帖【Lv=26,xxxPT】 mailto:sage [2011/05/19(木) 18:19:14.33 ]: CPU版と比較しても仕方が無い。
iccでチューニングした上で、更にその先の為にcudaを使うのだから。
まして、エミュレーションだなんてなんの悪夢だか。
670 名前：デフォルトの名無しさん mailto:sage [2011/05/19(木) 21:39:40.22 ]: >662
メモリ読み書きするなら32byte, 64byte, 128byte単位にしてください。
これだけ。
671 名前：デフォルトの名無しさん [2011/05/20(金) 14:19:17.96 ]: SDK,toolkitを再インストールしたところ
サンプルを実行すると
cudart32_31_9.dllが見つからなかったため、このアプリケーションを開始できませんでした
。cudartを間違えて消してしまったと思われるのですが、どうやったら解消されるのでしょうか？
ぐぐってみたのですが、わかりませんでした。
672 名前：デフォルトの名無しさん mailto:sage [2011/05/20(金) 14:31:37.09 ]: PCを再起動させてみたらどう
673 名前：デフォルトの名無しさん mailto:sage [2011/05/20(金) 15:41:59.99 ]: 最悪でも再インストールで済むだろ。
674 名前：デフォルトの名無しさん [2011/05/20(金) 17:07:36.68 ]: 再インストールしてみたのですが変わりませんでした泣
675 名前：デフォルトの名無しさん mailto:sage [2011/05/20(金) 17:37:59.88 ]: >>674
よかったな、cudartを間違って消してしまったわけじゃないことが判ったじゃないか。
676 名前：デフォルトの名無しさん mailto:sage [2011/05/27(金) 03:06:11.39 ]: CUDA 4.0
677 名前：デフォルトの名無しさん [2011/05/27(金) 11:09:49.93 ]: __CUDACC__ってもう定義されていないの?
678 名前：デフォルトの名無しさん mailto:sage [2011/05/28(土) 08:21:44.26 ]: みんな業務でCUDAつかってんの？
679 名前：デフォルトの名無しさん mailto:sage [2011/05/28(土) 08:56:58.21 ]: 一部使ってますが。
680 名前：デフォルトの名無しさん mailto:sage [2011/05/28(土) 21:17:04.74 ]: CUDA4.0、GTX480M + Notebook Developer Drivers for WinVista and Win7 (270.61)では
device query含め、CUDA Cのサンプルが全部動かない。
OpenCLとDirectComputeは動く。
なんでだろう。
681 名前：忍法帖【Lv=1,xxxP】 mailto:sage [2011/05/29(日) 11:17:41.19 ]: 坊やだからさ
682 名前：デフォルトの名無しさん mailto:sage [2011/05/29(日) 11:48:30.13 ]: VS2010とCUDA4.0の環境でrulesファイルを使えた方おりませんか？
$(CUDA_PATH)\extras\visual_studio_integration\rules
にあるファイルをコピーしたのですが、これを設定してプロジェクトを読み込みなおすと
「要素 <UsingTask> 内の属性 "AssemblyFile" の値 "$(CudaBuildTasksPath)" を評価した結果 "" は無効です」
といったエラーが出てプロジェクトが読み込めないようです…。
683 名前：デフォルトの名無しさん mailto:sage [2011/05/29(日) 21:12:48.05 ]: 完全な逐次処理（どう頑張っても並列化できない）のを、
CUDAで高速化してちょ、といわれたんだけど殴っていい？
684 名前：デフォルトの名無しさん mailto:sage [2011/05/29(日) 21:18:05.50 ]: >>683
いや無理にでもやれよ
それが仕事ってもんだ
685 名前：デフォルトの名無しさん mailto:sage [2011/05/29(日) 21:27:29.77 ]: ゼロコピーってキャッシュされないでしょ。

んで、ゼロコピーを使うと遅くなるケースが出てきたんで、
仕方なく使わないことにしたんだけど、
やっぱりcudaMallocやcudaMemcpyの時間がもったいないんで
どうにかしたいんだけど、なんか知恵ある？
686 名前：デフォルトの名無しさん mailto:sage [2011/05/29(日) 21:53:26.07 ]: >>683
どう頑張っても並列化できないを証明すれば良いじゃん．
なかなか難しいよ．
687 名前：デフォルトの名無しさん mailto:sage [2011/05/29(日) 22:08:28.81 ]: >>685
ゼロコピーって何よ。
何をしたいのか判らんけど、cudaMemset()も使えない内容なの?
688 名前：682 mailto:sage [2011/05/29(日) 22:53:30.43 ]: 解決しました。ちょうど公式フォーラムで同じ問題に遭遇してた人が居たようです。
先ほどあった投稿で解決することができました
forums.nvidia.com/index.php?showtopic=201433
689 名前：デフォルトの名無しさん [2011/05/29(日) 23:36:05.38 ]: もふ。
マジレスするとゼロコピーとは、cudaHostGetDevicePointer()を使うやつのことだす。
cudaMallocHost()でホスト側のメモリを確保しとかないとダメだけど。
690 名前：デフォルトの名無しさん mailto:sage [2011/05/30(月) 01:01:06.08 ]: ゼロコピーと言うのは知らんかった。
処で、勿体無いと言うけどcudaMalloc()って時間掛かる?
691 名前：デフォルトの名無しさん mailto:sage [2011/06/01(水) 01:01:48.66 ]: 演算速度について質問です。
プログラミングガイドによるとfloatにおけるmulとmadの速度は同一だとあるのですが、
for(i=0;i<N;i++)
a = b*c;

と

for(i=0;i<N;i++)
a = d*(b+c);

では明らかに前者のほうが短い時間で処理を行っています。
dをコンスタンとメモリや定数にしても同じでした。
いったいどういうことなんでしょうか？
助けてください。
692 名前：デフォルトの名無しさん mailto:sage [2011/06/01(水) 01:07:21.88 ]: >>691
どこのプログラミングガイドだよ教えろよ教えてくださいお願いします
693 名前：デフォルトの名無しさん mailto:sage [2011/06/01(水) 01:12:04.19 ]: madって(A*B)+Cじゃないの？
694 名前：デフォルトの名無しさん mailto:sage [2011/06/01(水) 01:14:43.17 ]: すみませんまったく環境とか書いてませんでした。

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.

だそうです。
695 名前：デフォルトの名無しさん mailto:sage [2011/06/01(水) 01:27:04.68 ]: >>693
ごめんなさい書き間違えました。

for(i=0;i<N;i++)
a = d*b+c;

です。
696 名前：忍法帖【Lv=2,xxxPT】 mailto:sage [2011/06/01(水) 02:09:02.07 ]: まぁ、ptxを貼ってみろ。
697 名前：デフォルトの名無しさん mailto:sage [2011/06/02(木) 01:22:02.76 ]: 積和演算なら
a = b*c+d*e
とかじゃないのん？畳込みとか行列積とかで頻出する計算です。
698 名前：デフォルトの名無しさん mailto:sage [2011/06/02(木) 01:30:56.84 ]: DSPの命令見てみたら、積和演算は
a = a + (b*c)

だったわ。
a+= (b*c)

でもええかもしれんが。
699 名前：デフォルトの名無しさん mailto:sage [2011/06/02(木) 03:02:35.87 ]: >>680
270.81でないとCUDAランタイムAPIが変な動作をするみたい
つまり現時点ではPC用のドライバしかない
700 名前：デフォルトの名無しさん mailto:sage [2011/06/02(木) 11:07:00.65 ]: 270.61に
「Unified Virtual Addressing (UVA) やGPUDirect? v2.0を特徴とするCUDA 4.0を使ったアプリケーション用サポートの追加」
と書いておきながら動かないなんてふざけた話だわ
701 名前：デフォルトの名無しさん mailto:sage [2011/06/02(木) 22:17:32.06 ]: CUDA勉強はじめました。
でも世の中に出てる日本語のCUDAの本って4冊ともFermi以前のアーキテクチャでの説明だからsharedメモリが16kbyteだったり、1ブロックあたりのスレッド数が512だったりするんで困ります。
なんとかして。
702 名前：691 mailto:sage [2011/06/02(木) 22:33:23.38 ]: コンスタントメモリや定数の値をレジスタに読み込む時間をまったく無視していました。
たぶんそのせいで遅くなっているんだと思います。
ご迷惑かけてすみません。
703 名前：デフォルトの名無しさん mailto:sage [2011/06/02(木) 22:59:17.85 ]: >>700
ちょうど新しいドライバが出たみたいです
704 名前：デフォルトの名無しさん mailto:sage [2011/06/02(木) 23:33:53.87 ]: >>701
英語も勉強汁
705 名前：デフォルトの名無しさん mailto:sage [2011/06/02(木) 23:44:11.15 ]: oh! i can't read english. :-P
706 名前：デフォルトの名無しさん mailto:sage [2011/06/03(金) 00:27:22.61 ]: >>701
Fermi用の正しい値が分かる人なら何も困らないような気が・・・・。
707 名前：デフォルトの名無しさん mailto:sage [2011/06/03(金) 22:26:32.06 ]: Compute Visual Profilerを使いこなしてる人っている？
あれ、結局、どう活用していかわからん。
708 名前：デフォルトの名無しさん mailto:sage [2011/06/04(土) 10:50:31.43 ]: >>701
Fermiの方が簡単になってるからいんじゃね

>>707
ローカルメモリ使われてたりバンクコンフリトあるかをチェックすんじゃね
709 名前：デフォルトの名無しさん mailto:sage [2011/06/04(土) 11:02:21.92 ]: 北川景子かわゆす
710 名前：デフォルトの名無しさん mailto:sage [2011/06/04(土) 11:18:35.39 ]: そうか？
711 名前：デフォルトの名無しさん mailto:sage [2011/06/04(土) 20:06:16.65 ]: あれ？4.0ってcutil64D.libなくなっちゃった？
712 名前：デフォルトの名無しさん mailto:sage [2011/06/04(土) 23:19:09.39 ]: 本当だなくなってる
713 名前：デフォルトの名無しさん mailto:sage [2011/06/05(日) 00:30:29.73 ]: 　
714 名前：デフォルトの名無しさん mailto:sage [2011/06/05(日) 00:56:21.39 ]: 自前でビルドしなさい
715 名前：デフォルトの名無しさん mailto:sage [2011/06/05(日) 04:04:47.93 ]: 今Cを勉強中ですが、CUDAと並行して勉強できますか？
716 名前：デフォルトの名無しさん mailto:sage [2011/06/05(日) 10:23:06.12 ]: >>715
少なくとも小中で算数（数学）や社会、英語などを並行して勉強してきたはず
717 名前：デフォルトの名無しさん mailto:sage [2011/06/05(日) 10:42:05.41 ]: >>715
CUDAのプログラム記法以外に知るべき事が多い。と言うかむしろそちらが大半じゃないか。並行してできる、がんばれ
718 名前：デフォルトの名無しさん mailto:sage [2011/06/05(日) 18:19:38.33 ]: どっちもものにならないと思う。どっちかだけやればものになるものでもないけど。
719 名前：デフォルトの名無しさん mailto:sage [2011/06/05(日) 18:43:41.04 ]: 4.0でコンパイルが通らないの結構多いね
CUDA BY EXAMPLEのサンプルをいくつか試したがエラーをはく。
まだ3.2使うわ
720 名前：デフォルトの名無しさん mailto:sage [2011/06/05(日) 19:14:22.50 ]: 4.0 から cutil32.lib とかって自分でビルドするようになったのね。
インストール後に検索しても見つからないから、変だな～と思った。
Direct Compute の Example が動かないのもよくわからん。
721 名前：デフォルトの名無しさん mailto:sage [2011/06/06(月) 09:41:55.58 ]: ある程度C(というかプログラミング)しってなきゃ百パー無理だろ
722 名前：デフォルトの名無しさん mailto:sage [2011/06/06(月) 13:37:59.32 ]: %を略すな、池沼
723 名前：デフォルトの名無しさん mailto:sage [2011/06/06(月) 13:54:50.09 ]: 。
724 名前：デフォルトの名無しさん mailto:sage [2011/06/06(月) 21:16:41.99 ]: >>721
それを知らなかったので、PATHの設定やmake実行さえ一苦労でした。
人によっては、ディレクトリの概念が難しい場合もあるかもしれません。
725 名前：デフォルトの名無しさん mailto:sage [2011/06/06(月) 21:26:48.74 ]: cudaはただでさえ難しいぞ
初心者に手にを得る代物じゃね～べよ
726 名前：デフォルトの名無しさん mailto:sage [2011/06/06(月) 21:32:48.38 ]: 資料も英語しかないしかなりきつかったわ
727 名前：デフォルトの名無しさん mailto:sage [2011/06/08(水) 14:58:17.39 ]: OpenCLとどっちが楽ですか？
728 名前：デフォルトの名無しさん mailto:sage [2011/06/08(水) 16:43:07.02 ]: Cudaの方が楽
OpenCLはOpenGLと同様、自助努力が基本になる
729 名前：デフォルトの名無しさん mailto:sage [2011/06/09(木) 05:38:12.86 ]: OpenCLとどっちが速いですか？
730 名前：デフォルトの名無しさん mailto:sage [2011/06/09(木) 06:45:26.86 ]: 自助努力って具体的に何？
731 名前：デフォルトの名無しさん mailto:sage [2011/06/09(木) 10:03:14.93 ]: >>423
亀だけど、cufftPlanは対象とするデータのサイズと同じサイズをバッファとして確保するみたい。だから、二倍以上メモリが空いてないとこける。
732 名前：デフォルトの名無しさん mailto:sage [2011/06/09(木) 17:23:54.07 ]: >>730
そういうのを自分で考えたり調べたりする姿勢や行動
733 名前：デフォルトの名無しさん mailto:sage [2011/06/09(木) 18:41:25.08 ]: OpenGL de プログラミングというサイトや
CUDAプログラミング入門とかフリーであるもの
見たほうがそこらの書籍よりかなり役立つよ

OpenCLは少し大変だけどCPUならdouble型
オプションあるからいいんだよね
ＣＵＤＡにもあるのかな?
734 名前：デフォルトの名無しさん mailto:sage [2011/06/09(木) 18:58:20.29 ]: あった、ごめん
735 名前：デフォルトの名無しさん mailto:sage [2011/06/11(土) 18:37:21.68 ]: visual studio pro, radeon 6000 台　で ati stream ないし　open cl
使っての　並列FPU高速化ってどのくらい大変ですか？
一般のプログラムと英語ドキュメント読解に問題はないです
736 名前：デフォルトの名無しさん mailto:sage [2011/06/11(土) 18:40:16.57 ]: スレタイ読む能力はなかったようだな
737 名前：デフォルトの名無しさん mailto:sage [2011/06/11(土) 19:04:07.23 ]: 間違えました。すいません。
738 名前：デフォルトの名無しさん mailto:sage [2011/06/12(日) 06:56:49.05 ]: 漫才かw
739 名前：デフォルトの名無しさん mailto:sage [2011/06/12(日) 19:39:36.69 ]: >>732
CUDAを勉強してみようかと少し取り組んでみましたが、
C言語の基本も知らぬプログラミング初心者なので
C言語から勉強しようと思ってます。
740 名前：デフォルトの名無しさん mailto:sage [2011/06/12(日) 19:49:59.75 ]: Rやoctaveに触れると、Cの世界に戻れない。
故にCUDAに触れることも出来なくなる。
741 名前：デフォルトの名無しさん mailto:sage [2011/06/12(日) 19:55:24.78 ]: Rでモンテカルロ出来るんだっけ？
742 名前：デフォルトの名無しさん mailto:sage [2011/06/12(日) 20:12:36.74 ]: >>739
C言語の文法を知ることは必須としても、まずプログラムがどういうもので何が出来るものなのか、って視点からも眺めた方がいいよ。
文法を覚えることに溺れないように
743 名前：デフォルトの名無しさん mailto:sage [2011/06/12(日) 20:32:53.39 ]: >>741
統計解析ソフトだから乱数の生成は得意だと思うけど、
ループが苦手だから本当にモンテカルロが必要な場面でが役立たずかも。

Rユーザから見ると、curandライブラリがRから呼び出せたら嬉しい。
744 名前：デフォルトの名無しさん [2011/06/13(月) 16:13:05.08 ]: カーネルを起動するのにかかるコストってNVIDIAの資料等に載ってるんでしょうか？
745 名前：デフォルトの名無しさん mailto:sage [2011/06/13(月) 17:09:26.42 ]: 資　料　読　め　よ
746 名前：デフォルトの名無しさん mailto:sage [2011/06/13(月) 20:52:40.31 ]: 　　　　　　　　　　　YES　→　【見つかった？】　─　YES → じゃあ聞くな死ね
　　　　　　　　　／　　　　　　　　　　　　　　　　　＼
【資料見た？】　　　　　　　　　　　　　　　　　　　　　　　 NO → なら、ねぇよ
　　　　　　　　　＼
　　　　　　　　　　　 NO → 死ね
747 名前：デフォルトの名無しさん mailto:sage [2011/06/13(月) 22:35:02.62 ]: 日本語の試料ないですか？
頭痛いです。
748 名前：デフォルトの名無しさん mailto:sage [2011/06/13(月) 22:46:41.93 ]: 　　　　　　　　　　　YES　→　【見つかった？】　─　YES → じゃあ聞くな死ね
　　　　　　　　　／　　　　　　　　　　　　　　　　　＼
【探した？】　　　　　　　　　　　　　　　　　　　　　　　 NO → それより僕と踊りませんか？
　　　　　　　　　＼
　　　　　　　　　　　 NO → 死ね
749 名前：デフォルトの名無しさん mailto:sage [2011/06/13(月) 22:56:54.15 ]: >>746
くだすれなんだし並列に書いてみてはどうか
750 名前：デフォルトの名無しさん mailto:sage [2011/06/13(月) 23:03:14.51 ]: >>749
ifが多すぎてパフォーマンスが落ちると思う
751 名前：デフォルトの名無しさん mailto:sage [2011/06/13(月) 23:07:55.03 ]: >>750
嫁>>1
752 名前：デフォルトの名無しさん mailto:sage [2011/06/14(火) 09:21:34.64 ]: >>1の序文が全く無視されているな
753 名前：デフォルトの名無しさん mailto:sage [2011/06/14(火) 09:45:27.25 ]: >>747
日本語の書籍購入するのが早いとおもうよ。
・ASCII.technologies (アスキードットテクノロジーズ) 2010年 08月号
・はじめてのCUDAプログラミング
・CUDA by Example 汎用GPUプログラミング入門

この辺がお勧め。自分はこの順で読んだ。
でもプログラミングガイドとか読むのは避けられないはず。

・CUDAプログラミング実践講座 - 超並列プロセッサにおけるプログラミング手法
これは買ってないけど良さそう。
754 名前：デフォルトの名無しさん mailto:sage [2011/06/14(火) 09:47:57.13 ]: >>753
上から三つ目は一つ目と二つ目が参考資料だっからな
アスキーは触りとしてはわかりやすかった
どれにしろFermiについて書いてないけどね
755 名前：デフォルトの名無しさん mailto:sage [2011/06/14(火) 09:55:59.01 ]: 4.0になると、もうどうしましょ
756 名前：デフォルトの名無しさん mailto:sage [2011/06/14(火) 20:59:49.06 ]: ■後藤弘茂の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の間での無駄なメモリコピーを排除するゼロコピーが実現できる。
757 名前：デフォルトの名無しさん mailto:sage [2011/06/14(火) 22:14:25.10 ]: あとはメインメモリがGDDR5になるのを待つだけだな．
758 名前：デフォルトの名無しさん mailto:sage [2011/06/14(火) 22:50:31.08 ]: また妙なことを始めたなAMDw
759 名前：デフォルトの名無しさん mailto:sage [2011/06/14(火) 22:59:19.26 ]: メモリコピーがネックにならない計算ってむしろおかしい？
計算の方が時間かかってるんだけど
もっと最適化できそうとか
760 名前：デフォルトの名無しさん mailto:sage [2011/06/15(水) 01:56:49.38 ]: 利用しているデータ量が少ないなら別におかしくはないと思うけど。
761 名前：デフォルトの名無しさん [2011/06/15(水) 14:58:40.10 ]: SDK4.0(x32)のtemplateを別フォルダにコピーしてVS2008でビルドしようとしたらcutil32D.libが開けませんって出ます。
助けてください。
762 名前：デフォルトの名無しさん mailto:sage [2011/06/15(水) 16:01:33.65 ]: >>761
720
763 名前：デフォルトの名無しさん mailto:soab_b_nb@yahoo.co.jp [2011/06/15(水) 17:10:58.62 ]: >>762
レスありがとうございます。
cutilとshrUtilsは既にビルドはしているんですが、開けないとでるんですよ。。
764 名前：デフォルトの名無しさん mailto:sage [2011/06/15(水) 21:41:20.65 ]: >759
メモリコピーの時間がかかってでも並列演算したいって問題がほとんどだから
別におかしくないんじゃない。
並列計算が早すぎてメモリコピーがネックのように見えるってことはあるけど。
765 名前：デフォルトの名無しさん mailto:sage [2011/06/15(水) 22:05:59.38 ]: メモリのコピー時間は0でも
同時にはアクセスできないよね、きっと。
でも大した問題でもないか。
766 名前：デフォルトの名無しさん mailto:sage [2011/06/16(木) 03:21:57.35 ]: Microsoft Going All-in on GPU Computing ≪ NVIDIA
blogs.nvidia.com/2011/06/microsoft-going-all-in-on-gpu-computing/?sf1642229=1
767 名前：デフォルトの名無しさん mailto:sage [2011/06/16(木) 06:03:44.90 ]: >>754
アスキーは2冊あって2010/08はFermiについても掲載されてるよ
そのほかにOpenCL、Direct compute, ATI streamについても載ってる
768 名前：デフォルトの名無しさん mailto:sage [2011/06/16(木) 15:34:19.10 ]: 物凄く基礎的な質問ですいません。
今大学でCUDAの勉強を始めた者です。
現在貸し出されたPCに入っていたのがGeForce7600GSで、
CUDAの公式を見る限りだとCUDAに対応していないようですが
どうにかして動かすことは出来ないでしょうか？

一応一番初期のVer1.0～最新の4.0までひと通り試してみたのですが
サンプルのdeviceQueryを動かしてもCUDA Driver Versionは0.0のままでした。

もし可能でしたら少しでも助言を・・・
無理なら無理と言ってくだされば大学側に申請して
可能なGPUの購入を考えてます。
769 名前：デフォルトの名無しさん mailto:sage [2011/06/16(木) 15:37:06.49 ]: CUDAの勉強するならCUDAに対応したグラボを用意して貰え
対応してない物が仮に動いたとしても、それが正確な動作かわからんから勉強にならんぞ
770 名前：デフォルトの名無しさん mailto:sage [2011/06/16(木) 15:54:45.30 ]: cuda x86
Amazon EC2

GPU買わなくてもおｋ
771 名前：デフォルトの名無しさん mailto:sage [2011/06/16(木) 16:13:17.23 ]: >>768
安いのなら1万切るから自分でグラボ買ってもいいんじゃない？
772 名前：デフォルトの名無しさん mailto:sage [2011/06/16(木) 16:48:57.04 ]: >>771
大学研究室ならキッチリ金出してくれて580くらいはポンと買ってくれるだろうから自分で買う意味ないだろう
773 名前：デフォルトの名無しさん mailto:sage [2011/06/16(木) 18:37:25.67 ]: ゲーム用と思われるのもあれだからTesla 20 Cシリーズ買ってもらいなよ
774 名前：デフォルトの名無しさん mailto:sage [2011/06/16(木) 22:59:07.51 ]: >>768
授業？独学？
研究のわけないよね
GPGPUの研究でDX9のGPUとか遠回しなアカハラだｗ
775 名前：デフォルトの名無しさん mailto:sage [2011/06/17(金) 04:29:44.26 ]: はじめての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
776 名前：デフォルトの名無しさん mailto:sage [2011/06/17(金) 07:42:07.25 ]: >>775
横からだが、これは面白そう読んでみるわ
777 名前：デフォルトの名無しさん mailto:sage [2011/06/17(金) 10:15:21.27 ]: 同じく横からだが、ブクマした
全部印刷しておいても良いかもしれん
778 名前：デフォルトの名無しさん mailto:sage [2011/06/19(日) 02:32:29.65 ]: >>775
参考になりそうだ。DLできないの？
779 名前：デフォルトの名無しさん mailto:sage [2011/06/19(日) 10:39:21.50 ]: できますん
780 名前：デフォルトの名無しさん mailto:sage [2011/06/19(日) 22:33:54.54 ]: 二つ質問があります
①felmiにはキャッシュが搭載されているみたいですがプログラム中で明示しなくてもキャッシュにより高速化するのですか？

②GPUでcuda化して高速化が見込めるアプリケーションが載ってるサイトとか知りませんか？ソースコードつきで。できればC言語だとうれしいです
781 名前：デフォルトの名無しさん mailto:sage [2011/06/19(日) 22:41:52.94 ]: めんどく
だれかこたえてあげれや
782 名前：デフォルトの名無しさん mailto:sage [2011/06/19(日) 22:54:32.14 ]: この程度の事を自分で調べられず、聞いてしまう人間には
使いこなすのは無理だよ。
783 名前：デフォルトの名無しさん mailto:sage [2011/06/19(日) 23:07:34.93 ]: >>782
一応は調べました
せめて①だけでも答えてもらえると助かります
②に関しては自分もいくつか見つけましたが他の方がどういったサイトを使ってるか気になったので聞いてみました
784 名前：デフォルトの名無しさん mailto:sage [2011/06/19(日) 23:08:08.80 ]: 機種依存文字-10
綴り間違い-10
sage+3
投稿時間+10

内容に言及しなくても-7点だ
もうちょいがんばれ
785 名前：デフォルトの名無しさん [2011/06/19(日) 23:31:02.35 ]: 教えるきないなら書き込むなよ、暇人
786 名前：デフォルトの名無しさん mailto:sage [2011/06/19(日) 23:52:40.95 ]: >>785
ﾌﾟｯ
787 名前：デフォルトの名無しさん mailto:sage [2011/06/19(日) 23:52:54.57 ]: キャッシュは勝手に使われるが、
キャッシュによって高速化(何と比べて？)するかはアプリによる。

ttp://journal.mycom.co.jp/articles/2010/07/21/fermi_cache/index.html
こんな話もあるから、キャッシュが如何使われるか位は
意識しないと駄目だし、究極を目指せば結局コアレスアクセスに
なってキャッシュの意味は無くなるかも知れない。
で、そこまで頑張れば微妙にteslaの方が速くなる。
788 名前：デフォルトの名無しさん mailto:sage [2011/06/20(月) 02:53:19.67 ]: ランダムアクセスでテストしてみたが（sm1.1でコンパイル）、コアレスと比較して、260では1/5まで速度が落ちたが、470は2/5だった
びっくらこいた
高速化できる例はCUDA SDKのサンプルみればよろし
789 名前：デフォルトの名無しさん mailto:sage [2011/06/20(月) 02:55:44.71 ]: 補足
G80/G92のように総レジスタ数の少ないGPUはもっと遅くなる。大体1/10
790 名前：デフォルトの名無しさん mailto:sage [2011/06/20(月) 07:37:03.23 ]: キャッシャが要らなきゃ共有メモリを増やせるし、Fermiは便利だね。
791 名前：デフォルトの名無しさん mailto:sage [2011/06/22(水) 16:10:37.96 ]: fermiで２ワープを２サイクルで処理していると記載を見かけるのですが、なぜ１サイクル１ワープでなく2サイクル２ワープなのでしょうか？
どういう原理なのかご存知の方がいらっしゃいましたら教えていただけませんか。
また、ご存知でしたらその記載がどこにあるか教えていただけませんか。
お願いします。
792 名前：デフォルトの名無しさん mailto:sage [2011/06/22(水) 18:40:23.89 ]: 「fermi half warp」でググればいくらでも出てくる．

[ 新着レスの取得/表示 (agate) ] / [ 携帯版 ]

read.cgi ver5.27 [feat.BBS2 +1.6] / e.0.2 (02/09/03) / eucaly.net products.
担当:undef