【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/ 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出た時点で一番コストパフォーマンスいいの選ぶわ。 311 名前：デフォルトの名無しさん mailto:sage [2009/12/06(日) 01:23:34 ] 試行錯誤するのにcompute capability 1.3に対応した ファンレスカードとかあれば良いんだけどな。 260や295を付けっぱなしってのは、なんか精神衛生に良くない。 312 名前：デフォルトの名無しさん mailto:sage [2009/12/06(日) 01:25:46 ] CUDA勉強中の者ですが共有メモリの利用で躓いてるところです。アドバイス頂けたら幸いです。 下記の二つのカーネルでCUDA_karnel_sの方が5倍時間がかかってしまうのですが原因がわかりません。 違いは読み込んだデータをグローバルメモリに保存するか共有メモリに保存するかです。 __global__ void CUDA_karnel_g(uchar4 *vram, int sx, int sy, uint1 *vram2) { 　　int i = threadIdx.x; 　　vram += blockIdx.y*sx; 　　vram2[i] = ((uint1 *)vram)[i]; 　　__syncthreads(); 　　uchar4 px; 　　*((uint1 *)&px) = vram2[i]; 　　unsigned char Y = (unsigned char)(0.2126f*px.z+0.7152f*px.y+0.0722f*px.x); 　　px.z = px.y = px.x = Y; 　　vram[i] = px; } __global__ void CUDA_karnel_s(uchar4 *vram, int sx, int sy, uint1 *vram2) { 　　int i = threadIdx.x; 　　vram += blockIdx.y*sx; 　　__shared__ uint1 shared[125*32]; 　　shared[i] = ((uint1 *)vram)[i]; 　　__syncthreads(); 　　uchar4 px; 　　*((uint1 *)&px) = shared[i]; 　　unsigned char Y = (unsigned char)(0.2126f*px.z+0.7152f*px.y+0.0722f*px.x); 　　px.z = px.y = px.x = Y; 　　vram[i] = px; } まずバンクコンフリクトを疑ったのですがCUDA Visual Profilerでみるとwarp serializeは0で発生していませんでした。 意図的にバンクコンフリクトを発生させると更に１０％程遅くなるのでバンクコンフリクトは原因ではなさそうです。 共有メモリは速いはずなのにグローバルメモリよりなぜ遅くなるのか悩んでいます。初歩的なミスだろうとは思うのですが。 313 名前：デフォルトの名無しさん mailto:sage [2009/12/06(日) 07:23:36 ] >>312 sharedが一定以上多いとOccupancyが下がるから、そこらへんじゃない？ Occupancyは実行効率にダイレクトに効いてくる。 Visual Profilerの実行ログにも出てくるし、SDKのtools/CUDA_Occupancy_calcurator.xlsで試算可能。 Shared Memory Per Block (bytes)のところに16000って入れると良い。(125*32*sizeof(uint1)) ちょっと計算してみると、 スレッドブロックのサイズが512ぐらいならまだマシ(67%)だけど、 64とかだと壊滅的に遅くなる(8%)。 あとpxがアドレス参照でローカルメモリ（VRAM上)に行ってないか心配だ。 314 名前：デフォルトの名無しさん [2009/12/06(日) 07:53:46 ] vram2[i] がレジスタのってたりしないかな。 親から vram3 として渡してみるとか、参照を i+1 にするとか。 315 名前：312 mailto:sage [2009/12/06(日) 10:18:05 ] >>313 仰る通りだったようです。shared[125*32]をshared[32]としたら劇的に速くなりました。 バンクコンフリクトを疑うあまりブロックあたりのスレッド数を32にしてたのも不味かったようですね。まだ未確認。 >あとpxがアドレス参照でローカルメモリ（VRAM上)に行ってないか心配だ。 *((uint1 *)&px) = shared[i]　を　px.x = shared[i].x; と書き換えたところ速度アップしたのでその可能性大です。最適化で同じコードになるんじゃね？と勝手に思ってました。 >>314 >vram2[i] がレジスタのってたりしないかな。 このコードは質問のために単純化したものだったのですが、単純化前にでptxファイルを出力して比較したときは ld.global.u32のところがld.shared.u32に変化しただけでしたのでレジスタにのったり最適化で消えたりはしてなさそうです。 とはいえptxの書式の資料を見つけられなかったので自信ありません。 参照を i+1の意味はちょっとわかりませんでした。すみません。 >>313-314 おかげさまでグローバルメモリより共有メモリが遅いという現象は解決しました。ありがとうございます。 ブロック専用なのにextern __shared__ の構文がなぜあるのか不思議に思っていたのですがこういう理由だったのかと。 しかし大きな共有メモリで遅くなるのは厄介な仕様ですね。 316 名前：313 mailto:sage [2009/12/06(日) 16:21:39 ] >>315 おつ。 勘が鈍ってなくて安心したｗ 共有メモリは多数のコアで共有される有限リソースだから、そこも一定以下に抑えないといけないって話ね。 あと本当にsharedを使う必要があるのか、一時変数（レジスタ）で済ませられないか？ sharedを使うべきなのは、同時に動くコア同士のやりとりがある場合と、 どう処理を分配してもcoarescedアクセスにできなくて、でもsharedを使えばできるという場合ぐらいかと思う。 あ、そういや、逆にレジスタが多くてOccupancyが上がらない場合の退避用とかにも使えるだろうか？ 317 名前：312 mailto:sage [2009/12/06(日) 17:53:52 ] 訂正　*((uint1 *)&px) = shared[i]　を　px.x = shared[i].x;の部分は勘違いでした。 １バイトしか書き換えてないので等価ではありません。無視してくださいです… >>315 実はどういう理屈で遅くなるのかよくわかってないのですがお蔭様で対処の方法がわかってひと安心です。 最終的にはいろいろな画像処理に使う予定なのでキャッシュ的な使い方をsharedでさせてみるテストでした。 比較のためテクスチャ版とshared版作ろうとしてはまってしまいました。 318 名前：デフォルトの名無しさん mailto:sage [2009/12/07(月) 12:02:53 ] 自宅でGPU4枚とかって人はなにに使うの？ 319 名前：デフォルトの名無しさん mailto:sage [2009/12/07(月) 12:36:53 ] エンコ 320 名前：デフォルトの名無しさん mailto:sage [2009/12/07(月) 13:11:16 ] ベンチ見てほくそえむ 321 名前：デフォルトの名無しさん mailto:sage [2009/12/07(月) 16:15:39 ] >>317 ptx出力を眺めれば判るけど、普通のCならできる最適化もGPU向けにはできないことが多いよ。 敢えて言えば、ポインタ変数に尽くvolatileがついているかのように振る舞うみたい。 例えば、 int function(int * foo) {int bar = 3; foo[0] = bar; return foo[0];} みたいなコードはCPU向けには {foo[0] = 3; return 3;} のように最適化されるのにGPU向けには {foo[0] = 3; return foo[0];} のように律儀に解釈される。 なので、ptx出力を読めるかどうかは割りと重要かも知れず。 # つーか、メモリアクセスの個数を数えるくらいのことは普通にやってる。 322 名前：デフォルトの名無しさん mailto:sage [2009/12/08(火) 01:19:36 ] すみません、かなり初心者です。 行ごとに要素数の違う配列をデバイス側に渡したいんですけど、 a=1; //ホスト側 float** mat=(float**)malloc(size); for(i=0;i<num;i++){ mat[i]=(float*)malloc(size / a); a*=2; } //デバイス側 float** mat_d; CUDA_SAFE_CALL(cudaMalloc((void**)&mat_d,size)); for(i=0;i<num;i++){ CUDA_SAFE_CALL(cudaMalloc((void**)&mat_d[i],size/a)); a*=2; } CUDA_SAFE_CALL(cudaMemcpy(mat_d,mat,size,cudaMemcpyHostToDevice)); for(i=0;i<M;i++){ CUDA_SAFE_CALL(cudaMemcpy(mat_d[i],mat[i],size/a,cudaMemcpyHostToDevice)); a*=2; } というようにはできないんでしょうか？ 323 名前：デフォルトの名無しさん mailto:sage [2009/12/08(火) 06:39:08 ] >>322 >CUDA_SAFE_CALL(cudaMemcpy(mat_d,mat,size,cudaMemcpyHostToDevice)); これがいらないんじゃない？ ホスト側のポインタ列をデバイスに渡しても使いようがない。メモリ空間が違う。 324 名前：デフォルトの名無しさん [2009/12/08(火) 09:12:09 ] 開発環境の入っていないマシンで動かすには、どのファイルを持って行けばよいのでしょうか？ 325 名前：デフォルトの名無しさん mailto:sage [2009/12/08(火) 10:25:10 ] >>291 個人的な興味なんですけど、CUDA使ってるcodecってなんで速いんでしょうね？ BYTEアクセスじゃ無いとしても、どの部分をGPUでやらせれば爆速になるのか不思議ﾁｬﾝﾃﾞｽ 326 名前：デフォルトの名無しさん mailto:sage [2009/12/08(火) 14:53:55 ] SDKのドキュメントの通りに，VS2005で サンプルプロジェクト template を元に使っているのですが このプロジェクトに新しい.cuファイルを追加した場合に そのcuファイルへの更新がビルド時に察知されません 毎回リビルドすれば一応反映できるのですが 通常のビルドで反映させるにはどうすればよいのでしょうか？ 327 名前：322 mailto:sage [2009/12/08(火) 15:52:18 ] >>323 すみません。 抜けてましたが、 >CUDA_SAFE_CALL(cudaMemcpy(mat_d,mat,size,cudaMemcpyHostToDevice)); は、 ホスト側であらかじめ計算を行い、 その結果を初期値としてポインタ列、matに与えた上で、 その値をデバイス側のポインタ列、mat_dにコピーし、 その値を使ってデバイス側で計算をする。 というつもりで書きました。 このようにホスト側で計算した値をデバイス側に渡す時は どのように記述するのがよいのでしょうか？ 328 名前：デフォルトの名無しさん mailto:sage [2009/12/08(火) 16:15:43 ] {1個、2個、4個、8個、16個、…}　みたいなデータを渡したいのかな?? 固めてまとめて渡しちゃったほうがオーバーヘッドが少ないと思います。 トータル何列あるよ、は別にパラメータで渡す。 (実際にCUDAのルーチン書く前に、コピー/戻しの時間を色々計ってみるといいです) cudaMallocしたデータにはホストからは触れないので、 ホストでmallocしたデータ(mat)に計算結果格納 →同じサイズでcudaMalloc(mat_d) →cudaMemcpyHostToDeviceで渡す なのでそれでいいです 329 名前：323 mailto:sage [2009/12/08(火) 17:35:58 ] >>327 >CUDA_SAFE_CALL(cudaMalloc((void**)&mat_d,size)); > >for(i=0;i<num;i++){ >CUDA_SAFE_CALL(cudaMalloc((void**)&mat_d[i],size/a)); >a*=2; >} 考えてみるとここも問題があって、cudaMallocということはデバイス側でポインタ列を確保しているんだけど、 そうすると&mat_d[i]というアドレスは、デバイス側にはバッファがあるが、 ホスト側には存在しないから、ここで例外になりそう。 やるならこんな感じかな？↓（未検証） float** mat_d; // GPU側に確保する（ポインタ列用）バッファ CUDA_SAFE_CALL(cudaMalloc((void**)&mat_d, count * sizeof(float*)); float** mat_d_tmp=(float**)malloc(count * sizeof(float*)); // ホスト側に確保する。内容はGPU側ポインタ列 for(i=0;i<num;i++){ // GPU側データバッファのポインタを格納 CUDA_SAFE_CALL(cudaMalloc((void**)&mat_d_tmp[i],size/a)); a*=2; } // GPU側ポインタ列をGPUに転送 CUDA_SAFE_CALL(cudaMemcpy(mat_d,mat_d_tmp,count * sizeof(float*),cudaMemcpyHostToDevice)); free(mat_d_tmp); てやっておいて、mat_dをカーネル呼び出しの引数にしてやるとか。 ここまで、バッファ作りしかしてないので、データ転送は別途必要。 たぶん>>322の最後４行のとおりで構わない。 ただデータ転送回数多いと多少とも時間かかるから、>>328の言うように 固めたほうがベターではある。 330 名前：323 mailto:sage [2009/12/08(火) 17:48:04 ] >>329 ちなみに、ポインタサイズが32bitか64bitかはホスト側に依存してることを確認した。 forums.nvidia.com/index.php?s=&showtopic=149501&view=findpost&p=943899 331 名前：デフォルトの名無しさん mailto:sage [2009/12/08(火) 17:59:16 ] >>325 あいまいな質問だけど答えてみる。 CPUよりGPUのほうが計算能力の総量はでかいから、本来GPUが速いほうが当然なんだけど、 以下のようなものは遅くなる場合もある。 ・計算負荷よりもメモリ転送がボトルネックになるもの ・処理を細かく分解して並列化することが難しいもの codecなどは、メモリ転送がボトルネックになりやすい傾向はあるものの、 だいたい画面の領域ごとに分解できる処理が多いから、多少とも速くはできるでしょう。 あと速くならない処理はCPU側でやればいい。 332 名前：デフォルトの名無しさん mailto:sage [2009/12/08(火) 20:37:12 ] >>328>>329 参考になりました。ありがとうございます！ まだ不慣れなので色々試行錯誤しながらやってみます！ 333 名前：デフォルトの名無しさん mailto:sage [2009/12/08(火) 21:48:05 ] >>325 CUDAだけ使って早いと思ってる？ｗ 334 名前：,,・´∀｀・,,）っ-○○○ mailto:sage [2009/12/08(火) 23:12:01 ] 動き検出なんかは当然CPUにやらせたほうがいいぞ SpursEngineなりCellなりがあるならそれでもいいが 335 名前：デフォルトの名無しさん mailto:sage [2009/12/09(水) 00:11:13 ] >>333-334 流石にGPUだけとは思ってないけど、atomでh264なんかエンコすると日が暮れる勢いだけど、ionプラットフォームだと実用範囲になる。 エンコ中でも完全にCPUを使い切ってないところを見ると、やっぱGPUをかなり使ってるんだなと勝手な想像 336 名前：デフォルトの名無しさん mailto:sage [2009/12/09(水) 06:34:40 ] IONはメインメモリをグローバルメモリとして使う、つまりマッピングしておけば CPUからもGPUからもフルスピードでアクセスできるメモリ空間を作ることができる。 ION2は買おうかな・・・。 337 名前：デフォルトの名無しさん mailto:sage [2009/12/09(水) 08:29:47 ] >>336 実際はそんなに甘くはなく、GPUからフルスピードでアクセスするためにはWriteCombinedモードに する必要があって、それを付けるとCPUからのアクセスが死ぬほど遅くなります。 通常は素直に転送（実際にはメモリコピーですが）したほうがマシです。 低価格で4GB近いデバイスメモリが使えるというメリットはありますが。 338 名前：デフォルトの名無しさん mailto:sage [2009/12/09(水) 09:47:42 ] そのためのMOVNTDQA/MOVTDQだろう。 まあ前者はSSE4.1でないと使えないがな。Atomは仕様不可。あしからず。 339 名前：デフォルトの名無しさん mailto:sage [2009/12/09(水) 10:11:54 ] それらを使っても十分に遅いと思うのですが。 340 名前： ◆0uxK91AxII mailto:sage [2009/12/09(水) 10:22:13 ] WCにしつつ、CPU->GPUはmovnt*でcopyすれば良さそうだね。 CUDAとか知らないけど:b 341 名前：デフォルトの名無しさん mailto:sage [2009/12/09(水) 10:26:15 ] >>339 だからCPU専用のワーク領域と分けて読み書きを最小限にする。 342 名前：デフォルトの名無しさん mailto:sage [2009/12/09(水) 19:25:43 ] >>337 そか、CPUのキャッシュが使えなかったね。 死ぬほどまで遅くなるとは思ってなかった。 343 名前：デフォルトの名無しさん mailto:sage [2009/12/10(木) 10:12:49 ] GTX295 Vista x64 CUDA2.3 の環境で、CUDA VS WIZARDで作成したプロジェクトを使っています。 質問1. ブロックあたりのスレッド数を512にすると計算結果が正しくなりません。256なら正しいです。なぜでしょう？ shared memory はグローバル変数として宣言して、各スレッドで共用（リードオンリー）しています。スレッド内で__shared__の宣言はありません。 質問2. 完成したので、SDKのsimpleMultiGPUを参考にGPU 2個使おうとしています。 とりあえずマルチスレッド用の空の関数を書いたところ、コンパイルはできるのですが、リンクできません。 cutStartThreadなどが外部で定義されてないと言われます。 simpleMultiGPUはビルドできるので、プロジェクトのプロパティを見比べてみましたが特に違いは見あたりません。 何が悪いんでしょうか？ どなたかお助けください。 344 名前：デフォルトの名無しさん mailto:sage [2009/12/10(木) 10:45:27 ] >>343 ・sharedにはCPUから書けませんが、各スレッドがリードオンリーならどこで書き込んでいるのでしょう。 ・cutで始まる名前の関数はcutilライブラリ内にあります。cutilライブラリもビルドし、あなたのプロジェクトからリンクできるようにしないといけません。 345 名前：デフォルトの名無しさん mailto:sage [2009/12/10(木) 12:13:48 ] >>344 上については、言葉足らずでした。カーネルの最初に if( threadIdx.x == 0 ) { /* デバイスメモリからコピー */ } __syncthreads(); でコピーしています。 将来的にはコピーも並列化する予定です。 下についてはありがとうございます。既存のlib(dll?)だけでは駄目だってことですね。 346 名前：デフォルトの名無しさん mailto:sage [2009/12/10(木) 12:40:10 ] threadIdx.xはスレッドの識別子で、実行順番とは無関係では？ 347 名前：デフォルトの名無しさん mailto:sage [2009/12/10(木) 13:15:43 ] 共有メモリのつくりに関する知識が思いっきり欠如していると思われ。 threadIdxが0の人だけが書こうとしたら、他の人はみんな何もできなくて時間が無駄。 おまけに、スレッド数が一定数を超えると(GTX295の場合は256か)、別の実行単位になってしまうから(表に見えない) 共有メモリ間転送が行なわれてしまってその点でも時間の無駄。 尤も、>345の様にthreadIdx0でしかコピーしない場合は転送は発生しない代わりに結果がご覧の通りなわけで。 どうせthreadIdx0がコピーしている間他はなにもできないんだから、一斉に同じものを書いてしまえばいいんでない? 348 名前：デフォルトの名無しさん mailto:sage [2009/12/10(木) 13:33:51 ] >>347 > threadIdxが0の人だけが書こうとしたら、他の人はみんな何もできなくて時間が無駄。 開発途中なので、とりあえずこうしているだけです。 将来的にはコアレスにします。 > スレッド数が一定数を超えると(GTX295の場合は256か)、別の実行単位になってしまうから(表に見えない) 512で上手くいかないのはこのあたりが原因のようですね。 349 名前：デフォルトの名無しさん mailto:sage [2009/12/10(木) 13:49:31 ] >>347 たびたびすみません。 「別の実行単位」の意味がよく分からないのですが、 要は1つのブロックが2つのSMに割り当てられるって意味ですか？ 350 名前：デフォルトの名無しさん mailto:sage [2009/12/10(木) 14:03:14 ] >>345 Sharedは文字通り、共有できるメモリなので、 各スレッドが１ワードずつ協力して作成→どこでも読み書き可能、 ですよ。0番スレッドだけに任せなくてもいい 351 名前：デフォルトの名無しさん mailto:sage [2009/12/10(木) 17:12:55 ] 自分は>>343じゃないけど256スレッドでOKで512スレッドで計算結果がおかしくなる理由がわからない。 >おまけに、スレッド数が一定数を超えると(GTX295の場合は256か)、別の実行単位になってしまうから(表に見えない) >共有メモリ間転送が行なわれてしまってその点でも時間の無駄。 というのは時間の無駄を指摘してるだけで結果がおかしくなる理由ではないという認識なのですがあってますか？ この時の共有メモリ間転送というのは別の実行単位に転送しているということ？ >>343 256スレッドでやればOKなのでもう興味はないかもしれないが、シェアードメモリの代わりにグローバルメモリで 共用して計算した場合は512スレッドの時も結果は正しくなるのですか？ 352 名前：デフォルトの名無しさん mailto:sage [2009/12/10(木) 19:27:33 ] >質問1. ブロックあたりのスレッド数を512にすると計算結果が正しくなりません。256なら正しいです。なぜでしょう？ カーネルコールの後でエラーキャッチのコード入れてる？エラー出てない？ 353 名前：デフォルトの名無しさん mailto:sage [2009/12/10(木) 21:05:48 ] 質問の傾向をみると、 CUDAって、普通のCのように自由度高く書けるけど、 実際は、サンプルソースからは読み取れない あれは駄目、こうしたら駄目、的な制約が多すぎるのかな？ DXAPI+HLSLのような、自由度の少ない環境の方が むしろ、良質なソースが書けるのかも。 354 名前：デフォルトの名無しさん mailto:sage [2009/12/10(木) 21:54:43 ] 参考資料 Device 0: "GeForce GTX 295" CUDA Driver Version: 3.0 CUDA Runtime Version: 2.30 CUDA Capability Major revision number: 1 CUDA Capability Minor revision number: 3 Total amount of global memory: 939524096 bytes Number of multiprocessors: 30 Number of cores: 240 Total amount of constant memory: 65536 bytes Total amount of shared memory per block: 16384 bytes Total number of registers available per block: 16384 Warp size: 32 Maximum number of threads per block: 512 Maximum sizes of each dimension of a block: 512 x 512 x 64 Maximum sizes of each dimension of a grid: 65535 x 65535 x 1 Maximum memory pitch: 262144 bytes Texture alignment: 256 bytes Clock rate: 1.24 GHz Concurrent copy and execution: Yes Run time limit on kernels: No Integrated: No Support host page-locked memory mapping: Yes Compute mode: Default (multiple host threads can use this device simultaneously) 355 名前：デフォルトの名無しさん mailto:sage [2009/12/10(木) 22:17:49 ] >>353 CUDAをただの言語だと思っていてGPUのアーキテクチャを理解していないからかと。 後半については同意。限られた記述で実装可能なアプリしか書かないからというのがより正確かも。 356 名前：343 mailto:sage [2009/12/11(金) 00:35:49 ] >>351 スレッド数512、デバイスメモリ使用でやってみましたが、やはり正しくありませんでした。 Emuではちゃんと動くので、原因はよく分かりません。 コンスタントも使って多のでそっちが原因かも知れません。 分からん。もういい。 >>343 の質問2については、multithreading.cppをプロジェクトに追加してないだけでした。 俺死ね。 357 名前：デフォルトの名無しさん mailto:sage [2009/12/11(金) 01:23:14 ] >353 テクスチャはまじで鬼門。 凝った使い方をしようとすると、コンパイル通っても動かない・動いても結果間違ってるがざらにある。 358 名前：デフォルトの名無しさん mailto:sage [2009/12/11(金) 01:42:00 ] >>357 あるある。 deviceemuともまた動作違うこと多いしね。 NEXUS出ればデバッグもうちょっと楽になりそう。 359 名前：デフォルトの名無しさん mailto:sage [2009/12/11(金) 09:51:22 ] >>343 歩留まり向上のため、重い処理をさせると計算が狂うことがあるらしい どこかのPDFで見た 360 名前：デフォルトの名無しさん mailto:sage [2009/12/11(金) 10:57:49 ] >>359 それちょっと表現が微妙に誤解を招くような・・・。 出荷されている状態で、計算がまともじゃないコアがたまにまぎれこんでいる。 グラフィック用途と違って、メーカーの想定より負荷が偏った計算をさせることがあるから、 やや耐性の低いコアの耐性がちょっと足りない部分がエラーを起こす。 高負荷プログラムなんかを使ってそういうコアというか製品をはじいてやれば、あとはほぼ安定稼働する、という話だと思う。 361 名前：デフォルトの名無しさん mailto:sage [2009/12/11(金) 11:17:51 ] はじめてのCUDAがいつまでたっても届かない 一体どこに発注したんだよ上の人 362 名前：デフォルトの名無しさん mailto:sage [2009/12/11(金) 12:42:26 ] >>360 CUDAのスパコンで1/6が不良品だったと言ってたが 363 名前：デフォルトの名無しさん mailto:sage [2009/12/11(金) 12:44:30 ] えっ　Teslaとかでもそうなのか？？

---

---

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

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