08/03/03 10:32:18
8600GT WIN2Kの俺に謝れ！！！
XPが無い…
176:デフォルトの名無しさん
08/03/03 10:36:43
>>175
Win2kで動かないのであれば、WinXP買うのも馬鹿馬鹿しいね。ゴメン、配慮が足りなかった。
177:デフォルトの名無しさん
08/03/03 18:37:47
XPが嫌いってだけのもあるんだけどなー
ムダに重たくて
178:デフォルトの名無しさん
08/03/03 18:50:25
>>177
テーマをクラシックにするか、システムプロパティでパフォーマンスを優先にすればWin2Kに較べて重くないと思いますよ。
179:デフォルトの名無しさん
08/03/03 18:58:12
構ってちゃんの相手するなよ・・・
180:デフォルトの名無しさん
08/03/04 02:14:10
そもそもWindowsが糞高いし、ソース付いて来ないし。
不自由過ぎてハックする気が萎える。
181:デフォルトの名無しさん
08/03/04 17:32:33
こんちわす。
.NET環境からCUDAとデータの受け渡しをする例、なんてのは
無いもんでしょうかね。既存のVB.NET、C#.NETの統計解析を
高速化致したく。
相関係数＝Σ(xj-xk)^2÷√Σ(xj-xjの平均)^2*(xk-xkの平均)
なんてのが多発し、j,kは1～1000、それぞれ1万点ずつ、x1～x1000の全ての
組み合わせで計算、という具合です。でも「総和を取る」とかはGPUは
苦手でしょうかね・・・。
182:デフォルトの名無しさん
08/03/04 17:44:36
総和は苦手なので、適当に分割して小計を求めてから最後に総計を求めることになると思う。
それはさておき、先ずはその相関係数の式自体を変形してみてはどうだろう。
分子も分母のそれぞれも平均値を使わない形にできるはずだが。
# ついでに言えば分母の最後は2乗が抜けている希ガス。
それをやると、平均を使う2パス処理ではなくΣxj, Σxj^2, Σxk, Σxk^2,...などを求める1パスですむ筈。
あ、肝腎のVB/C#からの呼び出し方はわかんね。CUDA側は"*.obj"になるからリンクさえできればなんとかなるのかな?
まぁ、計算量が充分大きいならCUDA側を"*.exe"にして起動するようにしても大差ない気もするけど。
183:デフォルトの名無しさん
08/03/05 12:43:43
どうもどすえ。自分なりにお勉強してみました。
.NET←→CUDAは、UnmanagedMemoryStream/共有メモリ/セマフォ、とかで
なんとかなりそうです。典型的な処理は以下のようにすればよいのかなと。
・データは、(128×8)センサー{測定データ64個単位×256｝のように
CPU側で整形する。平均0、分散1に正規化し、欠損・異常値の前処理も
しておく。これはNオーダーの処理なので十数秒ですむ。
・GPU側には、C(j,k)＝(A(j,k)－B(j,k))^2 (j,kは0-7)
つまり8×8画素の画像の、差の二乗？を繰り返し計算させる。
D(j,k)＝D(j,k)＋C(j,k)で集計する。
センサー(0～1023番)対センサー(0～1023番)なんで、N^2オーダー。
今はここに数時間掛かってます。
・終わったらホスト側でD(j,k)を取り出して、画素を全部足す。
数時間が数分になってくれると嬉しいのですがｗ
「GPU Gem」第三版を注文したので、もうちょっとお勉強してみるす。
184:デフォルトの名無しさん
08/03/05 13:27:36
だから、相関係数を求める式を変形しろって。
CUDA云々以前だぞ。
185:デフォルトの名無しさん
08/03/05 14:02:20
え、相関係数を求める式にはもう「平均」は無いすよ。
ごめんどこ怒られてるのかわかんね・・・
186:デフォルトの名無しさん
08/03/05 14:29:04
「相関係数」でぐぐれかす
187:デフォルトの名無しさん
08/03/05 14:33:41
お前ら両方とももういい加減にしろ
188:デフォルトの名無しさん
08/03/05 16:49:01
統計の基本:
相関係数は普通ピアソンの積率相関係数を指す。
これは、r = Sxy / √(Sxx * Syy)で求められる。
このとき、Sxyはx,yの共分散、Sxxはxの分散、Syyはyの分散。
分散はSxx = Σ(xi - x￣)^2。Syyについても同様。
# ここではxの平均を便宜上x￣で現わす。
この式は、Sxx = Σxx - Σx*Σx/nに変形できる。
このとき、Σxxはxiの平方和、Σxはxiの総和。
また、共分散はSxy = Σ(xi - x￣)*(yi - y￣)。
この式は、Sxy = Σxy - Σx*Σy/nに変形できる。
このとき、Σxyはxiとyiの積の総和。
つまり、相関係数はr = (Σxy - Σx*Σy/n)/√((Σxx - Σx*Σx/n)*(Σyy - Σy*Σy/n))。
従って、線形時間で求められる。また、1回のスキャンで充分なので、各要素を保存する必要がない。
189:デフォルトの名無しさん
08/03/06 01:21:05
上のほうにもあるけど、総和の計算って実際どーすればいいんだ？
スレッドブロックごとに小計を求めて、各ブロックの結果をまとめるとかいう
方法になりそうなのは予想できるが、実際どんな風に書けばいいのかわからん。
具体的にはベクトルの長さが出したいんだが、各要素の２乗の和がうまく出せない。
ヒントでもいいので分かる方教えてください。
190:デフォルトの名無しさん
08/03/06 02:29:35
1つのブロックを総和の計算に割り当てる…とか
そんなことしか思い浮かばん
191:上の方の人
08/03/06 07:58:00
後で暇があったら実際のコードの集計部分だけ晒してみるか。
192:デフォルトの名無しさん
08/03/06 19:32:29
CUDAを使ったjpegの展開ライブラリはありませんか？
193:デフォルトの名無しさん
08/03/06 19:55:22
IDCTはそこそこ効率よさそうだけど前段のビットストリーム分解が不向きっぽいぞ。
194:デフォルトの名無しさん
08/03/06 22:20:39
>>191
そうしてくれると助かります。
よろしくお願いします。
195:デフォルトの名無しさん
08/03/06 23:43:14
デイビッド・カーク氏来日，CUDAカンファレンス2008開催
URLﾘﾝｸ(www.4gamer.net)
196:デフォルトの名無しさん
08/03/07 04:38:20
やっとできたらしい日本語版のCUDAZONEとプログラミングガイドの日本語版。
URLﾘﾝｸ(www.nvidia.co.jp)
しかし、機械翻訳そのままなのかな。余りに訳が酷くて原文読まないと意味が判らんとは……
197:デフォルトの名無しさん
08/03/07 15:34:38
CUT_EXIT(argc, argv)はなんでargc, argvを与えなくちゃいけないのでしょうか？
198:デフォルトの名無しさん
08/03/07 15:54:16
>>197
引き数にnopromptが指定されていない場合は
"Press ENTER to exit..."のメッセージを出して入力待ちにするためです。
詳しくは、cutil.h参照で。
199:デフォルトの名無しさん
08/03/07 18:23:26
floatであった変数をいくつかまとめてfloat4を使うようにしました。
そうしたところプログラムの実行速度が10%ほど落ちたのですが、float4よりfloatの方が
実行効率がよいのでしょうか？
200:デフォルトの名無しさん
08/03/07 23:32:28
よく知らないが、float4ってSIMD用の型じゃないの？
201:デフォルトの名無しさん
08/03/08 00:20:55
単にまとめただけじゃ速くならないと思う。一番いいのは-ptxして出力眺めることなんだけどね。
202:デフォルトの名無しさん
08/03/08 09:24:11
なんで、nvidiaのフォーラムには韓国語はあるのに日本語はないの？
203:デフォルトの名無しさん
08/03/08 11:42:50
流石にそれはnvidiaに聞いてくれ。需要が少ないと言うか、要望が少ないのだろ。
半島人と違って日本人は奥床しいからw
204:デフォルトの名無しさん
08/03/08 11:57:24
アメリカの地下鉄とかの案内板や自販機にハングルはあるけど
日本語はないとかいう話を思い出しました w
東大キャンパスであった、CUDAカンファレンス2008の動画とかは公開されないのでしょうかね？
205:デフォルトの名無しさん
08/03/08 21:17:16
ずいぶん前の話だが、20-26あたりで議論されてる問題って結局どんな風に
並列処理するのが最適なんだ？
総和の計算の話が出てたから思い出したのだが。
前に試行錯誤したが、高速化できなかったり、値がおかしくなったりして
あきらめてしまっていた。
どうぞ皆様知恵をわけてください。
この前のカンファレンスでも言われていたがチューニングにかける時間を
惜しんではいけないのだね。
206:デフォルトの名無しさん
08/03/08 21:27:59
総和みたいな計算は向いてないとも言ってたけどね
207:デフォルトの名無しさん
08/03/08 21:59:03
>>205
結局>20は、>24で書いたようにインデックスで並列にした。
要はこんな感じ。
--
大きな配列tmpがあるとき、下図の+で区切られた範囲ごとに集計したい。
tmp+-----+-----+--------+-----+---+...
つまり、1番目のセクションは最初の(先頭の)+から次の+まで、2番目のセクションはその次の+まで……
ここで、こんな構造体を考える。
struct sections {short begin, end;}
こいつの配列を作って次のように値をセットする。
{{0, 6}, {6, 12}, {12, 21}, {21, 27}, {27, 31}, ...}
これをデバイス側に転送しておいて、一つのデータスレッドが一つのセクションを担当する形で集計した。
この方法の問題点：
・セクションのサイズが不均衡なので、ワープ内でも不均衡だと無駄なからループが発生してしまう。
# 1ワープ内では同じインストラクションが走ってしまうため。つまり、使用効率が落ちる状態。
・セクションサイズがワープ内では極力等しくなるようにソートすると、今度はアクセス場所がランダムになる。
# 先程の配列が、例えば{{0, 6}, {6, 12}, {21, 27}, {301, 307}, ...}のようになってしまう。
いずれにしても、綺麗に並べることができない。
但し、今回は前段のtmp配列への演算結果の格納が所要時間の大半を占めたために適当にチューニングして放置。
尚、予告した集計のロジックは、総和ではなく↑とも違う小計算出だったので条件が違うと言うことでパス。
総計は興味があるので、その内テストできたら改めて晒そうと思う。
208:デフォルトの名無しさん
08/03/09 13:44:37
いきなりですが質問です
xp64でVisualStudio2005と8800GTS(G92)がありCUDAドライバー、
ツールキット、SDK、と入れたんですが、結局ビルド方法とか
どのPDFに書いてあんだろ、というのがよくワカンネですよ。
VS2005のプロジェクト設定をx64に変えて、cutil32D.libを
拾ってきたZIPから解凍して入れたら、mandelbrotとか一応
動いたけど。他のmandelbrot描画ソフトより劇速なんで、
動いてはいるみたいです。
209:デフォルトの名無しさん
08/03/09 16:10:18
うーん、IDEは使ってないから判らないなぁw
pdfじゃなくて、ReleaseNoteか何かにビルド方法かいてなかったっけ?
後で見てみるわ。
210:205
08/03/09 23:15:54
>>206
確かに言っていた。まぁそれは言われなくとも自明なことなんだが。
ついでに、CUDAがあるからGPUに向いていない処理も容易に書けるようになってしまった
とも言っていたな。
だが、膨大なデータをホストに書き戻して、総和を計算し、その結果をまたデバイスに
送って使うということは出来れば避けたかった。
GPUが苦手な処理でもデバイス側でやらせるというのがいいのか、転送コストを覚悟の上で
ホストに戻して計算したほうがいいのか実際確かめるために試行錯誤してみている。
>>207
説明ありがとう。
おかげで24でかかれていたことが理解できたよ。
その考え方は思いつかなかったな。
俺ももうちょい考えてみていい感じにできたら晒そうかと思う。
211:デフォルトの名無しさん
08/03/10 02:56:03
ベクトルの総和を出すscanlarge.pdfを見てみると、
elementsが65536ならCPU－GPUでほぼ一緒、
100万elementsでGPUがCPUの5倍速い、ようですね。
なるほどGPU向きではないな・・・
212:デフォルトの名無しさん
08/03/10 03:54:26
HostがCore2Duo辺りなら未だいいんだけど、Woodcrest辺りだと厳しいんだよねw
超越関数を大量に使う演算だと、GPUが随分有利になるんだけど。
# SSEには超越関数がない→ベクタ化のためには関数テーブルが必要→メモリ消費→キャッシュミス
213:デフォルトの名無しさん
08/03/10 11:09:00
cudaでfloat4*float4ってできないの？
各要素ごとに4行に分けて書かないとだめ？
214:デフォルトの名無しさん
08/03/10 18:33:07
>>213
float4 * float4 はないみたいだ。
ちなみに、
--
__device__ float4 a, b, c;
__device__ static void func()
{
c = make_float4(a.x * b.x, a.y * b.y, a.z * b.z, a.w * b.w);
}
--
こんなコードを書いてみたけど、
--
.global .align 16 .b8 a[16];
.global .align 16 .b8 b[16];
.global .align 16 .b8 c[16];
:
:
ld.global.v4.f32 {$f1,$f4,$f7,$f10}, [a+0]; //
ld.global.v4.f32 {$f2,$f5,$f8,$f11}, [b+0]; //
.loc 4 12 0
mul.f32 $f3, $f1, $f2; //
mul.f32 $f6, $f4, $f5; //
mul.f32 $f9, $f7, $f8; //
mul.f32 $f12, $f10, $f11; //
st.global.v4.f32 [c+0], {$f3,$f6,$f9,$f12}; //
--
こうなった。
215:デフォルトの名無しさん
08/03/11 19:17:03
FreeBSDのLinuxエミュでCUDAを使っているんだけど、
Linuxと比較してどれくらい速度が落ちているのかな?
ほぼネィティブと同程度ならこのまま使っていこうかと
思っているんだけど。
216:デフォルトの名無しさん
08/03/12 05:34:03
Windows上でコンパイルしているのですが、nmakeは使えないのでしょうか？
cygwinのmakeコマンドであればコンパイルできるようなのですが、nmakeを使うと
NMAKE : fatal error U1077: 'C:\CUDA\bin\nvcc.EXE' : リターン コード '0xc0000005'
Stop.
となってコンパイルできません。このエラーは何を意味しているのでしょうか？
217:デフォルトの名無しさん
08/03/12 05:43:10
float4ってsimdで計算してくれないの？
218:デフォルトの名無しさん
08/03/12 07:57:11
>>217
そもそもsimdなんてありませんが何か。つーか、>214で回答ついてるっしょ。
>>216
nvccが終了ステータスに5を積んでいるようですが、
ちゃんと動いた結果ならnmake側で無視することができると思います。
そもそもちゃんと動いてますか?
>>215
4種類ほどボード別に速度を測ったデータがあるからそれでも載せてみましょうか。
219:216
08/03/12 08:30:07
>>217
早速のご回答ありがとうございます。
Makefileの内容ですが、非常に初歩的な
test.exe: test.cu test_kernel.cu
nvcc test.cu
です。直接、コマンドプロンプトから打ち込んだ場合、問題なくコンパイルできます。
220:デフォルトの名無しさん
08/03/12 08:49:24
>>219
レス番違いますぜ。
取り敢えず、nvccがエラーステータスを返しているかどうかと、
nmakeでそれを無視できないか調べてみてはいかがでしょう。
221:デフォルトの名無しさん
08/03/12 12:49:19
そういうわけで、各ボードの比較表を作ったので貼ってみる。
--前半
"name","Tesla C870","GeForce 8800 GTX","GeForce 8800 GTS","GeForce 8800 GT","GeForce 8600 GTS","GeForce 8600 GT","使用コマンド"
"totalGlobalMem","0x5ffc0000","0x2ff50000","0x27f50000","0x1ffb0000","0xffb0000","0xffb0000","deviceQuery他"
"[MiB]",1535.75,767.31,639.31,511.69,255.69,255.69,
"sharedMemPerBlock","0x4000","0x4000","0x4000","0x4000","0x4000","0x4000",
"[KiB]",16,16,16,16,16,16,
"regsPerBlock",8192,8192,8192,8192,8192,8192,
"warpSize",32,32,32,32,32,32,
"memPitch","0x40000","0x40000","0x40000","0x40000","0x40000","0x40000",
"[KiB]",256,256,256,256,256,256,
"maxThreadsPerBlock",512,512,512,512,512,512,
"maxThreadsDim",512,512,512,512,512,512,
,512,512,512,512,512,512,
,64,64,64,64,64,64,
"maxGridSize",65535,65535,65535,65535,65535,65535,
,65535,65535,65535,65535,65535,65535,
,1,1,1,1,1,1,
"totalConstMem","0x10000","0x10000","0x10000","0x10000","0x10000","0x10000",
"[KiB]",64,64,64,64,64,64,
"major:minor",1:00,1:00,1:00,1:01,1:01,1:01,
"clockRate","1350000[kHz]","1350000[kHz]","1188000[kHz]","1512000[kHz]","1458000[kHz]","1188000[kHz]",
"textureAlignment","0x100","0x100","0x100","0x100","1x100","0x100",
222:デフォルトの名無しさん
08/03/12 12:49:57
--後半
"with Xeon",,,,,,,"使用コマンド"
"convolutionFFt2d","145.05[MPix/s]","107.99[MPix/s]","99.57[MPix/s]","107.00[MPix/s]",,,"convolutionFFT2D"
"convolutionRowGPU","1100.00[MPix/s]","1071.34[MPix/s]","717.34[MPix/s]","1457.11[MPix/s]",,,"convolutionTexture"
"cudaMemcpyToArray","1209.69[MPix/s]","1141.85[MPix/s]","1049.30[MPix/s]","1483.92[MPix/s]",,,
"simpleTexture","1180.83[MPix/s]","1125.08[MPix/s]","799.22[MPix/s]","910.22[MPix/s]",,,"simpleTexture"
"clockTest",24766,24930,25346,22396,,,"clock"
"with Core2Duo",,,,,,,"使用コマンド"
"convolutionFFt2d",,,,"106.64[MPix/s]","41.21[MPix/s]","32.93[MPix/s]","convolutionFFT2D"
"convolutionRowGPU",,,,"1433.39[MPix/s]","467.43[MPix/s]","375.54[MPix/s]","convolutionTexture"
"cudaMemcpyToArray",,,,"1475.25[MPix/s]","893.47[MPix/s]","596.14[MPix/s]",
"simpleTexture",,,,"873.01[MPix/s]","370.00[MPix/s]","285.30[MPix/s]","simpleTexture"
"clockTest",,,,24066,65006,71736,"clock"
223:デフォルトの名無しさん
08/03/12 13:58:27
これってプログラミングガイドの最後にあるやつ？
224:デフォルトの名無しさん
08/03/12 14:29:37
CUT_EXITを実行せずにプログラムを終了した場合、どのような悪影響がありますか？
225:デフォルトの名無しさん
08/03/12 15:01:43
>>221
SharedMemoryって16KBしか割り当てられないのですか？？
226:デフォルトの名無しさん
08/03/12 15:10:28
幾ら何でも同じスレの中くらい検索してくれたっていいじゃないか(TT
>>223
いいえ、サンプルプログラムの出力です。
>>224
なーーんにも。どうしても気になるならcutil.hを眺めてください。
>>225
この表では、1block辺りのSharedMemoryのサイズを示しています。
227:デフォルトの名無しさん
08/03/12 18:15:17
Cのmallocの場合、NULLかどうかでメモリが確保できたか判別できますが、
cudaMallocでメモリが確保されたかどうかを調べる方法を教えてください
228:デフォルトの名無しさん
08/03/12 18:31:03
CUT_CHECK_ERROR
229:デフォルトの名無しさん
08/03/12 18:31:16
cudaMalloc()
cudaError_t cudaMalloc(void** devPtr, size_t count);
allocates count bytes of linear memory on the device and returns in *devPtr a pointer to the allocated memory.
The allocated memory is suitably aligned for any kind of variable.
The memory is not cleared. cudaMalloc() returns cudaErrorMemoryAllocation in case of failure.
230:228
08/03/12 18:42:05
ああ、あとデバッグモードじゃないとだめだよ
231:デフォルトの名無しさん
08/03/13 03:10:17
warpサイズって簡単に言うと何？
232:デフォルトの名無しさん
08/03/13 06:48:58
同一インストラクションが同時に走る、演算ユニットの集まり?
233:デフォルトの名無しさん
08/03/13 13:50:16
CUDA初心者です。
CUDA難しいよ
難しいよCUDA
・<<<Dg, Db, sizeof(float)*num_threads*4>>> て書いたら動かないんすかね。
int sharemem_size = sizeof(float)*num_threads*4;
<<<Dg, Db, sharemem_size >>> て書いたら動いた。
・畳み込み演算させたら、CPUでやるより1.7倍しか速くないの(´・ω・`)
・動いてる間、システム全部固まってる。最初焦った。
CUDAの道険しいﾅﾘ
234:デフォルトの名無しさん
08/03/13 15:37:14
初歩的な質問かとは思いますが、各ストリームプロセッサにそれぞれ1つずつスレッドが割り当てられるのですか？
235:デフォルトの名無しさん
08/03/13 18:22:14
__syncthreads()ではブロック内でしか同期できませんが、全ブロックを同期したいときは
カーネルを抜けていったんCPU側に戻すしかないのでしょうか？
236:デフォルトの名無しさん
08/03/14 02:24:24
cudaMallocHostとCのmallocの違いを教えてください
237:デフォルトの名無しさん
08/03/14 05:10:46
>>235
確かカーネルは非同期で実行されたはずだから、
カーネルを連続で呼んだら同期されない・・・と思う。
<<<カーネルの呼び出し>>>
cudaThreadSynchronize();
<<<カーネルの呼び出し>>>
ならいけるんじゃないの？自信ないけど。
238:デフォルトの名無しさん
08/03/14 11:06:04
>>236
cudaMallocHostでメモリ確保すると、ページングなしスワッピングなし
物理メモリに連続して常駐、のメモリ領域が確保されるはずなのだ！
と思ってcudaMallocHost使ってます。mallocとmlockの合せ技？
239:デフォルトの名無しさん
08/03/14 16:02:07
Windows, VC2005なんですが、PTXを出すためのオプションはどこに
書けばいいんどすか( ・ω・`)
苦労してPhenomの6倍速いまで持ってきたんですが、さすがに
この先はアセンブラ見てみないと何をどうすればいいか分からんどす。
モノは信号処理用の畳み込み(4096)どす。
240:デフォルトの名無しさん
08/03/14 16:41:32
>>239
プロパティでカスタムビルドであることをチェックできない?
オプションは言うまでもなく -ptx
241:デフォルトの名無しさん
08/03/14 16:56:28
グラフィックカードを2枚挿した状態で、メインメモリからGPU側に転送する際に
どのようにしてGPUを指定すればよいのでしょうか？
カーネルから見られるのは、1枚のGPUだけという認識でよろしいのでしょうか？
242:デフォルトの名無しさん
08/03/14 23:48:11
>>241
cudaGetDeviceCount()からcudaChooseDevice()あたりでは。プログラミング
ガイドのD.1。二枚刺しﾃﾗｳﾗﾔﾏｼｽ
>>240
プロジェクトのプロパティのカスタムビルドステップ、等にもコマンドライン
指定が何もなかとです。次はVSディレクトリに潜って.cuのRuleファイルを
探してみるです。
243:デフォルトの名無しさん
08/03/15 03:32:18
RuntimeAPIとDriverAPIの違いを教えてください。
244:デフォルトの名無しさん
08/03/15 12:17:00
CUDA素晴らしすぎです。ようやくグリッドとブロックとスレッドと
globalとsharedとconstantを理解して、今「Xeonより100倍速いぜ！」
な速度を叩き出しています。もっとチューニングがんばってみよう。
>>243
RuntimeよりDriverの方が低レベルとして書いてありますけど、
似たようなの一杯有るし、違い良く分からないですよね。
245:デフォルトの名無しさん
08/03/15 13:20:14
恐らくは、Driverの方はC++で使うことを想定してない。
実際問題として、nvccはC++としてコンパイルするのでRuntimeAPIの方が使いやすいと思う。
246:デフォルトの名無しさん
08/03/15 17:14:57
>>244
segdmm のC800では120GFlops程度と言われていますが
4coreのXeonでも75GFlopsです。
どのくらい違いますか教えてください。
247:244
08/03/16 00:22:24
>>246
やってみたのが4096単位データ×128個に対して4096のフィルタで
コンボリューション掛ける、なんですが、Xeon2GHzの1スレッド
だと250M操作/秒位しか出来ないですよ。キャッシュミス連発？
GeForce8800ですと、28G操作/秒で動きます。
248:244
08/03/16 04:08:51
あ、そんで、Xeonでやるとこうだってのはもれじゃ無くて既存の
なので、弄れないし、作りはよく分からんのです。とりあえず
「このルーチン速くするライブラリ作れんか？CUDAとか言う奴で」
と言われてやってみたら100倍速いので怒りが有頂天になった所です。
でも冷静に考えると、CPU側もカリカリに書けばあと10倍位速くなる
可能性はありますわな。たぶんSSEとか使ってないし。
でもそこから10倍にするには、つまりもう何台かパソコン追加しないと
ならんわけなので、やっぱりCUDA可愛いよ可愛いよCUDAｗ
249:デフォルトの名無しさん
08/03/16 08:32:33
私のところじゃ、Xeonの3-10倍だなぁ。単体テストで。
# 他の演算と組み合わせると、1.5-2倍が限度（TT
250:デフォルトの名無しさん
08/03/16 09:07:29
「Xeon」てみなさんが言ってるのは、DPマシン8コアのこと？それとも1コアのこと？
1コアに負けるんじゃ確かにちょっと悲しいよね。
251:デフォルトの名無しさん
08/03/16 09:14:14
1スッドレ、って書いてるじゃん
252:デフォルトの名無しさん
08/03/16 09:34:49
つまりこういう暑苦しい感じが最強だと。
喧嘩してる場合じゃねぇ！Quadコア×デュアルCPU＋8800GTXの2枚刺し、
の友情で最強を目指そうじゃないかお前ら！夕日に向かって海まで走れ！
敵は巨大行列のLU分解だ！手ごわいぞO(N^3)！
253:デフォルトの名無しさん
08/03/16 12:57:56
グリッド数、ブロック数、スレッド数の決め方がよくわかりません。
例えば独立な300個のスレッドがあった場合、
一つのブロックに300のスレッドを割り当てるのがよいのでしょうか？
それともスレッド数は1つで300のブロックを作成するのがよいのでしょうか？
ワープ数などはどのように考慮すればよいのでしょうか？
254:デフォルトの名無しさん
08/03/16 13:54:06
>>253
あくまでも漏れの場合ですが、
ブロック数＝16。これは使っているカードの「Multi Processor」に合せて。
違うグリッドに属するデータのやり取りは出来ない事に注意。
スレッド数＝512。目安で。
で、カーネル関数でループで処理を回す場合、最外では
for (i = 0; i < maxcalccount; i=i+スレッド数×ブロック数)
のようにします。これで8192単位での並列化になります。
ループ内部では、グローバルメモリからたとえば１ワード読むなら
sharedメモリ[スレッドID]
＝グローバルメモリ[スレッド数×ブロック数＋ブロックID×スレッド数＋スレッドID]
となります。
floatを使う場合、スレッドが512なら、各スレッドで最大8ワードもてますが、
sharedメモリは使い切らないほうが良いようです。
255:デフォルトの名無しさん
08/03/16 14:57:39
ブロック数は、どうせ多く割り当ててもCUDA内部で直列に並べるだけだから
非同期で少しでもCPUと並列にしたい場合を除けば大目に割り当ててOK。
スレッド数についても多い分はどうせ別のワープに割り当てられるから多めでOK。
但し、同期を取る場合には多過ぎるとダメ。
手元のデバイス関数の場合、ブロック数*スレッド数は少なくとも1024か2048以上必要(8800GTXの実測で)。
これらを踏まえると、スレッド数が32ならブロック数は64以上、スレッド数が64ならブロック数は32以上くらいか。
ブロック数の上限は、実測しながら適当に調整するとして、大体1024を超えるといくつでも変わらないと思う。
# これも、具体的なテスト用のサンプル用意したいところだね。
256:デフォルトの名無しさん
08/03/16 16:42:50
>>188みたいな話が載っている本で使っている本があれば教えてほしいなあ。
入門書の理屈を素直になぞっているような人間には、普通のcpuとgpuとの差を出せなさそう。
257:デフォルトの名無しさん
08/03/16 17:01:32
>>256
>188は統計学を知らんと自分で導けないし、統計学学ぶ香具師がプログラマになるとも限らんしね。
Webでも知識があることを前提としているか、律儀に平均値との差を求める方式しか載ってなかったり。
私の場合は、電卓の統計機能が個別のデータと平均値の差を使わずにどうやって標準偏差を得るのか
不思議に思って調べて以来の知識かな。
ってことで、統計学を知らない私も>188みたいな話の解説があるなら見たい希ガス。
258:デフォルトの名無しさん
08/03/16 17:05:10
CUDAというかGeforceってSIMDなの？
SSEは確かに1つの命令で4つとかの足し算が行われるけど、
CUDAの場合、どれがそれに当てはまるの？
259:253
08/03/16 17:17:24
>>254>>255
早速のご回答ありがとうございます。
基本的にはスレッド数は目一杯であとはブロックすうを調整すると言うことでいいみたいですね。
ところで、処理する数がスレッド数×ブロック数で割り切れない場合、最後のループはどうするのがいいのでしょうか？
やはりifでやるしかないのでしょうか？
260:255
08/03/16 17:38:35
>>259
んにゃ、私はブロック数もスレッド数も実測で決定している。
で、割り切れない場合はループ数で制御するのがいいみたい。
# やっぱり説明するためにはさんぷるが必要かw
>>258
GPUの場合は、SIMDじゃなくてMIMDということになるのかな。
CUDAで考えると、Warp単位で同じインストラクションを走らせて、
複数のデータスレッドを処理するイメージ。
だから、条件分岐をCUDAで書いても、Warp単位で同じインストラクションなので
データスレッドごとに条件が違うと無駄が生じてしまう罠。
詳細は手ごろなサンプルがないので割愛w
261:253
08/03/16 18:08:52
>>260
やっぱり基本は実測なのですね。
ところでループ数で制御するというのはどういうことなのでしょうか？
処理する数/(スレッド数×ブロック数)=ループ数
だと思うのですが・・・
262:255
08/03/16 18:27:15
例えばこんな感じ。
__global__ void func(float const * a1, float * a2, unsigned num)
{ // 実際には一時変数を使った方がいい希ガス
for (unsigned idx = blockIdx.x * blockDim.x + threadIdx.x; idx < num; idx += gridDim.x * blockDim.x) {
a2[idx] = a1[idx] * a1[idx];
}
num=ブロック数*スレッド数ならどのスレッドでもループは一回だけ回る。
num<ブロック数*スレッド数なら、暇なスレッドが発生する(から効率は宜しくない)。
num/(ブロック数*スレッド数)が11と12の中間なら一部のスレッドは12回回って残りは11回回る。
このループ回数が1や2じゃなければ、暇なスレッドの割合が相対的に少なくなる寸法。
# これが理由で、ブロック数は無闇に増やせばいいというわけでもないということになる。
あー説明が下手な私（：；
# サンプル作ったら誰か買わない?w
こういう回し方をすることによって、近くのスレッドが近くのメモリをアクセスする状態のままループが進行する。
つまり、ProgrammingGuideで言うところの"coalesced"。これをプログラミングガイドでは「結合した」と訳しているけど……
263:253
08/03/16 20:27:25
>>262
ありがとうございます。
なるほど非常にうまい手ですね。
プログラミングガイドはなんか日本語訳がめちゃくちゃで非常に読みにくいですが、
がんばって解読したいと思います。
264:デフォルトの名無しさん
08/03/16 20:55:07
shared memoryを使ったら、global memoryを直接使うより遅くなってしまいました。
どうも、カーネルを実行する際のshared memoryの確保に時間がかかっているようなのですが、
カーネルを実行するたびに毎回確保するのではなく、常時確保するような方法はないでしょうか？
265:255
08/03/16 21:20:30
>>264
その現象は確認していないので不明です。
が、保持する方法は多分ないでしょう。
実は一つ問題があって、shared memoryは16バイト同時アクセスができないので
その点でglobal memoryより遅いのは確認済みです。
# ptx出力を見れば見当がつきますが。
266:デフォルトの名無しさん
08/03/17 00:53:35
>>261
もれは、CUDAに投入する前に、CPU側で後詰めゼロとかやって、
「CUDA内では一切条件判断は不要」を原則にしています。
>>264
global、sharedへのアクセスがcoalescedじゃないと、めがっさ遅くなります。
それと、sharedに確保した後で、ブロック内の各スレッドでデータを
shareしない場合には、そもそもsharedを使う理由無いです。レジスタ
に順番にグローバルから読んで処理すればいいです。
267:デフォルトの名無しさん
08/03/17 06:10:01
ブロック間の同期をとる方法を教えてください
268:デフォルトの名無しさん
08/03/17 09:17:39
「ブロック」は、そもそも同時にブロック001とブロック002が動いているか
どうかも分からないんで、同期出来ないんでは？
CPU側で、カーネル関数その１呼び出して、cudaThreadSyncronizeして、
カーネル関数その２呼び出す。ではないかと。
269:255
08/03/17 09:22:38
一応、__syncthreads()という手はありますよ。
但し、全ブロックが同時に動いているわけではない点には注意。
つまり、上の方に書いた分割方法で書いたような大きなブロック数の時には
恐らく巧くいかないかかなり待たされることになるかも知れず。
私の手元のロジックでは、ブロック数が64の時には巧いこと同期が取れています。
# C870, 8800GTX, GTS, GTでしか確認してないけど。
270:デフォルトの名無しさん
08/03/17 14:41:06
>>269
それって各ブロック内のスレッドの同期じゃないの？
271:255
08/03/17 14:55:21
あー、そうそう。失敬、その通りです。
>269は全て、「ブロック」を「スレッド」に読み替えてください。
# どっかでブロック数に切り替わってしまった……_/￣|○
と言うことで、やっぱりブロック間の同期は取れないのでした。
272:デフォルトの名無しさん
08/03/17 15:29:11
>>268
カーネル2はカーネル1が終了してから実行されるんじゃないの？
273:デフォルトの名無しさん
08/03/17 15:56:55
>>272
そこが良く分からんのです。
サンプルではcudaThreadSyncronizeでカーネル呼び出しを挟んでるんで、
一応もれも呼び出すことにしておるのですが。
274:デフォルトの名無しさん
08/03/17 16:04:25
>>273
今、挟んだのと挟んでないのを比べてみたけど、全く同じ結果だった。
ifがあるので、各ブロックはばらばらのタイミングで終わるはずなのだが。
275:255
08/03/17 16:10:18
前にも書いたかもしれないけれど、カーネル終了を明示的に知る手段はないので注意。
つまり、<<<>>>による呼び出しから復帰しても、GPUの処理が終了したわけではないので。
なので、次のような呼び出しをした場合、k1()とk2()は並列に動く可能性があるわけ。
--
k1<<<1, 8>>>();
k2<<<1, 8>>>();
--
終了を保証するには、cudaThreadSynchronize()を呼ぶかcudaMemcopy*()を呼ぶかすればOK。
276:デフォルトの名無しさん
08/03/17 19:44:31
カーネルのループの段数を増やしたら、スレッド数が512だとカーネルの実行ができなくなった。
256ならできる。257はだめ。なんでだろう・・・
277:デフォルトの名無しさん
08/03/17 19:51:18
1BlockのThread数の制限512ってのと関係があるのでは？
278:デフォルトの名無しさん
08/03/17 19:59:35
ちょっと舌足らずだったかもしれませんが、カーネル内のループの段数を増やしたところ、そのカーネルが実行できなくなってしまいました。
そこで、スレッド数を512から256に減らしてみたところうまくいきました。
スタックか何かがあふれてしまったのではないかと疑っているのですが・・・調べる方法がわからず行き詰まってしまいました。
279:デフォルトの名無しさん
08/03/17 20:10:15
実行できないってことは何かエラーが出るの？
280:デフォルトの名無しさん
08/03/17 23:59:13
使っているレジスタの総数が8192を超えたとか。
unrollしすぎだとか。
281:デフォルトの名無しさん
08/03/18 05:17:35
カーネルを2つ同時に実行できますか？
2枚のカードでそれぞれ別のタスクを割り当てて、並列で実行したいのですが。
282:デフォルトの名無しさん
08/03/18 07:49:07
ブロック数が少なければ、勝手に同時実行するな。
283:デフォルトの名無しさん
08/03/18 08:14:14
Maximum memory pitchって何？
284:デフォルトの名無しさん
08/03/18 08:33:44
VCのnmakeって使えないの？
cudaをコンパイルできるサンプルMakefileがあったら教えてください。
285:デフォルトの名無しさん
08/03/18 11:56:28
>>281
サンプルのMultiGPU、8800GTS＋8400で試したら動いたよ。
単にこれ、WINAPIのCreateThreadして中でカーネル関数呼んでるだけですな。
でWaitForMultipleObjectsでWindowsスレッドの終了待っている。
>>283
2次元配列でGPUのメモリを使うときの幅じゃないすか。
cudaMemcpy2Dとか使うときに気にするんだと思われ。
ごめん二次元配列使ってないので良く分からない。
8800GTSも8400も262144bytesです。
286:デフォルトの名無しさん
08/03/18 14:42:22
演算性能がボトルネックになっているのか、バンド幅がボトルネックになっているのか知る良い方法はないものか
287:デフォルトの名無しさん
08/03/18 15:15:30
>>286
演算抜きでcudaMemcpyのHostToDevice、DeviceToHostだけやって
時間を測って見る、そして演算入りの時間と比べる、でどうだろう。
288:デフォルトの名無しさん
08/03/18 16:42:29
__constant__ floatとconst floatの違いって何？
__constant__の方が速いの？実測してみたけど、全然違いが出ないのだけど。
ちなみに何故か定数を直接、式の中に書いたら唯一5％ほど遅くなった。
289:デフォルトの名無しさん
08/03/18 18:11:38
プログラミングガイドの5章が全然わかんねー
結局のところcoalescedって何よ
Figure 5-2の右側は何でダメなの？？？？
教えて偉い人
290:デフォルトの名無しさん
08/03/18 23:54:19
>>288
__constant__は全部で64Kバイトある定数格納用メモリ、ホストから
memcpyすることが出来る。const floatは8192個のレジスタを使って
処理される。なので、普通のコード内定数はconst、ホストからどさ
っと書き込みたい固定データは__constant__でいいと思い。
>>289
5-2の右側は、132=33×4、だから？ 各スレッドは、
16×sizeof()×何か＋スレッドID、にアライメントされた
アドレスにアクセスしないとcoalescedしてくれないと。
291:289
08/03/19 04:35:23
>>290
coalescedを辞書で調べても「合体した」とかしか出てこないので、いまいち意味がよくわからないのですが、
結局0番スレッドは
a[0]
a[16]
a[32]
:
１番スレッドは
a[1]
a[17]
a[33]
:
以外のアドレスにアクセスしないようにすればよいと言うことでいいのでしょうか？
292:デフォルトの名無しさん
08/03/19 04:51:32
要は、配列のアクセスはarray[m * blockDim.x + threadIdx.x]のような形で
blockDim.xは16の倍数になればよろしって話なんじゃないかと。
293:デフォルトの名無しさん
08/03/19 11:08:25
cuMemGetInfoフリーメモリと合計メモリを取得したのだけど、４MBにしかならない。
使い方は
unsigned int f,m;
cuMemGetInfo(&f,&m)
でいいんだよね？
294:デフォルトの名無しさん
08/03/19 11:35:28
>>291
つまり「coalesced」は・・・もれの理解では・・・
・sharedメモリには入出力が32bit×8ポートずつあるかも。
でも開始アドレスは一個かも。
・各スレッドからのアクセス要求アドレスが「0,4,8,12,16,20,24,28」の場合、
GPU側で「ではsharedメモリの出力を綺麗にストリーミングプロセッサに
一個ずつ割り当てればいいのであるな、開始アドレスは0だな」と判断し
てくれるのかも。
>>293
もれの場合は、debugでは動くけどReleaseでは動かないのね。
debugでは FreeMem 465529088 TotalMem 536543232
と取れた。8800GTS/512MB。理由分かったら教えてください。
295:デフォルトの名無しさん
08/03/19 19:26:24
>>290
__constant__って書き換えられるの？
296:デフォルトの名無しさん
08/03/19 22:06:01
>>295
__constant__は定数メモリと考えればいい。
つまり、GPUからは書き換えられないのでCPUから書き換えると。
# constな変数は初期化はできるけど代入できないのと似てるかもしれない。
>>294>>293
>294のdebugでの数値は多分正しい値が出ているみたいだね。
>293の方は、もしかしたらエラーが返っていないかな?
私のところでも、整数にして201だったかのエラーが返ってきていたのだけれど、
その状態では値をセットしてくれないようだった。
つまり、もしかしたら>293の4MBとというのはfとmが未初期化で偶々入った不定値かも。
# 気になるなら、0で初期化してみればいい。
>>294>>291
coalescedなアクセスパターンにすることは、sharedだけでなくglobalでも重要なので
慣れておくことをお勧め。まぁ、例えば>262に書いたようにすればいいだけだけど。
297:デフォルトの名無しさん
08/03/19 23:54:01
>>295
cudaMemcpyToSymbol()でCPUからセット出来るよ！
速いし64KBあるんで結構便利です。
298:デフォルトの名無しさん
08/03/20 00:01:09
問題は、ベクタアクセスできる__global__と違って1データずつのアクセスになるから
逆に遅くなるケースもあることだな。
299:デフォルトの名無しさん
08/03/20 12:55:05
ややこしさはCellとあんまり変わらなそうだな。
300:デフォルトの名無しさん
08/03/20 14:25:54
>>299
とにかくメモリのコピーがヤヤコシイんだわ。ホスト側、GPUの
グローバルメモリ、GPUのチップ内メモリで転送しまくらないと
いかんで。やり方間違えると全然性能出ないし。
Cellはその辺どうなの？
301:デフォルトの名無しさん
08/03/20 14:28:32
256kしかないLSでやりくりするのが大変って聞いたな
302:デフォルトの名無しさん
08/03/20 15:10:07
両方試した私に言わせて貰えば、どっちもどっち。
確かにCBEは256KiBの壁がねぇ。GPUも64KiBの壁やcoalescedの沼があるけど。
超越関数を使える点ではCUDAが有利。ホストの性能でもPPEじゃ結構泣けるし。
303:デフォルトの名無しさん
08/03/20 16:02:10
PPEはひどいよな。ホステスにC2D使ってるけど、ifがいっぱいあるような場合はC2Dの方が速いしね。
304:デフォルトの名無しさん
08/03/20 16:14:13
>>298
それ実験してみたんだけど、変わらないみたい。
・global→shared(行列多数)と__constant__に置いた定行列で行列積
・global→shared(行列多数)とsharedの一部に置いた定行列で行列積
で、後者が1%遅いくらいだった。リードオンリーなだけで、コアとの距離や
所要クロック数はconstantもshared・レジスタも同じなのかもと？
305:デフォルトの名無しさん
08/03/20 16:28:58
CUDAで逆行列を求める方法を教えてください
306:デフォルトの名無しさん
08/03/20 16:49:55
page-lockedってどういう状態を表すのでしょうか？
307:298
08/03/20 17:42:11
>>304
-ptxでニモニックを出力してみれば違いが判るかと。
例えば、ld.global.v2.f32はあるけどld.const.v2.f32はないからld.const.f32が2回になってしまう。
coalescedなアクセスができるglobalは4クロックでアクセスできるからsharedやconstと変わらないわけで。
尤も、一旦そのパターンから外れるとglobalは数百クロックだそうだから途端に劇遅になるけど。
308:298
08/03/20 17:48:04
>>305
先ずは、逆行列を求めるプログラムをCで書いてみてください。
それの、ループの部分を分割する方向でGPU関数を作るのが第一歩になります。
つーかさぁ、一行質問する人達ってなんなの?
情報交換する気がなくて、単にくれくれの精神なんだったら勘弁して欲しいんだけど。
# なんか回答者が数人とそれより少し多いレス要員だけで持ってる気がするよ……
>>304
あーそうそう、書き忘れたけどレジスタアクセスは1クロックじゃなかったカナ。
確実にメモリよりも速いみたい。
309:デフォルトの名無しさん
08/03/20 19:43:17
nvccにはO2とかの最適化オプションは無いのでしょうか？
310:デフォルトの名無しさん
08/03/20 20:26:00
-Wallが使えないみたいね
たぶん無さそうな予感
311:298
08/03/20 21:19:57
一部のオプションはgccと共通ですよ。
例えば、-O3(恐らく-O2なども)や-pgは使えます。
312:デフォルトの名無しさん
08/03/20 23:15:10
>>299
どっちもどっちで、用途や慣れの問題だと思う。302がいうように、SFUの有無は
大きいが、それも用途しだい。
自分的には比較的情報がオープンなCellのほうがいろいろいじりやすいと思うが、
最近人気なくて心配。たぶんそのうちLarrabeeが全部かっさらってくんじゃないかな。
313:デフォルトの名無しさん
08/03/22 18:23:49
GeForce 9800 GX2で現バージョンのCUDAは動きますか？
314:デフォルトの名無しさん
08/03/22 18:28:28
ストリーム番号の同じカーネルやメモリコピーは順番に、違うものは並列に実行される？？の認識でいい？？？
315:デフォルトの名無しさん
08/03/22 18:43:27
テクスチャメモリの容量はどのようにすれば調べられるのでしょうか？
どこかに資料があるのでしょうか？それとも何かAPIがあるのでしょうか？
316:デフォルトの名無しさん
08/03/22 23:25:42
いつも適当に答えてるおじさんです。まことにもうしわけない。
>>313
もれ、8800GTS(G92)で動いてるから、×２で二個見えるんじゃないかと
予想しますけど、一応どっちも対応外だから。漢は度胸でひとつ人柱に。
>>314
ストリームが二つあっても、実行順は関係ないのではと予想。ストリーム
１の半分実行した後ストリーム２の半分実行、もありえるかと。
cudaStreamSynchronizeでストリームの実行完了を待たないといけないです。
>>315
テクスチャメモリはグローバルメモリのキャッシュですが、実サイズが
不明ですね。cudaGetDevicePropertiesでは「alignment=256Bytes」と
読めます。これ以下と以上で実行時間を調べてみると、キャッシュミス
してるかしてないかわかるのではないでしょうかと。
317:デフォルトの名無しさん
08/03/23 13:19:47
openSUSE 10.3でCUDAは動きますか？
318:デフォルトの名無しさん
08/03/23 13:33:44
動くらしい
319:デフォルトの名無しさん
08/03/23 17:02:07
次のバージョンっていつ頃出るの？
新しいカードも出ているわけだしそろそろ、SDKの方もアップデートしてほしいものなのだが
320:デフォルトの名無しさん
08/03/23 23:52:41
Vista64と9600,9800対応版が欲しいすね。
早く出してくれないとAMD/ATIに浮気したくなりますよ。
321:デフォルトの名無しさん
08/03/24 10:55:41
サンプル作ると言いつつなかなか暇が取れない私が来ましたよ。
>>313
ドライバが対応しているかどうかが問題。ドライバが認識さえすれば、動くとは思いますが。
>>314
ストリームのご利用は計画的に。ちなみに、CUDA1.0世代の(要はG8xかな)GPUだと
処理の隠蔽ができないらしく、全く速くならないという噂もあります。
>>317
>27-30
荒らしじゃないんだったら、スレくらい検索してほしいと思う今日子の頃。
# いや、荒らしなら検索如何に関わらずご遠慮願いますが。
>>319
取り敢えず、待つしか。下手すると夏の新製品までお預けの可能性も。
>>320
え～、待っていれば倍精度も来るのに。
322:デフォルトの名無しさん
08/03/26 23:04:49
英wikiでは、9600GTでもcudaできることになっている
悪い子と漢は居ね゛ーがー
323:デフォルトの名無しさん
08/03/26 23:38:52
>>322
もれ9600GT買おうと思ってる。
ARCTICのS1Rev2が付く、という報告待ち・・・。ファンレスで使いたくて。
324:デフォルトの名無しさん
08/03/31 21:17:49
URLﾘﾝｸ(www.nvidia.com)
9800GX2 CUDA OKだとでてるよん
325:デフォルトの名無しさん
08/04/01 00:04:20
おー、更新されている。情報THX。
ドライバは変わっていないから、チップ依存なんですかね。
326:デフォルトの名無しさん
08/04/04 22:15:16
ML115祭りに参加してOpetron使用をポチりました。
で、nbody.exeを実行すると・・・・ 160～170GFLOPSという値が。
うぅ・・速過ぎる。初めて見るそのパワーに感動ものですた。
327:デフォルトの名無しさん
08/04/04 22:15:54
忘れてた GPUは8800GTですぅ。
328:デフォルトの名無しさん
08/04/04 22:24:35
そろそろまじめに実用的実装に入らないのかな？
圧縮解凍ソフトに組み込むとか
暗号化ソフトに組み込むとか
一番簡単なのは画像処理ソフトだろうけど
329:デフォルトの名無しさん
08/04/05 02:08:09
大真面目に実用的な実装で悪戦苦闘していますが何か。
330:デフォルトの名無しさん
08/04/05 09:56:17
画面の解像度によって使えるメモリ量は変わりますか？
331:デフォルトの名無しさん
08/04/05 12:08:56
CUDA使うアプリをインストーラで配るにはどうしたらいいもんだろうか？
というところが。Vista対応してないし。
332:デフォルトの名無しさん
08/04/05 13:45:15
CALならドライバ要らずなんだけどなぁ。
なんで別途ドライバが必要なんだろ。
つかVista版ドライバは3月の終わりごろじゃなかったのかよ！
333:デフォルトの名無しさん
08/04/05 15:56:48
XP用のビデオドライバをインストールできるはずだから、その状態でXP用のCUDAドライバ入れられれば動くんじゃないかな
と無責任なことを言ってみる
334:デフォルトの名無しさん
08/04/05 16:12:49
>>332
たぶん時差
335:デフォルトの名無しさん
08/04/05 19:34:29
URLﾘﾝｸ(japanese.engadget.com)
cudaなんかやってる場合じゃないよな
336:デフォルトの名無しさん
08/04/05 19:52:25
戦犯はマイクロソフトだろ･･･
337:デフォルトの名無しさん
08/04/05 20:08:59
黒歴史確実なOSだからスルーでいいよ
338:デフォルトの名無しさん
08/04/05 20:15:04
え、俺パフォーマンス100倍UPにちょー期待してんだけど
339:デフォルトの名無しさん
08/04/05 22:09:08
普通にopenGL使ったオープンソースのキット使ったほうがいいだろ
340:デフォルトの名無しさん
08/04/05 23:52:39
今月リリース予定のCUDA 2.0ベータ版でVistaサポートだってさ
341:デフォルトの名無しさん
08/04/06 00:55:39
CUDAを有効に使うのなら、gccとの相性から考えてもLinuxで使うだろ。JK
342:デフォルトの名無しさん
08/04/06 01:32:22
仮想メモリが使ってみたいんだもの。
343:デフォルトの名無しさん
08/04/06 06:31:34
ローカル変数はどこに格納されるのですか？
shared memoryでしょうか？
344:デフォルトの名無しさん
08/04/06 08:35:00
スタックフレームはないようなので、全てレジスタに格納されると思って宜しいかと。
ローカルに大きな配列取ろうとしたらどうなるのかはしらんなぁ。
ptx出力して読んでみたら?
345:デフォルトの名無しさん
08/04/06 09:32:54
>>341
商売にならんじゃないか。建前はどうでもいいんですよ。
売れるかどうかだけです。
346:デフォルトの名無しさん
08/04/06 09:49:44
「商売にできない」の間違いですね。
347:デフォルトの名無しさん
08/04/06 11:51:20
エンコソフトでCUDA使ったやつ売れれば儲かりそうね。
動作検証が非常に大変だろうけど。
Vista対応不可って所で今の所どうにもならないですなぁ。
漏れがいま一番熱望してるのは、XilinxがCUDAに対応して、
FPGAのコンパイル・シミュレーションが劇速になること。
ああ夢のようだ・・・
348:デフォルトの名無しさん
08/04/06 16:11:58
GPUで計算させているのだけど、その間のCPU使用率って100%なんだけど、これって正常なの？
349:デフォルトの名無しさん
08/04/06 16:28:42
正常です。
ブロック数がPE数より多い場合は待たされることがありますし、
メモリ転送や同期を取るときには当然待たされます。
実は待っている間もしっかりCPU時間を消費するのです。
従って、GPUとCPUを巧く連携させて高パフォーマンスを狙うには
いかに待たずに済ませるかが鍵になるので、Streamの使用は必須になってくるかと。
350:デフォルトの名無しさん
08/04/06 16:37:26
くそ速いglReadPixelsとして使えるかと思ったが
CUDAで扱うデータしかWrite Readできないのね
ぐすん
351:デフォルトの名無しさん
08/04/06 16:53:37
CUDAにもGLサポート関数があるから、もしかしたら連携できるんじゃないの?
サンプルの、simpleGL辺りになんかない?
352:デフォルトの名無しさん
08/04/06 19:06:19
PBO使え
353:デフォルトの名無しさん
08/04/06 19:22:37
GT200と99GX2はどっちが高性能でしか？
354:デフォルトの名無しさん
08/04/07 10:28:11
このスレの住人なら知っていますね、あの糞開発ツールのことを
・自分のプログラムのバグなのかコンパイラのバグなのかわからない
・他の仕事に応用できない糞開発ツールの独自世界を必死に学習している
・テキストエディタで書いたほうが効率的なのに糞UIツールを懸命に使っている
・糞開発ツールを批判すると「性格が悪いから糞ツールを批判するんだ」と言われる
糞だけど、政治的な理由で無理やり使わされているんですよね。
もう、あんな厨の作った糞ツールを我慢して使うのはやめましょう。
・糞開発ツールを部下に押し付ける上司の命令は無視しましょう。
上司は糞開発ツールが使われる実績を作ることであの会社のごきげんをとっているのです。
・糞開発ツールを使わせる上司の下では働けません、と上司の上司に直訴しましょう。
・あの糞開発ツール提供会社には「おたくの糞開発ツールは話にならない」と突き放しましょう。
バグレポートなどしてはいけません。改善要求などもってのほかです。
あの会社はあなたたちのことをテスター/モルモットとしか思っていません。
・あの会議で「糞開発ツールを使ったら生産性がxx％アップしました」
なんて話が出たら力強く机を叩き、会議室を出ましょう。
あの人たちは糞開発ツールをマンセーすることで立場を確保しているのです。
糞な開発ツールを糞だと言える、そんな当たり前の環境をみんなの力で取り戻しましょう。
355:デフォルトの名無しさん
08/04/07 14:17:07
GX2はSLIと同じように2つのgridを使う必要があるのでしょうか？それとも、1つのgridで両方のGPUを使えるのでしょうか？
356:デフォルトの名無しさん
08/04/07 14:26:08
CUDAにSLIは関係ありません。
二枚挿しのことをSLIと呼ぶのは間違いです。
何故なら、SLI用コネクタを挿さなくてもCUDAではGPUを二つ使うことができるからです。
で、肝腎な9800GX2ですが予想ではGPUが二つ見える筈なので、一つの時のままだとダメな気がします。
357:デフォルトの名無しさん
08/04/07 15:24:49
2つのGPUを使うには、CPU側もマルチスレッドにする必要があるのですか？
358:デフォルトの名無しさん
08/04/07 16:17:21
Streamを使えばシングルスレッドでも何とかなるんじゃないかと思うので、是非とも試してみてください。
# 因みに、StreamはCUDA1.1からの機能なのでサンプルのmultiGPUでは使っていません。
359:デフォルトの名無しさん
08/04/08 18:48:35
ビデオカード2枚刺しの場合、2枚のビデオカード間のデータのやりとりはPCIe経由なのですか？
360:デフォルトの名無しさん
08/04/08 22:08:50
>>359
CUDAにはチップ間の転送なんてなかったと思いますが。
361:デフォルトの名無しさん
08/04/08 23:54:35
CUDA 2.0 in this month
URLﾘﾝｸ(forums.nvidia.com)
Vistaもサポートするよ。
362:デフォルトの名無しさん
08/04/09 00:10:04
一方のVistaはMSに見捨てられた。
363:デフォルトの名無しさん
08/04/09 00:27:55
PGがゴシップ記事鵜呑みにするなよ・・・
364:デフォルトの名無しさん
08/04/11 16:05:33
Linux でCUDA使っていますが、pthreadを使って以下のようにスレッド内でcudaMemcpyを呼ぶと
invalid device pointerでcudaMemcpyに失敗します。なぜでしょうか。
スレッドを作らずに、mainから直接testCopy(NULL)と呼び出すと、問題ありません。
#include <pthread.h>
#include <cutil.h>
#define NUM 512
int *h_buf=NULL;
int *d_buf=NULL;
pthread_t thread;
void *testCopy(void* args) {
CUDA_SAFE_CALL(cudaMemcpy(d_buf, h_buf, sizeof(int)*NUM, cudaMemcpyHostToDevice));
CUT_CHECK_ERROR("Copy failed");
}
int main(int argc, char **argv) {
CUDA_SAFE_CALL(cudaMalloc((void**)&d_buf, sizeof(int)*NUM));
CUDA_SAFE_CALL(cudaMallocHost((void**)&h_buf, sizeof(int)*NUM));
for(int i = 0; i < NUM; i++) h_buf[i] = i;
pthread_create(&thread, NULL, &testCopy, NULL);
pthread_join(thread, NULL);
}
365:デフォルトの名無しさん
08/04/11 18:24:09
Linuxよく判ってないけど次の点をチェック。
・cudaMallocHost()はメモリのページロックをしてしまうので、malloc()にしたらどうなるか。
・そもそもスレッドを分ける意図は何か。Stream系APIでは事が足りないのか。
366:デフォルトの名無しさん
08/04/11 19:16:30
CUBLASの使用においても他スレッドで確保したメモリは扱えないので注意が必要です
367:364
08/04/12 00:04:48
>>365
mallocも、newも試したのですが、だめでした。
CUDAの制御をするスレッドと、GUIの面倒を見るスレッドを分けたかったのです。
>>366
スレッド、またいじゃだめなんですか。それは、CUDAの仕様なのでしょうか。
368:デフォルトの名無しさん
08/04/12 00:23:19
>>367
CUDAの面倒を見るスレッドでメモリ確保すればいいだけじゃん。
369:デフォルトの名無しさん
08/04/16 17:03:58
URLﾘﾝｸ(pc.watch.impress.co.jp)
370:デフォルトの名無しさん
08/04/16 18:12:45
NVIDIA必死だなｗ
371:デフォルトの名無しさん
08/04/16 23:56:06
AdobeのPremierとか、先に取り込んだとこの勝ちになるんじゃね？
372:デフォルトの名無しさん
08/04/17 14:41:57
おまえら2.0（Beta版）きたぞ
URLﾘﾝｸ(forums.nvidia.com)
373:デフォルトの名無しさん
08/04/18 16:34:55
otu
374:デフォルトの名無しさん
08/04/23 01:01:57
CUDAプログラミングガイドに書かれているGPUの解説は
CUDAで使う場合限定の話ですか？CUDA対応GPUが行う処理全てについてですか？
例えば複数のグリッドで共有するメモリなど。
375:デフォルトの名無しさん
08/04/23 23:17:01
>>374
共通する点もあるとは思いますが、CUDAでは(利便性と引き換えに)GPUの利用に
制限が掛かっていることから類推して当てにならないと思った方がよさそうです。
376:デフォルトの名無しさん
08/04/24 23:35:24
夏ぐらいまでにCUDAで一本何か作ろうと思うけど
今から始めるなら2.0がいいの？というか何が違うんだ？
377:デフォルトの名無しさん
08/04/24 23:50:27
まずはEmuでつくれ
378:デフォルトの名無しさん
08/04/27 18:53:52
Visual Studio 2005じゃないとだめなのね
379:デフォルトの名無しさん
08/04/27 19:05:52
サンプルが2005じゃないとビルドできないですね
380:デフォルトの名無しさん
08/04/27 20:39:58
AMDの新しいbetaSDKって
Radeonじゃ動かんよね？
381:デフォルトの名無しさん
08/04/27 20:46:11
>>380
スレ違い。
>>379
Linux版はインストールされているgccがあれば大丈夫。
# cygwinのgccが使えるといいんだけどねぇ。
382:デフォルトの名無しさん
08/04/27 20:52:27
>>379
金がないってことか?
Express試用したら?
383:デフォルトの名無しさん
08/04/27 21:32:45
>>381
スレねーじゃん
ケチくせーこと言わないで教えろよ
384:デフォルトの名無しさん
08/04/27 21:34:29
GPGPUのスレあるだろ
385:デフォルトの名無しさん
08/04/27 22:38:28
>>381
逆にgccのバージョン違ってると上手く動かないけど。
個人的には*BSDで動いて欲しい。それも64bitで。
386:デフォルトの名無しさん
08/04/27 22:44:17
>>385
そんときは、違うバージョンのgccを入れて設定を変えておけばいいらしい。
なので、ドライバは兎も角nvcc(とemu)は*bsdでも動く可能性はあるかと。
あー、今確認のために見に行ったら2.0betaの案内がCUDAZoneにも出てますね。
CudaVisualProfilerも2.0betaが出ているらしい。
>>384
だってAMDのは知らないし～
387:デフォルトの名無しさん
08/05/04 02:02:07
かけ算にかかるクロック数や足し算にかかるクロック数を知りたいのですが、
そのような資料はないものでしょうか？
388:デフォルトの名無しさん
08/05/04 08:38:43
ガイドブック。
389:デフォルトの名無しさん
08/05/04 08:47:28
>>387
基本的に、掛け算も足し算も同クロック。
問題は、そこに至る過程なのでptxを出力してループ内の行数を数えるとか
メモリアクセスの個数を数えるとか。
そもそもプロファイラを使いこなせとか。
390:デフォルトの名無しさん
08/05/06 18:16:16
SLI環境で、GPUから別のGPUに直接データを転送することはできるのでしょうか？
もしできるのでしたら、やり方を教えてください
391:デフォルトの名無しさん
08/05/06 19:48:28
>>390
>355-360
392:デフォルトの名無しさん
08/05/13 14:13:18
CUDAってｎVidia限定でしょ？
サンプル動かして感動したけど
せっかく作っても使える環境が限られると意味ないね
393:デフォルトの名無しさん
08/05/13 22:32:10
確かに、仮令MacやLinuxで使える場合があったとしても、Windowsの半数以上で使えないからね。
まぁ、使いたい人が使うからいいのよ。
394:デフォルトの名無しさん
08/05/14 00:14:56
スレッドをいくつか使うときに
スレッド0、１，２の計算が終わってから
スレッド4の計算を始めるってどうやって書けばいいの？
待ち同期どうやって作ればいいか解らない
395:デフォルトの名無しさん
08/05/14 00:27:53
GPU側関数で同期を取るのは__syncthreads()でできるけど、
恐らくそれではパフォーマンスが出ないと思われ。
スレッドは数個だけ動かしている積もりでもWarp単位でしか動かないので、
分割の仕方を見直した方がよさそう。
396:デフォルトの名無しさん
08/05/26 20:14:25
ついに、gtx 280がでるね。
第２世代型の統合シェーダーアーキテクチャってどんなんかな。
倍精度サポートとか、10秒で計算打ち止めへの対処とかいろいろありそうだけど。
自分的には、分岐への弱さの改良とか、ブロック間の同期機構とか、手を入れてくるんじゃないかと思う。
あとは、各種メモリの配分をどうしてくるか注目。shared memoryを倍増とかあるかな。
397:デフォルトの名無しさん
08/06/01 23:08:01
取らぬ狸のなんだが
GTX 280/260ではかなりの高速化が見込まれるがどう？
398:デフォルトの名無しさん
08/06/01 23:16:46
消費電力がやばすぎると思う
399:デフォルトの名無しさん
08/06/01 23:17:36
>>398
お前そんなネガ発言すると
このスレをN房の溜まり場にすっぞコラ？
400:デフォルトの名無しさん
08/06/01 23:53:38
400Wゲトー
401:デフォルトの名無しさん
08/06/02 00:09:32
IU鯖で使えるくらいになってほしいとこだよなあ
402:188
08/06/03 12:50:08
今頃になって説明の間違いに気づいた……
>188の説明において、途中に出てくる分散と共分散の式が全て(n-1)で割られてないや。
相関係数の段階では分母子に出てくるから消えてなくなるんで忘れていた。
スレ違いだけど間違いを放置するのも難なんで念の為。
# Excelの分散は分母が(n-1)で、共分散は分母がnなんて仕様なのに気づいて、ここを思い出した。
403:デフォルトの名無しさん
08/06/09 19:12:05
CUDAを使ったプログラムは専用ドライバが入っていない環境では動かないですか？
配布は無理がありますか
404:デフォルトの名無しさん
08/06/09 23:32:27
っ[knoppix for CUDA]
405:デフォルトの名無しさん
08/06/09 23:55:59
普通に使うWindowsソフトの高速化として使うものではないと
406:デフォルトの名無しさん
08/06/12 15:49:19
最新のnVIDIAドライバにだったら入ってるでしょ
407:デフォルトの名無しさん
08/06/12 16:34:56
でも手軽にCUDAテクノロジを使えるのソフトというのも面白いね。
408:デフォルトの名無しさん
08/06/12 16:36:45
>>406
最新のnVIDIAドライバにはCUDAが使える機能が入ってるってこと？
409:デフォルトの名無しさん
08/06/13 00:12:22
スレチなんですけど、ほかに質問できそうなところがなくて・・
nVIDIA のグラボで水平スパンやDualViewの設定をアプリから制御する API みたいなのはあるのでしょうか？
アプリケーション起動時に自動でそのあたりを設定したいのですが。
410:デフォルトの名無しさん
08/06/13 00:55:07
あると思うけど、しらね。NVIDIAに聞いたら?
411:デフォルトの名無しさん
08/06/13 09:59:32
CUDAは別途ドライバが必要
Geforx8以上ならドライバ入れれば動くけど
一般ユーザーにドライバ入れるなんて作業が簡単に出来るわけないので
ソフトとして配布するのは無理じゃないかな
専用のインストーラーでも作るなら別だろうけど
412:409
08/06/13 10:11:50
>>410
以前 別件でnvidiaジャパンに問い合わせてみたけどシカトされたからなあ。
とりあえず ELSAジャパンにメールしてみました。
413:デフォルトの名無しさん
08/06/13 11:31:55
なんで標準ドライバにCUDAのドライバ入れとか無いんだろう・・・・
414:410
08/06/13 11:50:33
>>412
あ、やっぱり?w
1000枚単位の客じゃないと相手にしないって噂は本当だったか＜NVIDIA
ELSAもどうだろ。最近連絡くれなくなったからなぁ。
415:デフォルトの名無しさん
08/06/13 12:31:56
NVIDIAジャパンは営業ばっかで技術スタッフ皆無だからな
416:デフォルトの名無しさん
08/06/13 16:51:49
>>411
そうなんすかー
>>413
ですよねー
配るならDirectXやOpenGLでやるしかないということか。
417:409
08/06/13 17:33:23
>>410
エルザから返事がきましたよ！
>基本的には制御不可能かと思われます。
泣ける。あとは nVIDIA のデベロッパーサイトのリンク。号泣。
1000枚とはいわんけど、30枚くらいは買ってるんだがなあ。
418:デフォルトの名無しさん
08/06/13 22:17:30
だから、最初からドライバの入ってる
knoppix for CUDA と一緒に配布すれば
良いじゃん。
419:デフォルトの名無しさん
08/06/13 22:51:47
ドライバ内の公開向けには作ってない制御APIを、
公開しろとか騒いでも無駄だと思うぞ。
ドライバのバージョンでコロコロ変わってもおかしくないもんだし。
420:デフォルトの名無しさん
08/06/17 09:29:37
GPUはin-orderですかそれともout-of-orderですか？
421:デフォルトの名無しさん
08/06/17 10:41:22
ミクロではインオーダ
マクロではアウトオブオーダ
CUDAで見えるレベルだったらOoOと思っといてよい。
422:デフォルトの名無しさん
08/06/17 10:42:39
速攻でミスった。
× CUDAで見えるレベルだったらOoOと思っといてよい。
○ ～キホンインオーダ。ただしコンパイラはそれなりの最適化を施す。
423:デフォルトの名無しさん
08/06/17 15:03:34
>>417
今のバージョンで使えるかどうかはわからんけど、NVControlpanel（nvcpl.dll）のAPIはある。
URLﾘﾝｸ(developer.download.nvidia.com)
年代からして、Vistaはおそらく無理。
424:デフォルトの名無しさん
08/06/17 18:41:39
CUDAのサイトがリニューアルしてるぞ。
賛否両論あるとおもうが、とりあえず整理されている。
425:デフォルトの名無しさん
08/06/17 19:29:20
ROPユニットとストリームプロセッサの違いって何？
426:デフォルトの名無しさん
08/06/17 23:00:58
>>424
見た見た。
そして、例によってCUDA-Enabled ProductsにGTX280/260が載っていない罠。
427:デフォルトの名無しさん
08/06/17 23:37:37
GTX280 なんすかこれｗ
Vipperより糞じゃんｗ終わってるな
428:デフォルトの名無しさん
08/06/18 00:36:52
>>413
>>416
最新のドライバから統合するようになったみたいだよ
URLﾘﾝｸ(www.nvidia.co.uk)
提供されてるのが200シリーズ向けだけなので現在確かめようがないけど
429:デフォルトの名無しさん
08/06/18 00:39:29
ん、いや統合はvistaだけなんだろうか・・・？
URLﾘﾝｸ(www.nvidia.co.uk)
こっちにははっきり
Adds support for CUDA? Technology.
と記述されてる
430:デフォルトの名無しさん
08/06/18 04:25:53
PhysXとCUDAは同時に使えますか？
431:デフォルトの名無しさん
08/06/18 07:35:59
ゲロビディア終わったなｗ
432:デフォルトの名無しさん
08/06/18 19:53:32
だれかgtx280でcudaやったやつおらんのか。
433:デフォルトの名無しさん
08/06/18 22:38:05
ゲロビディア氏ね
434:デフォルトの名無しさん
08/06/19 10:43:41
>>432
やってみました。確かに速いが
435:デフォルトの名無しさん
08/06/20 13:31:41
>>434
確かに早いが？。。。その先が気になる。
すまないが、他のボードとの比較とかの情報があるととても嬉しい。
436:デフォルトの名無しさん
08/06/20 16:17:10
発熱が凄いってこったろう
437:ヽ・´∀｀・,,）っ━━━━━━┓
08/06/20 21:52:57
GPUとしては今はラデのほうが良いみたいね。
Brook+ってどうなのよ？
438:デフォルトの名無しさん
08/06/20 23:18:08
実用にはまだちょっと厳しい。
F@HはBrook+で書いてるし使えないこともないってレベル。
今のところOpenGLやDirectXと協調出来ないからそこも問題かな。
439:デフォルトの名無しさん
08/06/21 00:01:39
これから暑い夏を迎えて、発熱が凄いボードはどう評価されていくか興味深いね。
440:デフォルトの名無しさん
08/06/21 00:49:53
洞爺湖サミットでNvidia名指しで批判されるらしい
だから焦っていろいろアピールしてるらしい
環境もっとも悪い製品を作ってる会社の代表格て
声明が盛り込まれる予定
441:デフォルトの名無しさん
08/06/21 13:03:24
CUDA 2.0 Beta 2
URLﾘﾝｸ(forums.nvidia.com)
442:デフォルトの名無しさん
08/06/22 19:11:15
cudavideodecodeなんて今更何に使うんだと思ったけどDXVAがないOS向けか
443:デフォルトの名無しさん
08/06/27 21:18:32
倍精度計算したいどの変数使うの？
444:デフォルトの名無しさん
08/06/27 22:24:33
すいません。助けてください。
VC2005でサンプルをビルドしようとすると以下のエラーが出ます。
Visual Studio configuration file '(null)' could not be found for installation at 'C:/Program Files (x86)/Microsoft Visual Studio 8/VC/bin'
コンフィグファイルが必要なのかと思い、C:/Program Files (x86)/Microsoft Visual Studio 8/VC/binの下にnvcc.profileを
置いてみたのですが結果は変わりませんでした。
どうしたらいいですか。
ＯＳはwindows xp 64で、CUDAのtoolkitとSDKは1.1です。
445:デフォルトの名無しさん
08/06/27 23:14:43
>>444
cuda_build_rule.zip入れた？
倍精度で計算ってどうやって書けばいいの？
446:デフォルトの名無しさん
08/06/29 14:05:37
大学の課題でで使うことになったのでマニュアル読んでみたけど、日本語訳がファッキンなのは仕様ですか？
冗談抜きで英語の方が判りやすかったｗ
447:デフォルトの名無しさん
08/06/29 14:11:15
>>446
授業ページおしえてください
448:デフォルトの名無しさん
08/06/29 14:18:00
>>447
うちの大学の教授が俺用に30秒で考えて出した課題なので、特に授業ページはありません。 ごめんね。
449:デフォルトの名無しさん
08/06/29 14:18:03
>>446
授業ページおしえろｺﾞﾗｧｱｱｱｱｱｱ
450:デフォルトの名無しさん
08/06/29 14:18:13
>>446
NVIDIAジャパンには、まともな技術者がいないので仕方ありません。
尤も、営業にもまともなのがいるかどうか些か疑問ではありますが。
451:デフォルトの名無しさん
08/06/29 14:24:24
Gerovidia Japanには屑営業しかいねーじゃん
実際CUDAとか国内研究で利用していて
すごそうなのは、Gerovidia本社から来るし
452:デフォルトの名無しさん
08/06/29 14:46:57
cudaのおかげで英語力が向上しました
453:デフォルトの名無しさん
08/07/01 11:15:26
cudaのおかげで彼女ができました
454:デフォルトの名無しさん
08/07/01 15:54:31
double float はいつ???
455:デフォルトの名無しさん
08/07/07 14:14:38
いつって、2.0β使えばすぐできるんじゃないの? エミュで。
456:デフォルトの名無しさん
08/07/09 18:26:35
CUDAが.NETから使えるようになってるらしいので興味持ったのですが（今現在動作環境なし）
C#とかVSとかとの親和性とか、ネイティブと比べて処理性能どれくらいかとか、試した人いたら教えてください
CUDA.NET
URLﾘﾝｸ(www.gass-ltd.co.il)
457:デフォルトの名無しさん
08/07/09 22:07:47
ランタイムだけじゃね
458:デフォルトの名無しさん
08/07/10 22:56:51
アセンブリ読み込めば後はどの言語でも使えてしまうんですかね
PowerShellやIronPython（pyCUDAというのもあるみたいですが）から使えたりしたら面白そう、面白そうなだけですけど
459:デフォルトの名無しさん
08/07/11 15:12:02
超初心者ですが質問お願いします。
自分のパソコンにグラボが2枚刺さっているのですが
CUDAで使うグラボを選択するにはどうしたらいいんですか？
めっちゃ初歩的な質問ですが回答お願いします。
460:デフォルトの名無しさん
08/07/11 16:42:36
サンプルにあるだろ
461:デフォルトの名無しさん
08/07/13 02:02:23
さわりだけだけど日本語資料みっけた
URLﾘﾝｸ(www.easize.jp)
462:デフォルトの名無しさん
08/07/13 21:18:58
>>461
今月のインターフェイス８月号にも似たような資料があったYO!
誰かGpuCV使ったやついねーか…
makeできねーん、OpenCVで聞けばいいのかこっちなのかわかんね
463:デフォルトの名無しさん
08/07/13 21:55:32
>>462
環境書かけよくず
464:デフォルトの名無しさん
08/07/13 22:34:52
>>462
釣られてやったけどVS2005、VS2008
SuseLinux RedhatES 5.0 Ubuntu 8.04 CentOS 5.1で
ビルドできたんだがｗ
ビルドすらできないPGってカスもいいとこだろｗ
465:446
08/07/15 06:03:10
うわぁぁぁぁファッキンなのは日本語訳だけかと思っていたらSDKもだったぁぁぁぁ
誰だcutil_math.hの float / float[2,3,4]の演算子オーバーロード書いた馬鹿は
俺の3時間返しやがれ
inline __host__ __device__ float2 operator/(float s, float2 a)
{
float inv = 1.0f / s;
return a * inv;
}
ひょっとしてcuda SDKのcommon/inc/以下のヘッダーってわりと危険だったりする？
466:462
08/07/15 06:50:30
>>463-464
ちょっｗｗｗｗｗ
環境はCentOS 5.19 64bit
GpuCV 0.4.2 driverはcudaのやつ174.55
gpucv,resources,sugoitoolsをダウンロードして
URLﾘﾝｸ(picoforge.int-evry.fr)
を参考にファイルコピー
cp -Rf bin\gnu\ ..\resources\bin\gnu\
cp -Rf lib\gnu\ ..\resources\lib\gnu\
は sugoitools/bin/gnu,libが無いので実行できず無視して↓のみやった
cp -f include\SugoiTools\*.h ..\resources\include\SugoiTools\
cp -f include\SugoiTools\*.inl ..\resources\include\SugoiTools\
cp -f include\SugoiTracer\*.h ..\resources\include\SugoiTracer\
しかしgpucvに入って./createSolutions.shで
./createSolutions.sh
Generating GNU makefiles:
...GPUCVHardware
...GPUCVTexture
...GPUCVCore
...GPUCV
...GPUCVConsole
...GPUCVSimpleApp
を読み込めません: そのようなファイルやディレクトリはありません
を読み込めません: そのようなファイルやディレクトリはありません
sed: ./projects/gnu/example/Makefileを読み込めません: そのようなファイルやディレクトリはありません
467:462
08/07/15 06:58:15
無視してmakeしても
#make
==== Building CUDA plugin ====
==== Building GpuCV ====
==== Building GPUCVHardware ====
GLBuffer.cpp
<コマンドライン>:1:20: 警告: missing whitespace after the macro name
/usr/include/GL/glxew.h:150: error: ‘GLulong’ has not been declared
make[2]: *** [../../../Debug/gnu//GPUCVHardware/GLBuffer.o] エラー 1
make[1]: *** [GPUCVHardware] エラー 2
make: *** [SUB_DIR_ALL] エラー 2
468:デフォルトの名無しさん
08/07/15 07:35:41
>>465
cutilはサンプル集だと思ってた方がいい。
469:デフォルトの名無しさん
08/07/15 19:02:35
カーネル内で二次元配列を使うと
Advisory: Cannot tell what pointer points to, assuming global memory space
というwarningが出るのですが、どのようにすれば出なくなるのでしょうか？
470:デフォルトの名無しさん
08/07/16 00:38:28
>>469
「カーネル内」ってどこのこと?
いっそ再現する最小限のソースを貼ってくれた方が話が早そうな希ガス。
471:デフォルトの名無しさん
08/07/16 01:02:45
>>467
必要なもの
premake(これはビルドするかパッケ拾ってくるのじゃいいな
OpenCV
libglew1.5-dev
SugoiToolsとかいうふざけた名前のライブラリSVNから盗んでくる
ちなみにこのSugoiBakaToolsを作ったやつはUnixとLinuxの.soを作る流儀を
知らん白雉なので許してやれ
svn co URLﾘﾝｸ(sugoitools.svn.sourceforge.net) sugoitools
cd sugoitools
premake --file Premake.lua --target gnu --os linux
make
これで./libにlibSugoiTools.soが生成される。
次に、GPUCVをSVNから盗んできてくれ
cd gpucv
premake --file Premake.lua --target gnu --os linux
cp ../sugoitools/lib/gnu/*.so ./lib
ln -s ../sugoitools/include/gnu/SugoiTools SugoiTools
ln -s ../sugoitools/include/gnu/SugoiTracer SugoiTracer
make
後は必要なライブラリ入れるだけでうまくいく
いかなかったら>>470にゴルァしてくれたらまた何かかくぉ？
472:462
08/07/16 11:23:00
>>471
うひょーPremakeの使い方を始めて知ったお！
でもGPUCVTextureのコンパイルで止まる…
>必要なもの
>premake(これはビルドするかパッケ拾ってくるのじゃいいな
>OpenCV
>libglew1.5-dev
premake,OpenCVはインストしました。libglew1.5-devがCentOSで
rpmがなかったので
URLﾘﾝｸ(chihara.naist.jp)
を参考に導入しました
>svn co URLﾘﾝｸ(sugoitools.svn.sourceforge.net) sugoitools
>cd sugoitools
>premake --file Premake.lua --target gnu --os linux
>make
>これで./libにlibSugoiTools.soが生成される。
>次に、GPUCVをSVNから盗んできてくれ
>cd gpucv
>premake --file Premake.lua --target gnu --os linux
ここまで完璧です！
473:462
08/07/16 11:27:58
>cp ../sugoitools/lib/gnu/*.so ./lib
コピー先は../resources/lib/gnuでしょうか、
でもgpucv/lib/gnuのフォルダにGPUCVHardwared.soができてるのですが。。。
うーん、両方(/resourcesと/gpucv/lib)試しましたがうまくいきません。。。
>ln -s ../sugoitools/include/gnu/SugoiTools SugoiTools
>ln -s ../sugoitools/include/gnu/SugoiTracer SugoiTracer
コピー先は../resources/include/SugoiTools等ですよね。。
こっちは/gpucv/Sugoitoolsで試してません
コピー先は
URLﾘﾝｸ(picoforge.int-evry.fr)
を参考に決めました。。
でmakeすると
# make
==== Building CUDA plugin ====
==== Building GpuCV ====
==== Building GPUCVHardware ====
..中略..
Linking GPUCVHardware
==== Building GPUCVTexture ====
DataDsc_GLTex.cpp
..中略..
TextureGrp.cpp
<コマンドライン>:1:19: 警告: missing whitespace after the macro name
../../../src/lib/GPUCVTexture/TextureGrp.cpp:100: error: prototype for ‘bool TextureGrp::AddTextures(DataContainer**, unsigned int)’ does not match any in class ‘TextureGrp’
../../../include/GPUCVTexture/TextureGrp.h:118: error: candidate is: bool TextureGrp::AddTextures(DataContainer**, size_t)
make[2]: *** [../../../Debug/gnu//GPUCVTexture/TextureGrp.o] エラー 1
make[1]: *** [GPUCVTexture] エラー 2
make: *** [SUB_DIR_ALL] エラー 2
うわああああああああああんごｒｒｒｒｒっるうううううううううううああああああ
474:デフォルトの名無しさん
08/07/16 23:05:01
>>473
GPUCVの馬鹿どもはMakefileすら満足にかけないみたいだから
あれだけど
premake間違ってるからちょっと修正
premake --file Premake.lua --target gnu --os linux --cuda
あとコンパイル関係のログは一番上見て
○○.h No sucha file or directoryって出てるはず
きちんとログは最初から理解していきましょう。
以下の場所にGPUCVxxxx.makだかって糞Makefile入ってるから
ここで○○.hが足りないって言われたら -Iでパス足してやって
vim ./project/gnu/lib/
それでうまくいくはず、この作業はちなみに6回ぐらい繰り返すので
GPUCVプロジェクト市ねって500回唱えるのがいい
475:デフォルトの名無しさん
08/07/16 23:59:18
sugoitoolsってなんだよｗ馬鹿にしてんのか？
と思ったらマジで日本語の「凄い」から命名してんだな
URLﾘﾝｸ(sugoiaware.free.fr)
476:デフォルトの名無しさん
08/07/18 13:45:05
>>475
どっかのインターフェイスカードかとおもったよｗｗｗ
477:デフォルトの名無しさん
08/07/18 15:48:34
確かにｗ
478:デフォルトの名無しさん
08/07/18 16:00:58
>>475
カナダにはSUGOIっていうスポーツウェアブランドもある
それくらいで驚いちゃ駄目だｗ
479:デフォルトの名無しさん
08/07/23 17:22:30
Windowsなんだけど、出来上がったexeって、
他のPCで動かすためには、exe以外に必要なものって何？
cudart.dllが無いって言われるのは、CUDA対応カードじゃないから？
480:125
08/07/23 19:17:13
cudart.dllをsystem32に放り込むかpath通せ。
481:デフォルトの名無しさん
08/07/23 19:27:03
.cuファイルって外部のincludeファイルは読み込めないのでしょうか
template.cuに
#include <cv.h>
と1行書き加え、（OpenCVという画像処理用のライブラリです）
Makefileに
# OpenCVのためにパスを通す
NVCCFLAGS+=`pkg-config opencv --cflags`
LIB+=`pkg-config opencv --libs`
でMakeすると
/usr/lib/gcc/x86_64-redhat-linux/4.1.2/include/mmintrin.h(49): error:
identifier "__builtin_ia32_emms" is undefined
/usr/lib/gcc/x86_64-redhat-linux/4.1.2/include/mmintrin.h(62): error:
identifier "__builtin_ia32_vec_init_v2si" is undefined
....................というのがダラダラ続く…
Error limit reached.
100 errors detected in the compilation of "/tmp/tmpxft_000010b9_00000000-4_template.cpp1.ii".
Compilation terminated.
make: *** [obj/release/template.cu_o] error 255
環境は
OS:CentOS 5.2 64bit
CUDA driver: NVIDIA Driver for Redhat Enterprise Linux 5.x with CUDA Support (174.55)
cuda toolkit:CUDA Toolkit version 1.1 for Redhat Enterprise Linux 5.x
CUDA SDK: CUDA SDK version 1.1 for Linux
です。よろしくお願いします
482:デフォルトの名無しさん
08/07/24 00:32:59
>>481
普通になんでもインクルードできるよ。
だからこそ、cv.h経由か何かで読み込んでも解釈できないmmintrin.hまでインクルードしてしまっているわけで。
取り敢えず、cv.hが何をインクルードしているのか順に見てみたら?
483:デフォルトの名無しさん
08/07/24 00:33:42
どうみても、Intrinsicの関数が無いと言ってるだろ
484:デフォルトの名無しさん
08/07/24 00:34:38
>>479
ドライバ入れてもらわないとダメじゃないのか
485:デフォルトの名無しさん
08/07/24 14:55:26
普段Radeonを使ってるんだけど、CUDAを使うならやっぱり取り替えなきゃダメかな？
もう一枚追加して普段はRadeonでCUDA用にだけGeforceを使えたらと思ったんだけど…
486:デフォルトの名無しさん
08/07/24 15:56:37
最近CUDACUDA言ってるやつ多いけど、ShやBrookGPUも忘れないでね。
487:デフォルトの名無しさん
08/07/24 15:58:31
>>485
ドライバがバッティングしないなら、その方が楽だね。GPU一枚だと、CUDAで暴走したときに画面が崩れる場合がある。
# ついでに言えば、Radeonで3Dアプリを動かしていてもCUDAの処理速度への影響が出にくいと言うメリットもある。
488:デフォルトの名無しさん
08/07/24 16:26:48
SIGGRAPH2008でいろいろやるらしーよ
民主化だってさ
URLﾘﾝｸ(www.siggraph.org)
489:デフォルトの名無しさん
08/07/24 21:50:54
GPGPUやってるやつらの論文
精度甘くてうそばっかりなんだけど
クイックソートレベルで嘘あるって
なんなんだろう
SIGRAPHも悪の片棒担いでるし
国内の正会員も歯切れ悪い胡散臭い爺多いし
あいつら本当に計算機科学全般に害悪な
ことばっかりしはじめたなぁ
490:デフォルトの名無しさん
08/07/24 21:58:06
>>489
具体的に晒してください
491:デフォルトの名無しさん
08/07/24 22:24:38
>>490
さらせねーよｗ俺社会的に終わる
それだけはヤダ
492:デフォルトの名無しさん
08/07/24 22:26:40
>>491
つまり、いいがかりなんですね
493:デフォルトの名無しさん
08/07/24 22:27:06
CPUに比べて○○倍ってよく見るけどCPUの方は最適化してなかったり１コアしか使ってなかったりやけに古いCPU使ってたり
数字をよく見せるために胡散臭いのはあるな
494:デフォルトの名無しさん
08/07/24 22:27:10
>>487
ドライバの干渉が一番の不安なんだよなあ。
Geforce2枚刺して用途別にするのが一番いいんだろうけど完全に環境を変えなきゃいけないし…
495:デフォルトの名無しさん
08/07/24 22:32:13
>>493
海外のも見てみてよ
査読ありの論文でレベル的ほとんど査読無しと
変わらないヒドイ質の論文多いから
GPUへのフェッチ性能とかそれのオーバヘッド
0扱いとかいう凄まじい論文あるぞ
496:デフォルトの名無しさん
08/07/24 22:52:43
>>493
結構同意
CPUの数十倍早くなりました、とあるが
CPU，GPU間の転送時間を含んでいませんとかアホかと
497:デフォルトの名無しさん
08/07/24 23:17:06
>>489
SIGGRAPHはもう正しく論文を精査できてない。
で、年度によってはとりあえず載せちゃえってスタンス。
それでも多過ぎてSIGGRAPH ASIAやっちゃえって事に。
498:デフォルトの名無しさん
08/07/24 23:19:53
あのーここは学会加入してないとだめなのでしょうか
皆さん頭のよさそうな話ばかりで困ります
499:デフォルトの名無しさん
08/07/24 23:27:22
同意する点もあるけれど、言いがかりに近いと思う点もある。
転送時間を含まないケースでも、CUDAに関して言えば転送時間を隠蔽できる可能性もあるし
初回だけで中間に転送がいらないケースもあると思う。
実際のところ、意味があるかはさておき単純なロジックを組んでみるとちゃんと理想値に近い時間で
処理できるのは事実だし、最新でないCPUでもGPUを使って高速化するという運用はありだと思うのだけど。
まぁ、一般論で片付けられる問題ではないから個別にレスするのは控えるけどね。
500:デフォルトの名無しさん
08/07/24 23:32:06
>>499
>転送時間を含まないケースでも、CUDAに関して言えば転送時間を隠蔽できる可能性もあるし
>初回だけで中間に転送がいらないケースもあると思う。
CUDA内部の処理と
CUDA-バス-キャッシュ-CPU間の話どっちを
しているわけなの？厳密に答えてくれないかな？思うじゃなくて
そこ厳密に計測すると汎用的なアルゴリズムになるとせいぜい40GFぐらいしか
出てない。8CoreのXeon２台MPIするよりも全然遅くなっちゃうんだけど
501:499
08/07/24 23:36:14
ついでに、WoodcrestXeon3GHzで8800GTを使った場合の実運用での処理時間について一言。
行列演算みたいなGPU向きの演算と違ってGPUには不利な演算なのだけど、
単体演算ではCPUのみに対してGPUを使った場合は約2倍の処理能力が得られた。
尤も、ファイル処理なども含めた実運用状態での処理能力比は1.3倍。
まぁこんなもんでしょってことで。要は、如何にGPU向きの演算に落とし込むかだね。
502:デフォルトの名無しさん
08/07/24 23:43:36
スレチだしそろそろ自重しようぜと過疎スレで言ってみる
503:デフォルトの名無しさん
08/07/24 23:44:14
>>501
それは何コア使って1.3なの？
SSEも入れて、TBL使ってMathLib使った場合と
比較して1.3倍？
504:499
08/07/24 23:46:29
>>500
「汎用的なアルゴリズム」で40GFlops出て、8coreXeon2台のMPIと勝負できるならいい方なんでない?
8coreXeonでかりかりにチューニングしたら、GPUなんて使わない方が速くて当然だと思う。
そういうCPUと勝負するなら、GPU単体の性能じゃなくてCPU+GPUでCPUより「何割」速くなるかでしょ。
あー、書き忘れてた。>501に書いたのは2core*2CPUのシステムで、CPUのみとCPU+GPUの処理能力比ね。
CPUでも1coreだけ使うよりは4core使った方が当然4倍近く速いのだけど、そこにGPUを足すことで
更に1.3倍になったということ。1coreだけと較べてどのくらい速くなったかは失念している。
1coreのみを1とすると、2core2CPUで3倍、それにGPUを足して4倍位じゃなかったカナ。
505:デフォルトの名無しさん
08/07/25 07:53:05
具体的な論文名出せよ、そうしないから話がややこしくなる
506:デフォルトの名無しさん
08/07/25 08:59:48
>>482
>取り敢えず、cv.hが何をインクルードしているのか順に見てみたら?
すみません、こういう時どう対処していいのかわからなくて。
includeしていく順番を探せば何か見えてくるのでしょうか
cv.h - cxcore.h - cxtypes.h - emmintrin.h - xmmintrin.h - mmintrin.h
とつながっていました。mmintrin.hはMMX？xmmintrin.hはSSEでしょうか
私の直感だとこのくらいの解釈が限界です…
includeファイルを順にコメントアウトしてみましたがやはり通らず…
gccではコンパイルを通っているのにnvccではダメということは
MMX？が使えないようになっているのでしょうか
__builtin_ia32ほげほげが何者なのかさっぱりです…
ググるとWinXPではOpenCVが使えているっぽいので
XPのnvccではMMXが使えてLinuxのは使えないのでしょうか…
507:デフォルトの名無しさん
08/07/25 10:18:10
何か足りてないんだろうね。
508:デフォルトの名無しさん
08/07/25 12:31:41
コンパイルできるようになりました。
cxtypes.hの中でSSE2がオンになっている場所
#if defined WIN64 && defined EM64T && (defined _MSC_VER || defined CV_ICC) \
|| defined __SSE2__ || defined _MM_SHUFFLE2
#include <emmintrin.h>
#define CV_SSE2 1
#else
#define CV_SSE2 0
#endif
を見つけました。見てみるとWIN64のときだけオンになるみたいで…
ここで
#include <emmintrin.h>
#define CVSSE2 1
をコメントアウトして
#define CVSSE2 0
と書いたらコンパイルできるようになりました。
Vistaではこのようなことをしなくても
コンパイルできたのでは32bitからだったのでしょうか…
509:デフォルトの名無しさん
08/07/25 12:44:05
>>509
お疲れさん。そう言えば、Windowsでは64ビットでコンパイルできないってレスがあったからその辺りも関係するかもね。
510:デフォルトの名無しさん
08/07/26 01:31:05
SIGGRAPHではないがCanny Edge Detection on NVIDIA CUDAなんてのを読むとなかなか笑える。
・Matlab並に速いアセンブラで最適化されたOpenCVと比較したよ！
~~~~~~~~~~~~~~~pugya-
・使ったCPUはIntel Core2 CPU 6600 @ 2.40 GHzだよ！
＾ EなのかQなのかはっきりしろと。
ま、所詮OpenCVだから結果に影響ないけどｗ
・GPUはGeForce 8800 GTX, 768 MB、OCなんてしてないよ！
~~~~~~~~~~~~~~~~~~~~研究レベルでOCを考慮するッ
そこにシビれる！あこがれるゥ！
511:デフォルトの名無しさん
08/07/27 02:49:51
研究なら並列アルゴリズムだけ書いておけばいいんじゃね
512:デフォルトの名無しさん
08/07/28 11:33:26
opencvってアセンブラで最適化されてたっけ？
513:デフォルトの名無しさん
08/07/28 19:52:47
>>512
移植性考えてまったくされてないよ
デフォルトではシングルスレッドだからマルチコアも使われてない
514:デフォルトの名無しさん
08/07/29 12:41:15
今後GPUはえらい成長遂げるのに対しCPUはもう脚詰まりだから
いきなり8coreXeon2台とGPU一個が張り合えるという事実は大したことだ
後はプログラムの処理内容を最適化させれば良い
515:デフォルトの名無しさん
08/07/29 16:07:34
ハードgf98GX2 *2
ソフトwinXP VS2005
ドライバ cuda2.0β
上記の環境でC++とcudaを混在させることはできるのでしょうか？
キャプチャーカードのSDKがC++なので・・・
516:デフォルトの名無しさん
08/07/29 16:29:57
>>515
ドライバはCUDAじゃねぇぞ。
CUDAの開発では、*.cuのファイルがnvccでコンパイルしてデバイスモジュールとホストモジュールが出力される。
その後リンクするのはVSのリンカになるので、この段階でVSのオブジェクトモジュールとリンクできる。
但し、nvccはベースがgccなのでextern "C"を使うなどの工夫が必要かも知れず。
517:デフォルトの名無しさん
08/07/29 23:27:10
GCC使う以上、素直にLinuxで開発すれば良いのに。
なんでワザワザ苦労してWin糞使うのか？
518:デフォルトの名無しさん
08/07/29 23:49:17
大昔のように、
プログラム開発者＝ハード＆ソフト環境を自由に整えられる人
じゃないんだよ。
519:デフォルトの名無しさん
08/07/30 02:02:48
nvccってclのバージョンでエラー吐くから
Windowsだとclベースじゃねーの。
520:515
08/07/30 09:46:46
>516
ありがとうございます。自身で試行錯誤してみます。
>517
MFCを使ってwindowsアプリケーションを作るのにwinの方がいいかと思いまして
521:デフォルトの名無しさん
08/08/13 00:43:16
対応カードのドライバがXP～しか対応してないみたいだけど、Win2kでプログラミングしてる人はいますか？
余ったPCでCUDAをやってみようと思ってるんですが>>173あたりを見ても対応して無さそうで…
522:デフォルトの名無しさん
08/08/14 01:03:10
余ってないPCに二枚挿した方がよくね?
523:デフォルトの名無しさん
08/08/14 07:57:32
余ったPCにLinux載せちゃうのがよくね?
524:デフォルトの名無しさん
08/08/14 20:52:04
余ったPCを俺にくれればよくね？
525:デフォルトの名無しさん
08/08/19 14:01:19
いつのまにかにノート向けも対応リストに載ってるのね
526:デフォルトの名無しさん
08/08/19 21:25:04
2.0 北っぽいよ
URLﾘﾝｸ(www.nvidia.com)
527:デフォルトの名無しさん
08/08/20 08:54:18
>>515
.NET for CUDA つうのがあるみたいなんだわ
突撃してみてほしいです。
528:デフォルトの名無しさん
08/08/21 20:11:04
日本版でもBeta2取れたけど相変わらずアナウンスはないな
529:デフォルトの名無しさん
08/08/25 01:43:22
SDK2.0入れて見たんだが、threadMigrationってサンプルプログラムが
"cuCtxDestroy","cuCtxPopCurrent","cuCtxPushCurrent"が定義されてないってエラー吐いてる。
他のサンプルはコンパイル通るので、インストールが悪いのかSDKが悪いのか分からん。
530:デフォルトの名無しさん
08/08/26 00:27:10
バージョンアップの観察はいいから何か作れよお前らｗｗｗ
531:デフォルトの名無しさん
08/08/26 00:32:38
作ってるよ～
仕事だから詳細語れないけど。
532:デフォルトの名無しさん
08/08/26 00:40:43
作ってるけど研究なので言えない
けど上手くいかない
ああああああ
はぁ…
533:デフォルトの名無しさん
08/08/26 00:48:53
研究って学生の卒論ですか？
それとももっと高いレベル？
534:デフォルトの名無しさん
08/08/26 04:37:54
自宅研究員…
535:デフォルトの名無しさん
08/08/26 11:45:11
>>533
CUDAをどう使うかと、論文のレベルは関係ないだろｗ
536:デフォルトの名無しさん
08/08/26 15:24:13
URLﾘﾝｸ(gpu4vision.icg.tugraz.at)
CUDAを使った領域選択
537:デフォルトの名無しさん
08/08/27 14:26:05
CUDAのビデオ講義、リンクきれちゃってるね
せっかくiPod nano買ったから英語の勉強に聞きたかったのに…
538:デフォルトの名無しさん
08/08/27 22:19:54
２日でCUDAマスターしますた
すごい簡単だなこれ
なんか質問あれば聞いて良いよ
539:デフォルトの名無しさん
08/08/27 22:28:04
>>538
100万桁のπは何秒くらいで計算できますか？
1995年当時のスパコンで5秒らしいですが超えられますか？
540:デフォルトの名無しさん
08/08/27 22:40:52
>>539
やってみないとわかりません
541:デフォルトの名無しさん
08/08/27 23:03:22
CUDAでソートするアルゴリズム思いついた うひょ
542:デフォルトの名無しさん
08/08/27 23:40:53
だれか多倍長をCUDAで効率的に実現する方法教えてくださいー
543:デフォルトの名無しさん
08/08/28 01:21:54
>>541
URLﾘﾝｸ(www.nvidia.co.jp)
544:デフォルトの名無しさん
08/08/28 03:35:38
>>543
バイオニックソートってやつと同じｗｗｗ
車輪の再発明かｗ
545:デフォルトの名無しさん
08/08/29 19:20:37
当方 Mac なんだけど、CUDA 2.0 インスコしてサンプル make したのは
いいんだけど実行しようとすると
dyld: Library not loaded: @rpath/libcudart.dylib
Referenced from: /Developer/CUDA/bin/darwin/release/scalarProd
Reason: image not found
Trace/BPT trap
って出て先に進めない。どういうことなの。。。。。。
UNIX 詳しい人教えて下さい
546:デフォルトの名無しさん
08/08/29 19:43:28
並列に計算して計算結果を１つの変数に合計するってのをやってるんだけど
10回に1回くらい計算結果が狂うのはなぜだ
547:デフォルトの名無しさん
08/08/29 22:37:59
>>546
Ａｔｏｍｉｃ関数使ってみるとか・・・後、parallel reductionを読んでみると
良いかも
548:デフォルトの名無しさん
08/08/29 23:27:49
>>546
そういうのはCUDAの最も苦手とするところだ。
全スレッドで同期を取って、代表1スレッドが合計するのが手っ取り早いが遅い。
全スレッドで同期を取って、代表nスレッドがmスレッド分合計してからnスレッド分を合計するのが無難か。
或いは、n個になった時点でCPUに転送してしまう方がいいかもしれない。