- 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/
- 488 名前：487 [2010/01/13(水) 11:05:28 ]
- 追加報告です。
cudaMemcpyでcudaMemcpyDeviceToHostでコピーするところで異常に時間がかかっているようです。
この1行を消すだけで、100秒ほどかかっていたのが2秒にまでなりました。
逆方向のcudaMemcpyHostToDeviceは問題ありません。
- 489 名前：デフォルトの名無しさん mailto:sage [2010/01/13(水) 11:13:17 ]
- CUDAプログラムがGPUで計算等をしている間は画面表示の更新が遅くなります。
１回のループにどれくらいの時間がかかるのか、一定回数がどれくらいかによりますが、
経過の表示される間隔が短いとその影響を受ける可能性があります。
あるいは、マウスを動かすことでCPUに負荷をかけないとパワーセーブモードに入ってしまうとか・・・・。
- 490 名前：デフォルトの名無しさん mailto:sage [2010/01/13(水) 11:17:06 ]
- >>488
転送に１００秒かかるのは現実的にありえないので、恐らくcudaThreadSynchronize()等でカーネルの実行完了待ちを
していないのではないかと思います。カーネル呼び出しそのものは一瞬で終わりますが、
後続のcudaMemcpy(.....,...,...., cudaMemcpyDeviceToHost)時に暗黙的にそれまでに呼び出したカーネルの実行完了を待っています。
・・・・・ということかどうかは分かりませんが。
- 491 名前：デフォルトの名無しさん mailto:sage [2010/01/13(水) 11:29:27 ]
- >>489,490さん
早速の返信ありがとうございます。
cudaMemcpyでcudaMemcpyDeviceToHostを消したプログラムで、
プログラムの一番最後にだけ結果を取り出してみたところ、
正しく計算されていました。
次に、カーネルの実行部分のみをコメントアウトした場合、
やはり非常に時間がかかってしまいました。
やはり、GPUからCPUへのデータ転送に（というよりなにか同期の部分の様な気もしますが）
時間がかかっているようです。
計算用マシンなので省電力の設定はしていません。
- 492 名前：デフォルトの名無しさん mailto:sage [2010/01/13(水) 11:48:15 ]
- >>491
> cudaMemcpyでcudaMemcpyDeviceToHostを消したプログラムで、
> プログラムの一番最後にだけ結果を取り出してみたところ、
> 正しく計算されていました。
これを読むと、そもそも何度もDeviceToHostの転送をする必要がない処理という解釈で
よいのでしょうか？
DeviceToHostが遅くなる理由はハードウェア的なものから色々あります。
マザーボードを交換したという人もいました。
SDKに含まれているbandwidthTestの結果はどうなっていますか？
"--memory=pinned" を付けた場合と付けなかった場合をそれぞれ調べてみてください。
- 493 名前：デフォルトの名無しさん [2010/01/13(水) 11:59:12 ]
- Device 0: GeForce GTX 285
Quick Mode
Host to Device Bandwidth, 1 Device(s), Paged memory
Transfer Size (Bytes) Bandwidth(MB/s)
33554432 4340.2
Device to Host Bandwidth, 1 Device(s), Paged memory
Transfer Size (Bytes) Bandwidth(MB/s)
33554432 4193.7
Device to Device Bandwidth, 1 Device(s)
Transfer Size (Bytes) Bandwidth(MB/s)
33554432 119962.3
--memory=pinned
Device 0: GeForce GTX 285
Quick Mode
Host to Device Bandwidth, 1 Device(s), Pinned memory, Write-Combined Memory Enabled
Transfer Size (Bytes) Bandwidth(MB/s)
33554432 5758.1
Device to Host Bandwidth, 1 Device(s), Pinned memory, Write-Combined Memory Enabled
Transfer Size (Bytes) Bandwidth(MB/s)
33554432 5471.9
Device to Device Bandwidth, 1 Device(s)
Transfer Size (Bytes) Bandwidth(MB/s)
33554432 119852.3
- 494 名前：デフォルトの名無しさん mailto:sage [2010/01/13(水) 12:00:32 ]
- 結果はこんな感じです。
特に片方だけ遅いというわけでもないようです。
Linuxで起動した場合、遅くなるという問題は起きませんでした。
DeviceToHostが必要なのは結果を取り出すときだけで、
計算自体には必要ありません。
- 495 名前：デフォルトの名無しさん [2010/01/13(水) 12:03:47 ]
- 似たような話をnvidiaのforumで見つけました
forums.nvidia.com/lofiversion/index.php?t68705.html
forums.nvidia.com/lofiversion/index.php?t73047.html
返信は結構あるのですが、解決策はよくわかりません。
- 496 名前：デフォルトの名無しさん mailto:sage [2010/01/13(水) 14:46:26 ]
- ・CUDA内での計算処理、ループ数とループ内のステップ数、スレッド数ブロック数、だいたいどんなもんすか
・HostToDevice、DeviceToHostで転送するメモリのサイズはどのくらいですか
- 497 名前：デフォルトの名無しさん mailto:sage [2010/01/13(水) 15:02:38 ]
- とりあえず
タイマー1,2作成
cudaMemCopy(HostToDevice)
タイマー1,2起動
<<<>>>カーネル呼ぶ
cudaThreadSyncronize()
タイマー1停止
cudaMemCopy(DeviceToHost)
タイマー2停止
してほんとにMemCopyなのかカーネルなのか確かめてみる必要は。
- 498 名前：デフォルトの名無しさん mailto:sage [2010/01/13(水) 19:23:55 ]
- コピーが細切れでforループでやたら回数呼んで転送してたりしない？
１発で全部転送するのと1/1000ずつ1000回コピーするのとでは、
かかる時間に雲泥の差があるけど。
- 499 名前：487 mailto:sage [2010/01/14(木) 04:10:12 ]
- PCを再起動したら、上記の問題は出なくなりました。
お騒がせしました。
- 500 名前：デフォルトの名無しさん mailto:sage [2010/01/14(木) 09:16:24 ]
- おいｗｗ 終わりかよｗｗ
- 501 名前：デフォルトの名無しさん [2010/01/14(木) 11:44:30 ]
- cutil64.dllってどこに置けばいいの？
- 502 名前：デフォルトの名無しさん [2010/01/15(金) 05:05:43 ]
- GPU側で配列の合計を求める方法を教えてください
- 503 名前：デフォルトの名無しさん mailto:sage [2010/01/15(金) 07:59:04 ]
- サンプルのreduction見れ
- 504 名前：デフォルトの名無しさん [2010/01/15(金) 09:52:56 ]
- syncthreadsって異なるブロック間のスレッドも同期されるのでありましょうか？
- 505 名前：デフォルトの名無しさん mailto:sage [2010/01/15(金) 20:54:27 ]
- >>504
programing guide よめ
- 506 名前：デフォルトの名無しさん [2010/01/15(金) 23:10:51 ]
- Best Practices Guide
もいいこと盛りだくさん書いてあるよ
- 507 名前：デフォルトの名無しさん mailto:sage [2010/01/16(土) 19:28:55 ]
- cudaでは多次元配列の使用は推奨されてないの？
- 508 名前：デフォルトの名無しさん mailto:sage [2010/01/17(日) 12:17:08 ]
- pthreadとの組み合わせでうまく動作せずに困っています。
引数として渡す構造体に、float * device_Xを用意して
for(int k=0;k<NUM_THREADS;k++){
cudaMalloc((void**)&targ[k].device_X,
sizeof(float)*(hostX.nRow)*(hostX.nCol));
cudaMemcpy((void*)targ[k].device_X,
host_X,
sizeof(float)*(hostX.nRow)*(hostX.nCol),
cudaMemcpyHostToDevice);
pthread_create(&handle[k], NULL,
compute_thread_func,
(void *)&targ[k]);
}
と渡すと、pthread内で
cudaMemcpy(temp,argument->device_X, //argumentは渡されたvoid* argをキャストしたもの
sizeof(float)*(nRow)*(nCol),
cudaMemcpyDeviceToHost);
としてもうまく取り出せません。（中身がぐちゃぐちゃになる）
同様の処理をpthreadで呼び出す関数内で行うとうまく動作します。
スレッドごとにメモリを取るのだから、中でやってもいいのですが気になります。
何か原因になっていそうなことは無いでしょうか。
- 509 名前：デフォルトの名無しさん mailto:sage [2010/01/17(日) 12:59:47 ]
- >>508
＃1個のGPUに複数のスレッドからアクセスしたいアプリケーションなのかよく分かりませんが・・・・。
CUDAはスレッド毎にGPUのリソース（コンテキスト）を管理しています。
従ってこの例では
子スレッドのcompute_thread_func内の処理は
親スレッドがcudaMallocした領域にアクセスできません。
- 510 名前：デフォルトの名無しさん mailto:sage [2010/01/17(日) 13:32:14 ]
- >>509
ありがとうございます。
複数のスレッドから一台のGPUへのアクセスです。
（後々、スレッドごとに一個のGPUにしたいのですが）
スレッドごとに管理しているのははじめて知りました。
解決しました。
- 511 名前：デフォルトの名無しさん mailto:sage [2010/01/18(月) 20:51:33 ]
- 「Picture Motion Browser VAIO Edition」を追加。
超解像HDアップコンバートやノイズ除去、手ぶれ補正などの編集や動画作成を、NVIDIA CUDAによるGPU演算を使って行なえる。
- 512 名前：デフォルトの名無しさん mailto:sage [2010/01/18(月) 22:00:36 ]
- GT200 Fermi
トランジスタ数 14億 30億
倍精度浮動小数点 30FMA Oｐｓ/clock 256FMA Oｐｓ/clock
単精度浮動小数点 240 512
シェアドメモリ(SM) 16KB 64KB
L1メモリ 無し 48KB or 16KB
L2メモリ 無し 768KB
アドレス幅 32bit 64bit
複数のカーネル実行
無し 16個まで（GigaThredエンジン)
＊L1キャッシュ搭載
GT200では16KBのSMが32個あり、それぞれCUDAコアを8個割り当てられ、L1キャッシュが無し。
Fermiでは64KBのSMが16個。それぞれにCUDAコアが32個割り当てられ、SMから16KBか48KBのどちらかをL1キャッシュとして使用可能。
GT200に対して、3倍のSMと16KBのL1が使用可もしくは同じサイズのSMと48KBのL1が使用できるようになった。これにより、今までCUDA化できなかったプログラムの対応を増やし、さらに高速化もできる。
各CUDAコアに含まれるFPユニットは倍精度浮動少数演算を強化し、GT200に対し8倍の能力。
＊L2メモリの搭載 グローバルメモリに対する高速化。
＊C++をサポート
＊複数のカーネルの動作をサポート
SM内部のパイプラインを強化。SFUが複数に分けられたのでタスクとデータをより効率化。スレッドスケジューラを2個。
＊双方向PCIE通信
GT200ではPCIEバスの送受信をどちらか片方しか一度に実行できず、理論値8GB/s・実測4〜5GB/s程度だが
Fermiでは双方向通信が可能になり12GB/sを実現
＊新しいメモリコントローラ FermiよりGDDR5まで対応し、ECCにも対応する。
＊コア内部を各部でモジュール化
設定された価格帯や用途ごとにコアを設計しなおさず、機能をカットオフしたり追加したりできる。
SM単位でCUDAのコアを減らしたり、D3DやOpenGLなどの固定ハードウェアが不要なTeslaなどでオフになったりする可能性もある。
- 513 名前：,,・´∀｀・,,）っ-○○○ mailto:sage [2010/01/19(火) 02:32:35 ]
- GT200 理想のFermi 現実のFermi
トランジスタ数 14億 30億 30億
倍精度浮動小数点 30FMA Oｐｓ/clock 256FMA Oｐｓ/clock 224FMA Oｐｓ/clock
単精度浮動小数点 240 512 448FMA Oｐｓ/clock
シェアドメモリ(SM) 16KB 64KB 〃
L1メモリ 無し 48KB or 16KB 〃
L2メモリ 無し 768KB 〃
アドレス幅 32bit 64bit 〃
複数のカーネル実行 無し 16個まで（GigaThredエンジン) 14個まで（GigaThredエンジン)
- 514 名前：デフォルトの名無しさん mailto:sage [2010/01/19(火) 04:04:04 ]
- OpenCL待ちしてたけど、GPUの対応はCUDAだけっていう開発ばっかな
Fermiで本気出すから
- 515 名前：デフォルトの名無しさん mailto:sage [2010/01/20(水) 02:23:16 ]
- GT100ってなんかメモリアクセスのレイテンシがすごくでかくない？
それを隠蔽するために、L2キャッシュが128Kになっているけど、
コヒーレンシとかどうなのかな？
やっぱりGPUはデータ並列なアプリしか向いていないのかね。
- 516 名前：デフォルトの名無しさん mailto:sage [2010/01/20(水) 16:04:40 ]
- Fermiのグローバルメモリのレイテンシが遅いってのは何処から来た情報？GDDR5に対応するんだから、帯域は大きくなりそうだけど。
- 517 名前：デフォルトの名無しさん mailto:sage [2010/01/20(水) 17:08:55 ]
- >>512
SFUは今のGT200でも複数ある。
ただ、SFUが外部のパイプラインに分かれたからSFUを使っているときに違う種類の演算ができるようになったってことじゃね・。
>>515
最適化しなければいけない
- 518 名前：デフォルトの名無しさん mailto:sage [2010/01/21(木) 03:41:00 ]
- >>516
これをみて思った。
pc.watch.impress.co.jp/img/pcw/docs/343/352/html/kaigai-01.jpg.html
でもGT200と一緒なんだね。
pc.watch.impress.co.jp/docs/2008/0617/kaigai_10l.gif
いずれにしても、何でGPUがあんなにレイテンシが大きいのかがわかったよ。
だから、coalesced accessが必要なのね。
- 519 名前：デフォルトの名無しさん mailto:sage [2010/01/21(木) 04:03:50 ]
- >>518
>いずれにしても、何でGPUがあんなにレイテンシが大きいのかがわかったよ。
俺のために説明してくれ！
- 520 名前：デフォルトの名無しさん mailto:sage [2010/01/21(木) 22:43:04 ]
- GPUのメモリのレイテンシが大きいというよりも、CPUの場合と同じぐらいはかかるといった方がいい
- 521 名前：デフォルトの名無しさん mailto:sage [2010/01/22(金) 01:50:01 ]
- CPUだとメモリチャンネルが精々2〜3チャネル位しか無いが、GT100だと12チャネルもあるから、DRAMのチャネルのスイッチングに
それ以上かかると思っていい。あと、仮にGPC1?がDRAM0チャネル上にあるデータと、
DRAM11チャネル上にあるデータを同時にアクセスする場合、当然レイテンシがかかる。
もし、GPC1がDRAM0、DRAM1だけのデータアクセスですむなら、CPUと変わらない。
最適化が必要というのはそういったところだと思う。
- 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だったら行くんだけどな
|
 |
