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

1 名前：デフォルトの名無しさん mailto:sage [2009/10/08(木) 19:29:37 ] このスレッドは、他のスレッドでは書き込めない超低レベル、 もしくは質問者自身何が何だが分からない質問を勇気を持って書き込むスレッドです。 CUDA使いが優しくコメントを返しますが、 お礼はCUDAの布教と初心者の救済をお願いします。 CUDA・HomePage www.nvidia.com/cuda 関連スレ GPUで汎用コンピューティングを行うスレ pc11.2ch.net/test/read.cgi/tech/1167989627/ GPGPU#3 pc12.2ch.net/test/read.cgi/tech/1237630694/ 210 名前：197 mailto:sage [2009/11/29(日) 22:02:50 ] ベンチマークやったら電源が落ちた・・・ 780Wじゃ足りないのかな 211 名前：デフォルトの名無しさん mailto:sage [2009/11/29(日) 22:28:55 ] >>210 よっぽどの詐欺電源でも買ってない限りは、さすがに足りないってことは無いと思うけど 初期不良じゃないか？ 212 名前：デフォルトの名無しさん mailto:sage [2009/11/29(日) 22:35:29 ] 熱落ちじゃないの？ 213 名前：デフォルトの名無しさん mailto:sage [2009/11/29(日) 22:39:18 ] Geforceは80 PLUS シルバー以上の 電源じゃないとまともに高負荷に耐えられないぞ 214 名前：197 mailto:sage [2009/11/30(月) 06:48:15 ] 80 Plusって書いてないわ・・・ CORAZON ってやつです。 215 名前：デフォルトの名無しさん mailto:sage [2009/11/30(月) 08:07:19 ] >>214 www.keian.co.jp/products/products_info/kt_780as_sli/kt_780as_sli.html 12Vが25AだなHDDとかいっぱい着いてたら微妙だな 最小構成なら問題ない範囲だと思うけど てかこれ安いな、これ買えばよかったｗ 216 名前：197 mailto:sage [2009/11/30(月) 17:40:28 ] 25Aが2つ書いてあるのはどういう意味？ 217 名前：デフォルトの名無しさん mailto:sage [2009/11/30(月) 20:42:55 ] 定員25人のエレベーター2基と、定員50人のエレベーター1基は違うというのは分かるな？ 218 名前：デフォルトの名無しさん mailto:sage [2009/11/30(月) 20:46:53 ] ＨＤＤとかPCIE用の電源が２本電源の中にあるってことだ ちゃんと分割して接続すれば12x25=300WあるからGTX295のTDP300Wは支えられる HDDとかつけるともう足りない ちなみにうちの700W表記の電源は一本36Aある 219 名前：197 mailto:sage [2009/11/30(月) 21:20:06 ] IDE電源にハードディスクは接続しているけれども、 PCIEの電源は2本ともGTX295に挿してるぜ。 そういう意味じゃない？ ごめん詳しくないんだ。 大人しくこのあたり買っておくかな。 www.scythe.co.jp/images/energia/energia-label0800.jpg 220 名前：デフォルトの名無しさん mailto:sage [2009/11/30(月) 21:25:34 ] >>219 検索すりゃわかるが今の電源でGTX295を動かしてる人が居るから動くって ＨＤＤ１個にしてみて周辺機器もはずしてＯＣしてるならデフォルトにして 配線を入れ替えたりしてみてダメなら初期不良だろう 221 名前：デフォルトの名無しさん mailto:sage [2009/11/30(月) 21:30:13 ] つまり、安いもんじゃないし保障が切れたらもともこもないから 電源が原因だとしても保障があるうちに買った店に持っていって確認してもらったほうがいい お店の人が電源がだめだと言うなら電源を買えばいいし お店にあるちゃんとした電源でもダメだったら初期不良で交換してもらえるからね 222 名前：197 mailto:sage [2009/11/30(月) 21:50:02 ] そうしてみる。 電源も一緒に買ったものだし、持って行ってみるわ。 めっちゃ勉強になった。ありがとう。 223 名前：197 mailto:sage [2009/12/01(火) 06:40:06 ] 追記 HDを1つにしたら、落ちるまでの時間が長くなりました 224 名前：デフォルトの名無しさん mailto:sage [2009/12/01(火) 07:27:11 ] 排熱もやばいんだろw きちんとしたケースとFAN買えw 225 名前：デフォルトの名無しさん [2009/12/01(火) 17:12:58 ] 多次元配列の領域確保、コピーってcudaMallocとcudaMemcpyでできる？ 226 名前：デフォルトの名無しさん mailto:sage [2009/12/01(火) 18:48:14 ] できる。 てか普通のCみたいにコピーできる。 ちなみにCudaだと多次元配列だと面倒だから1次元配列として扱うことがおおい。 cudaMemcpyAsyncってのもある 227 名前：デフォルトの名無しさん mailto:sage [2009/12/01(火) 20:22:02 ] 面倒っていうか１次元しか扱えないし 228 名前：デフォルトの名無しさん mailto:sage [2009/12/01(火) 21:14:19 ] 1次元だけだったか 自分の技量が足りなくてできないのかと思ってた。 229 名前：デフォルトの名無しさん mailto:sage [2009/12/02(水) 19:12:31 ] 今日からCUDA触ってみたのですが、全然速くない… device側で1MB×2(dest, src)をアロケートして、hostからデータをコピー for (int n = 0; n < 1024*1024; n += 512) { CUDA_Func<<<1, 512, 0>>>(dest, src, nPos); } hostからdeviceへコピー ===== __global__ static void CUDA_Func(BYTE* pDest, BYTE* pSrc, int nPos) { int i = blockIdx.x * blockDim.x + threadIdx.x + nPos; pDest[i] = ((int)pDest[i] + (int)pSrc[i]) >> 1; } ===== なんて事をやっているのですが、CPUの方が速いです Visual Profilerを見ると、各CUDA_FuncのGPU Time は 8〜9us で終わってますが、CPU Timeが80〜150us になってます こんなものでしょうか？アドバイス頂けると嬉しいです Win7/GF8800/SDK 2.3 230 名前：デフォルトの名無しさん mailto:sage [2009/12/02(水) 19:56:18 ] >>229 CUDA_Func<<<1, 512, 0>>>(dest, src, nPos);を2ｋ回も呼んでるのがまず悪いんじゃね？ あとは詳しい人に任せた。 俺も勉強中。共有メモリのバンクコンフリクトがわけわからねえ。 231 名前：デフォルトの名無しさん mailto:sage [2009/12/02(水) 21:07:48 ] >>229 そりゃカーネルをキックするコストはかなりでかいから、ループで何回も呼んだらCPUにまけるだろ。。。。 >>230 shared memoryのbank conflictは ものすごーーーーく大雑把にいうと、thread_id順でshared memoryのアドレスにアクセスすると各バンクのチャネルがぶつからなくて、パラレルに出来るよってお話 232 名前：229 mailto:sage [2009/12/02(水) 21:33:41 ] レスありがとうございます。 リニアに1次元配列を処理するような事は意味がないと言う事でしょうか？ 例えば、ある程度の長さの、サンプリング単位のPCMの演算や、ピクセル単位の画像の演算とか… 233 名前：229 mailto:sage [2009/12/02(水) 21:40:38 ] 連投スマソ なんかレスを書いてて、やっとピンと来たんですが、 例えば各スレッドでさらにループで回して、CUDA_Funcを減らせば良い的な話だったりします？ >>229 のコードで言うと CUDA_Func内で1KB分ループさせて、各スレッドへは1KBのオフセットを渡す。 その分、CUDA_Funcの起動回数を減らす。 違う…？ 234 名前：初心者 mailto:sage [2009/12/02(水) 21:55:19 ] >>299 通常CPUなら、forで何回もやるような処理を CUDAのカーネルを一発たたくことによって処理させるっていうのが基本的な考え方じゃないの？ あと、メモリは一度になるべく大きくとって転送したほうが効率がいいらしいよん 235 名前：デフォルトの名無しさん mailto:sage [2009/12/02(水) 22:17:38 ] >>233 そのコードを見た感じ、1024*1024*512スレッドつかってることになってるけど、 何をしてるの？ 236 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 06:31:18 ] いや>>229のやり方は正しいよ 画面描画とＣＵＤＡは同期処理だから大きい単位でやると画面がタイムアウト起こす これ以上の最適化はＣＵＤＡでは不可能 これで遅いというならそれがそのカードの性能限界だと考えるしかないな 237 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 06:39:52 ] ちなみに>>229のCPUとカードの具体的な名前と周波数と PCIEの速度とx16 gen2とかね 遅いって実際にどれくらい遅かったのは知りたいね 上位のクアッドCPUと8400GSなんかじゃ勝負にならないのは当たり前だから 238 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 06:59:13 ] >>229 ごめん元のソースが何したいのかよく分からないんだけど、こんなことしたいわけ？ 1024ピクセル×1024ピクセルの二つの領域の明るさ平均を取るみたいな？ CPUなら、縦座標が外側ループで1024回×横座標が内側ループで1024回、回ると。 //device側で1MB×3(dest, src1, src2)をアロケートして、hostからsrc1,src2にデータをコピー //512スレッドを起動するなら、外側ループの回る回数はCPUだと1024、GPUだと1024÷スレッド数512 for (int n = 0; n < 1024 / 512; n ++) { 　CUDA_Func<<<1, 512, 0>>>(dest, src1, src2, nPos); } //deviceからHostへdestをコピー ===== __global__ static void CUDA_Func(BYTE* pDest, BYTE* pSrc1, BYTE* pSrc2, int nPos) { 　int i = nPos*512 + threadIdx.x * 1024;　//各GPUスレッドが動き出す起点、縦にずれてるわけ 　for (k = 0; k < 1024; k++)　//各GPUスレッドは呼ばれると横向きに1024回ループする。内側ループをGPUスレッドで処理。 　{ 　　pDest[i + k] = ((int)pSrc1[i + k] + (int)pSrc2[i + k]) >> 1; 　} } ===== ちなみにこれだといちいちグローバルメモリへのアクセスになるんで、 　//Sharedにスクラッチコピー 　//スクラッチコピー分だけループ処理 　//SharedからGlobalに書き出し するともっと速くなる 239 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 07:11:45 ] >>238 早けりゃいいってもんじゃないぞ そんなもん低クラスのカードで動かしたら一発で画面真っ暗だわ 240 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 07:16:32 ] あとsharedメモリはそんな使い方するもんじゃないだろう 毎回コピーしてたらそのコストの方がデカイわな 241 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 07:22:16 ] sharedメモリってあれだ １スレッドでやる計算が複雑な時に頻繁に変数の値を更新するだろ そういう時にグローバルメモリよりもアクセスが早い一時領域として利用するもんだ こういう計算自体が単純なケースでは効果はない 242 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 07:30:33 ] >>239 別人だけど8400GSくらいだとそうなの？経験上何msを超えるとハングする？ ググルとOSにより2秒や5秒でタイムアウトとあるがギリギリまでやるのはまずそうな気はする。 243 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 07:38:09 ] あ、失礼　coalescedになってなかった。こうかな?? //4回CUDA_Funcを呼び出す方向で。 for (int n = 0; n < 1024 / 256; n ++) { 　CUDA_Func<<<1, 512, 0>>>(dest, src1, src2, n); } //nは「縦」の分割数 //512スレッドが連続した512バイトを取り込む。二回動くと、1ピクセル×横に1024ピクセルを処理。 //上に向かって縦256回回る（k) __global__ static void CUDA_Func(BYTE* pDest, BYTE* pSrc1, BYTE* pSrc2, int n) { 　for (k=0; k < 256; k++) 　{ 　　//動き出す起点は各スレッドで1バイトずつずれてる。 　　int address = n*1024*256 + k*1024 + threadIDx.x; 　　//1024バイトを512スレッドで処理するので、二回。 　　pDest[address] = ((int)pSrc1[address]+(int)pSrc2[address])>>1; 　　pDest[address+512] = ((int)pSrc1[address+512]+(int)pSrc2[address+512])>>1; 　} } 実際書いて動かさないと良く分からないすな。グレーのビットマップ二つ用意してやってみる形かな。 244 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 08:17:42 ] >>242 タイムアウトが何秒とか議論することですらない マウスすら動かない状態が２、３秒も続くようなアプリはアプリ失格だろ 245 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 09:05:21 ] CPUより遅いくらいの8400GSで動かそうとしたなら2,3秒のフリーズ程度なら止む無し。 クラッシュしてデータを失わせるかもしれないリスクを犯すよりは 起動時に簡単なベンチ走らせて遅いGPUはハネちゃうのもありかな。 8400GSを考慮したせいでミドル以上のグラボの足を引っ張るとか馬鹿すぎる。 246 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 09:14:35 ] プログレスバー表示したら遅くなるから表示しないで画面を固まらせるなんて そんなものは個人で使うだけにするんだなｗ 247 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 10:20:49 ] ちょっと待て、みんな一台のGPUであれこれやろうとしているのか? それじゃ出るスピードも出ないぞ。 248 名前：229 mailto:sage [2009/12/03(木) 10:22:20 ] 皆様おはようございます。そして、レス感謝です。 朝イチで打ち合わせがあるので、結果だけ取り急ぎ報告します。 前のコード 0 memcpyHtoD 332.928 2155.51 1867.52 memcpyHtoD 332.512 1848.49 3403.26 CUDA_Func 10.624 1158.18 3588.86 CUDA_Func 8.864 119.289 (略) 767008 memcpyDtoH 289.504 997.333 CUDA_Funcでループ 0 memcpyHtoD 332.864 2149.65 1815.04 memcpyHtoD 332.512 1792.27 3264.26 CUDA_Func 11235.1 12351.3 28136.2 memcpyDtoH 286.368 1402.62 満足行く結果ではありませんが、速くはなりました。CPUでリニアに処理した方が速いです。AthlonX2 @1GHz〜3GHz あと、気づいたのですが、当方の環境ではRDP経由でCUDAが動きませんでした。ちょっとヤバイかも… 詳細は追ってフォローさせて下さい。 249 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 10:25:19 ] cudaMemcpyは同期を取ってから転送するから、結果の利用のタイミングぎりぎりまで実行を遅らせられれば 見掛け上の処理時間を短縮できるよ。 250 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 11:27:12 ] >>248 RDP経由でCUDAが動かないのは仕様です。 251 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 12:36:54 ] 俺もやってみたらこうなりましたが。 Using device 0: GeForce 9500 GT GPU threads : 512 Processing time GPU: 5.406832 (ms) Processing time CPU: 18.742046 (ms) Test PASSED GPU: 78.0000 87.0000 177.0000 1077.0000 CPU: 78.0000 87.0000 177.0000 1077.0000 Press ENTER to exit... カーネルはこう。 __global__ void testKernel( float* g_idata1, float* g_idata2, float* g_odata, int n) { 　// access thread id 　const unsigned int tid = threadIdx.x; 　// access number of threads in this block 　const unsigned int num_threads = blockDim.x; 　__syncthreads(); 　unsigned int startaddress = n * 1024 * num_threads; 　for (int j = 0; j < num_threads; j++) 　{ 　　for (int k = 0; k < 1024; k = k + num_threads) 　　{ 　　　unsigned int accessAddress = startaddress + k + tid; 　　　g_odata[accessAddress] = (g_idata1[accessAddress] + g_idata2[accessAddress]) / 2.0; 　　} 　　__syncthreads(); 　} } 252 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 12:40:39 ] //ホストでこうして、 for( unsigned int i = 0; i < 1024 * 1024; ++i) { h_idata1[i] = (float) (66.0 + i); h_idata2[i] = (float) (88.0 + i); } // こう実行。 for (int i = 0; i < (int)(1024 / num_threads) ; i++) { testKernel<<< grid, threads >>>( d_idata1, d_idata2, d_odata, i); } //CPUはこう。Athlon 2.3GHz。 computeCPU(float* idata1, float* idata2, float* reference) { for( unsigned int i = 0; i < 1024 * 1024; i++) { reference[i] = (idata1[i] + idata2[i]) / 2.0; } } これじゃサイズが小さすぎてあんまり比較にならないと思うっすよ。 景気よく4096×4096でやってみるといいかな？ 253 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 12:46:40 ] 4096x4096でやってみたらだいぶ差が出てきましたよ。 GPU threads : 2048 Processing time GPU: 33.380219 (ms) Processing time CPU: 260.214355 (ms) 254 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 13:12:34 ] accessAddressの計算おかしくありません？ 255 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 13:34:08 ] コードちゃんと見て無いけど、floatならGPUでパラで動かした方が速いに決まってる athlon系の浮動小数点演算のコストは、整数演算のざっと平均70倍 intel系は知らないにゃぁ 256 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 13:49:01 ] computeCPU(float* idata1, float* idata2, float* reference) { for( unsigned int i = 0; i < 1024 * 1024; i+=4) { static const __m128 div2 = { 0.5f, 0.5f, 0.5f, 0.5f }; __m128 tmp = _mm_load_ps(&idata1[i]); tmp = _mm_add_ps(tmp, _mm_load_ps(&idata2[i])); tmp = _mm_mul_ps(tmp, div2); _mm_store_ps(&reference[i], tmp); } } あとどっかでprefetchnta噛ませるといいかも。 CPU側は最低限SSE使おうや。 大学関係者も含めて比較用のCPU側コードが酷いのが多すぎる。 257 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 13:57:09 ] ぶたぎりですが、 geforce gtx285 と　tesla c1060 って どのくらい違います？ c2050がどうせでるから、c1060と285が 対してかわらないなら、285がいいかと思うのですけど。 ところで www.nvidia.co.jp/object/personal_supercomputing_jp.html いつのまにかnvidiaのページに 「Tesla GPU コンピューティングプロセッサは発売されています: - Tesla C2050/C2070 - Tesla C1060」 こんなことが書いてあるけど、まだ出てないですよね。 258 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 14:25:16 ] >>257 そんなこと言うと、NVIDIAの営業に「そんなアキバ的発想はダメですよ」って言われちゃいますよw 259 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 14:27:55 ] CPU側は、やっぱりCore i7で8スレッド並列とか動かしてあげないとだめじゃね？ そんでGTX285と勝負するみたいな。 260 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 14:29:52 ] 率直に言って、なんで値段が5倍のTeslaが売れるのか良く分からないのです＜アキバ的発想ですが Fermiアーキテクチャも、GeForceが先に出るんじゃないですか？　来年2月くらいでしょうかね。 261 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 14:33:23 ] >>260 そりゃぁ、GeForceが叩き台だモノ。NVIDIA曰く、GeForceよりも信頼性が10倍高いから5倍の価格差はペイできるって発想でしょ。 262 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 14:37:11 ] >>259 HPC用によく最適化されたコードだとHT切ったほうが速い >>260 サポートが手厚い 263 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 14:49:44 ] Engadgetに、米軍が旧型PS3を2000台買ったとか書いてあったよな $200ドルで256MB、100GFlops、て安いのかな?? 10倍して$2000ドル、2.5GB、1TFLOPS、て。どうなんだろ。GPUのほうが安そうな。びみょう？ 264 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 16:25:21 ] GTX285とか、240SPに、512とか1024とか2^nで放り込んでもうまく処理してくれるのですか？ 半端な部分の処理がどうなるのか心配していると脳みそかゆくなります。 265 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 16:37:48 ] >>264 cudaライブラリが「適当に」割り振ります。グラフィック兼用だと、そもそも全部使う保証さえありません。 266 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 19:13:14 ] 信頼性とか保証とかサポートとかの違いか 267 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 19:36:45 ] いや、メモリ搭載量全然違うし。 268 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 20:35:09 ] そうだね、GeForceの方が安いメモリを使っているしね。 269 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 22:47:05 ] はじめてのCUDAプログラミング買ったが テクスチャには触れてないも同然だな。 270 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 22:52:47 ] >>263 ゲーム機としても使うつもりだから 271 名前：デフォルトの名無しさん mailto:sage [2009/12/04(金) 01:08:38 ] >>251 実コード動かしたのは乙だけど、なんでデータがfloatになったの？ 整数演算だったら、結局ちゃんとコードが書かれていればボトルネックは転送速度であって、 DDR2-CPUとDDR2-(PCI-Ex16)-GPUの比較になりそうな気がする。 272 名前：デフォルトの名無しさん mailto:sage [2009/12/04(金) 01:42:05 ] CUDAの処理速度を計りたいならまずプロファイラにかけろ。話はそれからだ 273 名前：デフォルトの名無しさん mailto:sage [2009/12/04(金) 01:47:26 ] >>263 コードの資産の問題なんでしょ。 Cell用に書いていたので、GPUにポーティングする手間を考えたら、 PS3を大量に買った方が安いのでしょう。 もちろんアプリによるけど、PS3は実効性能で100Gflopsはでるけど、 GPUは実効性能で1TFLOPSも出ないでしょ。 自分の経験では、Cｅｌｌに対して、G200は大体2倍〜3弱倍という感じだった。 Cellはカリカリにチューンしたけどね。 274 名前：デフォルトの名無しさん mailto:sage [2009/12/04(金) 02:53:46 ] アキバ的発想ではやっぱり、GTX295を二枚というのが現時点ではさいきょうCUDA環境ですか 275 名前：デフォルトの名無しさん mailto:sage [2009/12/04(金) 05:25:55 ] >>273 GPU内部だけで完結させる処理なら１Tいくかもしれないが PCの場合はデータ転送がボトルネックになって１Gｆlopsとかになるが 276 名前：デフォルトの名無しさん mailto:sage [2009/12/04(金) 06:04:38 ] >>256 Type Zでやってみた。P9500(C2D　2.53GHz)、GeForce9300M GS(8SP) GPU Threads: 256, memsize = 16777216 Processing time GPU: 81.911575 (ms) Processing time CPU: 218.790421 (ms) Processing time SSE: 84.257759 (ms) Test PASSED Press ENTER to exit... ほう 277 名前：デフォルトの名無しさん mailto:sage [2009/12/04(金) 08:54:02 ] 結局さぁ、「ゲーム」だったらCPUで全部演算しようなんて考えるやついないんだし、 別にGPUでやるほどのもんじゃねぇ、ってだけなんじゃないの＞229の例 278 名前：デフォルトの名無しさん mailto:sage [2009/12/04(金) 11:40:15 ] >>274 アキバ的にはGTX295を４枚が最強だけど、 そこまでするならtesla買った方が幸せだわな。 279 名前：デフォルトの名無しさん mailto:sage [2009/12/04(金) 12:09:47 ] Larrabeeがなんかよさそう 280 名前：デフォルトの名無しさん mailto:sage [2009/12/04(金) 12:33:49 ] >>271 4096x4096でやってみたら(float) GPU threads : 2048 Processing time GPU: 33.380219 (ms) Processing time CPU: 260.214355 (ms) BYTEにしてみた。4096*4096で GPU threads : 2048 Processing time GPU: 27.527241 (ms) Processing time CPU: 345.093292 (ms) になった。BYTEのほうがCPU遅い…のはこれはこれでよいのかな 281 名前：デフォルトの名無しさん mailto:sage [2009/12/04(金) 14:40:50 ] >>278 電源がそれだけで1200Wくらい要るなぁ。 1000W電源二台って感じか。うちのブレーカーは間違いなく飛ぶだろう 282 名前：デフォルトの名無しさん mailto:sage [2009/12/04(金) 14:50:01 ] 電子レンジとかドライヤーつけなければいけるだろう たぶん 283 名前：デフォルトの名無しさん mailto:sage [2009/12/04(金) 15:05:13 ] 夏は地獄を見るだろう 284 名前：デフォルトの名無しさん mailto:sage [2009/12/04(金) 15:55:44 ] CUDAで倍精度演算をしたいと思っているのですが、 ・GeForce GTX 285 ・GeForce GTX 295 なら、どっちがいいですか？ また、もっといいものがあれば教えてください。 Teslaは高いらしいのでそれ以外で、お願いします。 285 名前：デフォルトの名無しさん mailto:sage [2009/12/04(金) 16:02:55 ] >>284 Fermiが出るまで待つ でなければGTX295でしょうねぇ。(メモリ・帯域幅は小さくても倍精度演算リソースが二倍) 286 名前：デフォルトの名無しさん mailto:sage [2009/12/04(金) 17:01:27 ] >>285 ありがとうございます。 Fermiがいつ出るかわからない状況なので、GTX295の方向で考えます。 いつ出るか発表されてませんよね？ 287 名前：デフォルトの名無しさん mailto:sage [2009/12/04(金) 19:46:16 ] cudaMallocHostってmallocのラッパーではない？ ライブラリつかってて、なぜかcudaMallocHostを指定していて mallocで大丈夫だろうと思ったら動作せず cudaMallocHostなら動作したので。 どう違うの？ 288 名前：デフォルトの名無しさん mailto:sage [2009/12/04(金) 19:58:37 ] >>287 つ　page lock 289 名前：デフォルトの名無しさん mailto:sage [2009/12/04(金) 20:54:29 ] >>285 １個づつ付いてるのを２倍って呼ぶのは詐欺くさいですな 290 名前：デフォルトの名無しさん mailto:sage [2009/12/04(金) 21:02:18 ] >>287 メモリコピーのハードウェアアクセラレートって決まった区切りで始まる決まったサイズのメモリ単位でしか扱えなくて そういうのを考慮してメモリを確保するんでしょう 291 名前：,,・´∀｀・,,）っ-○○○ mailto:sage [2009/12/04(金) 23:12:59 ] >>280 byteこそSSE使わないと。 最大16倍程度は速くなる 292 名前：,,・´∀｀・,,）っ-○○○ mailto:sage [2009/12/04(金) 23:25:37 ] 1byte単位のストアは部分書き込みになるからあまり性能的によろしくない。 アラインメント境界にあわせて128ビット単位のストアが一番パフォーマンスよろしい #include <emmintrin.h> computeCPU_SSE2(BYTE* idata1, BYTE* idata2, BYTE* reference) { for( unsigned int i = 0; i < 1024 * 1024; i+=16) { __m128 tmp = _mm_load_si128(&idata1[i]); tmp = _mm_avg_epu8(tmp, _mm_load_si128(&idata2[i])); _mm_store_si128(&reference[i], tmp); } } 293 名前：,,・´∀｀・,,）っ-○○○ mailto:sage [2009/12/04(金) 23:54:33 ] >>279 最終的にはメモリ帯域ネックなんで同等のメモリ帯域を持つGTXと変わらんと思うよ。 ただ2つの配列の平均取るだけならほぼ帯域ベンチ。 PCIeかGPUのL/S帯域で律速。 キャッシュメモリが256KB/coreほどあって16コアあれば、1024x1024ならキャッシュにほぼ収まるので それを有効に使って処理を連結していくなら、そこではじめてLarrabeeの旨味が出てくる。 個人的にはN-Bodyとかが良い勝負しそう。 294 名前：デフォルトの名無しさん mailto:sage [2009/12/05(土) 01:12:39 ] >>293 だんごさんてきにはGrape-DRが お勧めということですね 295 名前：デフォルトの名無しさん mailto:sage [2009/12/05(土) 01:17:38 ] 多体問題だけやりたいんならいいんじゃね？ 296 名前：デフォルトの名無しさん mailto:sage [2009/12/05(土) 02:09:51 ] おいおまいら、こんなのがあるぞ。 www.acceleware.com/index.cfm/cuda-training/2010sunnyvale/ リッチな会社にいたら旅行がてらいってみるといいぞ。 これ見るとＦｉｘｓｔａｒｓがいかに良心的かわかるぞ。 297 名前：,,・´∀｀・,,）っ-○○○ mailto:sage [2009/12/05(土) 12:49:04 ] >>295 Larrabeeは机上計算だと1コアあたり16並列で15サイクル（＋α）くらいで回るんで、かなり理想的なデバイスなんだけどね。 （ただしvrcppsの精度補完ありならもう少しかかる） 298 名前：デフォルトの名無しさん [2009/12/05(土) 13:54:20 ] GTX260 で決まりすよ。18000円なんだから4枚位買えばいい。 299 名前：デフォルトの名無しさん mailto:sage [2009/12/05(土) 14:08:17 ] ららびーはもう終わったみたいな話が出てますけど venturebeat.com/2009/12/04/intel-cancels-larrabee-consumer-graphics-chip/ 300 名前：,,・´∀｀・,,）っ-○○○ mailto:sage [2009/12/05(土) 15:09:19 ] 読んで字のごとくコンシューマ向けディスクリートはね。 何のためにスパコン向けカンファレンスで発表したと思う？ まあなんにせよ「パソコン」に刺さるカードは無くなったな。 PTX2.0の資料請求しとかなきゃ。 301 名前：284 mailto:sage [2009/12/05(土) 18:47:17 ] >>289 っていうことは、285の方がよさそうでしょうか？ メモリの帯域が大きく、こちらの方が使いやすそうな気がします。 デュアルGPUのコーディングも大変そうなので…… 302 名前：,,・´∀｀・,,）っ-○○○ mailto:sage [2009/12/05(土) 18:58:27 ] >>301 待てるならFermiを待った方がいいよ 303 名前：デフォルトの名無しさん mailto:sage [2009/12/05(土) 19:32:04 ] 倍精度性能はGT200世代はおまけで１個演算機がついてるだけだから期待するだけ無駄だよ Fermi世代から各コアに演算機付けるって言ってるけど、 一般向けには倍精度削ってGF100の名前で出すとか言ってるからどうなるか分からん CUDAに拘らないならRADEON HD5xxxって選択もある GT200世代より倍精度演算能力は圧倒的にHD5xxxの方が高いし ただRADEONはドライバが糞だしCUDAも動かないしいろいろ中途半端 Larrabeeがどうなるかってところか 現状実用的なものは無いから、実験的な目的以外では買わない方がいいし 実験的な目的なら安いやつでいいじゃないかという話になる 304 名前：,,・´∀｀・,,）っ-○○○ mailto:sage [2009/12/05(土) 20:29:53 ] Larrabeeはコンシューマ向けディスクリートGPUとしてはキャンセル。 HPC向けには続投しそうだから100万くらい出せば手に入るようになるかもよ？ 305 名前：デフォルトの名無しさん [2009/12/05(土) 23:36:41 ] RADEONはFFTを出せないところを見ると行列積が精一杯のようだよ。 GTX280は512bitのバンド幅がどうもよろしくないのでGTX260を奨めます。 306 名前：デフォルトの名無しさん mailto:sage [2009/12/05(土) 23:41:53 ] もしRADEONで遊ぶなら、現時点ではDXCSでSM5.0のスレッド制御を 使ってどこまでできるかだろうなぁ。 DirectX 11 SDKにFFTのサンプルコードなんかもあるから、持ってる人は色々 ベンチマーク取ってみて欲しいな。 307 名前：デフォルトの名無しさん mailto:sage [2009/12/05(土) 23:43:57 ] >>304 Tesla by Fermiの値段にぴったり張り付いて売るのがIntel式な気が しないでもない。そのためにもnVidiaには頑張ってもらわないと！ 308 名前：デフォルトの名無しさん mailto:sage [2009/12/05(土) 23:44:40 ] 295だろ、普通に考えて。 スレッド数を多く使える分だけ、高速化が容易 メモリのバンド幅とかよりも重要だと思うが？ 309 名前：デフォルトの名無しさん mailto:sage [2009/12/05(土) 23:51:04 ] 迷っている段階なら、とりあえずGTX260を買って試すのがオススメかな 310 名前：,,・´∀｀・,,）っ-○○○ mailto:sage [2009/12/06(日) 00:24:51 ] >>307 流石にIntelシンパの俺でも30万は出せない。 Fermi出た時点で一番コストパフォーマンスいいの選ぶわ。

---

---

[ 続きを読む ] / [ 携帯版 ]

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