【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/ 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とかでもそうなのか？？ 364 名前：デフォルトの名無しさん mailto:sage [2009/12/11(金) 13:49:03 ] Teslaは選別品 365 名前：デフォルトの名無しさん mailto:sage [2009/12/11(金) 14:52:34 ] 低価格帯で一番安定してるGPUってなに？ 366 名前：デフォルトの名無しさん mailto:sage [2009/12/11(金) 15:53:13 ] slidesha.re/5FtABc のP.26 長崎大の人に言って選別プログラムを貰うことだなｗ 367 名前：デフォルトの名無しさん mailto:sage [2009/12/11(金) 16:12:17 ] CUDAは保障外だからCUDAやりたい奴は動くのが当たるまで買えってことなんだろう 368 名前：デフォルトの名無しさん mailto:sage [2009/12/11(金) 16:17:50 ] 仮に2/3超の確率で完動しないとしても、自前で問題無く選別できるなら Tesla買うよりコンシューマGPU買うほうがまだ安い位だから、たくさん 買う人は自前で選別できるように頑張るのが正解だよなぁ。 369 名前：デフォルトの名無しさん mailto:sage [2009/12/11(金) 16:40:06 ] 多少クロック落とせば動くんじゃないのかな。CPUのオーバークロックと一緒で。 370 名前：デフォルトの名無しさん mailto:sage [2009/12/11(金) 20:54:29 ] >>361 自分は11/27楽天でぽちって 12/01に到着。 ぼちったあとすぐ納期の表示が延びたから焦ったけど滑り込みセーフだったらしい。 書いてあることはわかり易い。しかし意外なことが触れられていなかったりする。 371 名前：デフォルトの名無しさん mailto:sage [2009/12/11(金) 22:49:23 ] >>360 その状態が続く限り、コンシューマ用パッケージソフトじゃGPGPUを使えないね。 「動かない」サポートの爆発でたいへんなことになる。 372 名前：デフォルトの名無しさん mailto:sage [2009/12/11(金) 23:34:51 ] てか>>356の作ってるソフトが選別ソフト代わりになるんじゃね？ 373 名前：デフォルトの名無しさん mailto:sage [2009/12/12(土) 13:58:47 ] おれらまだまだ、「主メモリ側が主役」って固定観念なくね？ いかんな。どかどかデバイスメモリにロード、大量のブロック×スレッドを駆使、 CPUはたまにお手伝いをさせていただく、位に思わないとな。俺まだまだ 374 名前：デフォルトの名無しさん mailto:sage [2009/12/12(土) 15:52:58 ] メモリ転送と計算を非同期で出来るのかね？ 375 名前：デフォルトの名無しさん mailto:sage [2009/12/12(土) 15:58:25 ] >>374 マルチスレッドで、CPU側は「GPU様が計算し終わるまで寝とけや」くらいの扱いにするわけよ 376 名前：デフォルトの名無しさん mailto:sage [2009/12/12(土) 16:43:09 ] >>374 できる。 基本GPUとCPUでは非同期。 ストリームを設定すれば、GPUで演算中にデータ転送させることも可能。 ちゃんと同期させるための関数もある。 377 名前：デフォルトの名無しさん mailto:sage [2009/12/13(日) 15:51:31 ] > NEXUS いつ出るんだっけ VisualStudio2008でいいんだっけ 378 名前：デフォルトの名無しさん [2009/12/15(火) 06:00:58 ] カーネルで計算した結果をCPU側に返すにはどうしたらいいの？ return aのように簡単にはできないのですか？ 379 名前：デフォルトの名無しさん mailto:sage [2009/12/15(火) 06:22:49 ] >>378 ちょｗ カーネルfunctionの引数で　float * outdata とかやってカーネルからこれに書き込む　outdata[i] = result; ホスト側でcudaMemCopyのdeviceToHostで、ホスト側でcudamallocしたfloat * に戻す みたいになる。ややこしいぞ？ 380 名前：デフォルトの名無しさん mailto:sage [2009/12/15(火) 19:31:04 ] thrust便利だなしかし 381 名前：デフォルトの名無しさん mailto:sage [2009/12/16(水) 02:52:17 ] thrust、自分で書いたカーネル関数はどう使うのっすか 382 名前：デフォルトの名無しさん mailto:sage [2009/12/16(水) 15:26:45 ] __shared__ int sh[num*2]; /* numはスレッド数 */ という共有メモリの配列をソートし、最大or最小のデータのみグローバルメモリに書き込みたいのですが これをカーネルの中でやるいい方法を教えてください。 383 名前：デフォルトの名無しさん mailto:sage [2009/12/16(水) 17:36:07 ] ソートは不要で、reductionで半分を繰り返せばいいのでは 384 名前：デフォルトの名無しさん mailto:sage [2009/12/16(水) 23:18:35 ] >>383 言われてみればそうですね。 解決しました。thx 385 名前：デフォルトの名無しさん mailto:sage [2009/12/17(木) 03:27:35 ] どこかにfortranでの使用例って ありませんか 386 名前：デフォルトの名無しさん mailto:sage [2009/12/18(金) 16:42:19 ] そういえば見かけたことねえな 387 名前：デフォルトの名無しさん mailto:sage [2009/12/18(金) 17:17:27 ] CUDAはそもそも実行速度がシビアで 高級言語向けではないと思うけどな 388 名前：デフォルトの名無しさん mailto:sage [2009/12/20(日) 02:21:29 ] 共有メモリは16個のバンクに分割されて、半ワープのスレッドが同じバンクにアクセスしなければ競合は起こらない。 みたいな事の意味がイマイチ分かりません。 例えばブロックの共有メモリが8KBなら0.5KBづつに分けられて、半ワープのスレッドが0.5KBの中ならどこにアクセスしても良いって意味なの？ それとも、先頭から4ﾊﾞｲﾄとか8ﾊﾞｲﾄとかの固定の領域に連続して分けられてるの？ でもそれだと分けられる領域が4ﾊﾞｲﾄなのか8ﾊﾞｲﾄなのか分からないじゃまいかヽ(`Д´)ﾉ カーネルの引数にvoid*とデータサイズを渡して任意のデータ長で処理させたいのですが、 こういう事をしようとすると元からバンク競合は避けられないのですか？教えて、エロイ人(´・ω・`) 389 名前：デフォルトの名無しさん mailto:sage [2009/12/20(日) 02:33:00 ] 今あるMPIのプログラムをGPUに対応させようと考えているんだけど、 CUDAでMPIを使うとき、データの転送は一端ホスト側のメモリに転送してから 他のノードへ転送することになるの？ だとするとレイテンシがすごく大きそうだね。 それとも専用のライブラリなんかあるのかな？ GPU側のグローバルメモリがホスト側のメモリにマップされていれば、 GPUを意識しないで転送ができそうなんだが。 390 名前：デフォルトの名無しさん mailto:sage [2009/12/20(日) 08:02:33 ] >>388 おれの解釈だと、共有メモリには16個のアドレス入力と16個のデータ入出力が有る と思ってるんだけど　違うかも 391 名前：デフォルトの名無しさん mailto:sage [2009/12/20(日) 09:55:19 ] >>388 連続する64バイトの領域（先頭アドレスが４の倍数）について、4バイト(32bit)のバンク１６個に分かれています。 例えばバンク０に属するアドレスは０、６４、１２８、１９２、２５６、・・・・・から始まる４バイトの領域。 任意のデータ長が４バイトより大きいのか小さいのか分かりませんが、テンプレート等を使う場合には 型に合わせて場合分けをする必要があるでしょう。 >>390 「１個のアドレス入力と１個のデータ入出力ができるバンク」が１６個あるというのが適切な表現かと。 392 名前：デフォルトの名無しさん mailto:sage [2009/12/20(日) 10:00:29 ] >>389 device memory on node 0 -> host memory on node 0 -> host memory on node 1 -> device memory on node 1 となるのでレイテンシは大きくなります。今のところCUDAではデバイスメモリをメモリマップする手段はありません。 ただし十分大きなデータを転送する場合にはパイプラン化すれば問題なくなると思います。 むしろpinnedメモリとMPIライブラリの干渉が一番の問題・・・・。 393 名前：デフォルトの名無しさん mailto:sage [2009/12/20(日) 10:47:33 ] >>392 レスサンクス やはりレイテンシは大きいのだね。 目的のアプリは数10kbyteから数100kbyteの転送を頻繁に行うから、 せっかくGPUで速くできても、転送ボトルネックで詰まりそう。 転送するサイズも小さいので、page lockさせない方がよいのかも。 394 名前：デフォルトの名無しさん mailto:sage [2009/12/20(日) 14:34:19 ] GPUの購入で悩んでいるのだが、 Tesla C1060, GTX295,GTX285のうち結局どれが 速いんですか？ちなみに流体に使います。 GTX295ってGPU2基積んでるけど並列プログラミング しないと2基機能しないとか？ 素人質問で申し訳ない。 Teslaの保証とサポートも魅力的だが。 395 名前：デフォルトの名無しさん mailto:sage [2009/12/20(日) 14:44:32 ] >>394 メモリ量でC1060になっちゃう　なんてことないかな？ 295はちゃんとマルチスレッドでプログラムしないと二基使えないです。 396 名前：デフォルトの名無しさん mailto:sage [2009/12/20(日) 15:45:47 ] >>390 >>391 なるほど…よく分かりました。ﾚｽ感謝。 処理させたい任意のデータ長は5ﾊﾞｲﾄだったり、11ﾊﾞｲﾄだったり半端な数もきます。 任意のデータ長に対してコピーやビット演算を行なうんですが、データ長が4ﾊﾞｲﾄより大きいと もうバンク競合は避けられない感じなんですね。 プログラムも汎用にさせたかったけどここは我慢するしかないのか…(´・ω・`) 397 名前：デフォルトの名無しさん mailto:sage [2009/12/20(日) 16:09:09 ] VCでプロジェクト作るのが面倒なんだけど、 なんかウィザード的なものはないのかな

---

---

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

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