- 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/
- 522 名前：デフォルトの名無しさん mailto:sage [2010/01/22(金) 05:43:22 ]
- MTGPがコンパイルできません。
>>463の言う、4つのファイル使って試したところ100個近くコンパイルエラーが出てしまいます。
あと「末尾を適当に書き換える。」の意味がよく分からんどす…
エラー内容は
1>\略)\mtgp32-fast.h(117) : error C2054: 'inline' の後に '(' が必要です。
1>\略)\mtgp32-fast.h(120) : error C2057: 定数式が必要です。
1>\略)\mtgp32-fast.h(120) : error C2466: サイズが 0 の配列を割り当てまたは宣言しようとしました。
1>\略)\mtgp32-fast.h(120) : error C2085: 'mtgp32_do_recursion' : 仮パラメータ リスト内にありません。
1>\略)\mtgp32-fast.h(121) : error C2061: 構文エラー : 識別子 'inline'
1>\略)\mtgp32-fast.h(122) : error C2054: 'inline' の後に '(' が必要です。
...
WinXP32bit、VC++2005、CUDA SDK2.3
該当箇所のコード見てもどこが悪いのか分からない…助けて…。
- 523 名前：デフォルトの名無しさん mailto:sage [2010/01/22(金) 07:00:08 ]
- 単にパイプラインステージが何百もあるからレイテンシがでかいのだと思ってました
- 524 名前：デフォルトの名無しさん mailto:sage [2010/01/22(金) 08:15:33 ]
- >>522
あ。おれか。
元のソースと自分のソースでdiffしてみた。
・inline を全部削除してみて。
・末尾を適当に書き換える。は、元だとプリントして捨ててしまっているので、
利用したいように書き換えてねと。
おれはmake_uint32_random(d_status, num_data, hostbuff); として
hostbuffを自分で使えるようにしました。
- 525 名前：デフォルトの名無しさん mailto:sage [2010/01/22(金) 20:11:37 ]
- >>523
正解
メモリチャンネル云々は殆ど関係ない
CPUはメモリとのクロック差のため
GPUは長大なパイプラインのためレイテンシがデカイ
- 526 名前：デフォルトの名無しさん mailto:sage [2010/01/22(金) 20:30:44 ]
- 別にパイプラインが深いわけではなくてバッファが大きいだけなんだけど。
- 527 名前：デフォルトの名無しさん mailto:sage [2010/01/22(金) 20:51:46 ]
- 200-400段は十分に・・
- 528 名前：デフォルトの名無しさん mailto:sage [2010/01/23(土) 00:27:17 ]
- >>524
レスサンクス。
inlineを削除したら見事にエラーが亡くなりました(`・ω・´)
これから自分のソースに組み込もうか…と状況に入れます。
重ねて質問申し訳無いですが、hostbuffの名前からしてMTGPの乱数は一度バッファにぶち込んでから使うという事になるのですか？
>>461ようにカーネルの中で特に意識せず使いたいのだけれども…
MTGPがglobalﾒﾓﾘやﾃｸｽﾁｬを使ってるのなら、Cのrand()のように連続で呼び出して使えたら良いなと思ってるんですが無理ですかね？
- 529 名前：デフォルトの名無しさん mailto:sage [2010/01/23(土) 02:00:42 ]
- >>528
おれはホスト側で使いたいので、make_uint32_randomの末尾でホスト側のバッファにコピーして使ってる。
デバイス側で使いたいなら、make_xxxxx_randomにd_buffをmain()側でとって渡し、
make_xxxx_random()内ではfreeせずにそのまま。すると次にmain()内でデバイス側関数を呼び出すときに
そのまま渡して使える。 ※スレッドが変わってはいかんので注意
- 530 名前：デフォルトの名無しさん [2010/01/23(土) 05:30:39 ]
- CUDAを使って大量のデータのごく簡単な統計値（最小、最大、平均、標準偏差）を計算したいんですが、何かいいサンプルとかありますか？
- 531 名前：デフォルトの名無しさん mailto:sage [2010/01/23(土) 08:51:31 ]
- CUBLASにあるんじゃ。
- 532 名前：デフォルトの名無しさん mailto:sage [2010/01/23(土) 11:54:45 ]
- >>530
そういう、「最後が一個」な演算は、reductionで探すと有る気がする
並列性を有用に使うのがなかなか面倒らしい
- 533 名前：デフォルトの名無しさん mailto:sage [2010/01/23(土) 12:09:54 ]
- OpenGL glutで時系列で変化する数値計算の結果を表示しようとしています。
GPU 2個で計算する場合、コールバック関数内でスレッドを生成することになると思うのですが、
この度（毎フレーム）にホストからグローバルメモリへデータ（時間による変化なし）を
コピーしなくてはいけません。
glutMainLoop(); 呼び出し前に両GPUにコピーした場合は、
コールバック関数内で生成したスレッドにグローバルメモリのポインタを渡しても参照できません。
データのコピーを一度だけですます方法はないでしょうか？ご教示ください。
スレッド生成は ttp://tech.ckme.co.jp/cuda_multi.shtml を参考にしました。
- 534 名前：デフォルトの名無しさん mailto:sage [2010/01/23(土) 21:06:46 ]
- >530
それぐらいだったらたぶんCPUで実行した方が速いよ
- 535 名前：デフォルトの名無しさん mailto:sage [2010/01/24(日) 00:53:03 ]
- >>533
GPU0用、GPU1用のスレッドをglutMainLoop()呼び出し前に生成するべきかと。
- 536 名前：デフォルトの名無しさん mailto:sage [2010/01/25(月) 22:22:13 ]
- thread数 = N
Mat[N];
id = threadIdx.x;
if(id%2==0){
演算Ａ
Mat[id]にコピー
}else if(id%2!=0){
演算Ｂ
Mat[id]にコピー
}
のようなプログラムを組んでいるのですが、結果をみると最後の2つの要素と同じ計算結果が全体に入ってしまいます。
N=16なら
14の結果が0,2,4…12に
15の結果が1,3,5…13に入ってしまいます。
どこに問題があるのでしょうか
- 537 名前：デフォルトの名無しさん mailto:sage [2010/01/25(月) 22:49:59 ]
- >>536
とりあえずC言語から始めようか
- 538 名前：デフォルトの名無しさん [2010/01/26(火) 00:02:22 ]
- >>530
俺もこれ知りたい。
N社のSDKを見ても、画像処理のサンプルとかたくさんあっても、単純な総和計算とかないもんだね。やはり向いてないからか・。。。
ご丁寧にCのほうが早いとか教えてくださる人もいるがｗ
マルチパスの画像フィルターとか、
デバイスの中にある中間結果を作って、
その統計値を、次のパスの計算で使ったりするのが常套手段だから
ここでいったんホストのほうにコピーするとボトルネックになってしまう。
デバイスの中のデータの統計値を出すライブラリとか作ってくれると本にありがたいんだが
- 539 名前：デフォルトの名無しさん mailto:sage [2010/01/26(火) 00:30:52 ]
- >>530, 538
SDKのreductionサンプルが参考になると思う。 確かpdfのわかりやすいスライドが一緒に入っているはず。
CUDAのアーキテクチャって、総和とかのreduction系演算は
不向きとまでは言わないまでもちょっと頭をひねらなきゃいけない分野だよね。
- 540 名前：536 [2010/01/26(火) 01:19:31 ]
- 演算途中で計算領域が被ってただけでした。失礼しました。
- 541 名前：デフォルトの名無しさん mailto:sage [2010/01/26(火) 01:27:25 ]
- >>539
そうだね。
沢山のデータがあるなら、最終はCPUにやらせても良さそうだけど、
転送がボトルネックになるなあ。
それだったらGPU内でやらした方が速いか。
あとI/O系があるとだめだよね。
ダブルバッファやトリプルバッファが使えるようなアプリならいいんだけど。
そうなると画像、映像系ばかりになってしまうなあ。
- 542 名前：デフォルトの名無しさん [2010/01/26(火) 09:14:50 ]
- 処理に割り当てるmultiprocessorの数を指定とか出来ませんか？
出来ればgridサイズ変えずに
そもそもOSもGPU使う場合割り当てがどうなってるのか分からないんですが
- 543 名前：デフォルトの名無しさん mailto:sage [2010/01/26(火) 10:43:54 ]
- >>542
CUDAプログラムが実行中はOSはGPUを使えないので全てのSMを使用しても問題ありません。
- 544 名前：デフォルトの名無しさん mailto:sage [2010/01/26(火) 10:47:22 ]
- 5秒以上使ったらエラーが出たよ。
- 545 名前：デフォルトの名無しさん mailto:sage [2010/01/26(火) 10:50:22 ]
- SMを1個しか使用していなくても5秒以上１つのCUDAカーネルが実行されるとタイムアウト。
- 546 名前：542 [2010/01/26(火) 11:15:31 ]
- へー、なるほど。
あとgridサイズ小さくする以外に使うSM数を制限できますか？
- 547 名前：デフォルトの名無しさん mailto:sage [2010/01/26(火) 11:41:29 ]
- できませんし、その必要もないでしょう。
- 548 名前：542 mailto:sage [2010/01/26(火) 11:48:03 ]
- >>547
そうですか、ありがとうございます
並列計算の研究の一環でCUDA使ってるんで
並列数の変化でデータが取れたら嬉しいな、という理由です
- 549 名前：デフォルトの名無しさん mailto:sage [2010/01/26(火) 14:26:47 ]
- >>548
素直に描画用とは別にGPUを用意しましょう。
総和を取る処理は私も書いているけど、仮想的に二次元にデータを配置して、縦に集計するところまでをGPUでやっている。
残りはCPU側に吐き出してからCPUではSSEを使って集計。
- 550 名前：デフォルトの名無しさん mailto:sage [2010/01/26(火) 14:53:47 ]
- >>548
1番最初に実行が始まるCTAの配置は予測可能なので、
その中で使わないSMに割り当てられたCTAはコア内でダミーの演算を長時間繰り返すことでそのSMを占有し続ける。
こうすれば本来の計算は残りのSMでのみ行われるようになる。
通常時間計測できるのは全CTAが完了するまでの時間なので以下のどちらかを。
１）最後のCTAがdevice memoryにマップされたhost pinned memoryに何か書き込む。
２）ダミー演算の繰り返し回数を段々少なくしていき、計測される時間の減少率が変わるところを探す。
なお、全SMを使わないとメモリアクセス効率も落ちるのであまり面白くない。
- 551 名前：デフォルトの名無しさん mailto:sage [2010/01/26(火) 15:20:53 ]
- > 仮想的に二次元にデータを配置して、縦に集計するところまで
なるほど。ふつくしい。ｎ次元でCPUでは最後の一次元だけやらせれば、対称になるな
- 552 名前：542 [2010/01/26(火) 21:31:38 ]
- >>549
別に用意してこの場合メリットってありますか？
あと総和ならReductionで組んだけど今回はGPU単体の話だったんで最後までGPUでやりました
デバイスホスト間の転送時間って馬鹿にならないイメージあるんですが、CPUにやらせたほうが速いんですかね？
まあ最後のちょっとだから誤差の範囲な気がしないでもないｗ
>>550
結構シビアですね、
直接的な方法が無ければgridサイズ縮めてIdx周り弄ろうと思います
↑の方法で弊害無いですよね？；
- 553 名前：528 mailto:sage [2010/01/27(水) 17:26:29 ]
- 少し前にMTGPについて質問した者ですが、どうやら自分の要求とズレた感じでした。
thread 256、block 1でmake_uint32_random()するとバッファに768個のデータが生成されるが、でっていう…。
これはﾒﾙｾﾝﾇﾂｲｽﾀの高周期な乱数列の１部分って事で、本当にあの高周期な乱数を使いたいならその分のメモリが必要だということなのかな。
独自の乱数ジェネレータを作って、その768個の中から１つを取り出して…みたいな事をやり始めるとまた性質が変わってしまうし、本末転倒に。
結局、カーネルの中でシミュレーション目的の使用方法としては微妙だったので絶望。。。
スレッドID毎に使える線形合同法(遷移型)の乱数として使えるようになんとかできないものか…
- 554 名前：デフォルトの名無しさん mailto:sage [2010/01/28(木) 00:49:57 ]
- え、おれ100万個単位で作って使えてるよ。
int num_data, にちゃんとでかい数与えてるかな
- 555 名前：デフォルトの名無しさん mailto:sage [2010/01/28(木) 09:22:34 ]
- >>554
いや、num_dataの数を変えて生成される個数の事はあまり重要じゃないのよ…
問題はカーネルの中で使おうとした時、バッファに作成された乱数が並んでいる形態では微妙なのです。
例えば、100万個作ったとして256のスレ数で使うなら使用部分を分割する事になりますよね（thID==0は、バッファのindex0〜約4000、という感じ）
いや、各スレッドは100万個のうち初めのindexだけseedとして決めて、あとは順次indexを増やして使っていく感じでもいいけど、
両者とも乱数列の周期はバッファのサイズに依存してしまいます。
一方、よくある{x=x*1103515245+12345; return x&2147483647;}このような方法は、アルゴリズムが優秀だとxが4byteなら最大で2^32の周期を持ちますよね。
今の状態のﾒﾙｾﾝﾇﾂｲｽﾀで2^32の周期を得ようとしたら、どんだけ大きいバッファが必要かっていう…
精度の良い乱数という点では利点ありますが、globalﾒﾓﾘを物凄く使う割にはなんだかなぁ…という複雑な気持ち。
MTGPの形態を知らなかった自分が悪いんだけど、要はこれ乱数生成を並列化して高速にしたもので、
実際にパラレルな計算で使う用にはなりえない事が分かりました。
自分の要求としては、GPU Gems3の「CUDAによる効率的な乱数の生成と応用」がﾁﾗｯと見た感じでは合致してたので、今からこっちを参考にしてみます。
長文ｽﾏｿ。>>554にはとても感謝している。こんな結末で申し訳ない。
- 556 名前：デフォルトの名無しさん mailto:sage [2010/01/28(木) 09:38:33 ]
- ＞乱数列の周期はバッファのサイズに依存してしまいます
横から失礼しますが、
for( ; ; )
{
make_uint32_random(, , d_buff); //デバイスメモリに作らせて残す
my_kernel_func<<<>>>(d_buff, , , , ); //MTGPで作った乱数を消費
}
こんな感じとして、256スレッドが一回に8192個の乱数を使う
→make_uint32_randomは2097152個の乱数を生成する
で良いのでは? make_uint32_random() は複数回呼び出すごとに
前回のmtgp32_params_fast_tの続きから処理するわけで、周期は
2^11213-1 でしょう。÷256しても2^11205とかだ
- 557 名前：デフォルトの名無しさん mailto:sage [2010/01/29(金) 08:52:57 ]
- stupid questionですみませんが、VC++ 9.0で
1. .cppと.hのように、.cuファイル内でインクルードしているファイルが更新されたら.cuを再コンパイル対象にしたい。
2. ptxを出力したい。nvcc -ptxでは無理でした。
以上について教えてください。お願いします。
- 558 名前：デフォルトの名無しさん mailto:sage [2010/01/29(金) 19:27:06 ]
- >>557
コマンドラインから nvcc -ptx *.cu とやっても駄目？
- 559 名前：デフォルトの名無しさん mailto:sage [2010/01/29(金) 23:45:27 ]
- -keepでいいんじゃね
1.についてはSource Dependenciesを個別に設定すれば一応できる
- 560 名前：デフォルトの名無しさん mailto:sage [2010/01/30(土) 15:52:51 ]
- >>558
XP Pro 32bitとVista Ultimate 64bitの両環境で、
コマンドラインからnvcc と打つとcl.exeがないと怒られます。
プロジェクトのプロパティを参考にパスとincludeを指定してやっても
エラーは控えていませんがコンパイルできません。
VC使ってる人はどうしてるんでしょう？
- 561 名前：デフォルトの名無しさん mailto:sage [2010/01/30(土) 15:56:22 ]
- -ccbin で指定しても駄目ですか？
- 562 名前：デフォルトの名無しさん mailto:sage [2010/01/31(日) 02:17:49 ]
- もちろんVSコマンドプロンプトから打ってるよな
- 563 名前：デフォルトの名無しさん mailto:sage [2010/01/31(日) 16:27:27 ]
- >>562
ふつうのコマンドプロンプトを使っていました。
どうもお騒がせしました。
- 564 名前：デフォルトの名無しさん mailto:sage [2010/01/31(日) 16:30:07 ]
- 普通にC++のコードを書けるようになってからじゃないと、学習効率が悪くてどうしようもないぞ。
- 565 名前：デフォルトの名無しさん mailto:sage [2010/01/31(日) 18:38:21 ]
- >>559
俺は556じゃないが-keep知らなかった。ありがとう。
- 566 名前：デフォルトの名無しさん mailto:sage [2010/02/01(月) 21:17:18 ]
- GPU Challenge 2010
ttp://www.hpcc.jp/sacsis/2010/gpu/
自由課題もあるそうな
- 567 名前：デフォルトの名無しさん mailto:sage [2010/02/02(火) 11:27:25 ]
- まあ俺は学生だから規定課題でるけどな
しかしCellとかに比べてあんまり最適化するとこないな
- 568 名前：デフォルトの名無しさん mailto:sage [2010/02/02(火) 18:26:32 ]
- どのくらい参加するんだろう？
俺もとりあえずエントリーしようかな。
- 569 名前：デフォルトの名無しさん mailto:sage [2010/02/02(火) 22:14:16 ]
- 自由課題の方で、パターン形成する発展方程式とかの数値計算すると、絵的にも面白そうなの出来そうじゃない？
- 570 名前：デフォルトの名無しさん [2010/02/10(水) 10:54:42 ]
- jpegをデコードするライブラリもしくはCUDAのコードはどこかにありませんか？
- 571 名前：デフォルトの名無しさん mailto:sage [2010/02/10(水) 11:54:41 ]
- GPU Challengeの課題が増えた
メルセンヌツイスタと言われるとHack the Cellを思い出すな
- 572 名前：デフォルトの名無しさん [2010/02/10(水) 12:30:21 ]
- SLI環境で、プログラムを2つ実行した場合、それぞれ別のGPUを
利用する方法を教えてくれ
- 573 名前：デフォルトの名無しさん mailto:sage [2010/02/10(水) 14:52:21 ]
- > SLI環境で
えっ、ひとつにしか見えないんじゃないの?? 出来るの？
- 574 名前：デフォルトの名無しさん mailto:sage [2010/02/11(木) 00:45:35 ]
- >>572
cudaSetDevice()でそれぞれ0と1を指定する。
- 575 名前：デフォルトの名無しさん mailto:sage [2010/02/11(木) 13:21:15 ]
- >>573
東工大青木先生はGeForce4つ並べてた
OpenMPで並列化していたと思う
- 576 名前：デフォルトの名無しさん mailto:sage [2010/02/11(木) 13:39:01 ]
- >>575
CUDA版OpenMPてあるの？
- 577 名前：デフォルトの名無しさん mailto:sage [2010/02/11(木) 14:25:56 ]
- CUDA版と言えるOpenMPはない
OpenMPのスレッド指定とCUDAのdevice指定を組み合わせただけのこと
- 578 名前：デフォルトの名無しさん mailto:sage [2010/02/11(木) 14:53:48 ]
- そういうことか、今度挑戦してみようかな
- 579 名前：デフォルトの名無しさん mailto:sage [2010/02/12(金) 00:56:59 ]
- > GeForce4つ並べ
と
> SLI
は違うんじゃね？ ケーブルで繋ぐのがSLI･･･かな
- 580 名前：デフォルトの名無しさん [2010/02/12(金) 10:07:29 ]
- OpenMPを使えば複数のGPUを簡単に使えるのですか？
やりかたをおしえてくれろ
- 581 名前：デフォルトの名無しさん [2010/02/12(金) 11:11:35 ]
- CUDAで顧客向けプログラムを作成しています。
CUDAプログラムの配布先には本体以外に何を配布すればよろしいのでしょうか？
- 582 名前：デフォルトの名無しさん mailto:sage [2010/02/12(金) 11:46:23 ]
- >>581
GeForce人数分
- 583 名前：デフォルトの名無しさん mailto:sage [2010/02/12(金) 11:49:25 ]
- TeslaかQuadroを配っておけば良いと思うよ
- 584 名前：デフォルトの名無しさん mailto:sage [2010/02/12(金) 11:56:09 ]
- cutil使わなければcudart.dllだけでよろしよ
- 585 名前：デフォルトの名無しさん mailto:sage [2010/02/12(金) 17:57:29 ]
- >>579
内部でケーブルでつながってても、
デバイスメモリが共有されるわけじゃないから
CUDA的には関係ない。
- 586 名前：デフォルトの名無しさん [2010/02/13(土) 04:48:44 ]
- >>580
SDKにサンプルがある
- 587 名前：デフォルトの名無しさん mailto:sage [2010/02/14(日) 12:30:31 ]
- CUDA FORTRANのセミナーが青木先生のとこで開催されるらしいが、おまいら行く？
- 588 名前：デフォルトの名無しさん mailto:sage [2010/02/14(日) 12:42:11 ]
- >>587
青木先生か、Cだったら行くんだけどな
- 589 名前：デフォルトの名無しさん mailto:sage [2010/02/15(月) 17:44:43 ]
- デバイスエミュレーション時の速度って、実際のCPUとの目安で考えたら
どのくらいスケールして考えればいいですか？
初めてエミュレーションモードを使ってみたんですが、3000倍以上の差が付いて明らかにおかしいと思うんです…
CPU: Core i7 Q720@1.6GHz、 GPU:Tesla C1060
- 590 名前：デフォルトの名無しさん mailto:sage [2010/02/15(月) 17:53:40 ]
- Q720って720QMのこと？
ノートPCにTeslaが搭載されているとか、聞いたことがないんだけど
- 591 名前：589 mailto:sage [2010/02/15(月) 17:56:06 ]
- 追記。
grid(2,1,1)、block(256,1,1)でのカーネルで、3000倍になります。
これからgridを増やすと、さらに差が広がっていきます。
カーネルで実装した内容を、CPU版で実装したくないけど速度比較はしたい。
･･･無理でつか？
- 592 名前：589 mailto:sage [2010/02/15(月) 18:06:39 ]
- >>590
Teslaは別のデスクトップPCので、エミュを動かしたのはノートPCでの方です。
紛らわしくて申し訳ない。
どちらもPCもCUDA使えるんですが、CPU自体はノートの方が性能良かったのでこちらを使いました。
ノートPCのGPU： GeForce GTX 260M
- 593 名前：デフォルトの名無しさん mailto:sage [2010/02/15(月) 18:14:45 ]
- そもそも、エミューレーションモードって非CUDA環境でも
CUDAカーネルのデバッグが出来るようにしたものでしょ
あくまでテスト用のもの
- 594 名前：デフォルトの名無しさん mailto:sage [2010/02/15(月) 18:15:25 ]
- 誤)エミューレーションモード
訂)エミュレーションモード
- 595 名前：デフォルトの名無しさん [2010/02/15(月) 19:10:55 ]
- 3000倍？そんなもんでしょ
- 596 名前：デフォルトの名無しさん mailto:sage [2010/02/15(月) 19:28:11 ]
- 効率の良いブロック分けの仕方？がわからず困っています．
実行時にＮ個のデータ系列が与えられて，
それぞれの系列へ施す処理内容自体は同一なのですが，
その処理に用いるパラメタ値が異なります．
例えばＮ＝３の場合，パラメタもp[0]〜p[2]の３個があって，
データ系列0の全データ { D[0,0], D[0,1], D[0,2], ..., D[0,m0] } にはp[0]を加算，
データ系列1の全データ { D[1,0], .... , D[1,m1] } にはp[1]を加算，
データ系列2の全データ { D[2,0], .... , D[2,m2] } にはp[2]を加算
という具合です．
全系列のデータ数が同じ(m0=m1=m2)ならば
グリッドの次元の１方向を系列(0〜N-1)に割り当てれば良いかと思うのですが，
系列毎にデータ個数がかなり異なる場合はどうすればいいのでしょうか？
データ個数は系列ごとに少なくとも数千個くらいになります．
同じような割り振り方だと何もしないブロックが大量にできてしまいそうです．
- 597 名前：デフォルトの名無しさん [2010/02/15(月) 20:36:27 ]
- CにD[0,0]というものはないのでよくわからないけど、
いったん長い配列にまとめて処理して、あとでCPUでばらせばいいのでは。
D[i,j]のjについて、自分はどのpに属するのか覚えさせて。
- 598 名前：デフォルトの名無しさん mailto:sage [2010/02/15(月) 23:36:05 ]
- >>592
Nvidia Nexus使えば？
- 599 名前：デフォルトの名無しさん mailto:sage [2010/02/15(月) 23:37:13 ]
- 追記
Nexus使うとネットワーク経由で、
コード書く用のPCとデバッグするPCを分けられるよ、ってことね。
- 600 名前：デフォルトの名無しさん mailto:sage [2010/02/16(火) 14:48:16 ]
- >>598
うーん、デバッグというよりは単にCPUとGPUで速度比較をしたいだけなんです。
うまく並列化して普通は、1〜50倍くらいの成果になると思うんですが･･･
目安でいいからエミュレーションモードから大体の速度が分からないのかなと。
- 601 名前：デフォルトの名無しさん [2010/02/16(火) 23:24:57 ]
- いまいち意味がわかんないけどCPUコードとGPUコードをデバイスエミュで実行したらCPUコードのが3000倍早いって事？
それだったらそんなもんかと。デバイスエミュは重いし。
違うんだったらごめん。
CPUとGPUで速度比較したいなら普通にCPUとGPUそれぞれ向けのコード書いて実行したらいいんじゃない？？
- 602 名前：デフォルトの名無しさん mailto:sage [2010/02/17(水) 00:50:55 ]
- >>601
あ、あれ？ 自分のデバイスエミュの認識自体が間違ってたかな…？
言いたかったのは、実行するハードの方での両者の比較です。
CPUコードと言うのはありません。
カーネルや、その内部で呼ぶ__device__の関数らがGPUコードだよね？それを普通に「GPU」が実行した時の速度と、
デバイスエミュを使ってCPUが実行した場合（内部では逐次計算？）の速度では、普通に「GPU」で実行した方が3000倍速いということです。
>>591の通り、GPUコードが多くなりすぎて、同じ事をするCPUコードを実装するのが面倒なのです。
デバイスエミュはCPUが実行するとの事で淡い期待を抱いてましたが無理そうな感じなんですな…
- 603 名前：デフォルトの名無しさん mailto:sage [2010/02/17(水) 01:32:32 ]
- >>602
比較する目的はなんでしょうか？
研究目的であれば面倒であろうがCPU用も実装しなければなりません。
そうでないなら比較なぞしなくてよいのではないかと。
ちなみにGPU：エミュが3000：１程度であればCPU用に普通に実装した方がGPUより速い可能性が十分にあります。
- 604 名前：デフォルトの名無しさん mailto:sage [2010/02/17(水) 15:10:51 ]
- 面倒でも計算結果の比較しろよｗ
nvidiaのサンプルコードでも結果の比較してるだろ
- 605 名前：デフォルトの名無しさん mailto:sage [2010/02/17(水) 15:41:47 ]
- >>602
emulation modeは，名前の通りGPUでの動作を模擬しているだけで，
その計算速度とGPUの計算速度を比較することに意味はない．
emulation modeがある理由は，カーネル内にprintfとかのコードを書いて
debugしたり出来るようにするため．
従って，CPUとGPUの計算速度の比較を行いたいなら，それぞれに最適化した
２つのコードを書いて比較する必要がある．
関係ないけど，
CPUとの比較しているときにCPUのコア一つとの比較が多い気がするけど，
それフェアな比較じゃないよね．せめてOpenMPくらい使おうよと．
まぁ使うとGPUの存在感が薄れるのだけれども…．
- 606 名前：デフォルトの名無しさん mailto:sage [2010/02/17(水) 17:32:51 ]
- >>603
研究なんて言えないようなものです。目的としては自己満足になりますね。
ただ、目安程度であれど比較できないとGPUとCUDAを使う意義に関わってきます。
早ければSUGEEEE!!ってなって満足し、遅ければそこできっぱり止めるという選択ができる。
そして3000：1ならまだCPUの方が早そうだというのは参考になりました。
>>605
おっしゃる通りですが、厳密に比較するまでは求めていないんです。
今自分がやってることは無意味なのか？（つまりCPUの方が普通に早い）が分かればいいんです。
grid(2,1,1)でフェアじゃないのは、ｺｰﾃﾞｨﾝｸﾞが糞なので2以上じゃないとGPUで動作しないんです('A`)･･･(ﾒﾓﾘ周りの関係で)
我侭な要求でスマン。
- 607 名前：デフォルトの名無しさん mailto:sage [2010/02/17(水) 18:07:50 ]
- 逆に考えるんだ
CPU側で動作をきっちり検証したプログラムを、
GPUに移植して、速度を比べる。
GPUに本当に適した問題なら、数十倍出る場合もあるし。
- 608 名前：デフォルトの名無しさん mailto:sage [2010/02/17(水) 18:47:11 ]
- >>605
GPUの優位性をうたうような文脈で
比較対象のCPUコードが1スレッドだったら
それを評価する人間は目が余程の節穴でない限り、
CPUのコア数倍した結果と比べるでしょ。
物凄く差がある場合はそれでも十分優位性をうたえるから。
- 609 名前：デフォルトの名無しさん [2010/02/17(水) 21:05:25 ]
- >>605
１コアの40倍とあれば4コアの10倍と読み替えればいい訳で。
- 610 名前：デフォルトの名無しさん mailto:sage [2010/02/17(水) 22:05:37 ]
- ４コアと比較したらどうなるかということではなく、
１コアと比較してる人がCPU版もまともにチューニングしてるとは思えない、ということかと。
まぁ、皆GPU版ばっかりチューニングしてますから。
- 611 名前：デフォルトの名無しさん mailto:sage [2010/02/18(木) 06:16:20 ]
- たぶんCPU版はSSEすら使っていないんだろうね。
メモリ帯域がものを言うコードでなくて、CPUがNehalemだったら、
安いGPUなんかじゃ全く優位ないからね。
- 612 名前：デフォルトの名無しさん mailto:sage [2010/02/19(金) 02:31:33 ]
- 俺はNvidiaちゃんを信じるよ
twitter.com/NVIDIAGeForce/status/9265680539
- 613 名前：デフォルトの名無しさん mailto:sage [2010/02/19(金) 02:32:33 ]
- 誤爆った／(^o^)＼
- 614 名前：デフォルトの名無しさん mailto:sage [2010/02/20(土) 04:09:26 ]
- うちは理論で「***手法より*%高速化して最速!」とかやってないってのもあるけど
GPUで組んだ手法と既存の手法を比べる場合、既存のほうはベーシックにしろと指導された。
複数CPUだとかSSEを使ってガチガチに最適化した手法と比べちゃうと基準が分からなくなるからだと。
他の高速化との差を知りたければその論文と比較しろということだと思う。
CPU最適化して無いなんて糞というのも分かるけど、こういうところもあるということで。
- 615 名前：デフォルトの名無しさん mailto:sage [2010/02/20(土) 08:30:23 ]
- コードの比較もいろいろだよな。
同じアルゴリズムを採用しても、CPUでも書き手によってGPUでも明らかに差が出てくる。
でもGPUを使う場合、多くの場合はCPUよりも速くなりました。というのが目的な訳で、
CPUの方が速いならあえてGPUを使う必要はないからね。
基準が曖昧になるのもわかるけど、そもそも基準が曖昧な気がするなあ。
場合によってはかなり恣意的になることもあるし・・・・。
- 616 名前：デフォルトの名無しさん mailto:sage [2010/02/20(土) 10:07:52 ]
- Femiやばいまた延期確定かも
- 617 名前：デフォルトの名無しさん mailto:sage [2010/02/20(土) 11:37:52 ]
- 一般人が入手できるのは１年後になる可能性もあるらしいね
- 618 名前：デフォルトの名無しさん mailto:sage [2010/02/20(土) 11:44:04 ]
- なんでそんな度々延期になるの
- 619 名前：デフォルトの名無しさん mailto:sage [2010/02/20(土) 12:14:45 ]
- >>618
ペーパーロンチで実際開発が
行われていないからだよ
- 620 名前：デフォルトの名無しさん mailto:sage [2010/02/20(土) 13:17:36 ]
- 今回のケースは大きな欠陥があることを知りながら、小手先の改良でなんとかしようとして
「完成品」を大量生産をして、まとにチップが取れなかったのが原因だろ
１％程度とされる歩留まり率で、１チップ当たり5000ドルの原価
これでは商売にならないね
- 621 名前：デフォルトの名無しさん mailto:sage [2010/02/20(土) 13:39:05 ]
- 3/19に東工大青木先生がCUDA Fortranのセミナやるんだって
- 622 名前：デフォルトの名無しさん mailto:sage [2010/02/20(土) 13:40:57 ]
- 関係者の宣伝おつ
|
 |
