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

[表示 : 全て最新50 1-99 101- 201- 301- 401- 501- 601- 701- 801- 901- 2chのread.cgiへ]
Update time : 02/21 05:22 / Filesize : 250 KB / Number-of Response : 931
[このスレッドの書き込みを削除する]
[＋板最近立ったスレ＆熱いスレ一覧 : ＋板最近立ったスレ／記者別一覧] [類似スレッド一覧]


↑キャッシュ検索、類似スレ動作を修正しました、ご迷惑をお掛けしました

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/
453 名前：デフォルトの名無しさん mailto:sage [2010/01/04(月) 16:05:01 ]: ⇒を何らかの演算として　0⇒256 256⇒512 512⇒0　といった感じにしたいのですが

a = (a+256)%513;

とする以外思いつきません。けど剰余は遅いので違う計算にしたいんです。
いっそのことテーブル作ったりif文にしたほうが早いのでしょうか？
馬鹿すぎてわかりません。誰か助けて
454 名前：デフォルトの名無しさん mailto:sage [2010/01/04(月) 16:10:45 ]: >>453
Cudaだと四則演算はほとんど処理時間には影響しないし、その式をつかうといいんじゃない？
っていおうとおもったけど、a=512のとき0にならなくね？
455 名前：デフォルトの名無しさん mailto:sage [2010/01/04(月) 16:13:31 ]: 257⇒0　258⇒1　…　でいいの？
でもそうすると512⇒0がヘンだけどな。

int val;
val = a + 256;
if (val > 512) val = val - 513;

でいいんじゃないの？
というかC言語の疑問をCUDAスレで聞くのはまちがっとる、ここは人口が少ないぞ
456 名前：デフォルトの名無しさん mailto:sage [2010/01/04(月) 16:26:48 ]: じゃあCUDAっぽく

const int TABLE[512] = { 256, ... , 512, ... , 0 }

__device__ int f(int a){
return TABLE[a];
}
457 名前：デフォルトの名無しさん mailto:sage [2010/01/04(月) 17:09:36 ]: レスありがとうございます
わけわかんないこと書いてすみません
テーブルとif文試して出直してきます
458 名前：デフォルトの名無しさん mailto:sage [2010/01/04(月) 17:37:31 ]: >>453
b⇒aとして、
a = (b+256)%768;
でいいのかな？
0⇒256
256⇒512
511⇒767
512⇒0
459 名前：デフォルトの名無しさん mailto:sage [2010/01/04(月) 21:45:29 ]: 釣りだよな？
460 名前：デフォルトの名無しさん mailto:sage [2010/01/05(火) 01:06:37 ]: 申し訳ないっす
461 名前：デフォルトの名無しさん mailto:sage [2010/01/07(木) 00:21:36 ]: ランダム関数の作成で悩んでます。
・CPUで作成した乱数リストをGPUに転送しない
・1スレッドの中で何度も呼ばれる
・グローバルメモリは使わない。スレッドIDやclock()をseedとして使う

この制約で良い乱数生成アルゴリズムを教えて下さい。
MTGPっていうのを参考にしようとしましたがムズ過ぎて挫折…まずコンパイルが通らず(´・ω・`)
462 名前：デフォルトの名無しさん mailto:sage [2010/01/07(木) 01:01:35 ]: おれMTGPコンパイルできたよ。GPU側で作成してそのままGPUで使えるし。
（シンボルコピー）
463 名前：デフォルトの名無しさん mailto:sage [2010/01/07(木) 08:15:53 ]: 使ったファイルはこれだけ
mtgp32-fast.h
mtgp32-cuda.cu
mtgp32-fast.c
mtgp32-param-fast.c
make_single_random()とかの末尾を適当に書き換える。
464 名前：デフォルトの名無しさん [2010/01/07(木) 12:09:10 ]: 誰かＯｐｔｉＸやってる人いませんかー？
465 名前：デフォルトの名無しさん mailto:sage [2010/01/08(金) 09:15:11 ]: CES 2010: NVIDIA Shows GeForce GF100 Fermi Video Card - Legit Reviews
www.legitreviews.com/news/7122/

3/2に本当に出るのかしらん
466 名前：デフォルトの名無しさん mailto:sage [2010/01/08(金) 09:22:19 ]: 気にするな、出たら買う、出るまではCUDAの修行をする。
単純でｲｲじゃないか。
467 名前：デフォルトの名無しさん mailto:sage [2010/01/08(金) 09:41:48 ]: >>461
Perlinノイズやってみてよ
ttp://mrl.nyu.edu/~perlin/doc/oscar.html
Cのソースが有る
468 名前：デフォルトの名無しさん mailto:sage [2010/01/08(金) 11:02:05 ]: Fermiもダメそうってどこかのブログで見た
NVIDIAまじやばい
469 名前：デフォルトの名無しさん mailto:sage [2010/01/08(金) 11:05:40 ]: ようやくCUDAに慣れてきたところなのに、StreamSDKとかOpenCLとか
DirectComputeとかまた勉強しないといけないの？　
辛いのうｗｗプログラマは辛いのうｗｗｗ　　　orz
470 名前：デフォルトの名無しさん mailto:sage [2010/01/08(金) 13:36:24 ]: windows7 64bit、Visual C++ 2008 Pro環境に、CUDA　toolkit、sdk 2.3をインストールして、
C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK\C\srcにある、
_vc90.slnファイルを開いて、問題なくビルドできます。
exeファイルの出力先は、デフォルトでC:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK\C\bin\win64\ですが、
これ以外の出力先に変更する方法を教えてください。
471 名前：デフォルトの名無しさん mailto:sage [2010/01/08(金) 13:42:23 ]: プロジェクトのプロパティで変更できるよ
$(OutDir)\$(ProjectName).exe
これを丸ごとフルパスで指定してしまいなさい。
472 名前：デフォルトの名無しさん mailto:sage [2010/01/08(金) 17:14:12 ]: XP,32bitでコンパイルしたプログラムを7,64bitで動かそうとしたら動かなかった･･･

こういうものなのか
473 名前：デフォルトの名無しさん [2010/01/08(金) 17:23:41 ]: 32/64ビットのcudart.dllが必要
474 名前：デフォルトの名無しさん mailto:sage [2010/01/08(金) 19:20:38 ]: >>471
thanks
475 名前：デフォルトの名無しさん mailto:sage [2010/01/10(日) 12:33:02 ]: 現在CUDA SDKを導入しVisual C++ 2005 Express Editionを使用しCPUでのエミュレーションを行っています。
いくつかのサンプルは動かせたのですがparticlesなど複数のサンプルでコンパイルはできているようですが実行すると以下のようなエラーが出ます

device emulation mode and device execution mode cannot be mixed.

エラーが出ている行は

cutilCheckMsg("cudaMalloc failed");
cutilSafeCall(cudaGLRegisterBufferObject(vbo));
cutilSafeCall( cudaMalloc3DArray(&noiseArray, &channelDesc, size) );

のように「cutil～」になってます。オプションから設定しているのでデバイスエミュレーション自体はうまく行ってるようなのですが
エミュレーションを行う際はソースを書き換える必要があるのでしょうか?
476 名前：デフォルトの名無しさん mailto:sage [2010/01/10(日) 12:55:52 ]: メッセージのとおりエミュONOFFまぜてる
またはリビルドしてないんじゃ？

プロジェクト設定じゃなくてビルド構成のほうでエミュレーションにかえられないっけ？
477 名前：デフォルトの名無しさん mailto:sage [2010/01/10(日) 13:05:00 ]: メッセージの内容は把握できています。
リビルドも多分できてると思います。
エミュレーションの設定は
プロジェクトのプロパティ→CUDA Build Rule→Emulation Mode はい
にしてます。ほかの方法はよくわからないのですがほかの方法があるのでしょうか?
478 名前：デフォルトの名無しさん mailto:sage [2010/01/10(日) 22:13:13 ]: ビルド構成のほうからエミュレーションの設定ができました。問題なくビルド&実行できました。ありがとうございます。

ただFPSが0.1なんでほぼ無意味でした。
479 名前：デフォルトの名無しさん [2010/01/10(日) 23:24:17 ]: エミュレーションなんてHello world以上のことをやるもんじゃないな
480 名前：デフォルトの名無しさん mailto:sage [2010/01/11(月) 00:20:59 ]: 処理する画面領域を１０００分の１ぐらいに限定して、バグを調査するくらいのことはできる。
481 名前：デフォルトの名無しさん mailto:sage [2010/01/11(月) 01:26:19 ]: なんだっけ、エミュじゃなくて実機で実デバッグできるやつ
早く来てほしいな
482 名前：デフォルトの名無しさん mailto:sage [2010/01/12(火) 14:25:34 ]: CPU版アプリの移植の際に、エミュはデータ検証に使える。
483 名前：デフォルトの名無しさん [2010/01/12(火) 20:21:38 ]: -G　オプションって使える？WinXPで
「'assembler' は、内部コマンドまたは外部コマンド、
操作可能なプログラムまたはバッチファイルとして認識されていません。」
と出てビルドできないけど。

--deviceemuもつけないとだめなの？
484 名前：デフォルトの名無しさん [2010/01/13(水) 09:38:31 ]: CUDAのバージョンを書き出しておきたいのですが、
CUDAのバージョンを調べる関数か何かありませんか？
485 名前：デフォルトの名無しさん mailto:sage [2010/01/13(水) 09:41:29 ]: >>484
>>436　でどう？
486 名前：デフォルトの名無しさん mailto:sage [2010/01/13(水) 09:50:27 ]: こっちを要求している気がする。
cudaRuntimeGetVersion()
cudaDriverGetVersion()
487 名前：デフォルトの名無しさん [2010/01/13(水) 09:56:08 ]: WindowsでCUDAを使っています。
ループが一定回数超える度に、coutで経過を表示しています。
Linuxに比べて遅かったので、いろいろ試してみたところ、
DOSウインドウの前で、マウスを動かすと、速くなります。
垂直同期か何かで引っかかっているのかと思い、設定で切ってみましたが
状況が変わりません。
どなたか似たような状況になって解決された方はいないでしょうか？
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だったら行くんだけどな
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 ]: 関係者の宣伝おつ
623 名前：デフォルトの名無しさん mailto:sage [2010/02/23(火) 08:21:59 ]: 青木先生に集客されたくねえなあ正直
624 名前：デフォルトの名無しさん mailto:sage [2010/02/23(火) 22:06:44 ]: nexusをリモートで動かそうとしたが、ブレークポイントでとまらねぇ・・・。
色々試したがどうにも解決しないので教えてください。

状況としては、nexusのユーザーズガイドに沿って設定。
ためしにnexusサンプル動かそうとしたら、
ランタイムAPIプロジェクトはGetDeviceで引数に０が。
で、次の行で落ちる。
DriveAPIは落ちない。んでホスト側の画面右下に青いポップアップ出て、
ターゲットマシンにコンソール画面出てるのでプログラムは正常に動いてるっぽい。
でもカーネル関数内にブレークポイント置いても止まらず。
自分でSDKサンプルのプロジェクトの設定変えて試しても同じ。

マシン環境はこんな感じ。
ホストマシン
Vista 64bit SP2
.Net3.5 SP1
Host nexus1.0(jan 64bit)
GPU FX570
VC++ 2008 SP1
DirectX10 August 09
CUDA SDK2.3 32bit
CUDA ToolKit2.3 32bit
625 名前：続き mailto:sage [2010/02/23(火) 22:08:57 ]: ターゲットマシン
Vista 64bit SP2
.Net3.5 SP1
Target nexus1.0(jan 64bit)
GTX285
VC++ 2008
DirectX 10 August 09
CUDA SDK2.3 32bit

他に設定としてはEnable Security Serverをfalse
nexus→option→Enable secure connectionをfalseにしてます。
使い方は、ホストマシンでVC起動→プロジェクト読み込み→nexusデバッグ
ターゲットマシンはデバッグモニタ起動のみ。
どこがおかしいのだろう？
626 名前：デフォルトの名無しさん mailto:sage [2010/02/24(水) 22:31:24 ]: Fermi終了したらこのスレも終了するんかなあ
627 名前：デフォルトの名無しさん mailto:sage [2010/02/24(水) 23:16:54 ]: the launch date for GeForce GTX 480 and GTX 470 is March 26
628 名前：デフォルトの名無しさん mailto:sage [2010/02/25(木) 12:42:26 ]: >>627
ペーパーリリースで全世界で１万枚以下の出荷といううわさだけどね
一般人が手に入れられるようになるのは、下手をすると来年
良くて年末という予想がある
629 名前：デフォルトの名無しさん mailto:sage [2010/02/25(木) 13:15:26 ]: やっぱ、シリコン丸ごと改良しないとだめなんか。
半分ダメで256コアでもいいんだけどｗｗ
630 名前：デフォルトの名無しさん mailto:sage [2010/02/25(木) 20:23:58 ]: 俺はNexusインストールすらできなかった
631 名前：デフォルトの名無しさん mailto:sage [2010/02/25(木) 22:02:09 ]: 消費電力280Wだっけ。。。GTX480
632 名前：デフォルトの名無しさん mailto:sage [2010/02/25(木) 23:43:58 ]: アム虫キモ
633 名前：625 mailto:sage [2010/02/26(金) 23:13:31 ]: もうよくわからんからnexusのエミュモードでやることにした。
カーネル関数内でブレークするし、値もちゃんと表示されてそう。

>>630
OS対応してないとか？Vistaか7しかできない。
あとはOSとnexusのbitが違うとか。
634 名前：デフォルトの名無しさん mailto:sage [2010/02/27(土) 19:37:59 ]: >>633
インストーラがVS2008 SP1入れてあるのに読み取ってくれなくて、
インストールができない状態だった。

Microsoftからダウンロードしたばっかりのイメージを使ってインストールしたから
当然SP1はあたっているものかと思っていたらあたっていなかった、っていう初歩的なミスだった。

ちょっくら遊んでくる
635 名前：デフォルトの名無しさん mailto:sage [2010/03/05(金) 11:16:51 ]: 　　　　　　　　 ,. -‐'''''""¨¨¨ヽ
　　　　　　 (.＿＿_,,,... -ｧァﾌ|　　　　　　　　　　あ…ありのまま今　起こった事を話すぜ！
　　　　　 |i i|　　 }!　}} /／|
　　　　　　 |l､{　　j}　/,,ｨ//｜　　　　　　　『おれはNvidiaにARE YOU READY?と言われて
　　　　　　　 i|:!ヾ､_ﾉ／ u {:}//ﾍ　　　　　　　　　　準備していたら準備しているのはNvidiaの方だった』
　　　　　　　 |ﾘ u' }　 ,ﾉ　_,!V,ﾊ |
　　　　／´fト､_{ﾙ{,ィ'ｅﾗ　, ﾀ人　　　　　　　　な…　何を言ってるのか　わからねーと思うが
　　　　 /' 　ヾ|宀| {´,)⌒`/ |<ヽﾄiゝ　　　　　　　　おれも何をされたのかわからなかった
　　　　,ﾞ　／ )ヽ iLﾚ　u' |　| ヾｌﾄﾊ〉
　　　 |／_／　ﾊ !ニ⊇　'／:} 　V:::::ヽ　　　　　　　　頭がどうにかなりそうだった…
　　　 /／二二二7'T'' ／u'　__ /:::::::/｀ヽ
　　　/'´r　-―一ｧ‐ﾞＴ´　'"´ ／::::／-‐ 　＼　　　　128bitメモリバスだとかリネームテクノロジーだとか
　　 / // 　广¨´ 　/'　　／:::::／´￣｀ヽ ⌒ヽ　　　　そんなチャチなもんじゃあ　断じてねえ
　　ﾉ ' /　ノ:::::`ー-､___／:::::／/ 　　　ヽ　　}
_／｀丶　/::::::::::::::::::::::::::￣`ー-{:::...　　　　　　ｲ　もっと恐ろしいものの片鱗を味わったぜ…
636 名前：デフォルトの名無しさん mailto:sage [2010/03/06(土) 00:09:36 ]: GTX480がそこそこ出回るらしい（core数は当初想定よりも少なくなる可能性があるが）
Fermiアーキテクチャを思っていたよりも早い時期に体験出来る可能性が出てきた
637 名前：デフォルトの名無しさん mailto:sage [2010/03/06(土) 20:59:14 ]: Fermi火事出すだろうな
電気食いすぎだ
638 名前：デフォルトの名無しさん mailto:sage [2010/03/07(日) 10:23:33 ]: GF100（GTX480, 470）は高電力でお値段も高めになるだろうから、
Fermiアーキテクチャをとにかく早く試したい人以外はその次のGF104がいいと思う
639 名前：デフォルトの名無しさん mailto:sage [2010/03/07(日) 11:05:46 ]: 半分だけの460とか出ないかな？
640 名前：デフォルトの名無しさん mailto:sage [2010/03/07(日) 11:26:00 ]: この辺りの情報がそれかも知れん

The AMD's Cafe:ローコストなFermiやQuadroについて - livedoor Blog（ブログ）
blog.livedoor.jp/amd646464/archives/51525107.html

まだまだ先のようだ
641 名前：デフォルトの名無しさん mailto:sage [2010/03/07(日) 15:31:52 ]: 売り物にならない奴はQuadro逝きか。
OpenGL市場ってほんと舐められてるな。
ゲイツに縛られない自由なAPIだったはずなのに。
642 名前：デフォルトの名無しさん mailto:sage [2010/03/09(火) 00:29:10 ]: www8.plala.or.jp/b4zabeat/
643 名前：デフォルトの名無しさん mailto:sage [2010/03/09(火) 23:19:41 ]: このレビューは何が言いたいのかよくわからんかった
pc.watch.impress.co.jp/docs/column/nishikawa/20090518_168541.html
644 名前：デフォルトの名無しさん mailto:sage [2010/03/16(火) 00:53:46 ]: CUDA上の命令がどれくらいのクロックで動くかまとめられていませんか？
整数の乗算やら三角関数はプログラミングガイドに載っていたのですが,ほかの命令も知りたいです
645 名前：デフォルトの名無しさん mailto:sage [2010/03/16(火) 08:46:48 ]: 「どのくらい」でいいなら実測すればいいかと
646 名前：デフォルトの名無しさん mailto:sage [2010/03/16(火) 16:42:16 ]: エミュレーションモードでやるとうごくのですが、GPUをつかうと動きません。

const int c = border + (blockDim.x * blockIdx.x + threadIdx.x) * step * 2;
const int r = border + (blockDim.y * blockIdx.y + threadIdx.y) * step * 2;
const int i = 1;
// atomicAdd(&count[0], 1);//ここでは動くのでatomicAddの問題ではない
if(c >= i_width - border || r >= i_height - border) return;
int i_max = -1, r_max = -1, c_max = -1;
float max_val = 0;
for (int ii = i; ii < min(i+2, intervals-1); ii += 1)
for (int rr = r; rr < min(r+2*step, i_height - border); rr += step)
for (int cc = c; cc < min(c+2*step, i_width - border); cc += step) {
float val = getVal(d_m_det, o, ii, cc, rr, i_width, i_height);
if (val > max_val) {
max_val = val, i_max = ii, r_max = rr, c_max = cc;
}
} // Check the block extremum is an extremum across boundaries.
/***********ここでd_iptsにiptを加えてもうごく*********/
// float4 ipt ;
// d_ipts[atomicAdd(&counter, 1)] = ipt;
if (max_val > 0.0004f && i_max != -1 && isExtremum(d_m_det,o, i_max, c_max, r_max, i_width, i_height, intervals)) {
float4 ipt = interpolateExtremum(d_m_det, o, i_max, r_max, c_max, i_width, i_height);
if(ipt.x >= 1)
{
d_ipts[atomicAdd(&counter, 1)] = ipt;//ここの行をコメントアウトすると動く
ipt.x += 1;
}}
最後のif文の中でd_iptsやcountにアクセスするのがだめっぽいのですが・・・
なにかif文を書いたときに同じような症状になった方や、これを見ただけでわかる方いらっしゃったら、教えてください。
よろしくおねがいします。
647 名前：デフォルトの名無しさん mailto:sage [2010/03/16(火) 23:39:20 ]: とりあえず問題になってるatomicAdd(&counter, 1)の戻り値調べようか
648 名前：デフォルトの名無しさん mailto:sage [2010/03/17(水) 00:13:02 ]: counterの宣言にちゃんと__device__は付いているのだろうか・・・
649 名前：646 mailto:sage [2010/03/17(水) 00:40:32 ]: >>647
エミュレーションで確認したところ戻り値はcounterと同じ値になっていました。
実際にGPUでうごかすと、
counterをデバイスからホストへ送るところでunspecific launchとなるか、
永久ループに入るか何かで画面が固まり、ブルースクリーンになって落ちます。

>>648
_device_をつけてグローバル変数(っていうのかわからないけど)として定義しています。
650 名前：デフォルトの名無しさん mailto:sage [2010/03/17(水) 09:55:20 ]: その反応だとアクセス違反のときが多いのだがそんなことないよな？
651 名前：デフォルトの名無しさん [2010/03/17(水) 13:42:58 ]: CUDAでつかうlong doubleってWindowsとLinuxでサイズは違うの？
652 名前：デフォルトの名無しさん [2010/03/18(木) 00:14:29 ]: >>650
アクセス違反ってcounterの位置がおかしいってことですか？
このプログラムをいれずに単にcounterをデバイスからホストに送ることは可能でした(初期値として０をおくっていたので０が帰ってきていました)。

つまり、このプログラムのようにatomicAddをif文とか分岐が多くなるような文章内でつかうと、
今回でいえばcounterのアドレスがかわるということですか？？
653 名前：デフォルトの名無しさん [2010/03/19(金) 07:46:47 ]: パスワードクラックでTeslaがHD5970に完敗

www.geeks3d.com/20100316/radeon-hd-5970-the-ultimate-password-cracking-hardware/
www.geeks3d.com/public/jegx/201003/elcomsoft-wpa-psk-password-test.jpg
654 名前：デフォルトの名無しさん mailto:sage [2010/03/19(金) 10:12:59 ]: Dual-GPUとSingle-GPUを比べてる時点でアウト
655 名前：デフォルトの名無しさん [2010/03/19(金) 10:34:24 ]: １スロットどうしの比較だから問題ない
656 名前：デフォルトの名無しさん mailto:sage [2010/03/19(金) 11:09:41 ]: >>655
それを言い出すと、TeslaDが出てくるぞ。ブリッジ自体は1スロットだからなw
657 名前：デフォルトの名無しさん [2010/03/19(金) 11:22:52 ]: はいはい、負け惜しみ
単純な計算性能では圧倒的にラデオンのほうが優れているんだよ
658 名前：デフォルトの名無しさん mailto:sage [2010/03/19(金) 11:32:44 ]: 比較対象にHD5870が入っていない時点で
なんか違和感があるんだが
659 名前：デフォルトの名無しさん mailto:sage [2010/03/19(金) 11:37:34 ]: Tesla C1070ってなに？Fermi？
660 名前：デフォルトの名無しさん mailto:sage [2010/03/19(金) 11:42:09 ]: www.elsa-jp.co.jp/products/hpc/tesla_s1070/index.html
これ
661 名前：デフォルトの名無しさん mailto:sage [2010/03/19(金) 11:42:32 ]: GT200じゃなかった?
662 名前：デフォルトの名無しさん mailto:sage [2010/03/19(金) 14:33:49 ]: ttp://developer.nvidia.com/object/gpucomputing.html

CUDA Programming Guideがリンク切れしているんですが
誰か正しいアドレスを知りませんか
663 名前：デフォルトの名無しさん mailto:sage [2010/03/19(金) 18:21:58 ]: 高い・遅い・熱い
664 名前：デフォルトの名無しさん mailto:sage [2010/03/19(金) 18:52:10 ]: S1070ってGT200×4の奴だろ。
665 名前：デフォルトの名無しさん [2010/03/20(土) 17:28:11 ]: CUDA 3.0 Downloads
developer.nvidia.com/object/cuda_3_0_downloads.html
666 名前：デフォルトの名無しさん mailto:sage [2010/03/20(土) 18:02:54 ]: CPUでの計算にインテルコンパイラを使いたいのですが、どのようにすればいいのでありまするか？
667 名前：662 mailto:sage [2010/03/20(土) 20:32:56 ]: >>665
ありがとうございます
668 名前：デフォルトの名無しさん mailto:sage [2010/03/20(土) 20:52:01 ]: >>666
nvcc -cで*.cuをコンパイルしたら、できた*.oをiccでリンクすればいい。
Windowsの場合はnvccもiccも別々にオブジェクトを作ることになるからそれをVCでリンク。
669 名前：デフォルトの名無しさん mailto:sage [2010/03/20(土) 21:52:07 ]: 利用者はCUDAなんて独占的技術を求めてませんよね
670 名前：,,・´∀｀・,,）っ-○○○ mailto:sage [2010/03/20(土) 22:52:58 ]: まして永久β版のSDKなんて論外
671 名前：デフォルトの名無しさん mailto:sage [2010/03/20(土) 23:30:36 ]: まあ利用者だけが決めるわけでもないのも悲しいけど現実なのよね
672 名前：デフォルトの名無しさん mailto:sage [2010/03/21(日) 04:48:10 ]: なんか面白いことに使えないかな
もったいない
673 名前：,,・´∀｀・,,）っ-○○○ mailto:sage [2010/03/21(日) 08:27:51 ]: 「何かには使える」って言ってるうちは何にも使えないまま終わるんですけどね。
ターゲットアプリケーションがあってはじめて、そのニーズに合わせてハードの機能・性能の拡充が行われうるわけで
今までであればゲームがそうだった。

ウン十並列のデータを同時処理するような用途のニーズが仮に高まってるとしても
それはCPUのSIMD拡張によってもカバーできるでしょ
674 名前：デフォルトの名無しさん [2010/03/21(日) 08:51:43 ]: Linuxでドライバのバージョンを調べる方法を教えてください
675 名前：デフォルトの名無しさん mailto:sage [2010/03/21(日) 08:58:16 ]: もう少し具体的に聞かないと・・・・。
676 名前：デフォルトの名無しさん [2010/03/21(日) 09:00:06 ]: ドライバのバージョンによって、振る舞いを変えたいので、
UbuntuでGPUのドライバのバージョンをプログラム上から調べる方法を教えてください
677 名前：デフォルトの名無しさん mailto:sage [2010/03/21(日) 09:22:25 ]: Ubuntu限定な必要があるかどうかはともかく、ドライバ自体のバージョンを知るAPIは用意されていません。
/usr/lib64のディレクトリでlibcuda.so.* のレギュラーファイルを探すのが確実かと思います。
678 名前：,,・´∀｀・,,）っ-○○○ mailto:sage [2010/03/21(日) 09:33:40 ]: CUDA実行環境がインストールされてるかどうかを調べてライブラリを遅延ロードできるような仕組みを
標準で用意して欲しいかな
679 名前：デフォルトの名無しさん mailto:sage [2010/03/21(日) 20:24:40 ]: CUDA3.0
ttp://developer.nvidia.com/object/cuda_3_0_downloads.html
680 名前：デフォルトの名無しさん mailto:sage [2010/03/21(日) 21:52:46 ]: Fermi対応版か
肝心のブツが手元にまわってくるかも怪しいのに
681 名前：デフォルトの名無しさん mailto:sage [2010/03/21(日) 22:30:11 ]: deviceQueryがあるじゃん
682 名前：デフォルトの名無しさん mailto:sage [2010/03/22(月) 02:31:14 ]: Fermiは、入手しやすくなるまで松わ。
初物は爆熱で卒倒価格だろうし。
683 名前：デフォルトの名無しさん mailto:sage [2010/03/22(月) 14:26:17 ]: 俺は特攻する
470か480かが問題
そもそも手に入るのか、という話もあるが
684 名前：デフォルトの名無しさん mailto:sage [2010/03/22(月) 15:02:55 ]: そうか。取り合えず貼っておかねばなるまい
::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::。:::::::::::::::::::::::::::::::::::::::::::::
:::::::::::::::::::::::::::::::::。::::::...... ...　　 --─-　　:::::::::::::::::::: ..::::: . ..::::::::
:::::::::::::::::...... ....:::::::゜::::::::::..　　（___ ）（___ ）　::::。:::::::::::::::::　゜.::::::::::::
:. .:::::。:::........　.　.::::::::::::::::: _　i/ ＝　＝ヽi　:::::::::::::。::::::::::: . . . ..::::
:::: :::::::::.....:☆彡::::　　　／/[||　　」　　||]　>>683 ::::::::::゜:::::::::: ...:: :::::
　:::::::::::::::::: . . . ..: ::::　/　ﾍ　| |　 ____,ヽ | |　:::::::::::.... .... ..　.::::::::::::::
::::::...゜　.　.:::::::::　　／ヽﾉ　　ヽ＿＿/　　.......　.　.::::::::::::........ ..::::
:.... .... ..　.　　　　く　 /　　　　三三三∠⌒＞:.... .... ..　.:.... .... ..
:.... .... ..:.... .... ..... .... ..　.:.... .... ..　..... .... ..　..... ............. ..　.　........　......
:.... . ∧∧　　∧∧　　∧∧　　∧∧　.... .... ..　.:.... .... ..... .... ..　.
... ..:（　　）ゝ（　　）ゝ（　　）ゝ（　　）ゝ無茶しやがって… ..........
....　 i⌒　/ 　 i⌒　/ 　i⌒　/ 　 i⌒　/　..　..... ................... ..　.　...
.. 　三　 | 　三　 | 　三　 | 　三　|　　... ............. ...........　.　.....
...　 ∪ ∪ 　 ∪ ∪ 　 ∪ ∪ 　∪ ∪　............. ............. ..　........　...
　　三三　　三三　　三三　　　三三
　三三　　三三　　三三　　　三三
685 名前：デフォルトの名無しさん mailto:sage [2010/03/24(水) 18:34:15 ]: 東工大、気象庁の次世代気象モデルのフルGPU化に成功
pc.watch.impress.co.jp/docs/news/20100324_356466.html

120GPUで3.22TFLOPSか。大変なんだろうけど微妙な数字だな。
686 名前：デフォルトの名無しさん mailto:sage [2010/03/24(水) 19:27:10 ]: どう微妙なんだい
687 名前：デフォルトの名無しさん mailto:sage [2010/03/24(水) 20:00:16 ]: なんか膨大な演算能力は殆ど遊んでいるな。
こういう用途なら演算機減らした方が、電力効率的にはマシになるんじゃ。
688 名前：デフォルトの名無しさん mailto:sage [2010/03/24(水) 20:06:08 ]: ha?
689 名前：デフォルトの名無しさん mailto:sage [2010/03/24(水) 20:13:50 ]: え？
690 名前：デフォルトの名無しさん mailto:sage [2010/03/24(水) 20:16:26 ]: ま、近い将来fermiに置き換わるんだろ
691 名前：デフォルトの名無しさん mailto:sage [2010/03/24(水) 20:34:22 ]: >>685
1GPUあたり44.3GFlopsかぁ。維持で対応したって感じだ。
692 名前：デフォルトの名無しさん mailto:sage [2010/03/24(水) 20:48:15 ]: 浮動小数点演算速度が3.22TFLOPSなのか？
693 名前：デフォルトの名無しさん mailto:sage [2010/03/24(水) 20:48:57 ]: それでもCPUより80倍の実行性能なんだね
694 名前：デフォルトの名無しさん mailto:sage [2010/03/24(水) 20:52:14 ]: 6時間の気象モデルが70分で終了するなら
同じ気象モデルを使えばリアルタイム予測が可能？
695 名前：デフォルトの名無しさん mailto:sage [2010/03/24(水) 23:28:24 ]: 偏微分方程式の数値解法をやる大学院生の演習課題みたいなもんですなw
696 名前：デフォルトの名無しさん mailto:sage [2010/03/25(木) 00:41:07 ]: 気象問題って、ノード間の影響はどうなの？
TESLAって結局PCI Expressで繋がっているから、
レイテンシが大きそうだな。
697 名前：デフォルトの名無しさん mailto:sage [2010/03/25(木) 01:59:47 ]: とりあえず運用してノウハウ貯めて、
28nmのFermi2で一気にパワーアップってのがいいんじゃないか？
明確なハズレ世代を大量導入するのはちとどうかと。
698 名前：デフォルトの名無しさん mailto:sage [2010/03/25(木) 07:32:00 ]: >>687
そういうのは GPU 以外に期待した方がいいんじゃないの？
電力効率なんて気にしたせいでピーク性能が落ちてゲームユーザが買わなくなったら、GPU の市場自体が崩壊するよ。
699 名前：デフォルトの名無しさん mailto:sage [2010/03/25(木) 14:56:59 ]: >>696
レイテンシを隠蔽するようにcode組んだらしい
700 名前：デフォルトの名無しさん mailto:sage [2010/03/26(金) 00:39:01 ]: >>699
へー、そーなんだ。
701 名前：デフォルトの名無しさん mailto:sage [2010/03/26(金) 00:59:41 ]: >>693
これがよくわからんな。
なんで80倍なんだろう？
G200って倍精度の理論値って80Gflops位じゃなかった？
それに対してCPUが10Gflopsだろ。
8倍の間違いじゃないのか？
702 名前：デフォルトの名無しさん mailto:sage [2010/03/26(金) 03:03:04 ]: techon.nikkeibp.co.jp/article/NEWS/20100324/181319/?ST=lsi
> 「最新のマイクロプロセサとの比較ではなく，チューニングの程度の
> 差もあるため，80倍という数字自体は重要ではない。GPUの活用で
> ケタ違いの性能が得られることを確認できた点に意味がある」（青木氏）

だそうで
703 名前：デフォルトの名無しさん mailto:sage [2010/03/26(金) 03:12:28 ]: 倍精度もハードウェアで本格対応して500GFLOPSになったんじゃないの
と思ったけどFermi世代からで、G200世代のものは90GFLOPS程度みたいだね。

これで倍精度で44.3GFLOPSなら効率50%近くということになるし
アルゴリズムとかかなり頑張っただろうね。

>>701
シェアードメモリをうまく使ったり、大量のスレッドでノード間のレイテンシを隠蔽したり
帯域がボトルネックになりにくいように出来たからじゃないの。

スカラープロセッサは、流体力学とか多体問題とかの
計算結果を相互に利用しながら並列計算する場合には
帯域がボトルネックになって効率がかなり低くなりがちらしいし、
比較対象のCPUを使ったシステムは効率が10%を下回っていたとかかも。

完全に並列化できて帯域がボトルネックになりにくい演算の場合は
10～20倍という話だったかな。
704 名前：デフォルトの名無しさん mailto:sage [2010/03/26(金) 05:40:03 ]: >>703
メモリ帯域が問題になるのはよくわかるんだけど、
それも今のNehalemあたりは30GB/sぐらいあるから、
精々teslsaとは5倍くらいだろう。
それだとメモリ転送時間は1/5、計算時間が1/8になるとして、
CPUのメモリ転送時間をA、計算時間をBとすれば、
実行時間は単純にするとC=A+Bになる。
GPUの場合はCg=0.2A+0.125Bになる。
A=10Bとしたとしても5倍程度にしかならない様な気がするなあ。
CPUの場合巨大なキャッシュがあるし、プリフェッチもあるから
その差が縮まるはずなんだけどな。
東工大のクラスのだからCPUのコードがくそだということもなさそうだから、
なんかOpteronのCPU1コアと比較してそうだなあ。

と思っていたら、>>702が書いてくれているね。
705 名前：デフォルトの名無しさん mailto:sage [2010/03/26(金) 05:49:43 ]: 比較するCPUが遅すぎたわけか
706 名前：デフォルトの名無しさん mailto:sage [2010/03/26(金) 07:17:15 ]: このcpuの速度は、普通にべたべたfpu演算を書いたときの数字だね。
理研の姫野ベンチ並みのバカコードと対照させても意味がない。

Linpackで最近のcpuをベンチマークすると、理論値の8掛け程度の数字は出る。
TUBAMEのopteronも1コアあたり10GFlops前後。しかも倍精度で。
Linpackに比べて気象エミュは速度が上がらんのは間違いないが、
この青木とやらの記事は全然ダメだ。
nVIDIAが成果としてレファできないレベル。
707 名前：デフォルトの名無しさん [2010/03/26(金) 08:52:49 ]: CPUというのは全然チューニングしてないレファレンスコードだろ。
よく使う手。しかもCPU名すら書いてない。

ベンチマークには２種類ある。嘘か大嘘。
708 名前：デフォルトの名無しさん mailto:sage [2010/03/26(金) 17:23:46 ]: PPTを見ると、44.3GFLOPSは単精度、倍精度では15GFLOPSだね。
比較のCPUは倍精度で0.5GFLOPS以下で、単精度の44.3GFLOPSと比較し、80倍以上といってるんだね。
青木いい加減にしろよｗ
709 名前：デフォルトの名無しさん mailto:sage [2010/03/26(金) 21:19:29 ]: これはひどい。
710 名前：デフォルトの名無しさん mailto:sage [2010/03/26(金) 22:05:03 ]: >>704
この手の計算の場合はCPUのキャッシュやプリフェッチはあまり効果が無いと思うけど
Nehalemだと30GB/s程度メモリ帯域があるのか。

ただCUDAでは単なるメモリ帯域の比較だけでなく、シェアードメモリの利用や
大量のスレッドによるメモリアクセスレイテンシの隠蔽が重要になってくると思う。

>702の記事見たけど、デュアルコア2.4GHzのOpteronの1コアと
TESLA S1070の中の1基の比較なのかな・・・

>>708
倍精度だと思ったら、単精度だったのか・・・
そのpptってどこで公開されているんだろ？
711 名前：デフォルトの名無しさん mailto:sage [2010/03/26(金) 22:33:58 ]: 倍精度でも30ばいかい？
712 名前：デフォルトの名無しさん mailto:sage [2010/03/27(土) 01:43:50 ]: なんでみんなプレスリリースくらい読まないの？
不思議
713 名前：デフォルトの名無しさん mailto:sage [2010/03/27(土) 03:02:10 ]: >>708
おいおい・・・まじかよ。
かなり酷いなあ。まだ騙すことは出来るかもしれないが、
これがばれてくると偉いことになるぞ。

そういえばNVIDIAのGPUカンファレンスでCPUの2000倍速くなりましたってのがあったが、
Tesrax4、Opteron 2.4GHz 1Coreのものだった。
かなりつっこまれていたよ。
で、GPUは最適化しましたけど、実はCPUは最適化してませんだった。

>>710
この手の計算はどちらかというとストリーミングに近いから、
キャッシュやプリフェッチは十分効くよ。
特にハードウェアプリフェッチはかなり効果的だよ。
ただ、SSEを使ってしまったりすると、計算時間を隠蔽できなくなって、
今度は転送時間を隠蔽する方法を考えなくてはならなくなる。

Nehalemの中でも2000MHzのメモリに対応したものであれば、
48GB/sにもなるよ。このあたりはオーバークロック気味になるので、
コンシューマレベルでしかないけどね。

GPUを使うのを否定はしないが、いい加減嘘に近い誇張は止めてもらいたいものだ。
数倍でもいいだろう。2倍でも2日かかったものが1日で終わるんだぜ。
あんまり速いと仕事が増えるじゃねぇか！
714 名前：デフォルトの名無しさん mailto:sage [2010/03/27(土) 03:47:59 ]: 嘘は言ってないんだよな
比較対象が微妙すぎるだけで
CPUでも効率を出すのが難しい問題はあるから、両方ともきちんとチューニングした上で比較するのが理想ではある
まあXX倍という数字が一人歩きするのはよくないね
715 名前：デフォルトの名無しさん mailto:sage [2010/03/27(土) 04:43:39 ]: >>712
そのプレスリリースってどこにあるの？

>>713
この手の計算って常に帯域を使い切っているわけじゃないの？
それならキャッシュの効果は限定的になりそうだけど、プリフェッチはかなり効果ありそう。

journal.mycom.co.jp/news/2010/03/24/055/index.html
> 単一のGPU(Tesla S1070)を用いた際の性能は
> 「単精度ではTSUBAMEに搭載されているOpteron比で最大で100倍を超すレベルを達成、
> 平均でも 44,3GFlopsを達成している。
> 性能が落ちる倍精度でも15GFlops程度を達成しており、
> CPU比では相当高い値を実現した」(青木教授)と説明する。

www.gsic.titech.ac.jp/contents/press_release0324.pdf
> スパコンTSUBAMEの単一GPUを使った計算で
> 44.3GFlops（CPUの1コアに対して約80倍）

これってやっぱり、Tesla S1070の4基を単一GPUと言って
Opteronの1コアと比較している同じパターン・・・？

あまりに酷い誇張はやめてもらわないと、有能な開発者がGPUの性能をそれなりに引き出しても
能力が低いとか言われたり、色々と深刻な事態になりかねない・・・
716 名前：デフォルトの名無しさん mailto:sage [2010/03/27(土) 04:54:00 ]: 実際、某社でなりかけている。

私の現場では、寧ろ旧世代のCPUを積んだサーバ機をGPUで延命させるってシナリオだから
処理能力は1.5倍でも御の字さw まぁ、実際のところはXeon1core対比で5倍程度は出たから
サーバ機一台で換算してなんとか2倍(5+1*3 vs 4)の数字は出たけど。
717 名前：デフォルトの名無しさん mailto:sage [2010/03/27(土) 05:08:13 ]: >>715
結局実装アルゴリズムによるけど、
プリフェッチはかなり効くよ。

最後の2行はすごく同意だな。
わかっていない人からすると、GPU使えば数10倍になるはずと思ったりするので、
有能な奴が5倍速い書いたコードを書いたとしても相対的に低いと評価されそうだ。

>>716
たぶんこういう使い方が一番いいのかもね。
718 名前：デフォルトの名無しさん mailto:sage [2010/03/27(土) 07:21:48 ]: >>715
CPU比ではそうかもしれんが
AMD比ではかなりしょぼくないか？ｗ
719 名前：デフォルトの名無しさん mailto:sage [2010/03/27(土) 09:10:23 ]: nttxstore.jp/_II_EA13119354
NVIDIA Tesla C2070　ETS2070-C6ER　475,545円（税込)　発売開始日 2010/9/30
720 名前：デフォルトの名無しさん mailto:sage [2010/03/27(土) 12:09:51 ]: GTX 470欲しいな。でもどうせ瞬殺なんでしょ？
いいよな秋葉原に住んでるおまえらは
721 名前：デフォルトの名無しさん mailto:sage [2010/03/27(土) 12:47:08 ]: >>716
GPUで延命→鯖予算獲得→新鯖にGPU移植
の流れが理想
722 名前：デフォルトの名無しさん mailto:sage [2010/03/27(土) 13:51:52 ]: GPGPUでの性能評価の指針って明確に定められないのかね
CPU単一コア，レファレンスコードと比較して何倍速くなったと言っても，
そんな条件で計算すること自体なさそうなんで，意味無い気がするんだが

GPUでチューニングするなら，CPUでもチューニングするべきだろう
723 名前：デフォルトの名無しさん mailto:sage [2010/03/27(土) 14:00:31 ]: AMDかintelかでめんどくさいからじゃない
724 名前：デフォルトの名無しさん mailto:sage [2010/03/27(土) 14:34:25 ]: >>722 CPUまでチューニングする必要はないと思うけど、
気象コードがMPIで並列化されてないとも思えないし、1GPU内でも相当の並列化を行ってるんだから、
CPUの1コアのみと比較するのは、ちょっとやりすぎ。
で実際Xeon/X5570と倍精度で比較したら、単一コアで10倍程度、4コアで3倍程度のアドバンテージしかないんじゃない。
ノード単位だと逆に遅くなるとか。120GPUの3.2TFLOPSも単精度だし…、比較自体を都合のいいように、いいとこどりしすぎ。
ベンチマークだけでなく、せっかく実アプリをフルGPU化したのは十分意味があるのに、世間受けを狙いすぎた発表で逆に良識を疑ってしまう。
725 名前：デフォルトの名無しさん mailto:sage [2010/03/27(土) 15:23:41 ]: 第三者（？）が書いた記事よりも
NVIDIAの発表のほうがまだ信頼できるな。
726 名前：デフォルトの名無しさん mailto:sage [2010/03/27(土) 16:40:26 ]: 詐欺師の言葉のほうが信頼があるとはこれ如何に
727 名前：デフォルトの名無しさん mailto:sage [2010/03/27(土) 16:46:56 ]: 詐欺師はお客様に信用されなければいけない。
だから最低限の嘘しかつかない。

大道芸人は嘘をつくたびに金をもらえる。
だからいくらでも嘘をつく。
728 名前：デフォルトの名無しさん [2010/03/27(土) 18:47:44 ]: モックアップＮＶＩＤＩＡとうそつきが多い野心的な学者連中の
たわ言なんて信用できないよな
729 名前：デフォルトの名無しさん mailto:sage [2010/03/27(土) 19:13:00 ]: ま俺は数字しか信じない
ttp://www.anandtech.com/video/showdoc.aspx?i=3783&p=6
730 名前：デフォルトの名無しさん mailto:sage [2010/03/27(土) 19:26:53 ]: 数字なんていくらでもいじれるんですよ！
信用のあるデータが欲しいなら自分で実測するしかない
731 名前：デフォルトの名無しさん mailto:sage [2010/03/27(土) 19:36:26 ]: つまりそのデータを公開しても
誰も信じないってわけね
732 名前：デフォルトの名無しさん mailto:sage [2010/03/27(土) 19:44:22 ]: その通り！
よく分かってるじゃん
論文だって他人の実験の結果なんて基本誰も信用しないよ
とりあえず自分で再現実験してみるのは常識
733 名前：デフォルトの名無しさん mailto:sage [2010/03/27(土) 20:09:19 ]: そしてその論文も誰も信用しない
734 名前：デフォルトの名無しさん mailto:sage [2010/03/27(土) 21:06:02 ]: 利害を共にしない数十人の人が
方法の妥当性を検証し、追試し、同じ結果が出て初めて
ある程度の信用を得る
それでも「ある程度」なのが学問の厳しい所だ
735 名前：デフォルトの名無しさん [2010/03/27(土) 21:40:16 ]: GPUは1コアだと思いますよ。TESRA内でもGPU同士はPCI越しなんで
遅くなると思います。
3000x3000x50って1GPUにおさまるサイズと思うし
736 名前：デフォルトの名無しさん [2010/03/27(土) 21:41:34 ]: 誰も追試できないと思っているから舐めているんだろ。
今後は青木の言うことはハイハイワロスだな。
737 名前：デフォルトの名無しさん [2010/03/27(土) 21:43:30 ]: >>735
それだったら120GPUの意味がないだろ。
あと、Teslaな
738 名前：デフォルトの名無しさん [2010/03/27(土) 21:49:48 ]: いくつかの別な測定を同じ測定かもしれない感じで書くのはテクニックじゃないですか
１GPUで44.3GFLOPSなのに120GPUでなんで3.22TFLOPSなんでしょうか。
739 名前：デフォルトの名無しさん mailto:sage [2010/03/28(日) 09:52:45 ]: Zotac GeForce GTX 480
Amazon.comでPre Order　$499だったから
申し込んでしまったｗ　(米→日転送業者使用)
740 名前：デフォルトの名無しさん mailto:sage [2010/03/28(日) 13:50:46 ]: 結局512spじゃなかったねGTX480
741 名前：デフォルトの名無しさん mailto:sage [2010/03/28(日) 15:30:51 ]: Ultraがくる
742 名前：デフォルトの名無しさん mailto:sage [2010/03/29(月) 03:53:35 ]: >>740
最初から、576にしておけばよかったのにね。
32x18で。2ブロックダメでも512で出せるのに。
743 名前：デフォルトの名無しさん mailto:sage [2010/03/29(月) 08:15:18 ]: そんなことしたらますますイールドが悪くなるわけで。
744 名前：デフォルトの名無しさん [2010/03/29(月) 11:43:02 ]: 1CUDAコア破損しただけで３２個分が台無しになるFermi
512すべて無事なのはほとんどないんだろうな
745 名前：デフォルトの名無しさん mailto:sage [2010/03/29(月) 11:50:30 ]: PS3と同じ作戦でござる。

　　　　＿＿＿＿＿＿＿＿＿＿＿＿
　　　　ヾミ || || || || || || || ,l,,l,,l 川〃彡|
　　　　　V~~''-山┴''''""~ 　　ヾニﾆ彡|　　　　　　　512SPは存在する・・・・・・！
　　　　 /　二ー―''二　　　　　ヾﾆニ┤　　　　　　存在するが・・・
　　　 <'-.,　　￣￣　　 _,,,..-‐､　〉ニﾆ|　　　　　　　今回　まだ　全部有効にするとの
　　　/"''-ﾆ,‐l 　　l`__ニ-‐'''""` /ニ二|　　　　　　　指定まではしていない
　　　|　==＝､!　　`＝====､　　l =ｌべ=|
.　　　| `ーﾟ‐'/　　 `ー‐ﾟ―' 　　l.=lへ|~|　　　　　　　そのことを
　　　 |`ー‐/　　　　`ー――　　H<,〉|=|　　　　　　　どうか諸君らも
　　　 |　　/　　　､　　　　　　 l|__ノｰ|　　　　　　　思い出していただきたい
.　　　| ／`ー　~　′　　＼　　　.|ヾ.ニ|ヽ
　　　 |ｌ　下王l王l王l王lヲ｜　　|　ヾ_,|　＼　　　　　つまり・・・・
. 　　 |　　　　≡　　　　　　　　　|　　 `l　　＼__　　　我々がその気になれば
　　　　!、　　　　　　　　　　 _,,..-'′ ／l　　　 |　~'''　 FermiのSP数は
‐''"￣|　`iｰ-..,,,＿,,,,,....-‐'''"　　／　 |　　　 |　　　 320SP　384SP　ということも
　-―|　　|＼　　　　　　　　　／　　 |　　　　 |　　　可能だろう・・・・・・・・・・ということ・・・・！
　　 |　　 |　＼　　　　　　／　　　 |　　　　　 |
746 名前：デフォルトの名無しさん mailto:sage [2010/03/29(月) 14:07:53 ]: 　　　　　＿＿＿_
　　　　　　　|<三`'ヨ′
　　　　　　_/6|ｰ廿┤
　　　 /l￣　ＫL.三.｣￣h
.　　 /　|　　レ兮y′/ l
　　〈　く　　 ∨ l/　,ｲ　|
　　　＼_,.>、　　　 /,L..｣_
.　0ﾆﾆﾆ)而}ﾆﾆﾆﾆニ),ﾘﾘﾆﾆ)
.　　Ｌ| |＿＿＿__|＿＿__| |
　　 l | |.＿＿＿＿＿＿_| |　　,:
　,　　l ＼ヽ　l　　|　　 , '／　　;'
　:, ____l＿|_|_;_|＿|＿__|_|＿＿　　　;
　　|＼ﾞ;三三ﾞ';三三三,;ﾞ三三＼　;'
　　|＼＼三三ﾞ三ｼﾞｼﾞ三三,''三;'＼,;'　　;'
　　|、＼＼三ﾞ;三三ｼﾞｼﾞ･'三三三;＼　;
　　0ﾄ､＼＼＼;'三三;'三三三;''三三,;'＼
　　　　＼＼＼| 炎炎炎炎炎炎炎炎炎 |
　　　　　＼＼| 二I二二I二二I二二I二 |
　　　　　　　＼ＬI二二I二二I二二I二二｣
　　　　　　　　0｣　　　　　　　　　　　　0｣
747 名前：デフォルトの名無しさん [2010/03/30(火) 10:14:08 ]: NVDIAフォーラムでGTX480の倍精度性能はTeslaの1/4という発言が
ありますがどうなんでしょう
748 名前：デフォルトの名無しさん mailto:sage [2010/03/30(火) 10:30:27 ]: >>747
おれ、人柱としてぽちったから待っててくれ。
749 名前：デフォルトの名無しさん mailto:sage [2010/03/30(火) 11:15:36 ]: >>748
俺はお前を待っているぞ
+　　　+
　 ∧＿∧ 　+
　（0ﾟ・∀・）　　　ﾜｸﾜｸﾃｶﾃｶ
　（0ﾟ∪ ∪ +
　と＿_）__）　+
750 名前：デフォルトの名無しさん mailto:sage [2010/03/30(火) 11:36:18 ]: Quadroと同じ戦略か。
妥当っちゃ妥当だな
751 名前：デフォルトの名無しさん [2010/03/30(火) 12:00:52 ]: EECだけでなく倍精度も性能を劣化させたのか
GPGPUの利点がどんどんなくなってるな
752 名前：デフォルトの名無しさん mailto:sage [2010/03/30(火) 12:20:38 ]: >>749
あ、748なんだけど、さすがにTesla買う金は無いのよ。
なんか、GTX480で動かして「明らかに倍精度の性能落としてやがる！！！１１」と
分かるプログラムはどこかにあるかな。
753 名前：デフォルトの名無しさん mailto:sage [2010/03/30(火) 12:21:31 ]: Teslaもないと比べようがなくね？

visual profilerの関数の実行時間を比較くらいしか厳密な計測はできなさそう。
754 名前：753 mailto:sage [2010/03/30(火) 12:22:12 ]: リーロードしてなかった、悪気はないんだｗ
755 名前：デフォルトの名無しさん [2010/03/30(火) 12:34:45 ]: NVIDIA_GPU_Computing_SDK/C/bin/linux/release$ ./matrixMul
Processing time: 0.120000 (ms)
Test PASSED

Press ENTER to exit...
あたりを倍精度化してもらえばいいかと
756 名前：デフォルトの名無しさん mailto:sage [2010/03/30(火) 12:36:50 ]: >>753
かぶって申し訳ないｗｗ
おれ>>739なんだけど、失敗かな。日本で入手可能になるのを
待ったほうが早くて安かったかもしんないよなorz
757 名前：デフォルトの名無しさん mailto:sage [2010/03/30(火) 12:41:40 ]: >>755
おけ、matrixMulでGTX280とGTX480の比較ならやってさしあげられる。
floatのままと、doubleに全部置換した版で。
758 名前：デフォルトの名無しさん mailto:sage [2010/03/30(火) 12:50:39 ]: floatは内部で型変換してるからdoubleの方が早いって聞いた事あるんだけど…
759 名前：デフォルトの名無しさん mailto:sage [2010/03/30(火) 12:59:07 ]: 倍精度のFMAをひたすら繰り返す感じのカーネルで計測するのがよいかと。
760 名前：デフォルトの名無しさん mailto:sage [2010/03/30(火) 13:09:51 ]: >>758
どこかのCPUでintをdoubleに変換してたという話かと。
761 名前：デフォルトの名無しさん mailto:sage [2010/03/30(火) 13:10:18 ]: しかし、480SPとか、AMDはリアル12コアだとか、
時代の進歩は速いものだな。
762 名前：デフォルトの名無しさん [2010/03/30(火) 14:06:01 ]: しかし、性能はGTX295から毛の生えた程度
ていうか、もし倍精度の性能がいまいちだったらGTX480/470より
GTX295を買ったほうが安くていいかもね
763 名前：デフォルトの名無しさん [2010/03/30(火) 14:14:19 ]: 長崎大のようなことをやられたらNVDIAが東工大から呼出をうけて
「２度目はないからな」とか言われても不思議じゃないよね
764 名前：デフォルトの名無しさん mailto:sage [2010/03/30(火) 15:00:14 ]: >>762
GTX295、5万円くらいだからなぁ。倍精度の性能がGTX280の２倍程度ならば…
GTX295でヨシってことになると。>>757よろしく！
765 名前：デフォルトの名無しさん mailto:sage [2010/03/30(火) 15:19:12 ]: matrixMulはメインメモリの転送が時間に含まれてるしサイズが小さめ
プログラムが書けるなら1000x1000あたりの性能をお願いします
766 名前：デフォルトの名無しさん mailto:sage [2010/03/30(火) 15:43:23 ]: >>760
そうなのか？
普通のCPUの構造上浮動小数点演算は、
64bitでやるから型変換をしてると聞いたんだけど
767 名前：デフォルトの名無しさん mailto:sage [2010/03/30(火) 15:48:57 ]: もしかしてGPUだと話は別なのかな…
768 名前：デフォルトの名無しさん mailto:sage [2010/03/30(火) 18:34:02 ]: >>766
x86アーキテクチャでは、普通は変数型に関わらず浮動小数点演算は80bitの拡張倍精度で行う。
メモリロード/ストアの際に、変数がfloat型ならfloat型に変換されはするが、
別に大した処理でもないので速度にそんな影響は出なかったと思うが。
769 名前：デフォルトの名無しさん mailto:sage [2010/03/30(火) 18:46:24 ]: R3000かなんかで64bitint乗算がなくてとかなかったっけ
770 名前：デフォルトの名無しさん mailto:sage [2010/03/30(火) 22:56:03 ]: >>767
GPU だとどころか、倍精度演算器が載ってない CPU だってあるから。
処理系によって話が別。
771 名前：デフォルトの名無しさん mailto:sage [2010/03/31(水) 02:58:13 ]: device emulationってなくなっちゃうのね
772 名前：デフォルトの名無しさん mailto:sage [2010/03/31(水) 05:57:59 ]: ちょっと整理。

・x86(fpu)
浮動小数点レジスタが80bitなので、floatを突っ込んでも80bitで演算する。
従って、floatとの変換処理が入るのでdoubleの方が速いことがしばしば。

・x86(sse)
MMXレジスタが汎用なので、floatはfloatのまま演算する。
従って、定数や標準関数でdoubleに汚染されないように気をつければfloatの方が多少速くなる。
また、MMXレジスタにfloatの方が2倍詰め込めるのでベクタ化した場合に2倍速くなる可能性もある。
更に、キャッシュ効率もよくなるのでより速くなるかもしれない。

・GPU(cuda)
単精度レジスタと倍精度レジスタが分かれているんだっけ? 倍精度についてよく知らんのでフォローお願い。
演算器の特性上、単精度の方がずっと高速に演算するし、転送量も当然半分にできる。
応用にも拠るけれど、CPUとのI/Fは単精度で内部だけ倍精度にできると最強かも。
773 名前：デフォルトの名無しさん mailto:sage [2010/03/31(水) 07:48:43 ]: 64bit版gccだと、-m32オプションをつけない限りもはやx87のコードは吐かないなあ。

>CPUとのI/Fは単精度で内部だけ倍精度にできると最強かも。
これって何回丸めが発生するかによるけど、精度は単精度+αだね。
でもこういうことはよくやるなあ。
データ量を減らしたくて少し精度が欲しい場合は、最終的に欲しいデータはfloatで持っておいて、
計算するときにすべてdouble型のデータにコピーして最後に結果をfloat型に戻す。
この方法をとれば、丸め誤差の発生は一回で済む。
計算の中身が超越関数使ったり、複雑であったりすると結構有効だよ。
その分速度は犠牲になるけど、メモリ帯域が支配的であると余り影響がないな。
G200系だとまだ効果はないだろうけど、Fermiなら結構役に立つんでない？
774 名前：デフォルトの名無しさん mailto:sage [2010/03/31(水) 11:35:13 ]: >・x86(fpu)
>浮動小数点レジスタが80bitなので、floatを突っ込んでも80bitで演算する。
>従って、floatとの変換処理が入るのでdoubleの方が速いことがしばしば。
doubleは64bitなんだから結局まるめが入るのは同じじゃないの？
775 名前：デフォルトの名無しさん mailto:sage [2010/03/31(水) 13:02:05 ]: おいおいあれだけCPU側のコードが糞だといっておいて、CPUのこと知らなさすぎだろｗ

floatが遅くなるのは、MSVCが、floatのときは毎回メモリに書き戻すことで精度を32bitに落としてIEEE互換にするコードを出すから
レジスタ間ならfloat/double/long double関係なく80bit
776 名前：デフォルトの名無しさん [2010/03/31(水) 14:32:17 ]: GPUの種類を取得する関数か何かはありませんか？
777 名前：デフォルトの名無しさん mailto:sage [2010/03/31(水) 14:42:10 ]: >>776
サンプルでついてくるdeviceQueryのソース読んでみれば
778 名前：デフォルトの名無しさん mailto:sage [2010/03/31(水) 14:44:59 ]: >>776
CUDA APIガイドに書いてあるだろ
779 名前：デフォルトの名無しさん [2010/03/31(水) 15:12:11 ]: そうじゃなくて製品名です
780 名前：デフォルトの名無しさん mailto:sage [2010/03/31(水) 15:20:59 ]: 製品名って、ASUSかEVGAか判別したいってこと？？
781 名前：デフォルトの名無しさん mailto:sage [2010/03/31(水) 15:26:11 ]: メーカーの判別はどうがんばっても無理だろ
782 名前：デフォルトの名無しさん [2010/03/31(水) 15:32:47 ]: じゃあせめてWindowsかMacかInaxかぐらいわかりませんか？
783 名前：デフォルトの名無しさん mailto:sage [2010/03/31(水) 15:33:23 ]: 製品名って言うくらいだから、ELSA GLADIAC 998 GTX Plus V2 512MBみたいなのじゃないのか。
784 名前：デフォルトの名無しさん mailto:sage [2010/03/31(水) 15:34:04 ]: >>782
CUDA以前の問題。てか、そのレベルじゃ絶対無理だろ。
785 名前：デフォルトの名無しさん mailto:sage [2010/03/31(水) 15:53:38 ]: #ifdef _WIN32
786 名前：デフォルトの名無しさん mailto:sage [2010/03/31(水) 15:55:08 ]: まさかOpenCLの話か？
#ifdef _apple
とかやった覚えがある。
787 名前：デフォルトの名無しさん [2010/03/31(水) 16:15:17 ]: #ifdef _appleは豆知識だな
788 名前：デフォルトの名無しさん mailto:sage [2010/03/31(水) 16:44:37 ]: Inax は釣りだろ。782 は偽者じゃないか？
TOTO 向けと別の最適化するのかなｗ
789 名前：デフォルトの名無しさん mailto:sage [2010/03/31(水) 19:31:24 ]: 水流をGPUでシミュレーションして最適化するのか
790 名前：デフォルトの名無しさん mailto:sage [2010/03/31(水) 21:04:19 ]: 流体シミュレーションはGPGPUのメインテーマだから、まさにうってつけだろう
791 名前：デフォルトの名無しさん mailto:sage [2010/03/31(水) 21:06:11 ]: OpenCL使えよ馬鹿ども
792 名前：デフォルトの名無しさん mailto:sage [2010/03/31(水) 21:08:09 ]: 日本の便器メーカーは水量削減に血道を上げているからな
いかに少ない水量で、効率良く、かつきっちり排泄物を流しきるか
793 名前：デフォルトの名無しさん mailto:sage [2010/03/31(水) 21:11:31 ]: 便器開発での計算に使用するだけでなく、
便器自体にTeslaを搭載し、
排泄物を画像認識して最適な水流を計算する。

排熱も有効利用できそうだ
794 名前：デフォルトの名無しさん mailto:sage [2010/03/31(水) 22:39:35 ]: >>793
画像認識のコード書くやつは大変だな
795 名前：デフォルトの名無しさん mailto:sage [2010/03/31(水) 22:42:09 ]: 確かに、今は節水ということで一般家庭向けの便器は
流れが悪いといくことを感じるな。

現場では流体演算とかして設計しているのだろうか？
796 名前：デフォルトの名無しさん mailto:sage [2010/04/01(木) 01:09:36 ]: TOTO 節水シミュレーションでググってみた。
www.toto.co.jp/saiyo/new/techno/person/person_05_2.htm
797 名前：デフォルトの名無しさん mailto:sage [2010/04/01(木) 01:17:06 ]: 流体どころか三相全てシミュレートしてるんだな
798 名前：デフォルトの名無しさん mailto:sage [2010/04/01(木) 07:04:31 ]: GTX470が先に発売されるみたいですね。
購入される方いますか。
799 名前：デフォルトの名無しさん mailto:sage [2010/04/01(木) 12:43:41 ]: どうしてcudaはosと密接なのか？
最新のubuntuをいれたくてもいれられへん
理由を教えてくれろ
800 名前：デフォルトの名無しさん mailto:sage [2010/04/01(木) 14:28:01 ]: >>799
ドライバレベルで提供されているから。
801 名前：デフォルトの名無しさん [2010/04/01(木) 19:55:02 ]: >>797
固体ってのが生々しいな。シミュレーションするためにウンコの物性とか
測定したりしたんだろうな。
802 名前：デフォルトの名無しさん mailto:sage [2010/04/01(木) 20:10:39 ]: 壊れ方とかがリアルな模型があるとか聞いたことがある。
でも物体の測定データのファイル名とかはなまなましいだろうなｗ

シミュレーションしているときのSSとかあればいいのに。
803 名前：デフォルトの名無しさん mailto:sage [2010/04/01(木) 21:27:09 ]: たしか法律で、流す水の量は決まってるんだよね。
804 名前：デフォルトの名無しさん mailto:sage [2010/04/02(金) 07:57:34 ]: >>799
たぶん、nvccがgccの進化について行けていないだけだと思う。
バイナリだけなら、最新のUbuntuでも動くよ。
805 名前：デフォルトの名無しさん [2010/04/02(金) 16:58:24 ]: 最近のディストリビューションって大抵1年かそこらでサポート終わりじゃん。
1つまえのバージョンにしか入れられないと、半年程度でいれかえなきゃいけないんだよね。
そこらへんを早く何とかしてくれよ。
CentOSとかつかえばいいのかもしれんけど
806 名前：デフォルトの名無しさん mailto:sage [2010/04/03(土) 03:06:19 ]: そこを何とかしてもらいたいね。
俺は古いバージョンのLinuxをVirtualBoxで新しいバージョンのLinux上で動かして、
クロスコンパイルライクなことをしているよ。コンパイルだけなら仮想化環境でも通るからな。
807 名前：デフォルトの名無しさん mailto:sage [2010/04/03(土) 12:32:49 ]: GTX480で倍精度削られたのは本当らしい…

Quadroではどうなるんだろう
808 名前：デフォルトの名無しさん mailto:sage [2010/04/03(土) 14:22:26 ]: 今すぐ計算して論文書かなきゃって人以外は↓ここ聞いてからにした方がいいんじゃないかと。
ttp://www.hardocp.com/article/2010/03/26/nvidia_fermi_gtx_470_480_sli_review/7

「当機はまもなく離陸しますw」
809 名前：デフォルトの名無しさん mailto:sage [2010/04/03(土) 14:24:49 ]: >>807
ドイツ語読めないけど、これ？
ttp://www2.hardware-infos.com/news.php?news=3497
810 名前：デフォルトの名無しさん mailto:sage [2010/04/03(土) 22:55:58 ]: Q1.同じGT-240を二枚挿せばCUDAも倍近く早く処理できるようになりますか？
Q2.GDDR5とGDDR3とDDR3でCUDA動画エンコードの処理スピードはどれくらい
違いますか？GDDR3を100として。
811 名前：デフォルトの名無しさん mailto:sage [2010/04/04(日) 00:22:50 ]: >>810
もう少し具体的に知りたいことを書いた方が答えやすい。

例えば動画エンコード用途に限ると「複数枚挿しはどうよ？」とか、「FermiとかATIとかの中でどれが一番速いか」とか。
812 名前：デフォルトの名無しさん mailto:sage [2010/04/04(日) 00:38:10 ]: 十分具体的だろ･･･。

ソフトウェアによって、複数挿に対応しているかどうかは変わるので、
使いたいソフトについて調べよう。

メモリの速さは重要だけど、GPGPUでは、メモリにアクセスするときの遅延のほうが問題となっている。
ハードウェアの構造的にもGT200系のほうがメモリアクセスが柔軟なので、GTX260あたりを検討してはどうだろう。
813 名前：デフォルトの名無しさん mailto:sage [2010/04/04(日) 01:19:30 ]: 具体的だろと言っておきながらソフトによって変わるとか意味不明
814 名前：デフォルトの名無しさん mailto:sage [2010/04/04(日) 01:38:58 ]: >>811-812
レスありがとうございます。

CUDAの使用目的：動画エンコのみ。使用するソフト：MediaCoder
使用するかもしれないソフト：TMPGEncKarmaPlus
現状：AVIUTLでロゴ消しとインタレ解除のプロジェクト→
TMPGEnc4で色γクロップ・リサイズしてHUFFYUVで出力→MediaCoderのx264で。
課題：Q6600でVGAでx264エンコが22fpsしか出ない。1080pだと4fps！orz
MediaCoderにCUDAでH.264エンコできる機能があるので使いたい。
GT240がGDDR3で6000円、GDDR5で7500円～なので、二台組むより二枚入れた
ほうがいいのかも？
GDDR3とGDDR5でＣＵＤＡエンコに殆ど差がないならGDDR3のほうにしたい。
複数枚挿しはどうです？
現状G43/G41/G31なので新しくマザー買わなきゃできないけど…。

>>812
GPGPUやメモリアクセスについて仕組みとか全く知らないのですが、
GTX260はGT240の倍以上の値段だけど、倍の性能あるんですか？
wikiに書いてあるSPとCUDAコア数がどう違うのか分からないので…
もし倍の性能（エンコが倍早くなりそう）ならGT240よりGTX260を選ぶかも
しれません。PXI-EX⑯が1つですむし。
815 名前：デフォルトの名無しさん mailto:sage [2010/04/04(日) 01:42:45 ]: >>810
素直にCorei7にしておけ。
GT240じゃ2枚さそうがCorei7の方がマシだ。
816 名前：デフォルトの名無しさん mailto:sage [2010/04/04(日) 02:30:20 ]: 現状CPUを強化したほうがメリット多い
CUDA使うならGTX260以上じゃないとCPUの足を引張る可能性がある
817 名前：デフォルトの名無しさん mailto:sage [2010/04/04(日) 04:58:16 ]: aviutlでの処理時間考えたらCPU強化の方が妥当だな
818 名前：巻添規制中(810=814) mailto:sage [2010/04/04(日) 06:04:59 ]: >>815
www.techarp.com/showarticle.aspx?artno=520&pgno=7
でみるとi7はQ6600の1.25倍くらい早い
www.katch.ne.jp/~kakonacl/douga/mediacoder/v0.71_cuda.html
でみるとCUDAはQ6600の3.7倍くらい早い
∴CUDAはi7より3倍近く早い…？？？　CUDAは実写向き？

X58+i7買いたいけど、LGA1155待とうかと。
819 名前：デフォルトの名無しさん mailto:sage [2010/04/04(日) 09:47:48 ]: >>818
あ～、俺もこれ使ってcudaエンコしてるけど、Bフレームが4までだったり2passできなかったり
画質はいまいちですよん。たしかにCPUよりは速いとはおもうけど・・・。

新しいバージョンだと改善されてる可能性はあるけどね。
820 名前：デフォルトの名無しさん mailto:sage [2010/04/04(日) 11:55:03 ]: >>814
GT240の3倍くらいの値段で売られているが、3倍速くなるとはいえないけど、
MediaCoderは複数差対応していなし、GT240を2枚買っても無駄になる。
あたらしくマザー買う予算が削れるならGTX260でもいいかと。

画質もとめるなら、i7ってのは合意。
821 名前：810 mailto:sage [2010/04/04(日) 14:18:33 ]: みなさん色々有難うございます。ググりながら考えてるんですが難しいですね。

>>816 >>817
GTX260だと電源交換、i7はM/B&DDR3全部揃えないと…先立つものが…。(ToT)
www.katch.ne.jp/~kakonacl/douga/mediacoder/v0.71_cuda.html の人
は　GeForce9600GT(VRAM 512MB)Vista HomePremium(SP2)で実行、CUDAエンコ
でQ6600使用率83％と書いてあるけど、GT240のほうが高速ですよね？
pc.watch.impress.co.jp/docs/column/tawada/20091117_329556.html によると
GeForce 9600 GTの1,800MHz/256bit メモリ帯域幅は57.6GB/sec、
GeForce GT 240のGDDR5/3,400MHz/128bit 54.4GB/sec、
GeForce GT 240のDDR3/2,000MHz/128bit 32GB/sec。
DDR2-667（5.3GB/s） DDR2-800（6.4GB/s）デュアルだと倍。
もしかしてDDR2がボトルネックになる？　
GT240 GDDR3とGDDR5でCUDAエンコの速さの違いどうなんでしょう？
そこらへんが一番気になります。
822 名前：810 mailto:sage [2010/04/04(日) 14:19:49 ]: >>819
最新ではBフレームが16まで、Average/Variable/CBR/2pass/3pass できるみたいだよ。
x264ではAverage/Variableは何故かコマ落ちする。

>>820
SP216&DDR3のGTX260とCUDAコア数96のGT240だと、エンコ速度二倍差が出ます？
PHARAOH 500W電源だとGTX260+Q6600はギリギリかなぁ。

(Q6600+GT240)複数台にするほうが経済的かな? G41とDDR2が1組余ってるし…
823 名前：デフォルトの名無しさん mailto:sage [2010/04/04(日) 15:38:15 ]: hothardware.com/articleimages/Item1477/san.png
doubleはfloatの1/2*9ってとこか。

Teslaでやって1/2とかだったら、ロックなんだろうな。
824 名前：デフォルトの名無しさん mailto:sage [2010/04/04(日) 17:43:24 ]: これは低すぎるのでOpenCLでDoubleがHWサポート
されてないのではないかと
825 名前：デフォルトの名無しさん mailto:sage [2010/04/05(月) 00:43:36 ]: >>821
GDDR3とGDDR5とでは、メモリ帯域が倍違うので、ストリーム系のアプリでは大きく変わります。
ていうか、このクラスのボードで2枚挿しとか意味がない。
826 名前：デフォルトの名無しさん mailto:sage [2010/04/05(月) 01:09:55 ]: >823－824
むしろ理論値で単精度の5分の1になるはずのRADEONの倍精度が半分程度で済んでる方が気になる

>822
GT240なんてゴミ買うぐらいならいっそ中古のQ9xxxのCPUでも買った方が良くね？
827 名前：デフォルトの名無しさん mailto:sage [2010/04/05(月) 03:39:12 ]: 安物買いの銭失いとはこのこと

CUDAに大きな期待を抱かないほうがいい
828 名前：デフォルトの名無しさん mailto:sage [2010/04/05(月) 19:17:53 ]: >>823で、GTX295が285よりもスコア低いのはなぜ？
829 名前：デフォルトの名無しさん mailto:sage [2010/04/06(火) 01:17:45 ]: >>828
単純にGPUを一個しか使っていないからだと思う。
単一GPUなら285の方が速いからね。
830 名前：デフォルトの名無しさん mailto:sage [2010/04/06(火) 08:54:53 ]: >>825
CUDAエンコはストリームと違ってGDDR3もGDDR5殆ど関係無いってさ。
SP数(CUDAコア数)でほとんど性能が決まる。

>>827
最新のMediaCoderでVBRでやってみたらどう？
831 名前：デフォルトの名無しさん mailto:sage [2010/04/06(火) 19:53:29 ]: MLB オバマ始球式
832 名前：デフォルトの名無しさん [2010/04/07(水) 09:15:06 ]: HPC向けGPGPU終わりつつあるな

102 ：Socket774 ：sage ：2010/04/06(火) 23:24:12 ID:n4owrnuu

>>98
> HPC向けはどうなるんだろう？

マキーノの話だとこんなのが。
grape.mtk.nao.ac.jp/~makino/journal/journal-2010-04.html#3
> 一枚5万とかで買うのでない限り GPU は価格性能比では CPU に勝てなくなってしまった

以前はGPUの方が同コストのCPU比で10倍↑とか軽く叩き出してたけど、
CPUはマルチコア化が進みまくり値段下がりまくり…
ハイエンドGPGPU買うよりCPUの方がコスパが良くなってしまった。
基本直線番長のGPGPUよりCPUの方が扱いやすいし、プログラミングも先行きわからない
CUDAやらなくても、今までやってきた事そのままで走るし…
833 名前：デフォルトの名無しさん mailto:sage [2010/04/07(水) 09:26:05 ]: 確かに、OpenMPとかでそれなりに性能出るならそっちの方が超簡単だもんな。
834 名前：デフォルトの名無しさん mailto:sage [2010/04/07(水) 12:13:09 ]: 精度を削って性能を出したマシンで有名になったのにこういう時は倍精度の話だけか
835 名前：デフォルトの名無しさん mailto:sage [2010/04/07(水) 12:54:34 ]: >>832
マルチコア化が進みまくりって、↓か？当面CUDAエンコの半分にもならんだろう
akiba-pc.watch.impress.co.jp/hotline/20100403/ni_c6176se.html
akiba-pc.watch.impress.co.jp/hotline/20100327/ni_cw3680.html
836 名前：デフォルトの名無しさん [2010/04/07(水) 15:29:03 ]: >>835
CUDAエンコは実用性ゼロと言われているだろ
HPC以外の分野でCUDAを利用している人はいないだろ
837 名前：デフォルトの名無しさん mailto:sage [2010/04/07(水) 15:31:47 ]: >>834
そこ大事だよな。CPUと違って、倍精度・単精度の使い分けで性能違うもんな。
て、マキーノは昔、計算パスの場所によって計算精度が違う計算機を作ってたもんな
838 名前：デフォルトの名無しさん mailto:sage [2010/04/07(水) 15:33:36 ]: >>836
エンコの中の人が面倒くさがっているだけなんじゃないの
839 名前：デフォルトの名無しさん mailto:sage [2010/04/07(水) 15:38:47 ]: >>838
環境を作るのがめんどくさい。
コードを書き換えるのがめんどくさい。
最適化するのがめんどくさい。
840 名前：デフォルトの名無しさん [2010/04/07(水) 15:45:24 ]: >>838
まともなエンコを利用できないという事実が重要
841 名前：デフォルトの名無しさん [2010/04/07(水) 18:57:43 ]: たしかにOpteron12コアX4がFermiと同じ値段ならGPU終わるな。
既存のコードがそのまま動くし。
CUDAでこれ以上の性能が出るアプリは限られている。
842 名前：デフォルトの名無しさん mailto:sage [2010/04/07(水) 19:05:29 ]: fermiの場合ボッタクリなだけのような。
最終的にはfusionみたいな物に落ち着くだろうけど。
GPUのアーキテクチャとしては変に汎用に振るより
コンパクトな割に暗号解読みたいに得意なものが速い
と言う方が良いだろう。
843 名前：デフォルトの名無しさん mailto:sage [2010/04/07(水) 20:15:16 ]: >>838

正直エンコにCUDAを適応出来る処理が少なすぎる
844 名前：デフォルトの名無しさん mailto:sage [2010/04/07(水) 21:02:20 ]: x264の開発者が全員RADE愛好者だったとかいうオチなら面白い
845 名前：デフォルトの名無しさん mailto:sage [2010/04/07(水) 21:05:12 ]: >>839
ペガシスがKarmaPlusに導入したCUDAエンコをTXP4になかなか導入
しないのはそれが原因かｗ
846 名前：デフォルトの名無しさん [2010/04/08(木) 00:28:04 ]: >>841
メニーコア化が進展すればGPUは終わるよね
実際のところ、全然進んでないからNVは助かっているけど
847 名前：デフォルトの名無しさん mailto:sage [2010/04/08(木) 00:33:02 ]: メニーコアのメニーの次元が、GPUとCPUじゃ、全然違うしな
848 名前：デフォルトの名無しさん mailto:sage [2010/04/08(木) 00:52:10 ]: 新しいものが出てくるときはチャンスだと思うんだが、このスレではそんな気配かけらもないな。
やっぱり日本人てダメなのかね
849 名前：デフォルトの名無しさん [2010/04/08(木) 00:53:43 ]: だって、HPC向けはコストパフォーマンス悪いし
GeForceは機能削られまくりだし、いまいちなんだよね
850 名前：デフォルトの名無しさん mailto:sage [2010/04/08(木) 00:56:38 ]: ATOM+IONチップセットで
エンコ爆速になったりしないか
851 名前：デフォルトの名無しさん mailto:sage [2010/04/08(木) 01:25:40 ]: >>850
動画データの転送には最低でもPenDは必要。

>>848
バカンスの概念が無い日本人はここぞって時に余力が無い。

>>841
4万のマザーに32,480円の8コア載せるより
5千円のマザーに1.6万円の4コア載せて数万のビデオカード挿す方が数倍早い
んだろうし、どっちもムーアの法則どうり進化すればGPUは当分優勢では？
852 名前：デフォルトの名無しさん [2010/04/08(木) 06:07:48 ]: GPUで縁故するとなんであんなに汚いの?
853 名前：デフォルトの名無しさん mailto:sage [2010/04/08(木) 06:24:22 ]: ソースみたいとわからん。
854 名前：デフォルトの名無しさん mailto:sage [2010/04/08(木) 06:27:58 ]: SpursEngineでエンコして汚いのはハードのせい
CUDAでエンコして汚い場合はソフトが成熟してないから
855 名前：デフォルトの名無しさん mailto:sage [2010/04/08(木) 07:03:51 ]: >>851
適材適所ってことだよなー
y=a*x+b を100万個×10万回　みたいな計算には巨大コアはいらない、
小さいコアがたくさんある方が速い　みたいな。
856 名前：デフォルトの名無しさん mailto:sage [2010/04/08(木) 19:19:24 ]: >>851
なぜ今頃ムーアの法則？
成り立たなくなってかなり経つんだが
857 名前：デフォルトの名無しさん mailto:sage [2010/04/08(木) 19:46:55 ]: え？
858 名前：デフォルトの名無しさん mailto:sage [2010/04/08(木) 19:53:19 ]: >>856

え？
859 名前：デフォルトの名無しさん mailto:sage [2010/04/08(木) 21:00:18 ]: >>856
　　　　　　　　 _......_
　　　 __　　／.::::::::::｀:.､
　　／, - ｒ, /::::i::::ﾊ:i:::::;::',
　 r-'ヽ./ｲ i::::|ｲ/'　'-ﾍl:::i
　｀ｰ､i　{ l::::l '"´　￣ l:::l　＜またまたｗご冗談を
　　　　l!　l l::::ﾄ､ r_っ ,ィ:::l
　　　 l　ﾄヽ::l弋ニ＜l::::l!
　　　　ゝ- イ`　　　ｲ^ｲ
　　　　　　　 | 　　　/-{′

wiredvision.jp/news/200709/2007092021.html
www.itmedia.co.jp/news/articles/0906/18/news006.html
860 名前：デフォルトの名無しさん mailto:sage [2010/04/08(木) 21:01:30 ]: あと10年ぐらいはどうにかなりそうな気がするけど
861 名前：デフォルトの名無しさん mailto:sage [2010/04/08(木) 22:27:03 ]: >>856はfreelunchは終わった発言と混同しているに1000ペソ
862 名前：デフォルトの名無しさん mailto:sage [2010/04/08(木) 22:27:35 ]: >>860
どんだけ低クロックなんだよ、電算機系分野の１０年ってものすげぇ進歩するんだぞ
863 名前：デフォルトの名無しさん mailto:sage [2010/04/08(木) 22:58:17 ]: クロック？
ムーアの法則って集積密度の話じゃなかったっけ
864 名前：デフォルトの名無しさん mailto:sage [2010/04/08(木) 23:03:27 ]: ムーアの法則なんて、明らかに無理だと分かった時点で
定義の方を変えて無理やり存続させているだけだろ。
865 名前：デフォルトの名無しさん mailto:sage [2010/04/08(木) 23:14:23 ]: IntelはAMDの様子見して出し惜しみ＆殿様商売
「半導体の集積密度は18～24ヶ月で倍増する」ゆえ
CPUの性能は2年で倍近くになる。1年だと√2倍弱
GPUも同様。
866 名前：デフォルトの名無しさん mailto:sage [2010/04/08(木) 23:19:44 ]: 集積密度≠性能ということをきちんと理解しましょう
867 名前：デフォルトの名無しさん mailto:sage [2010/04/08(木) 23:53:15 ]: GPUで無理矢理あれこれするより、
計算専用のユニットを別途開発した方がいいんじゃないの？と思う
GPUより効率よくできる部分もあるだろうし
868 名前：デフォルトの名無しさん mailto:sage [2010/04/09(金) 10:42:33 ]: 　　∩＿＿＿∩　　　　　　　　｜
　　 | ノ＼　　　　ヽ　　　　　　　|
　　/　　●゛　　● |　　　　　　　|
　 |　∪　　( _●_)　ミ　　　　　　j
　彡､　　　|∪|　　 |　　　　　　Ｊ>>867
/　　　　 ∩ノ ⊃　ヽ
(　＼　／＿ノ　|　 |
.＼　“　　／＿＿|　 |
　　＼／＿＿＿／
869 名前：デフォルトの名無しさん [2010/04/09(金) 12:48:00 ]: x264がCUDA対応してないことと、AVIUTLでCUDA使えないこと、
MediaCoderのcudaH264Enc.exeの画質をx264並みにするにはビットレートを
何割増しにすればいいのか不明なこと、
PowerDirectorは設定がゴミなこと、

が問題かな
870 名前：デフォルトの名無しさん mailto:sage [2010/04/09(金) 15:40:20 ]: forums.nvidia.com/index.php?showtopic=165055

Double precision is 1/2 of single precision for Tesla 20-series, whereas double precision
is 1/8th of single precision for GeForce GTX 470/480
871 名前：デフォルトの名無しさん [2010/04/09(金) 17:18:50 ]: 倍精度だとteslaの1/4だけと、値段は1/5なんだよね
872 名前：デフォルトの名無しさん mailto:sage [2010/04/09(金) 17:46:49 ]: ノードの数がGTX480は480でteslaは442なんだよね
873 名前：デフォルトの名無しさん mailto:sage [2010/04/09(金) 18:47:55 ]: あれ、HD5870でよくね？
874 名前：デフォルトの名無しさん mailto:sage [2010/04/09(金) 20:10:18 ]: pc.watch.impress.co.jp/docs/news/20100405_359261.html

AMDは単精度しかなかったみたいね
875 名前：デフォルトの名無しさん mailto:sage [2010/04/09(金) 20:58:04 ]: IEEE754準拠ではない64bit double floatはR7xxの頃からあった。
俺は使ったことないからわからないけど。
少なくともドキュメントには、Radeon HD4xxx を除外するような文言はなかった。

R8xxはIEEE754準拠の命令が結構揃ってるよ。
876 名前：デフォルトの名無しさん mailto:sage [2010/04/09(金) 21:12:53 ]: なんでラデはFFTのライブラリを出さないのか理由がわかりますか
877 名前：デフォルトの名無しさん mailto:sage [2010/04/09(金) 21:15:33 ]: つくれば？
878 名前：デフォルトの名無しさん mailto:sage [2010/04/09(金) 21:23:09 ]: ですよね
879 名前：デフォルトの名無しさん mailto:sage [2010/04/09(金) 22:31:50 ]: ascii.jp/elem/000/000/513/513385/
秋葉に出たみたいです。どうしようかな。
880 名前：デフォルトの名無しさん mailto:sage [2010/04/09(金) 22:38:50 ]: 熱的にやめとけ。
881 名前：デフォルトの名無しさん mailto:sage [2010/04/09(金) 22:53:30 ]: 売り切れたそうです。GTX480、3枚買った方がいるとか。
882 名前：デフォルトの名無しさん mailto:sage [2010/04/09(金) 23:04:45 ]: >>875
IEEE754準拠の精度になったのがRV770で
RV670のころからdoubleはサポートされているよ。
883 名前：デフォルトの名無しさん mailto:sage [2010/04/09(金) 23:45:38 ]: 3枚とかアホとしかww
884 名前：デフォルトの名無しさん mailto:sage [2010/04/10(土) 00:08:11 ]: 470なら3台行けるかなー
885 名前：デフォルトの名無しさん [2010/04/10(土) 01:56:40 ]: オークションで売るんじゃないか
886 名前：デフォルトの名無しさん mailto:sage [2010/04/10(土) 02:23:45 ]: PCI-EXｘ16　3つあるマザーでPCIと交互にあるマザーでCUDAエンコに
使うんだろう。
887 名前：デフォルトの名無しさん mailto:sage [2010/04/10(土) 05:14:47 ]: 結局、倍精度は削られているのかね？
削られていなければ買いたいのだが。
888 名前：デフォルトの名無しさん mailto:sage [2010/04/10(土) 06:32:49 ]: まだ実測した人はいないみたいですね。スレの人が手に入れるのをまちましょう。
889 名前：デフォルトの名無しさん mailto:sage [2010/04/10(土) 07:09:42 ]: GTX480じゃ倍精度無効になっているだけで
sandraのベンチのやつはEMUですが
890 名前：デフォルトの名無しさん mailto:sage [2010/04/10(土) 07:42:43 ]: 無効じゃなくて1/4ですよね
ベンチはOpenGLがそうなってるということで
891 名前：デフォルトの名無しさん mailto:sage [2010/04/10(土) 07:47:12 ]: OpenGLのなにが？
892 名前：デフォルトの名無しさん [2010/04/10(土) 07:59:10 ]: sandraのあれはOpenCLだろ>倍精度無効
893 名前：デフォルトの名無しさん mailto:sage [2010/04/10(土) 08:04:30 ]: sandraはOpenCLだろうがGLだろうがCUDAだろうがCSだろうがvideorenderingだろうが
ハードで使えなきゃエミュでだすよ
894 名前：デフォルトの名無しさん mailto:sage [2010/04/10(土) 08:14:32 ]: pc11.2ch.net/test/read.cgi/jisaku/1269857880/740
購入者
895 名前：デフォルトの名無しさん [2010/04/10(土) 08:54:56 ]: >>894
ゲーム目的だったらHD5970を買ったほうがよかったのにね
896 名前：デフォルトの名無しさん mailto:sage [2010/04/10(土) 10:53:48 ]: M/Bとか貧弱なんでゲーマでもないみたいでなんで２枚も買ったのか
897 名前：デフォルトの名無しさん mailto:sage [2010/04/10(土) 11:38:52 ]: 480/470の倍精度演算削られたのか…
倍精度演算やる人はぼったくり価格のC2050/2070買えってか
898 名前：デフォルトの名無しさん mailto:sage [2010/04/10(土) 11:45:47 ]: >>897
そのソースはどこ？
899 名前：デフォルトの名無しさん mailto:sage [2010/04/10(土) 13:34:59 ]: >>897
まぁ、このスレ住人が実際にCUDAで試すのを待とうや。
900 名前：デフォルトの名無しさん mailto:sage [2010/04/10(土) 13:57:02 ]: 米アマゾンはまだ発売前になってるね
901 名前：デフォルトの名無しさん mailto:sage [2010/04/11(日) 10:27:23 ]: これだけ出回ってるのに倍精度はおろか単精度の演算を流す人もいないとは
902 名前：デフォルトの名無しさん mailto:sage [2010/04/11(日) 13:02:22 ]: 出来る人はこんな所に来ない、つまりここは無能の衆が集う所だから
903 名前：デフォルトの名無しさん mailto:sage [2010/04/11(日) 13:11:46 ]: みんな科研費で買うんじゃない？５月まで待たないと．
904 名前：デフォルトの名無しさん mailto:sage [2010/04/11(日) 13:37:56 ]: 無駄遣いはやめてもらうようにこのスレのことも仕分け人に伝えとかないといかんね
905 名前：デフォルトの名無しさん [2010/04/11(日) 13:48:15 ]: 科研費が無駄に物価を高騰させてるな
さっさと仕分けされろ
906 名前：デフォルトの名無しさん [2010/04/11(日) 18:10:27 ]: 無駄な科学者・技術者は農業や林業にまわって効率を上げてやってほしい。
907 名前：デフォルトの名無しさん mailto:sage [2010/04/12(月) 06:21:08 ]: 470ポチった。ついでにCUDAの入門書もポチった。
両方到着は14日予定。おまえらよろしく

>>897
倍精度のテストプログラムか何かあれば提示してくれないか？
手元にGeForce系列VGAが一切ないので一切やったことがないんだ
908 名前：デフォルトの名無しさん mailto:sage [2010/04/12(月) 07:42:29 ]: 14日、期待してますよ。
909 名前：デフォルトの名無しさん mailto:sage [2010/04/12(月) 09:50:33 ]: 人柱になってくれるのに情報薄くて申し訳ないが
このスレの上の方にあるCUDA公式での行列かけ算コードを倍精度化したものを走らせるのが一番簡単かな？
手持ちがあれば提供したいが
910 名前：時々書いている人 mailto:sage [2010/04/12(月) 13:40:47 ]: 未だELSAからボードが来ない……
911 名前：デフォルトの名無しさん [2010/04/12(月) 14:48:20 ]: 480のCUDAのベンチマーク結果はどかでみられませんか？
ゲームのベンチだといくらでもみつかるのですが。
912 名前：デフォルトの名無しさん mailto:sage [2010/04/12(月) 15:15:57 ]: 　　　　　　　　　　　　　_(こ＾)､_
　　　　　　　　　　　〃､__ﾉノ､__,ヽ
　　　　　　　　　 {.っ＞　＜っﾄ､
　　　　　　　 (⌒i　（千于`ｰ┴'─────┐
　　　　　　　　　（O人　 `ー|　　　　　　　　 |
　　　　　　　　　　　/⌒ヽ(＾う　見せられ.　　　 |
　　　　　　　　　　　`ｧー─イ　　　　ないよ！　 |
　　　　　　　　　　　/　　（0::|＿＿＿＿＿_＿_＿_|
　　　　　　　　 /＼＿＿＿_/
　　　　　　　　／　　　/　　⌒ヽ
　　　　　　___／　／￣￣｀)　　ﾉ
　　　　　(__r＿__ノ　　　　 (.＿＿つ
913 名前：デフォルトの名無しさん mailto:sage [2010/04/12(月) 18:54:31 ]: www.anandtech.com/show/2977/nvidia-s-geforce-gtx-480-and-gtx-470-6-months-late-was-it-worth-the-wait-/6
914 名前：デフォルトの名無しさん mailto:sage [2010/04/12(月) 19:10:50 ]: おｋ、14日・・・はWin7のセットアップなどもあるから無理として、15日か16日にはテストする
915 名前：デフォルトの名無しさん mailto:sage [2010/04/12(月) 19:23:43 ]: あいかわらず倍精度のベンチはないもののNVIDIAがコンフォームしたって書いてあるね
916 名前：デフォルトの名無しさん mailto:sage [2010/04/13(火) 08:17:19 ]: マーケティングの理由から倍精度つぶすとかやってくれるぜ全く
くあどろも同じだったら本当に萎えるわ
917 名前：デフォルトの名無しさん mailto:sage [2010/04/13(火) 08:20:01 ]: やっぱり倍精度は1/4になっているのか。
Teslaだと高いしな。
Nehalemが6コア、8コアになってきて、Opteronは12コアになってきてしまったので、
CUDAの優位性がかなり下がってしまったな。
918 名前：デフォルトの名無しさん mailto:sage [2010/04/13(火) 08:27:05 ]: nvidiaは1世代分戦略を間違えた希ガス
社運をHPCに賭けるなら、Fermiは倍精度つぶさずにバーゲンするべきだった
CUDAがある程度スタンダードになったことを確認した上で、Fermiの次をぼったくり価格にすれば良かったはず
PCIExpress3.0対応にしてさ

いまならCUDA捨てるの間に合うしなあ　ユーザ側が

リネーム商法といい、nvidiaはほんと強気だ
919 名前：デフォルトの名無しさん mailto:sage [2010/04/13(火) 10:20:40 ]: まじか。書いてあるな。萎えたな。
「NVIDIA has confirmed it - the GTX 400 series' FP64 performance is capped at
1/8th (12.5%) of its FP32 performance, as opposed to what the hardware
natively can do of 1/2 (50%) FP32. 」
920 名前：デフォルトの名無しさん mailto:sage [2010/04/13(火) 10:21:59 ]: これまでの何倍速くなりました！というのがかなり限定されるのがわかって来たからね。メモリ帯域に関しても確かにCUDAが始まった頃は、CPUの10倍以上あって早かったけど、今はDual CPUだと大差無くなって来ているので、高価なTeslaを入れる意味も無いなあ。
前に誰かが書いていたけど、古いPCの延命のために使うのはありだと思うけど、ちょっと高いよなあ。
歩留まりが悪いのはわかるが。
921 名前：デフォルトの名無しさん mailto:sage [2010/04/13(火) 10:29:35 ]: 結局、Fermiのほどまりは何％くらいで、普通は何％くらいなんですか？
922 名前：デフォルトの名無しさん [2010/04/13(火) 11:55:56 ]: >>921
不明
今後も確実なソースから数値が出る可能性はほぼゼロ
923 名前：デフォルトの名無しさん mailto:sage [2010/04/13(火) 13:23:11 ]: konozamaだったよ

orz
924 名前：デフォルトの名無しさん mailto:sage [2010/04/13(火) 13:32:21 ]: >>923
ｲ㌔
925 名前：デフォルトの名無しさん mailto:sage [2010/04/13(火) 16:23:45 ]: >>921
ほどまり？？？
926 名前：822 [2010/04/13(火) 17:12:13 ]: MediaCoder
CUDAエンコ爆速ｗｗｗ
Q6600の6倍早いｗｗｗ
画質x264と変わらんｗｗｗ
GTX260でCPU45～51％使用ｗｗｗ
927 名前：デフォルトの名無しさん mailto:sage [2010/04/13(火) 17:29:23 ]: ふどまり
928 名前：デフォルトの名無しさん mailto:sage [2010/04/14(水) 12:24:12 ]: どうやらこのスレが世界初GTX480/470CUDA倍精度性能実測報告スレとなりそうですね
929 名前：907 mailto:sage [2010/04/15(木) 04:57:37 ]: とりあえず470は明日到着予定。
けど、搭載予定のケースが明後日到着予定・・・
うきいいいいい

変な時間に起きてすることがない・・・
930 名前：デフォルトの名無しさん mailto:sage [2010/04/15(木) 05:48:41 ]: 瞑想しろ

[ 新着レスの取得/表示 (agate) ] / [ 携帯版 ]

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