【GPGPU】NVIDIA CUDA ..
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が二つ見える筈なので、一つの時のままだとダメな気がします。
次ページ最新レス表示スレッドの検索類似スレ一覧話題のニュースおまかせリスト▼オプションを表示暇つぶし2ch
4326日前に更新/252 KB
担当:undef