1 名前:デフォルトの名無しさん mailto:sage [2012/09/23(日) 23:17:47.58 ] このスレッドは、他のスレッドでは書き込めない超低レベル、 もしくは質問者自身何が何だが分からない質問を勇気を持って書き込むスレッドです。 CUDA使いが優しくコメントを返しますが、 お礼はCUDAの布教と初心者の救済をお願いします。 CUDA・HomePage ttp://developer.nvidia.com/category/zone/cuda-zone 関連スレ GPGPU#5 ttp://hibari.2ch.net/test/read.cgi/tech/1281876470/ 前スレ 【GPGPU】くだすれCUDAスレ【NVIDIA】 ttp://pc12.2ch.net/test/read.cgi/tech/1206152032/ 【GPGPU】くだすれCUDAスレ pert2【NVIDIA】 ttp://pc12.2ch.net/test/read.cgi/tech/1254997777/ 【GPGPU】くだすれCUDAスレ pert3【NVIDIA】 ttp://hibari.2ch.net/test/read.cgi/tech/1271587710/ 【GPGPU】くだすれCUDAスレ pert4【NVIDIA】 ttp://hibari.2ch.net/test/read.cgi/tech/1291467433/ 【GPGPU】くだすれCUDAスレ part5【NVIDIA】 toro.2ch.net/test/read.cgi/tech/1314104886/
403 名前:デフォルトの名無しさん mailto:sage [2013/07/04(木) NY:AN:NY.AN ] >>400 X79の最新のチップセット使うとUbuntu 12以上じゃないと 動かなくてですね、そいつのデフォのバージョンのgccだと CUDAが対応しないんですわ。
404 名前:デフォルトの名無しさん mailto:sage [2013/07/06(土) NY:AN:NY.AN ] >>403 単にnvccのベースがgcc4.4までだからだろ。 Ubuntuならソフトウェアセンターでインストールしてalternativeで切り替えればいいだけ。 これはCUDAに限らず、インテルコンパイラでも必要。
405 名前:デフォルトの名無しさん mailto:sage [2013/07/06(土) NY:AN:NY.AN ] >>404 はいはい情弱ですみませんねえ。 みんながみんなGetting Startedだけ読んでインストールできたら 「Ubuntu 12.**でCUDA 5.0入れてみた」系のブログを書く人も読む人いないですよ。 ふーんだ。
406 名前:デフォルトの名無しさん mailto:sage [2013/07/06(土) NY:AN:NY.AN ] CUDA 5.0がだめなら、5.5 RCを試せばいいじゃない。
407 名前:デフォルトの名無しさん [2013/07/07(日) NY:AN:NY.AN ] cuda5.0のgccは4.6だろ? それよりnVidiaはFedora16のサポートが切れてることについてどう思ってるんだろう。
408 名前:デフォルトの名無しさん mailto:sage [2013/07/07(日) NY:AN:NY.AN ] Ubuntuにせよ、FedoraにせよNVIDIAは最近Linuxに対してあんまりやる気ないな。 リーナスに中指立てられて批判されたからかな?
409 名前:デフォルトの名無しさん mailto:sage [2013/07/08(月) NY:AN:NY.AN ] windowsが最高の開発環境だし
410 名前:デフォルトの名無しさん mailto:sage [2013/07/13(土) NY:AN:NY.AN ] GPGPUはAMDになってしまったから、Nvはやる気でないだろ
411 名前:デフォルトの名無しさん mailto:sage [2013/07/14(日) NY:AN:NY.AN ] はい?
412 名前:デフォルトの名無しさん mailto:sage [2013/07/15(月) NY:AN:NY.AN ] 適切なスレが分からなかったので、ここで質問します。 今のCUDAはCUDA CとOpenCLでバックエンドが共通になっていると聞きましたが、 今もしくは将来のCUDAで、HSA(Heterogeneous System Architecture)を 共通のバックエンドで動かすことは技術的に可能ですか?
413 名前:デフォルトの名無しさん mailto:sage [2013/07/17(水) NY:AN:NY.AN ] nvidiaに聞け 公開資料にない事の予定問われても スレの住人はnvidia関係者な訳じゃないし 関係者が居たとしても、2chで非公開の予定情報の可否なんか答える訳ないだろ
414 名前:デフォルトの名無しさん mailto:sage [2013/07/17(水) NY:AN:NY.AN ] メジャーになればそれなりの対応もあるだろうが 影も形もないものを・・
415 名前:デフォルトの名無しさん mailto:sage [2013/07/17(水) NY:AN:NY.AN ] 技術的に可能かどうかと言われれば可能でしょ。 メモリ空間が共通化されれば、GPUの演算器がCPUのSIMD演算器のように扱えるわけだし。 ただCUDAである必要があるかどうかはNVIDIAが判断するんじゃないか?
416 名前:デフォルトの名無しさん mailto:sage [2013/07/29(月) NY:AN:NY.AN ] N×1行列とM×N行列を計算して結果をテクスチャに書き込むという単純な処理で これを合計512スレッド(Mに関して並列化)で実行しているんだけど(N=3000 M=512) 各ブロックを16×16スレッドの2ブロックよりも 各ブロックは16×1スレッドの32ブロックのほうが2〜3%速度が速いという不可解な結果が出てしまっている 何でこんなことが起こるんだろう
417 名前:デフォルトの名無しさん mailto:sage [2013/07/30(火) NY:AN:NY.AN ] 適度に粒度下がってスケジューラの効率上がったとか?
418 名前:デフォルトの名無しさん mailto:sage [2013/07/30(火) NY:AN:NY.AN ] local memoryにレジスタが溢れているとか
419 名前:デフォルトの名無しさん mailto:sage [2013/07/30(火) NY:AN:NY.AN ] こういう予測しにくい挙動こそがGPGPUのクソなところ。
420 名前:デフォルトの名無しさん mailto:sage [2013/07/30(火) NY:AN:NY.AN ] そういえばNVIDIAはPGIを買収したらしいね ここに書くことでもないかも知れないが
421 名前:デフォルトの名無しさん mailto:sage [2013/07/31(水) NY:AN:NY.AN ] >>416 Nを並列化せずに本当に512スレッドしか使ってないんだったら、16warp*32threadより 32warp*16threadの方がわずかに効率が良いというだけの話じゃないのか?分岐やなんかで。
422 名前:デフォルトの名無しさん [2013/07/31(水) NY:AN:NY.AN ] >>416 メモリの配置次第で不可解でも何でもないと思うが。
423 名前:デフォルトの名無しさん [2013/08/07(水) NY:AN:NY.AN ] 【AMD涙目】デファクトスタンダードへの道を突き進むCUDA・・IBMやGoogleも陣営に加わる engawa.2ch.net/test/read.cgi/poverty/1375883942/
424 名前:デフォルトの名無しさん mailto:sage [2013/08/08(木) NY:AN:NY.AN ] デファクトスタンダードへの道を突き進むCUDA IBMやGoogleも陣営に加わる kohada.2ch.net/test/read.cgi/pcnews/1375893221/
425 名前:デフォルトの名無しさん mailto:sage [2013/08/08(木) NY:AN:NY.AN ] 性能の低さを政治力でカバーするのか?
426 名前:デフォルトの名無しさん mailto:sage [2013/08/08(木) NY:AN:NY.AN ] まぁニュースサイトへの宣伝広告費でカバーしてるところより良いと思うよ
427 名前:デフォルトの名無しさん mailto:sage [2013/08/08(木) NY:AN:NY.AN ] カーネル関数内で,乱数が生成されたD_c配列を用いて計算したくて, D_c配列のポインタを渡していくのが面倒なので以下としたけど,上手くいかない. __device__ double D_c[is_110][is_110][is_110]; curandGenerateUniformDouble(generator, (double*)&D_c, count); cutilSafeCall(cudaMemcpyFromSymbol(H_a, D_c, size_a)); //ここで30:Unknown error cudaMallocの場合には動くから,乱数生成場所がおかしいんだと思うけど. でも,curandGenerateの第2引数はdouble*型だけど,&D_cで配列の先頭を示してるから, 実質同じ事だと考えてたんだけど,違うの?
428 名前:デフォルトの名無しさん mailto:sage [2013/08/08(木) NY:AN:NY.AN ] >>427 今手元にないから記憶で書くけど、D_cの型がdouble []じゃないからだと思う。 D_cの定義をdouble D_c[is_110 * is_110 * is_110]にしてみたらどうなる?
429 名前:デフォルトの名無しさん mailto:sage [2013/08/08(木) NY:AN:NY.AN ] >>428 __device__ double D_c[is_110* is_110* is_110]; curandGenerateUniformDouble(generator, D_c, count); にしても,同じエラーだった. (double*)&D_cでdouble*型に変換してるからイケると思ってたんだけど… __device__ double *D_c;でcudamallocしたら乱数生成はできたけど, カーネル関数内でD_c[id_array]がアクセスエラーっぽい. __device__でやるのは無理があるのかなぁ?
430 名前:デフォルトの名無しさん mailto:sage [2013/08/08(木) NY:AN:NY.AN ] >>429 うーん、よく判らんな。 少なくとも、cudaMemcpyFromSymbol()はCのコードとしては特殊(トリッキー)な仕様だから 使い難いんだよね。まるでプリプロセッサマクロのように。 だから私のところではリファクタリングの結果、使うのをやめてしまっている。 変に悩むくらいなら、代替手段を考えたほうがいいかもよ。
431 名前:デフォルトの名無しさん mailto:sage [2013/08/08(木) NY:AN:NY.AN ] >>430 特殊なのかー. アドバイス通り,普通にホストでMallocして,カーネル関数にポインタ渡す方式にします. ありがとう!
432 名前:デフォルトの名無しさん mailto:sage [2013/08/22(木) NY:AN:NY.AN ] ドライバをアップデートすると演算性能上がりますか?落ちますか?
433 名前:デフォルトの名無しさん mailto:sage [2013/08/23(金) NY:AN:NY.AN ] __global__の関数内でグローバルメモリの内容を一気にコピーしたいんだけどそういう方法ってある? 一要素ずつやった方が無難?
434 名前:デフォルトの名無しさん mailto:sage [2013/08/30(金) NY:AN:NY.AN ] >>433 コピーの手前でカーネル切って、新たに __global__ copy(double *a, double *b) { int tid = blockIdx * blockDim.x + threadIdx.x; b[tid] = a[tid] } を実行するのじゃだめ? あるいはダイナミックパラレリズムで何とかなるのかな。
435 名前:デフォルトの名無しさん mailto:sage [2013/08/30(金) NY:AN:NY.AN ] カーネル分けんでもindexの生成を工夫すれば同じことでしょ。 それがDtoDなら。DtoHやHtoDなら設計を見直すべき。
436 名前:デフォルトの名無しさん mailto:sage [2013/08/31(土) NY:AN:NY.AN ] ツールキット5.5プロダクションリリースってどういう意味ですか? ベータ版とかじゃなくて正式に使えるってこと?
437 名前:デフォルトの名無しさん mailto:sage [2013/08/31(土) NY:AN:NY.AN ] curandGenerateUniformって1.0を含むけど1.0を含まない乱数生成ってないの?
438 名前:デフォルトの名無しさん mailto:sage [2013/08/31(土) NY:AN:NY.AN ] それ使う所の式変形で対処できないの?
439 名前:デフォルトの名無しさん mailto:sage [2013/09/03(火) 08:02:37.87 ] >>435 ブロックまたいで同期する必要がある場合は?
440 名前:デフォルトの名無しさん mailto:sage [2013/09/03(火) 13:46:17.67 ] .cubin ファイルをエディターで・・・と 公式スライドにあるのですが、本当ですか?
441 名前:デフォルトの名無しさん mailto:sage [2013/09/04(水) 00:17:50.20 ] cudaErrorInvalidValueってカーネルがエラーはいたんだけどどういう状況でなる?
442 名前:デフォルトの名無しさん mailto:sage [2013/09/04(水) 08:47:55.31 ] >>441 詳細はAPIマニュアル見てねだけど、 ホントにカーネルで出てる? APIの引数を間違ってるんじゃない?
443 名前:デフォルトの名無しさん mailto:sage [2013/09/05(木) 06:10:12.96 ] 実行はできるし結果もそれっぽいけど画面が不安定になるなあ
444 名前:デフォルトの名無しさん mailto:sage [2013/09/08(日) 14:31:18.78 ] >>443 カーネル実行時間が長すぎると画面が真っ暗になる場合がある。 制限時間はレジストリをいじって変えられたはず。
445 名前:デフォルトの名無しさん mailto:sage [2013/09/09(月) 21:57:16.85 ] >>444 レジストリはいじったあとなんだけどそうなる
446 名前:デフォルトの名無しさん mailto:sage [2013/09/10(火) 12:17:17.37 ] >>445 メモリアクセスが間違ってても落ちることあるですよ。
447 名前:デフォルトの名無しさん mailto:sage [2013/09/10(火) 23:29:52.04 ] Geforce GT 530 で CUDAインストーラー(5.5)が「対応デバイスが無い」とかで失敗するので 古いドライバに変えてみたり再起動繰り返したりしたのですが、 developer.nvidia.com/cuda-gpus に、GT530が載ってませんでした・・・。 マシンに「nVidia GEFORCE with CUDA」のシールあるし、GPU対応のソフトも動いているのですが 開発用としては使えないってことでしょうか?
448 名前:デフォルトの名無しさん mailto:sage [2013/09/10(火) 23:41:30.09 ] セットアップがミスってんでしょ。
449 名前:447 mailto:sage [2013/09/11(水) 00:36:33.98 ] セットアップのミスの原因ってなにかありますか? 展開後は高速インストールかカスタムか選ぶだけだし、 どちらを選んでも失敗します・・・。 プログレスバー見てる感じ、Toolkitのインストール中10%くらいでエラーが出ます。 c:\NVIDIA\CUDAへの展開でかなり時間かかりますが、 インストール失敗すると、この下のインストーラー各種全部削除されて 最初からやり直す仕様なんですね・・・。(成功しても消えるのかもしれませんが)
450 名前:デフォルトの名無しさん mailto:sage [2013/09/11(水) 00:52:29.92 ] ちゃんとリリースノートやインストールノート読んでやってないところ あとはほんとに530のせいかどうかを他のグラボ(nvidiaね)に変えて切り分けして原因しぼっていくしかないだろ。 可能性だけなら グラボ、PCパーツの不良 DLしたソフトウエアの不良 システム不良 などたくさんあんだからさ
451 名前:デフォルトの名無しさん mailto:sage [2013/09/11(水) 02:12:53.82 ] 俺も今5.5のToolkit のインストールでずっこけてる。 Windows XP service pack 3 Quadro FX 580 Toolkit と sample以外はカスタムでインストールできたんだけど、 Toolkitが8割ぐらい行ったところで失敗しました、てなる。 インスコディレクトリにいくらかコピーできてるみたいなんだけど、 環境変数なんかは設定されてない。 リリースノート見てもGUI使うかmsiをシェルで実行しろとしか書いてなくて 困ったぽよ
452 名前:447 mailto:sage [2013/09/11(水) 04:56:20.56 ] もう朝だお。。。 Toolkitのインストーラーが失敗するから ログ取ろうとしたらなぜか成功したっぽい?? C:\NVIDIA\CUDA\CUDAToolkit> msiexec /i "NVIDIA (略).msi" /L*v install.log でも、サンプルのtemplateとか開いてビルドしようとしても 「error : The CUDA Toolkit v5.5 directory '' does not exist.〜」てなる。 環境変数(CudaToolkitDir?)が設定されてないのか、VisualStudioよく分かってないのか・・・。 スレチなら他行きますので・・・。
453 名前:447 mailto:sage [2013/09/11(水) 05:42:32.92 ] ここは俺の日記かお・・・。 VS2012でプロジェクト→プロパティ→構成プロパティ→デバッグ→環境 を選んで編集モードへ。 ここで「マクロ>>」をクリックすると設定されたCuda用環境変数もちらほらありますが $(CudaToolkitDir)の値がからっぽでした。 普通に動かせてる方、この辺の弄り方教えてください。 設定するパスはこれですかね?→C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.5
454 名前:デフォルトの名無しさん mailto:sage [2013/09/11(水) 10:14:42.53 ] >>453 うちでは、NVIDIAのGPUが無くてもインストールとビルドできてるよ。(当然、このPCでは実行はできないけど) コンパイラはVS2008 Standard SP1, VS2010 Professional, VS2012 Professional UP3の3つ。 マクロの値は>>453 でok。 OSの環境変数は: CUDA_PATH=C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.5 CUDA_PATH_V5_5=C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.5 NVCUDASAMPLES5_5_ROOT=C:\ProgramData\NVIDIA Corporation\CUDA Samples\v5.5 NVCUDASAMPLES_ROOT=C:\ProgramData\NVIDIA Corporation\CUDA Samples\v5.5 NVTOOLSEXT_PATH=C:\Program Files\NVIDIA Corporation\NvToolsExt\
455 名前:447 mailto:sage [2013/09/11(水) 16:32:58.66 ] >>454 ありがとうございます、template動きました。 他のサンプルも手作業コピーで動きました。 # C:\NVIDIA\CUDA\CUDASamples\の中身を # C:\ProgramData\NVIDIA Corporation\CUDA Samples\v5.5 # にコピーして、\v5.5\Samples_vs2012.sln開いて全部ビルド成功。 VolumeRenderのfpsが2.1〜2.2の貧弱環境ですが ひと通り習得できてきちんと開発できるようになったら 新しいグラフィクスカード買わないとな・・・。
456 名前:447 mailto:sage [2013/09/11(水) 16:47:38.48 ] 追加質問です。 ビルド中に警告が大量に出ますが、手動インストールによるものなのか判断できません。 > \include\math_functions.h : > warning C4819: ファイルは、現在のコード ページ (932) で表示できない文字を > 含んでいます。データの損失を防ぐために、ファイルを Unicode 形式で保存してください。」 > というのが、たくさん出ます。 普通にインストーラーが成功した方も同じようになりますか? エディタで開いたところ、 math_functions.hはUTF-8N、 cuda.hはSJISと表示されました。
457 名前:デフォルトの名無しさん mailto:sage [2013/09/11(水) 17:16:10.44 ] >>456 > > warning C4819: ファイルは、現在のコード ページ (932) で表示できない文字を > > 含んでいます。データの損失を防ぐために、ファイルを Unicode 形式で保存してください。」 > > というのが、たくさん出ます。 > > 普通にインストーラーが成功した方も同じようになりますか? CUDAにかぎらず、海外のコードは、よくその警告がでる。 自分は、プロパティマネージャで、C4819を無効にしちゃった。 あと、自分が書いたコードはUTF-8(BOM無し)にしてる。
458 名前:デフォルトの名無しさん [2013/09/12(木) 01:13:44.56 ] すみません。 まだまだ初めたばかりの初心者なのですが、 質問させてください。 Tesla C2075を使っていて、DeviceQuery.exeでスペックを見たところ Total amount of global memoryが1.2GB 程度しかありませんでした。 仕様では5GB程度あるはずなのですが・・・。 OSが32bit(win7)であることが関係してたりしますか?
459 名前:デフォルトの名無しさん mailto:sage [2013/09/12(木) 02:48:14.69 ] >>458 そうかも。 GPUのメモリはPCのメモリマップ上に割り付けられるけど、32bitでは4GBしかないので。 メモリマップは、Windowsのデバイスマネージャで「表示」→「リソース (種類別)」の「メモリ」で判る。
460 名前:デフォルトの名無しさん mailto:sage [2013/09/12(木) 10:32:10.68 ] 64bitでもVRAM全部を割り付けたりしないよ 多分、互換性の為だろうけど
461 名前:デフォルトの名無しさん mailto:sage [2013/09/12(木) 23:42:16.47 ] Windows 7 service pack 1 64bit Geforce TITAN の構成でも試してみたが5.5インストールできなかった。 Windowsでインストーラ使って5.5入れられた人いますか?
462 名前:デフォルトの名無しさん mailto:sage [2013/09/12(木) 23:54:10.36 ] >>454 さんはするっとインストールできちゃった感じなんでしょうか。
463 名前:デフォルトの名無しさん mailto:sage [2013/09/13(金) 00:17:54.99 ] 連投すまん。 Visual Studioが2008 Express しか入っていないのが原因かな?
464 名前:やんやん ◆yanyan/Pails mailto:sage [2013/09/13(金) 05:10:33.37 ] >>461-463 何の問題もなくインストールできたよ
465 名前:デフォルトの名無しさん mailto:sage [2013/09/13(金) 09:34:15.52 ] >>462 >>463 toolkit 5.0, 5.5 はVS StdかProが必要だよ。 4.0, 3.0 は忘れた。
466 名前:デフォルトの名無しさん mailto:sage [2013/09/13(金) 23:20:29.29 ] おかげさまでどうにか自作ソースを5.5でコンパイルするところまでこぎつけました。 インストーラでToolkitとSample 以外を先にインストールして、 Toolkitを単独でインストール、ファイルのコピーが終わったところでこけるので、 >>454 の手順で環境変数を設定して、 Sampleはインストーラを立ち上げて展開されたものを所定の位置にこぴぺしました。
467 名前:デフォルトの名無しさん mailto:sage [2013/09/14(土) 03:24:37.09 ] CUDAでホストのmain関数(グローバル関数の呼び出し込)で時間計測にclock()ってつかって問題あったっけ?
468 名前:デフォルトの名無しさん mailto:sage [2013/09/14(土) 18:37:57.64 ] >>467 OpenMPと一緒に使うとCPUタイムがでちゃって正確な経過時間を計れない場合があります。 精度が1msecというのも場合によっては足りないかもしれません。
469 名前:デフォルトの名無しさん mailto:sage [2013/09/14(土) 22:44:07.90 ] カーネル内でmallocとかnewとかってできる? できるとしたらどのメモリが確保されるん? グローバル?ローカル?
470 名前:デフォルトの名無しさん [2013/09/14(土) 22:52:42.35 ] グローカル
471 名前:デフォルトの名無しさん mailto:sage [2013/09/15(日) 01:00:15.49 ] ワロタw
472 名前:デフォルトの名無しさん mailto:sage [2013/09/15(日) 05:52:52.68 ] CUDA C Programming Guide の B.18. Dynamic Global Memory Allocation and Operations にcompute capability 2.0以上でカーネル内のmallocが使える みたいな事が書いてあるね。GPU側のグローバルメモリじゃないかな。
473 名前:デフォルトの名無しさん mailto:sage [2013/09/15(日) 10:04:46.08 ] >>472 サンクス! GPUのグローバルメモリのヒープ領域(デフォルト6MB)に確保されるのね
474 名前:デフォルトの名無しさん mailto:sage [2013/09/15(日) 13:48:19.08 ] >>468 CPUの方でも処理(マルチスレッドではなく)してるからそういう方式では問題ないってことでいいのかな?
475 名前:デフォルトの名無しさん mailto:sage [2013/09/15(日) 14:55:48.31 ] >>474 だったらたぶん。大丈夫
476 名前:デフォルトの名無しさん mailto:sage [2013/09/16(月) 19:37:26.22 ] カーネル関数でグローバル変数を使う方法ありますか? 渡したい値がたくさんあるのでカーネル関数の引数がかなり多くなってしまい面倒なのですが
477 名前:デフォルトの名無しさん mailto:sage [2013/09/16(月) 21:00:31.24 ] __device__付けたらいいんですね自己解決しました
478 名前:デフォルトの名無しさん mailto:sage [2013/09/20(金) 03:32:59.30 ] グローバル化するより構造体作って引数減らしたら駄目なん?
479 名前:デフォルトの名無しさん mailto:sage [2013/09/20(金) 22:20:48.09 ] >>476 構造体作ってポインタで渡す。
480 名前:デフォルトの名無しさん mailto:sage [2013/09/21(土) 08:08:30.99 ] 構造体渡しはつまり全メンバーコピーだからね
481 名前:デフォルトの名無しさん [2013/10/03(木) 18:57:43.41 ] cudaするのに、11インチくらいのディスプレイのノートPCってあるの? kakaku.comで調べても、14インチくらいのしかないけど、だれか教えてください。
482 名前:デフォルトの名無しさん mailto:sage [2013/10/03(木) 22:45:25.45 ] なぜそこでノート限定なのか
483 名前:デフォルトの名無しさん mailto:sage [2013/10/04(金) 00:18:17.93 ] ちっこいノートはインテルHDグラフィックスなんでねーの? と適当に書いてみる。
484 名前:デフォルトの名無しさん mailto:sage [2013/10/04(金) 07:25:58.57 ] いい加減、グラボを換装できるノートが欲しいよな
485 名前:!ninja mailto:sage [2013/10/04(金) 07:49:10.28 ] 普通にあるけど割高だよ。
486 名前:デフォルトの名無しさん [2013/10/04(金) 08:42:04.48 ] 以前は、IONの載ったちいさいノートPCでCUDAできたのに・・・
487 名前:デフォルトの名無しさん mailto:sage [2013/10/04(金) 21:53:58.77 ] どこかが展示してた外付けグラボが欲しいんだが、発売する気配なしw
488 名前:デフォルトの名無しさん mailto:sage [2013/10/04(金) 23:18:05.53 ] 外付けグラボは過去に何製品か出たけど あまりにもニッチすぎる、発熱や性能の問題でどれもヒットしなかった
489 名前:487 mailto:sage [2013/10/04(金) 23:20:13.46 ] >>488 そうなのかぁ・・・。 残念。
490 名前:デフォルトの名無しさん mailto:sage [2013/10/05(土) 00:05:23.57 ] くだらないしつもんですが、 1デバイス内のグローバルメモリ内でデータをコピーするのに、 cudaMemcpyDeviceToDeviceを使うのと、 カーネルを書くのではどちらが速いですか?
491 名前:デフォルトの名無しさん mailto:sage [2013/10/07(月) 18:54:19.99 ] 簡単なんだから、実測してみたら? コピー専用のカーネル作ってたらロードで不利だけど、メインの処理に抱き合わせるならきっと速いよ。
492 名前:デフォルトの名無しさん mailto:sage [2013/10/08(火) 10:22:41.50 ] >>481 一台GPU積んだlinux機作ってsshできるようにすると快適だよ
493 名前:デフォルトの名無しさん mailto:sage [2013/10/08(火) 17:50:51.74 ] 最近cudaを触りはじめた者なのですが質問があります 多次元配列(2とか3次元)の中で指定した値に一番近い値を抽出する方法をcudaで処理するにはなにが一番適しているのですか? 色々ggってみて、とりあえずソートして絶対値をとって比較すれば良いのではと思い バイトニックソートなどで試そうとしましたがソースコード付の物はどれも一次元配列の物ばかりでどうすればいいのかわかりません それにその方法をとっても総当たりより速度が余計遅くなるのではないかとも考えましたがどうなのでしょうか? 抽象的な質問になってすみませんが誰か回答をお願いします
494 名前:デフォルトの名無しさん mailto:sage [2013/10/08(火) 18:32:54.64 ] >>493 kd木構築してNNでいいんでないの どう並列化するかは細かい状況による
495 名前:デフォルトの名無しさん mailto:sage [2013/10/08(火) 18:39:54.30 ] 総当りを(並列で)やるしかないと思うよ。あれをreductionって言うのかわからないけど。 たとえばMaxを求めるのにソートは必要なくてO(n)でできるのと同様に、 一番近い値を求める(だけ)ならばソートする必要はないはず。
496 名前:495 mailto:sage [2013/10/08(火) 18:43:15.91 ] ああ、多次元配列は変化せずに指定する値をコロコロ変えるなら最初にソートしておけば その後はO(logn)でいけるね。条件次第。
497 名前:デフォルトの名無しさん mailto:sage [2013/10/08(火) 19:37:01.90 ] >>494-496 返信ありがとうございます とりあえず自分の知識が全く足りてない事が分かったのでもう少し色々調べてみます あと、できれば496あたりの詳しい解説を知りたいのですがO(long)とはなんでしょうか?
498 名前:デフォルトの名無しさん mailto:sage [2013/10/08(火) 19:46:50.79 ] >>497 O(long)じゃなく、O(logn)な ロガリズムnのオーダー
499 名前:495 mailto:sage [2013/10/08(火) 19:56:25.36 ] >>497 O(n)とかについては「ランダウの記号」でググるかウィキペディると説明がみつかるはずです。 (ウィキペディアの記事は「一般的なオーダー」よりも前のところは正直俺はよくわからないけど…)
500 名前:デフォルトの名無しさん mailto:sage [2013/10/09(水) 08:46:32.40 ] >>499 了解です 回答ありがとうございました また分からなくなったら来ます
501 名前:デフォルトの名無しさん mailto:sage [2013/10/09(水) 10:13:27.99 ] その多次元配列は、一次元配列に投影できないのだろうか。
502 名前:デフォルトの名無しさん mailto:sage [2013/10/09(水) 23:03:59.00 ] >>501 (i, j, k)の三次元配列だったら index = k * N^2 + j * N + i で一次元になるね。
503 名前:デフォルトの名無しさん mailto:sage [2013/10/12(土) 23:33:36.04 ] Visual Studioでcompute_20,sm_20,-maxrregcount=128でコンパイルしたら, 16 bytes stack frame, 24 bytes spill stores, 12 bytes spill loads ptxas : info : Used 63 registers, 32 bytes cmem[0], 16 bytes cmem[16] ってスピルしちゃうんですが. どこを確認すればいいんでしょうか
504 名前:デフォルトの名無しさん mailto:sage [2013/10/13(日) 04:07:02.74 ] >>503 2.xと3.0は最大レジスタ63個までだから 処理少なくするかコード見直して不要冗長な部分を削るとか あるいはレジスタ制限がゆるい1.xか3.5にする
505 名前:デフォルトの名無しさん mailto:sage [2013/10/13(日) 18:48:01.71 ] >>504 63個までだったなんて知らなかった…. 18K Shared使ってるのと,三次元グリッドなんでCP3.5にしたら,ちゃんとコンパイル通って動作しました. ありがとうございます!
506 名前:デフォルトの名無しさん [2013/10/14(月) 04:34:33.91 ] カーネル呼び出しの行で実行がストップしてしまうのですが(次の行に進まない)、 原因がわかりません。 しかし、nvccの「-G」オプションを付けると、この事象が発生せず、 普通に動作します。 原因のわかる方、あるいは原因の特定方法がわかる方、 教えて下さい。
507 名前:デフォルトの名無しさん mailto:sage [2013/10/14(月) 10:57:27.12 ] >>506 カーネル内で不正なメモリアドレスにアクセスすると 問答無用で落ちる場合はありますね。 -Gをつけた上で、カーネル実行の後ろに printf("%s\n", cudaGetErrorString(cudaGetLastError()); と入れてみてはどうでしょうか。
508 名前:デフォルトの名無しさん mailto:sage [2013/10/14(月) 16:11:07.38 ] >>506 >>507 さんが言ってるように,まずはエラー原因を拾ってみると良いのでは? 30: Unknownが出そうな気はするけど. デバッグ情報付与するとメモリ確保位置が異なるので, デバッグ版では偶然に計算に影響を与えない領域にアクセスしていたという可能性も. まず思いつくのは,ある配列の要素Nに対してthreadを生成しているならば, そのthread数がNを超えて生成されるので(BlockDims*GridDims>N), N以上のIdxのthreadが配列外アクセスしてるとか. あとは,カーネル関数内でPrintf()して確認していくとか.
509 名前:デフォルトの名無しさん mailto:sage [2013/10/14(月) 16:15:22.22 ] >>506 追記 CudaGetLastError()は,カーネル関数がエラー起こしてても拾わない事があるので, そのまえのカーネルが原因という可能性もありますよ. 前に起動したカーネルがあるなら,その後に主要変数をmemcpyして調べるといいのでは?
510 名前:506 [2013/10/14(月) 16:45:39.95 ] 回答ありがとうございます。 >>507 なるほど、メモリアクセスが怪しいのですね。 確認してみましたが、エラーとななっていない事がわかりました。 「cudaGetLastError()」の代わりに、「cudaPeekAtLastError()」でも 試したのですが、エラーとなりませんでした。 >>508 メモリ確保位置が異なるということで、「-G」を付けずに試してみましたが、 やはりカーネル呼び出しの行で実行がストップしてしまい、 エラー取得の行まで進みません。 (ログを入れて確認しました。) 配列のインデックスのバグの可能性があるのですね。 細かく確認してみます。 >>509 問題の起きているカーネルが、最初に実行されるカーネルになります。 カーネル実行の前では、「cudaMemsetAsync」やら「cudaMemcpyAsync」等を 実行していたりしますが、戻り値を確認しましたが、正常値でした。 (ストリームを使っています。) 非同期関数なので、すぐにはエラーがでないそうですが。 この辺も怪しいですね。確認します。
511 名前:デフォルトの名無しさん mailto:sage [2013/10/14(月) 17:57:20.43 ] >>510 >配列のインデックスのバグの可能性 それはもちろんあるんだけど,ThreadIdxが計算したいNを超えていないか?という事を言いたかったわけで. 例えば8,000個の計算をしたかったときに,1,024Threads/block,8Blockでカーネル起動したとき, 192個余分にThreadを起動しているので,その処理は大丈夫?って意味だったんだけど >カーネル呼び出しの行で実行がストップしてしまい、エラー取得の行まで進みません ホストからカーネル起動したらすぐに次の行にステップして, カーネルの動作は投げっぱなしだったと思うんだけど…違ったっけ? なんか,メモリアクセスエラーよりもカーネル起動すら出来ていない様な気がするので, 簡単なテストカーネル作って起動させて確認するといいのでは?
512 名前:デフォルトの名無しさん mailto:sage [2013/10/14(月) 18:14:52.48 ] 意外と凡ミスでHW制限越えた値で起動させようとしてるとか パラメータの指定のような気がする
513 名前:デフォルトの名無しさん mailto:sage [2013/10/14(月) 18:15:28.26 ] c/パラメータの指定/パラメータの指定ミス
514 名前:506 [2013/10/15(火) 01:01:28.01 ] おかしな動作をしている箇所を特定しました。 (「cuda-gdb」で再現しないと、デバッグに時間がかかりますね。。。) ワープ内で32回ループ(ワープ内スレッド数分のループ)している箇所があり、 その中で1つ1つのスレッドが順番に「とあるシェアードメモリ変数」を 書き換えているのですが、別のスレッドが書き換えたはずの値が、 正しく取れていない動作をしていました。 追跡はこれからですが、その後の処理でその変数を使っているところで、 何かしらカーネルが止まる動作になっているものと思います。 とりあえず、その変数に「volatile」を付けたところ、 カーネルが動く様になりました。 まだまだ、これで直ったかどうかは、これからじっくり試験しないとダメですが、 変な箇所が発見できたので、大きく前進しました。 ありがとうございました。 511>> threadIdx/blockIdxについて確認してみました。 大丈夫でした。 カーネル呼び出しは、通常であればご指摘の通りの動作をしますが、 今回のバグでは、カーネル呼び出しで処理がストップしていました。
515 名前:デフォルトの名無しさん mailto:sage [2013/10/15(火) 08:45:34.52 ] >>514 別のスレッドが書き換えたシェアードメモリの内容を参照しょうとしてるの? 排他処理とかどうなってんの?
516 名前:506 [2013/10/15(火) 13:31:30.83 ] >>515 排他処理は、ワープ内の1番目のワープが代表してロックを取っています。 その後、514で書いた32回ループの処理を行い、 再び1番目のワープが代表してロックを解放しています。
517 名前:506 [2013/10/15(火) 21:40:36.51 ] すいません、 ワープ内の「1番目のワープ」とか書いてますね。 誤記です。 「1番目のスレッド」です。
518 名前:デフォルトの名無しさん mailto:sage [2013/10/19(土) 15:35:37.45 ] あるプログラムでビジュアルプロファイラー使ってみたら、 カーネルの実行時間が数マイクロ秒で、 カーネルの起動とcudaThreadSynchronizeの オーバーヘッドが数百マイクロ秒だた。 FermiとKeplerを比べると、 Keplerの方がカーネル実行時間は短くなっているのに、 オーバーヘッドがでかくなって、トータルで遅くなっている。 カーネルの実行時間が数十から数百ミリ秒のプログラムなら 問題にならないんだけど、カーネルちっこいと効率悪い。 エヌヴィディアさん何とかしてください。
519 名前:デフォルトの名無しさん mailto:sage [2013/10/19(土) 20:38:48.02 ] >>518 それはもうGPUの宿命だね。 カーネル起動やデータ転送のオーバーヘッドがゴミみたいなレベルの ど〜んとデカい並列処理をやらすのがGPU。 一時期、AMDがHSAで細粒度並列処理に振ることも検討していたが、 結局、粗粒度での効率を追求するアーキテクチャに絞ったようだ。
520 名前:デフォルトの名無しさん mailto:sage [2013/10/21(月) 21:24:19.54 ] 将来の解決策の一つがハイブリッドアーキテクチャなんだろうなあ レイテンシコアがオーバヘッドのでかい処理を担当するという
521 名前:493 mailto:sage [2013/10/25(金) 16:43:42.75 ] 以前(>>493 )の質問の続きというか似たような質問なのですが誰か教えてください あれから少し調べて配列から二分探索で値を探索し、バイトニックでソートしようと思いcでプログラムを組みました(>>494 に関してはよくわからなったので止めておきました) そこで思ったのですが1*nの配列ならその方法で出来ても目標であるn*nの配列から複数回の入力に近似した値を抽出するという作業を行うにはどうすればいいのかわかりません 自分が考えた方法としては無理矢理配列を1*nに置き換えて探索するという方法ですがそれは処理時間的に有用な方法なのでしょうか? また根本的に探索法の選出に誤りがあるなどの点があれば教えてください 最後にc言語からcudaを適用する際に現状で理解できたのはソートの配列入れ替えを並列化するという点なのですがそこをプログラムとして書くときに注意すべき点があれば教えてください 理解力の乏しい素人の質問ですみませんが誰か回答をお願いします
522 名前:デフォルトの名無しさん mailto:sage [2013/10/25(金) 17:31:46.90 ] ゲームだと処理された画像がすぐに表示されるのに、 GPGPUだとオーバーヘッドが大きいのは何でなん?
523 名前:デフォルトの名無しさん mailto:sage [2013/10/25(金) 18:06:42.51 ] JITしてるから
524 名前:デフォルトの名無しさん mailto:sage [2013/10/25(金) 19:56:16.96 ] >>522 ゲームの場合は、 初期化時にモデルやテクスチャなどの容量大きめのデータをCPUからGPUのVRAMに転送。 ランタイム時はGPU⇔VRAMの広帯域を利用してでモデルやテクスチャなどの容量大きめのデータをやりとりし、 CPU⇔GPUの間(比較的帯域の狭いPCIE)では行列や描画コマンドといった小容量のデータ転送しか行わない。 また、処理結果はCPUに戻さず、GPUからビデオ出力するだけ。 大して、GPGPUでは、 ランタイム時にもCPUからGPUへ処理対象となるデータを転送するため、 帯域の狭いPCIEに頻繁に大量のデータを流すことになる上、 処理された結果をCPUに戻すために往復分の帯域を必要とする。 結局ゲームなどのグラフィック処理がGPU自体はもちろん、 周辺システムに関しても相性の良いアプリケーションとなっている。
525 名前:デフォルトの名無しさん mailto:sage [2013/10/26(土) 01:04:42.99 ] >>521 何次元配列で各次元の長さはどのぐらいあるん?
526 名前:493 mailto:sage [2013/10/26(土) 03:06:54.12 ] >>525 二次元配列で大きさとしては512*512ぐらいを条件と考えています
527 名前:デフォルトの名無しさん mailto:sage [2013/10/26(土) 11:10:24.64 ] >>521 まずさ、1次元配列で近似値抽出するコード書きなよ。 それで2次元(というか多次元)配列を1次元として処理するイディオムがあるからそれ使って処理する。 そのイディオムを知らないなら入門書買って読みな。
528 名前:デフォルトの名無しさん mailto:sage [2013/10/26(土) 14:36:23.97 ] >>526 CUDAのスレッドの生成はブロックもグリッドも 三次元構造までサポートしている。 「CUDA by Example 汎用GPUプログラミング入門」 に画像(2次元配列)処理の例が載ってる。 今扱ってる問題が二次元だったら、 2次元構造にスレッドを起動すればいいんじゃないかな。 配列サイズもその規模なら比較的容易にプログラムできそう。 配列中の一番近い値の検索って要するに一番近い値を 持っているインデックス(i,j)の検索だから、 まず2次元のグリッドで入力値と配列値の差の絶対値を 全要素に対して計算して、 次に、j方向に差の絶対値とインデックス(i,j)をセットでリダクションして、 各iの中の最適なjを見つける。ここの段階でリダクションされたデータは一次元になる。 つぎにiに関してリダクションすれば入力値に一番近い値を持っている インデックスが見つかるはず。 あと、多次元配列と言えどメモリ上では一次元で管理されているから、 コアレスアクセスになるように気をつける必要がある。 まあこれがパフォーマンス的にいいアルゴリズムなのかは責任持てんけどw
529 名前:デフォルトの名無しさん mailto:sage [2013/10/26(土) 15:11:16.10 ] ふと思ったんだが、C++で (value, i, j)の構造体作って、 そいつの512*512長のvector作って配列の値とインデックス放り込んで valueでsortしてlower_boundとかupper_bound使えば CUDAいらなくね?
530 名前:デフォルトの名無しさん mailto:sage [2013/10/26(土) 15:47:59.74 ] ふと思ったのなら検証してくれ
531 名前:デフォルトの名無しさん mailto:sage [2013/10/26(土) 16:53:08.04 ] 速度は知らんけどできると思うよ。検証は知りたい人がやればいいと思う。 >あれから少し調べて配列から二分探索で値を探索し、 >バイトニックでソートしようと思いcでプログラムを組みました とあるけどc(非CUDA)で組んでちゃんと結果はでたのかな? さすがに、実は課題は各要素がスカラーではなくベクトルで、指定したベクトルに (ノルム的に)一番近いものを抽出したかったなんてオチはないと思うけど。 そしたら問題の根本が変わるし。>>493 で「絶対値をとって比較」の絶対値がふと気になりました。
532 名前:493 mailto:sage [2013/10/26(土) 22:19:06.13 ] >>527 回答ありがとうございます お答えいただいた通りまずは一次元でcudaを適用しその後そのイディオムとやらのやり方を調べてみようと思います
533 名前:493 mailto:sage [2013/10/26(土) 22:23:53.50 ] >>528-531 回答ありがとうございます 皆さんの意見を参考にもう少し自分で考えてみます ありがとうございます!
534 名前:デフォルトの名無しさん mailto:sage [2013/10/26(土) 22:36:50.51 ] そう、自分で考えることが大事だ。 がんばれよ。
535 名前:デフォルトの名無しさん mailto:sage [2013/10/26(土) 22:47:24.53 ] c(非CUDA)で組めてなかったのか
536 名前:デフォルトの名無しさん mailto:sage [2013/10/27(日) 06:17:23.68 ] 非cudaで組まないで、どうやってデバッグするつもりだったんだろう……
537 名前:デフォルトの名無しさん mailto:sage [2013/10/27(日) 22:48:39.68 ] アーキテクチャーが変わるたびに 「既存のプログラムが遅くなった」という悲鳴を聞くのだが、 どうよ?
538 名前:デフォルトの名無しさん mailto:sage [2013/10/27(日) 23:01:38.83 ] どうもこうもない。 各アーキテクチャごとに最適な記述が変わるというだけだ。 そのままのコードで速くなるx86が特殊だと思ったほうがイイ。
539 名前:デフォルトの名無しさん mailto:sage [2013/10/27(日) 23:12:01.87 ] 同じハードウェア・同じソースでSDKのverを上げてビルドすると速度が落ちたというケースも…
540 名前:デフォルトの名無しさん mailto:sage [2013/10/27(日) 23:17:21.49 ] 新しいアーキテクチャーが出てきた時にそれに対応する コードを注ぎ足していければいいのだが、 マクロCUDA_ARCHって使えるの? 定義されてないとコンパイラに怒られるんだけど、 nvccが勝手に定義してくれるもんじゃないの?
541 名前:デフォルトの名無しさん mailto:sage [2013/10/28(月) 22:40:19.72 ] ダイナミックパラレリズムのカーネル起動コストは ホストからの起動よりも軽いですか?
542 名前:デフォルトの名無しさん mailto:sage [2013/10/28(月) 22:42:50.37 ] どうして自分で検証して確認しようとしないのですか?
543 名前:デフォルトの名無しさん mailto:sage [2013/10/29(火) 01:13:24.21 ] 既に誰かが試したことを再度試すのは時間の無駄
544 名前:デフォルトの名無しさん mailto:sage [2013/10/29(火) 01:18:03.10 ] 回答を待つ時間は無駄ではないと
545 名前:デフォルトの名無しさん mailto:sage [2013/10/29(火) 01:55:32.53 ] そいつの試験方法が正しいということを証明する方法を述べよ
546 名前:デフォルトの名無しさん mailto:sage [2013/10/29(火) 08:20:10.29 ] くだすれやんw
547 名前:デフォルトの名無しさん mailto:sage [2013/10/30(水) 01:19:15.18 ] 重いとは何度か聞いているが自分で試したことはないな 今のところD-Pを使う機会がないというのもあるが
548 名前:デフォルトの名無しさん mailto:sage [2013/11/02(土) 11:53:56.20 ] D-P試してみたよ。 とりま、孫カーネルまでネストしちゃうとオーバーヘッドが かなりきつそう。 子世代ぐらいまでにして、うまくグリッドのサイズを調整すれば、 普通にストリーム使う以上に大量のカーネルを同時実行できるから いいかも。
549 名前:デフォルトの名無しさん mailto:sage [2013/11/02(土) 12:00:15.51 ] 今のところD-Pを使える機械を持っていません><
550 名前:デフォルトの名無しさん [2013/11/07(木) 21:03:34.40 ] CUDAは一般用とだとどういう分野が得意なの? フォトショみたいに単純な画像処理だけなん? レイテンシだとかオーバーヘッドみたいな文章を読んでると、 CPUで処理したほうがいいからGPGPUが普及しないんじゃないかと勘ぐってしまうんだが・・
551 名前:デフォルトの名無しさん mailto:sage [2013/11/07(木) 21:05:52.75 ] >>550 GPUのアーキテクチャにうま〜くハマる処理ならイイんだよ。 んで、結局グラフィック処理が一番イイというオチになるw
552 名前:デフォルトの名無しさん mailto:sage [2013/11/07(木) 21:22:35.38 ] TSUBAME2は一般用ではないのか?
553 名前:デフォルトの名無しさん mailto:sage [2013/11/07(木) 21:27:03.92 ] CUDAで囲碁を強くすることはできんのかな
554 名前:デフォルトの名無しさん mailto:sage [2013/11/07(木) 21:54:40.05 ] 浮動小数点演算をたくさん使うならGPGPUが優位になる可能性はある。単精度で済むならなおよし。 整数演算はKeplerからぱっとしなくなってるんだよなぁ。おそらくMaxwellも。
555 名前:デフォルトの名無しさん mailto:sage [2013/11/08(金) 00:34:39.03 ] >>553 囲碁は評価関数がネックになってるからなあ 着手可能手の膨大さに関しては上手くはまれば期待できるけど、盤状態の評価はifだらけのCPU向きなコードにならざるを得ないから無理目に感じる 積み手もないから枝刈り判定が出来なくてDPの出番もなさげ
556 名前:デフォルトの名無しさん mailto:sage [2013/11/08(金) 08:21:22.96 ] モンテカルロ木探索とかどうよ
557 名前:デフォルトの名無しさん mailto:sage [2013/11/09(土) 00:13:30.26 ] 木探索とか一番苦手なんじゃないのか あとメモリ大量に使う処理も苦手
558 名前:デフォルトの名無しさん mailto:sage [2013/11/09(土) 00:59:47.96 ] 上界/下界と探索済/未探索空間の共有以外は依存性がないから 並列計算に向かないわけじゃないと思うがなぁ。
559 名前:デフォルトの名無しさん mailto:sage [2013/11/09(土) 01:51:28.08 ] 並列計算にむかないんじゃなくてCUDAにむかないということだろう。 ワープ内のスレッドがあるifーelse文でそれぞれちがうステートメントを実行するような場合、 両方実行されるためにペナルティが発生するから。
560 名前:デフォルトの名無しさん mailto:sage [2013/11/09(土) 13:07:58.43 ] 完全な単純全手探索の速度ならCPUぶっちぎるだろうけど、10の360乗ともなるとたとえ千倍速かろうが焼け石に水だなw
561 名前:デフォルトの名無しさん mailto:sage [2013/11/09(土) 13:24:22.72 ] test
562 名前:デフォルトの名無しさん mailto:sage [2013/11/09(土) 15:03:01.44 ] 780tiってCuda5.5フルに使える? 公式いってもCuda対応としかかいてなかったんだけど
563 名前:デフォルトの名無しさん mailto:sage [2013/11/09(土) 17:37:43.86 ] >>562 Tesla向け機能以外は使えるよ
564 名前:デフォルトの名無しさん mailto:sage [2013/11/09(土) 17:42:40.41 ] >>557 モンテカルロ木は普通の木探索とは違うよ
565 名前:デフォルトの名無しさん mailto:sage [2013/11/09(土) 19:06:37.72 ] >>563 ダイナミックなんちゃらは不可?
566 名前:デフォルトの名無しさん mailto:sage [2013/11/09(土) 22:40:17.21 ] >>565 可
567 名前:デフォルトの名無しさん mailto:sage [2013/11/09(土) 23:23:01.04 ] 現役CUDA使いの方々,教えて下さい。 久々に趣味で学生時代のコードを触ろうと思って調べ始めたら, Kepler世代になって色々と変わってて,たった3年で浦島状態に… そこでお聞きしたいんですが, チューニングのコツとかでFermi迄と大幅に変わってたりします? 特に,リダクションとか,キャッシュ回りとか, メモリパディングとかのアドバイスを頂きたいです。
568 名前:デフォルトの名無しさん mailto:sage [2013/11/10(日) 13:17:24.64 ] ワープ内のリダクションは共有メモリを介さずシャッフル関数で できるようになりましたね。
569 名前:デフォルトの名無しさん mailto:sage [2013/11/11(月) 06:02:24.91 ] >>557 そこそこメモリバンド幅あるんでねがったか?
570 名前:デフォルトの名無しさん mailto:sage [2013/11/15(金) 08:08:45.50 ] www.anandtech.com/show/7515/nvidia-announces-cuda-6-unified-memory-for-cuda NVIDIA Announces CUDA 6: Unified Memory for CUDA
571 名前:デフォルトの名無しさん mailto:sage [2013/11/17(日) 07:54:34.91 ] Unified Memoryはホストデバイス間の通信コストがなくなるのかね。
572 名前:デフォルトの名無しさん mailto:sage [2013/11/17(日) 16:29:21.30 ] 何その魔法の技術は?
573 名前:デフォルトの名無しさん mailto:sage [2013/11/18(月) 11:27:47.63 ] >>571 news.mynavi.jp/news/2013/11/15/539/ CUDA6で出るみたい。 ソースコードレベルでMemcpyを消せるだけなのか、 実際にPCI Expressの転送が無くなるのかは分かんないね。
574 名前:デフォルトの名無しさん mailto:sage [2013/11/18(月) 17:47:33.99 ] 有るとしたら ・仮想アドレスレベルでのポインタ共有 ・CPUのメモリ内容がGPUから直接参照可能(今でも可能)なだけでなく、 GPUのメモリ内容がCPUから直接参照可能になる。 ・GPUから他のGPUのメモリへの直接参照(GPU-DirectはCPUメモリを介した間接コピー) てなところじゃね。 明示的コピーをしなくてもオンデマンドでデータ転送が 行われるようになるだけで、PCIeの転送は無くならないだろう。 アクセスパターンによっては、明示的コピーをした方が 速いとかは十分にあり得る。
575 名前:デフォルトの名無しさん mailto:sage [2013/11/18(月) 21:04:51.43 ] てかどう考えたらPCIeの転送が無くなるという発想に至るんだ? APUみたいにVRAMとメインメモリが完全に共有されない限り無理だろ。
576 名前:デフォルトの名無しさん mailto:sage [2013/11/19(火) 00:23:27.89 ] Intelがスパコン用かなんかで似たようなことしてなかったっけ?
577 名前:デフォルトの名無しさん mailto:sage [2013/11/19(火) 01:54:00.83 ] x86なAPUを持ってないNVIDIAがPCIe転送を回避する方法はないよな HSA構想を足蹴にできる現実的な新手法でも閃いたんだろうか >>576 そんなんあった?今時IA64でもないだろし、 Phi上で独立したLinux走らせるという荒技っぷりしか記憶にないのだけど あ、IrisのSRAMがCPU共有のキャッシュになってるとか読んだような。あれは何か野心感じたな、よく知らんけど
578 名前:デフォルトの名無しさん mailto:sage [2013/11/19(火) 08:03:32.80 ] 次世代GPUではarm組み込むんでGPU側だけで完結とか
579 名前:デフォルトの名無しさん mailto:sage [2013/11/19(火) 08:56:58.75 ] >>578 ますますソースコードの可搬性が無くなりそうな悪寒。
580 名前:デフォルトの名無しさん mailto:sage [2013/11/19(火) 09:46:43.72 ] >>578 それはそれで、組み込み用に特化しそうだw
581 名前:デフォルトの名無しさん mailto:sage [2013/11/19(火) 09:59:50.15 ] それなんてcell?
582 名前:デフォルトの名無しさん mailto:sage [2013/11/19(火) 10:10:47.36 ] phiに近くなるだけでは www.atmarkit.co.jp/ait/articles/1309/19/news073.html Xeon Phiは、PCI Express(PCIe)バスに接続して使う数値演算用の拡張ボードである。 「コプロセッサ」と呼ばれているが、実際にはボードコンピュータに近い。 ホストとなるPCとは独立したメモリ空間を持ち、機種によって異なるが6G/8G/16Gバイトの独立した主記憶を備える。 専用のLinuxが稼働し、ホストPCとは異なるIPアドレスが割り当てられる。 ホストPCからはSSHなどでログインでき、プログラムはホストPCと独立して稼働する。 このような演算性能を備えるXeon Phiで稼働させるプログラムをIntel C++ Composerで開発する場合、 大きく分けて2つの実行モデルがある。「オフロード実行モデル」と「ネイティブ実行モデル」だ。 いずれのモデルでも、Intel C++ Composerを利用すれば、Xeon Phi向けの特別なコードを記述する必要はない。 変数piの値は、ホストPCとXeon Phiの両方にあらかじめインストールしてあるミドルウェアによって、自動的にやりとりされる。 ソースコード上に、ホストPCとXeon Phiとの間でデータをやりとりするためのコードを記述する必要はない。 もしも、このプログラムを起動したホストPCにXeon Phiが実装されていない場合には、 プラグマで指定したオフロード部分もホストPCで実行される。 つまり、オフロード部分は、Xeon Phi向けのオブジェクトとともにホストPC向けのオブジェクトも生成されており、 実行時にオフロード部分に入る直前でXeon Phiの有無を検知し、 存在すればXeon Phiにオブジェクトを転送して実行、なければホストPCで同等の処理を実行する。
583 名前:デフォルトの名無しさん mailto:sage [2013/11/19(火) 10:21:30.90 ] シリコンだけでその構造とったら意味が変わるでしょ
584 名前:デフォルトの名無しさん mailto:sage [2013/11/19(火) 10:39:06.45 ] ?
585 名前:デフォルトの名無しさん mailto:sage [2013/11/19(火) 22:37:35.50 ] devblogs.nvidia.com/parallelforall/unified-memory-in-cuda-6/ サンプル 簡単な説明あり
586 名前:デフォルトの名無しさん mailto:sage [2013/11/20(水) 08:39:13.50 ] CufftでMKLであるDftiSetValeに該当すものってある?
587 名前:デフォルトの名無しさん mailto:sage [2013/11/20(水) 09:38:20.39 ] DftiSetValeなんてもんはない。
588 名前:デフォルトの名無しさん mailto:sage [2013/11/20(水) 18:02:46.08 ] 単なるミス。 status = DftiSetValue(desc_handle,config_param, config_val) だよ。
589 名前:デフォルトの名無しさん mailto:sage [2013/11/20(水) 23:48:10.48 ] だから該当すものなんてないってば。
590 名前:デフォルトの名無しさん mailto:sage [2013/11/21(木) 00:36:01.29 ] なんだないのか。
591 名前:493 mailto:sage [2013/11/25(月) 15:27:38.56 ] またよくわからない状態になったので再度質問に来ました。 よろしくお願いします。(過去の質問>>493 >>521 ) 以前の状態からいろいろやってみて二次元配列においてCUDAを用いて 数値探索をできるようになったのですがどうも領域の確保というか ブロックやスレッドの使い方が理解できずスレッドの最大数(?)である 512個以上のデータを扱おうとした場合に正しくない結果が出てしまう状態になってしました。 そもそも使っているのがthreadIdx.xだけでそこに代用として blockDim.x*blockIdx.x + threadIdx.xなどを入れてみてやったところ 一定周期でソートされていて全体ではソートされていないという出力になってしました。 まだよくわかったないことが多いですが512以上のデータを扱う場合に どのようブロックなどを扱えばよいのでしょうか?
592 名前:デフォルトの名無しさん mailto:sage [2013/11/25(月) 15:30:37.09 ] スレッドの最大数が512??、
593 名前:493 mailto:sage [2013/11/25(月) 15:42:32.50 ] >>592 すみません言い方が悪かったです。 www.gdep.jp/page/view/253 を見た感じ1ブロックの最大スレッド数が512らしいので たぶん1ブロック分しかとれていないのだろうと思ったので512だと書きました。
594 名前:デフォルトの名無しさん mailto:sage [2013/11/25(月) 22:56:17.42 ] スレッドが1ブロックでのスレッドが512を超えるとスルーされるな。 GTXだからかもしれんが。
595 名前:デフォルトの名無しさん mailto:sage [2013/11/26(火) 00:07:26.34 ] >>593 いや、プログラミングガイド見ろよ。 CCによってその制限はちがうよ
596 名前:デフォルトの名無しさん mailto:sage [2013/11/26(火) 00:37:36.99 ] もでなぜか512を超えるカーネルが起動しないケースがある。 レジスタが足りんのかわからんが。
597 名前:デフォルトの名無しさん mailto:sage [2013/11/26(火) 12:33:17.31 ] >>591 まず blockDim.x*blockIdx.x + threadIdx.x の意味は分かってる? 左の項は自分がどのblockにいるかを計算していて、右は自分がそのblockの何番目のthreadなのかを示してる つまりこれは、CUDAを一次元的に使うなら、それぞれのスレッドに固有の値になるってわけ
598 名前:デフォルトの名無しさん mailto:sage [2013/12/06(金) 06:19:21.24 ] Tesla k40。。。NVIDIAは日本で売る気無いんだな。たかすぎw。
599 名前:デフォルトの名無しさん mailto:sage [2013/12/06(金) 11:13:24.60 ] 最近CUDAに興味持ちました。 識者の方教えてください。 CUDAやる時はGPUは画面表示用と計算に使うGPUの2つが必要になるのですか?
600 名前:デフォルトの名無しさん mailto:sage [2013/12/06(金) 12:08:08.33 ] >>599 ひとつでいいけど、GPGPUの負荷が高いと画面表示がかなり重くなるよ。
601 名前:デフォルトの名無しさん mailto:sage [2013/12/06(金) 12:08:36.62 ] 一つでもできるけど、画面表示が重いと計算速度が落ちるよ。 3dゲームなんか動かしてたら、メモリもろくに使えなくなるし。
602 名前:デフォルトの名無しさん mailto:sage [2013/12/06(金) 16:16:48.41 ] 表示のリアルタイム性が重要ならおなじGPU。 そうでない処理ならば表示のコストなんて屁みたいなもん。 複数GPU使うのは計算自体を分散したくなってからでいい。
603 名前:デフォルトの名無しさん mailto:sage [2013/12/06(金) 22:19:53.38 ] >>601 CUDAで計算しながら3Dゲームでもするのかと思ったw
604 名前:デフォルトの名無しさん mailto:sage [2013/12/07(土) 08:13:35.39 ] 何かのゲームでDirect3Dで通常のゲーム描画しつつ、 CUDAを演算に使ってたぞ。
605 名前:デフォルトの名無しさん mailto:sage [2013/12/07(土) 10:18:58.23 ] CUDA使う3Dゲームなら一つのGPUでCUDAとDirectXのバッファ共有だろうな。
606 名前:デフォルトの名無しさん mailto:sage [2013/12/07(土) 12:32:22.59 ] SDKにCUDAとDirectX同時に使うサンプルいっぱいあるよ
607 名前:599 mailto:sage [2013/12/07(土) 14:57:55.56 ] 皆様ありがとうございました。 早速買いに行ってきます!
608 名前:デフォルトの名無しさん mailto:sage [2013/12/12(木) 06:44:43.27 ] 開発中にOSごとクラッシュしたりディスプレイが落ちたりすることがあるから 面倒なんだけどGPU2付けるのも面倒なんだよな
609 名前:デフォルトの名無しさん mailto:sage [2013/12/12(木) 06:54:23.18 ] >>608 クラウドを使え
610 名前:デフォルトの名無しさん mailto:sage [2013/12/12(木) 23:59:49.49 ] 教えて頂きたいのですが、エンコード等でCUDAを使う場合、 グラフィックカードに搭載されているCUDAコアプロセッサの 搭載数の多い少ないだけで、性能を判断してよいのでしょうか? 例えばGTX Titanのように、倍精度に制限がない場合、 CUDAでの処理性能に違いが出たりしますか?
611 名前:デフォルトの名無しさん mailto:sage [2013/12/13(金) 08:37:16.00 ] いいえ GeForceでは3Dグラフィックスに使われる浮動小数点演算が最も重視され エンコード等に使われる整数演算は重視されません GPUのアーキテクチャにより整数演算性能は大きく左右されるので単純にCUDAコアの多少で決まるわけではありません
612 名前:610 mailto:sage [2013/12/13(金) 22:43:14.72 ] >>611 わかり易く教えて頂いて、ありがとう御座います。
613 名前:デフォルトの名無しさん [2013/12/29(日) 22:40:27.99 ] バイトニックソートに関して質問があります。 カーネル関数にバイトニックソートを実装してみたのですが、複数ブロックを立てた時のソート結果が正しくなりません。 ちなみにソースコードは「CUDA高速GPUプログラミング入門」に掲載されているものを使用しています。 (こちらのHPにも同じものがあります→www.shuwasystem.co.jp/support/7980html/2578.html) 過去ログを読んでみたところ、>>591 さんと似たような症状みたいです。 bitonicSort<<<2, 512>>>(...)みたいな感じで関数呼び出ししているのですが、 何かプログラムに手を加える必要があるのでしょうか?
614 名前:デフォルトの名無しさん mailto:sage [2014/01/04(土) 20:48:34.47 ] >>>613 #define BLOCK_SIZE 256 以外にいじったとこある?
615 名前:613 [2014/01/05(日) 21:21:55.05 ] >>614 レスありがとうございます! 底以外は特に何も手を加えていない状態なんですが、 実行したところ、512スレッド毎にバイトニックソートがかかっているみたいで、全体のソートが出来ていません。 <<<1, 512>>>のように1ブロックしか立てなかった場合はしっかりとソートがかかるのですが、 複数ブロックを立てるとやはり1ブロック中のスレッド毎にしかソートがかからないようです。 何かハードの制約でもあるのでしょうか?
616 名前:デフォルトの名無しさん mailto:sage [2014/01/05(日) 22:02:49.87 ] それは恐らくハードの制約ではなくブロック内でソートすると言うサンプルの制限なのではないだろうか。 つーか、サンプルを理解できないなら使うなよ。
617 名前:613 [2014/01/05(日) 22:41:59.96 ] >>616 レスありがとうございます。 自分なりにいろいろと勉強はしているのですが、サンプルが掲載されている書籍がもうひとつ説明が少なくて苦労してました。 ブロック内でソートするとかいう説明もなく、ひたすらデバッグやトレースを繰り返してはいるものの、 解決策が見つからなかったのでこちらで質問させていただきました。 CUDAのバイトニックソートのプログラムはほとんどが1ブロックしか立てていないものばかりなので、 参考になる情報も少ないので・・・
618 名前:デフォルトの名無しさん mailto:sage [2014/01/05(日) 22:48:01.57 ] >>615 書き換えない状態では256スレッドのブロックが四つという設定の筈ですが、 サンプルコードがそのままでも動かないということですか?
619 名前:デフォルトの名無しさん mailto:sage [2014/01/05(日) 22:51:43.74 ] >>615 というか、 #define SIZE 1024 に相当する部分をいじってないですか? カーネルにも SIZEが使われてるので、 下手に数字をべた書きすると整合性が 取れなくなりそうな希ガスるですよ。
620 名前:613 [2014/01/05(日) 23:06:57.98 ] >>618 迅速なレスありがとうございます。 今サンプルコードのままで実行して結果を調べてみたところ、 ちゃんとソートされたり、>>615 のようになってソートされなかったりします。 ああ、余計に訳が分からなくなってきた・・・
621 名前:613 [2014/01/05(日) 23:09:22.98 ] >>619 レスありがとうございます。 defineの部分はいじってないです。 SIZEはサンプル通りの1024のままですね。
622 名前:デフォルトの名無しさん mailto:sage [2014/01/06(月) 00:30:08.60 ] その本読むよりも、他の資料、例えば CUDA by Example を読んだほうが、基礎が掴めると思うよ。
623 名前:613 [2014/01/06(月) 01:46:43.97 ] >>622 レスありがとうございます。 そうですね、自分の勉強不足もあると思います。 もう少し勉強して頑張ってみます。
624 名前:デフォルトの名無しさん mailto:sage [2014/01/06(月) 01:57:37.95 ] 勉強不足もあると思うけど、理解してから次に進むっていう意識が欠如してるよな blockDim.x*blockIdx.x + threadIdx.x の意味すら理解せずにアルゴリズムがどうこう言ったって理解できるわけないだろ。
625 名前:デフォルトの名無しさん mailto:sage [2014/01/07(火) 09:26:42.38 ] まあくだすれだし。 またーりしようよ。 >>620 まずサンプルコードをいじらずにそのまま使って 何度も実行して、問題が再現するかどうかちぇっくしませう。
626 名前:デフォルトの名無しさん mailto:sage [2014/01/15(水) 08:04:25.37 ] on-demand.gputechconf.com/supercomputing/2013/video/SC3108-New-Features-CUDA%206%20-GPU-Acceleration.mp4
627 名前:デフォルトの名無しさん [2014/01/16(木) 13:45:32.09 ] 関数の呼び出しについて質問があります 回答をお願いします c言語でつくられたプログラムの一部をCUDAに適用しそれを 本体(cppで作られたプログラム)とは別にCUDA用プログラムとして(拡張子cuとして別に作っておいて) 作り本体から呼び出そうとしたところ未解決の外部シンボルとしてエラーが吐かれてしまうのですが どのように処置したらC言語のプログラムからCUDAのプログラムを呼び出すことができるのでしょうか?
628 名前:デフォルトの名無しさん mailto:sage [2014/01/16(木) 14:01:34.82 ] CUDA SDKのサンプルはビルド&実行できた?
629 名前:デフォルトの名無しさん mailto:sage [2014/01/16(木) 14:45:40.63 ] >>627 Cはc++またはg++とかでコンパイルしている? ccで生成した中間生製オブジェクトとはリンクできなかったような
630 名前:627 [2014/01/16(木) 14:54:30.04 ] >>628 回答ありがとうございます サンプルとは「〜\NVIDIA Corporation\CUDA Samples\v5.0」の中にあるサンプルのことでしょうか? もしそのサンプルなら全てではありませんが一部をやってみたところ問題なく動きました。 また外部シンボル等でいろいろ調べてみたところライブラリのリンクが等々と書いていたので 構成プロパティでいくつか追加してみても特にエラーは直りませんでした。
631 名前:627 [2014/01/16(木) 15:01:57.65 ] >>629 回答ありがとうございます コンパイルについては問題ないと思います。 ネットに書いてあったような設定をしてやったところ.cuのプログラム一つだけの場合 問題なくコンパイルされ期待通りの実行結果が出てくれたのでCのプログラム内で CUDAのプログラムの関数を呼び出す際に問題が発生したのではないかと思います。 というかそのやり方がわかりません・・・
632 名前:628 mailto:sage [2014/01/16(木) 15:18:23.50 ] .cuファイルを複数使ってるのかな?…と思ったらCUDA5からそれでもOKっぽいんだね NVIDIA,「CUDA 5」を正式発表。第2世代Kepler「GK110」に向けた準備が整う ニコニコニュース news.nicovideo.jp/watch/nw401167
633 名前:628 mailto:sage [2014/01/16(木) 15:19:46.16 ] cuファイルが1つだけならビルドできるサンプルにコピペすれば どうにかなる(問題の切り分けを始められる)と思ったけど
634 名前:628 mailto:sage [2014/01/16(木) 15:21:38.59 ] 1つ抜けてた、サンプルというのは>>630 の言うそれのことで合ってますです。
635 名前:627 [2014/01/16(木) 15:32:40.44 ] >>632-634 回答ありがとうございます cuファイルは一つだけなので該当サンプルと内容をすり替えれば良いということでしょうか? よろしければそのサンプルはどのサンプルなのかお教えください。 よろしくお願いします
636 名前:628 mailto:sage [2014/01/16(木) 15:39:00.82 ] そうです>内容をすり替えればいい ごめん、今使ってるPCにCUDA SDKを入れてない(非Geforce機)ので何があるか わからんけど初級の短めのやつにすればいいと思う。 あと(少なくともcuda sdk3か4の頃だと)プロジェクトに相対pathか絶対pathかが使われてて、 フォルダの場所を移動させるとそれに合わせて設定変更しない限りビルドか実行かが うまくいかなかった覚えがあるので注意してくださーい
637 名前:デフォルトの名無しさん [2014/01/16(木) 15:45:39.71 ] コペンバローナ「ンーwwwwwwwwwwwwwwwwwwwwwwww」
638 名前:628 mailto:sage [2014/01/16(木) 15:46:41.43 ] ごめん、問題のポイントを誤解してたかも。 (CとC++とcuの入ったサンプルは見た覚えがない) とりあえず俺の発言は忘れてください。ごめんなさい。
639 名前:デフォルトの名無しさん [2014/01/16(木) 15:47:58.43 ] ペッコンバローナーwwwwwwwwwwwwwwwwwwwwwwwww
640 名前:デフォルトの名無しさん mailto:sage [2014/01/16(木) 17:03:38.03 ] >>631 未解決の外部シンボルのエラーって、リンク(コンパイル)の時じゃなくて、実行時に出るの?
641 名前:デフォルトの名無しさん [2014/01/16(木) 18:25:52.58 ] >>628 ペッコンペッコンペッコンバローナーwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwww
642 名前:デフォルトの名無しさん [2014/01/16(木) 18:29:18.81 ] ロ・・・ロバwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwww
643 名前:デフォルトの名無しさん [2014/01/16(木) 18:32:40.24 ] コペンハーゲ「ンーwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwww」
644 名前:デフォルトの名無しさん [2014/01/16(木) 18:33:44.44 ] コペンバロー「ナwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwww」
645 名前:デフォルトの名無しさん [2014/01/16(木) 18:34:43.10 ] バコナロ「バコーンwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwwww」
646 名前:デフォルトの名無しさん [2014/01/16(木) 18:58:08.66 ] 、 ′ 、 ’、 ′ ’ ;. `',. ’ ’、 ′ ’ . ・ 、′・. ’ ; ’、 ’、′‘ .・” ’、′・ ’、.・”; ” ’ ’、 (;;ノ;; (′‘ ・. ’、′”; ’、′・ ( (´;^`⌒)∴⌒`.・ ” ; ′・ , '´`ヽ.-──-,'´`ヽ. 、 ’、 ’・ 、´⌒,;y'⌒((´;;;;;ノ、"'人 / ゙i::::::::::::::゙i ゙: _、(⌒ ;;;:;´'从 ;' ;:;;) ;⌒ ;; :) )、 ミ / ;゙:;::::::;:::::::! ! ( ´;`ヾ,;⌒)´ 从⌒ ;) `⌒ )⌒:`.・ ,.;゙ r'^ー、;゙:;ィ:::ハ::λi,r'ヽ, i ‘: ;゜+° ′、:::::. :::( ::;; ノ ´⌒(,ゞ、⌒) ;;:::)::ノ i::i | i゙ノiノレ レ' ゙!!i | ! ....................`:::、 ノ ...;:;_) ...::ノ ソ ...::ノ ミ ハ::゙、 i ! > <!| !; :::::::::`- ´:::::::::::::::::::::::::::::::::::::::::::::::::: ノ:::λ゙ー| |〃 __ 〃i l' 〈ノ::イ::ノ::::! ! '、 ノ ノ'. ; 〈/!::;イ::/,r'´`゙゙i>‐_-_t.´r゙´ヾ ミ V^レ´i ,.〉`'゙'i/゙`i,_,..ノ ┏━━━━━━━━━━━━━━━━━━━┓ `ーi゙´ /> i ┃ 「キャー♪」 . . ....... ┃ | ´ | ┃ てゐちゃんは きょうも たのしそうだ!! ┃ミ ,.へ. ,ノ , ,. ! ┗━━━━━━━━━━━━━━━━━━━┛ /___゙ニ=-‐´ , / 〉 ゝ. /´ ̄ _ ノ / / ノ 〉
647 名前:デフォルトの名無しさん mailto:sage [2014/01/18(土) 13:44:16.61 ] denver(maxwellも含む)がかなり面白そうだ Instruction-optimizing processor with branch-count table in hardware https://www.google.com/patents/US20130311752 の特許関係者の前歴 Ben Hertzberg - intel Madhu Swarna - intel Ross Segelken - intel Rupert Brauch - ?Hewlett-Packard David Dunn - transmeta
648 名前:デフォルトの名無しさん mailto:sage [2014/01/18(土) 21:26:58.60 ] お、Transmetaの人が関わってるのか。
649 名前:デフォルトの名無しさん mailto:sage [2014/01/19(日) 14:04:30.61 ] プログラム実行中のGPUの温度をモニターしたいのでNVMLを試してみようと思っています。 ここで、nvmlDevice_tとCUdeviceの対応はどのようにとればいいんでしょう? CUDAのデバイスindexとNVMLのindexは必ずしも一致しないという記述はあったのですが、 じゃあどうすればいいのか、というところを見つけられませんでした。
650 名前:デフォルトの名無しさん [2014/01/19(日) 19:36:52.66 ] GPU2枚差して、CPU介さずにデータ共有ってできる?
651 名前:デフォルトの名無しさん mailto:sage [2014/01/19(日) 20:22:54.84 ] Teslaならできるらしい。持ってないんで試したことないが。 https://developer.nvidia.com/gpudirect
652 名前:デフォルトの名無しさん [2014/01/19(日) 21:48:00.18 ] 試してみたら、GeForceでもできました。 ありがとうございます
653 名前:デフォルトの名無しさん mailto:sage [2014/01/19(日) 23:40:42.70 ] >>652 詳細お願いします。
654 名前:デフォルトの名無しさん mailto:sage [2014/01/20(月) 00:48:56.84 ] >>649 普通にインデックスでいいよ 汎用に作るならデバイス数を取得して、それぞれのnvmlDeviceをインデックスで取得して、いろんな情報とればいい
655 名前:デフォルトの名無しさん mailto:sage [2014/01/26(日) 23:20:39.25 ] GeForce,Quadroはメインメモリ→ボードのDMACしか持ってないよね?
656 名前:デフォルトの名無しさん mailto:sage [2014/01/26(日) 23:25:59.06 ] なんでそう思ったのかが気になる。
657 名前:デフォルトの名無しさん mailto:sage [2014/01/30(木) 23:59:53.57 ] NNみたいなモロにメモリ律速な計算だとろくに速度出ないな 帯域80GB/s使って160GTlopsとかになる
658 名前:デフォルトの名無しさん mailto:sage [2014/01/31(金) 04:06:11.75 ] 結局どういう問題なら高速化できるんだ
659 名前:デフォルトの名無しさん mailto:sage [2014/01/31(金) 22:45:44.77 ] メモリへのアクセスが少ない、扱うデータサイズが小さい、分岐がない 最低数万スレッド以上で並列計算可能な問題であること
660 名前:デフォルトの名無しさん mailto:sage [2014/01/31(金) 23:40:00.25 ] メモリ量と計算量が比例する問題しか普段扱ってないんだよなあ 暗号解読とか?
661 名前:デフォルトの名無しさん mailto:sage [2014/02/01(土) 08:00:55.65 ] Geforce GT520(VRAM: DDR3 1GB)でもCore2Duo E4300に比べたらFFTを高速化できるかな?
662 名前:デフォルトの名無しさん mailto:sage [2014/02/01(土) 10:22:49.10 ] マンデルブロが超得意 データ量Nに対して計算量がN^1より大きいオーダーで 増えていくような処理 巨大な元データが必要でも、それ自体は変えずに 少量のパラメータを与えて再計算を繰り返すような処理 しかも結果をグラフィックス表示すればOKな用途
663 名前:デフォルトの名無しさん mailto:sage [2014/02/01(土) 10:55:21.73 ] >>662 >しかも結果をグラフィックス表示すればOKな用途 GPU⇒CPUが入ると途端にスループット落ちることになるもんね・・・。
664 名前:デフォルトの名無しさん mailto:sage [2014/02/01(土) 11:40:24.53 ] 人工ニューラルネットワークなんかは、 データ量N、i段目のニューロン数n_iとすると、 計算量=NΠ_i n_i だから実はあんまり適してねえのか
665 名前:デフォルトの名無しさん mailto:sage [2014/02/01(土) 13:38:05.35 ] GPUもCPUも足回りが全然ついていかないんだよな NvidiaもAMDもFlops値ばかり競ってるけどメモリ帯域はこの数年で1割程度しか増えていない 完全に頭打ちの傾向
666 名前:デフォルトの名無しさん mailto:sage [2014/02/01(土) 14:03:29.95 ] そして効率的な演算とデータアクセスの比率は高まるばかり・・・
667 名前:デフォルトの名無しさん mailto:sage [2014/02/02(日) 14:27:36.24 ] 石の性能が良くなっても仕方ないよな。 プロセッサの性能が無駄になってる。
668 名前:デフォルトの名無しさん mailto:sage [2014/02/02(日) 14:40:31.96 ] まぁ、VoltaでスタックドDRAM使うみたいだから、いくらか改善されるかもね。
669 名前:デフォルトの名無しさん mailto:sage [2014/02/02(日) 14:43:19.21 ] ☆ チン マチクタビレタ〜 マチクタビレタ〜 ☆ チン 〃 ∧_∧ / ̄ ̄ ̄ ̄ ̄ ̄ ̄ ̄ ̄ ̄ ̄ ̄ ̄ ヽ ___\(\・∀・) < データまだ〜? \_/⊂ ⊂_ ) \_____________ / ̄ ̄ ̄ ̄ ̄ ̄ /| | ̄ ̄ ̄ ̄ ̄ ̄ ̄| | | CPU・GPU |/
670 名前:デフォルトの名無しさん mailto:sage [2014/02/02(日) 14:56:23.47 ] GPUはバス幅を狭くすることでコストダウンを図ってるんだから仕方ないな。 それこそ、バス幅求めるならベクトル計算機でも使えと。全レジスタに対して本当の同時操作が出来るぞ。
671 名前:デフォルトの名無しさん mailto:sage [2014/02/02(日) 15:07:43.13 ] >>669 現状を表す最適なAA乙w
672 名前:デフォルトの名無しさん mailto:sage [2014/02/02(日) 15:12:44.27 ] このアンバランスな状態を解消できるのはプロセスルールが物理的限界に到達した後だろうな。
673 名前:デフォルトの名無しさん mailto:sage [2014/02/02(日) 15:20:27.01 ] しかしその頃には光コンピュータが実用化されていたのだった…… 速さが足りない!!
674 名前:デフォルトの名無しさん mailto:sage [2014/02/02(日) 15:40:48.83 ] 俺が遅い・・・ 俺がスロウリィ?!
675 名前:デフォルトの名無しさん [2014/02/03(月) 04:25:07.11 ] HOLY隊員のクーダーです
676 名前:デフォルトの名無しさん mailto:sage [2014/02/03(月) 15:22:21.17 ] FFTぐらいしか応用が思いつかねぇ。
677 名前:デフォルトの名無しさん mailto:sage [2014/02/03(月) 18:10:49.48 ] >>676 FFTに向いているなら自動的に円周率計算もバリバリなはずだが、ググっても 「円周率の小数点以下8000兆桁めをGeForceで求める方法」 (www.4gamer.net/games/120/G012093/20130323002/ ) といった話ぐらいしか出てこねぇ……
678 名前:デフォルトの名無しさん mailto:sage [2014/02/03(月) 18:28:29.72 ] 音声処理におけるFIRフィルタを想定してるぜ・・・。
679 名前:デフォルトの名無しさん mailto:sage [2014/02/03(月) 21:38:24.39 ] 世間が持てはやすのがFLOPS値ばかりだから一向に帯域増える方向にいかんな
680 名前:デフォルトの名無しさん mailto:sage [2014/02/03(月) 22:19:44.58 ] 帯域はコストが高く付くからな。 バランス取ろうと思ったら、途端に価格が跳ね上がる。 一般人じゃ手の届かない価格になるよ。
681 名前:デフォルトの名無しさん mailto:sage [2014/02/03(月) 22:47:08.60 ] >>680 別にHPC用なら値段高くても買う奴いるじゃんか……
682 名前:デフォルトの名無しさん mailto:sage [2014/02/04(火) 07:24:40.49 ] 普及してて値段が安いからGPGPUがもてはやされてるわけでさ。 値段が高くなればベクトル計算機のプロセッサをPC向けに販売して使ったほうが良いって。
683 名前:デフォルトの名無しさん mailto:sage [2014/02/04(火) 21:02:06.21 ] >>669 わらった。 GPGPUの一般用途での最大の問題点はCPU<=>GPU間データ転送。一般用途ではそれを解消したAMDのAPUでHSAする方が良いからな いくらGPUがすごくても、メモリ転送に時間掛かってはお手軽に使えないからな
684 名前:デフォルトの名無しさん mailto:sage [2014/02/04(火) 21:43:04.69 ] kaveri出たらHSA酷使した絶賛ベンチが次々と出てnvidia叩きレスで溢れかえると思ったら思いのほか静かで不思議
685 名前:デフォルトの名無しさん mailto:sage [2014/02/04(火) 23:20:40.89 ] >>683 データ転送せず極力内部で計算するようにしても結局GPU側の帯域で足引っ張られる 780Tiで単精度5.76Tflopsに対して330GB/sだから足回りが70倍も遅い
686 名前:デフォルトの名無しさん mailto:sage [2014/02/04(火) 23:27:45.03 ] >>684 言い出しっぺの法則 >>685 だが待ってほしい 70倍遅いなら70倍転送せずに計算すればトントンではないだろうか
687 名前:デフォルトの名無しさん mailto:sage [2014/02/05(水) 00:00:22.86 ] HSA使ってみたいんだけど、具体的にどうすればいいの? VisualStudioで始められる??
688 名前:デフォルトの名無しさん mailto:sage [2014/02/05(水) 02:10:41.38 ] CPUGPU間の転送が足を引っ張ってるってイメージはないな シェアードメモリやキャッシュ以外のVRAM・GPU間がただただ遅いのだ
689 名前:デフォルトの名無しさん mailto:sage [2014/02/05(水) 11:28:09.07 ] レイテンシ?
690 名前:デフォルトの名無しさん mailto:sage [2014/02/05(水) 17:37:53.61 ] 基本的にI/Oが遅いんだよ。 これが何とかなったらいいけど、何とかするとコストがかさむから一般向けでは無理。
691 名前:デフォルトの名無しさん mailto:sage [2014/02/05(水) 21:02:36.68 ] 一般向=>一般向CUDA用途==スパコン
692 名前:デフォルトの名無しさん mailto:sage [2014/02/05(水) 21:49:04.57 ] なぜそうなる。数十万でも買うのかよ。
693 名前:デフォルトの名無しさん mailto:sage [2014/02/06(木) 10:48:02.20 ] 重ーい超越関数をバリバリ使う計算ならメモリ転送はさほど器にしなくて良いのでは。
694 名前:デフォルトの名無しさん mailto:sage [2014/02/06(木) 11:35:48.47 ] 三角関数がそこそこ速いから最初に三角関数テーブルを作っておいて纏めて計算するんだけど、 キャッシュに乗らないとべらぼうに遅くなるw。
695 名前:デフォルトの名無しさん mailto:sage [2014/02/06(木) 23:35:46.17 ] 今や、テーブルにしてメモリから読み出すよりも、 手前で計算で作ったほうが速いからなw
696 名前:デフォルトの名無しさん mailto:sage [2014/02/06(木) 23:44:44.46 ] 昔「計算が遅いからメモリでなんとかしよう」 今「メモリが遅いから計算でなんとかしよう」 将来「???」
697 名前:デフォルトの名無しさん mailto:sage [2014/02/07(金) 00:13:26.20 ] PS3もちょうどその技術トレンドを読んで企画されたけど、ちょっと早漏すぎたな。
698 名前:デフォルトの名無しさん mailto:sage [2014/02/07(金) 01:03:30.43 ] >>689 VRAMのレイテンシは数百クロックもある上にピーク速度でも計算速度より何百倍も遅い
699 名前:デフォルトの名無しさん mailto:sage [2014/02/07(金) 01:19:34.56 ] >>697 早漏てw そこは先駆者として評価してやっていいんじゃないの。十分出回ったしハード的にもソフト的にも注目されて、長めのゲーム機サイクルの中で研究されたんだからアーキテクチャとしては幸せな方でしょ ソニーさんのビジネス的にどうだったのかは知らんけど
700 名前:デフォルトの名無しさん mailto:sage [2014/02/07(金) 02:55:08.39 ] >>698 え、マジで? >>699 さすがに逆ザヤはNG
701 名前:デフォルトの名無しさん mailto:sage [2014/02/07(金) 08:04:52.20 ] サブプロセッサの性能は兎も角、メイン側が遅過ぎ。 メインとサブの間のメモリ空間も狭いし。 あれでよくゲームに活かせたと思うよ。
702 名前:デフォルトの名無しさん mailto:sage [2014/02/07(金) 23:50:49.70 ] >>694 三角関数テーブルって精度的にはどうなん? 多項式補間とかするの?
703 名前:デフォルトの名無しさん mailto:sage [2014/02/07(金) 23:54:19.34 ] 用途によるだろう
704 名前:デフォルトの名無しさん mailto:sage [2014/02/07(金) 23:57:35.92 ] 多項式補間といっても奥が深くてだな…… 単なるテイラー展開(途中打ち切り)とよく練られた多項式との差はダンチ 例: cosx≒1-0.5x^2+0.04166666667x^4-0.00138888889x^6とすると 誤差はx=±1までで2.4528×10^-5(テイラー展開)。しかし、 cosx≒0.99999981155-0.49999395279x^2 +0.04166666667x^4-0.00138888889x^6とすると 誤差はx=±1までで2.4528×10^-5(テイラー展開)。しかし、
705 名前:704 mailto:sage [2014/02/08(土) 00:01:16.23 ] 途中送信してしまったorz テイラー展開→cosx≒1-0.5x^2+0.04166666667x^4-0.00138888889x^6で、x=±1までの最大誤差2.4528×10^-5 最良近似式→cosx≒0.99999981155-0.49999395279x^2+0.04163632912x^4-0.00134007047x^6で、x=±1までの最大誤差1.8845×10^-7 (出典:www.amazon.co.jp/dp/456301382X )
706 名前:デフォルトの名無しさん mailto:sage [2014/02/08(土) 01:28:11.34 ] >>702 私(>694)のところで使うのは周波数空間像の畳み込みだから、三角関数の引き数は格子上の点の距離。 なので、補間の必要もないの。ついでに、cufft相当も自前で実装した。
707 名前:デフォルトの名無しさん mailto:sage [2014/02/08(土) 04:18:08.80 ] テイラー展開とか教科書に載ってるだけで、 関数近似の方法としては、ほぼ実用されてねえよ
708 名前:デフォルトの名無しさん mailto:sage [2014/02/08(土) 04:33:47.29 ] >>705 URLが見つかりません
709 名前:デフォルトの名無しさん mailto:sage [2014/02/08(土) 11:00:08.81 ] 最大誤差が小さくても、cos(0)が0.99999981155になる関数なんて使いたくないな。 0みたいな重要点でおかしな値が出ると致命傷になることが多い。
710 名前:デフォルトの名無しさん mailto:sage [2014/02/08(土) 11:17:17.06 ] >>707 テイラー展開の誤差範囲の理論値が明確であるメリットは結構大きい
711 名前:デフォルトの名無しさん mailto:sage [2014/02/08(土) 12:45:17.63 ] >>708 URLの最後の)がいらない。
712 名前:デフォルトの名無しさん mailto:sage [2014/02/08(土) 12:48:31.76 ] スレ違いかもしらんが、 gccとかのソースを見れば超越関数の実装が分かったりするのかな。 >>708 URLの最後の)がいらない。
713 名前:デフォルトの名無しさん mailto:sage [2014/02/08(土) 16:15:49.24 ] >>712 FPUがサポートしている超越関数はソースがないかもね。 iccならSSE版の並列演算用の超越関数が実装されているんだけど。
714 名前:デフォルトの名無しさん [2014/02/13(木) 18:24:19.21 ] 質問です。CUDAを初めて使おうと思うのですが何を買っていいのかわかりません。 当方プログラマです。整数演算主体の力学シミュレータを自作しています。 その中にある絶望的に激重な評価関数が高速化できたらなぁと夢見ています。 その関数は同じデータセット(200キロバイトくらい)を、さまざまな初期値で評価するのですが、 条件分岐が殆ど発生しないアルゴリズムを発見しました。CUDA 向けなんじゃないかと使ったこともないのに妄想しております。 1回の評価計算そのものがめちゃくちゃ重い(単純に100万回くらいループさせているだけ)で、 ループさせるプログラムそのものは数キロバイトも無いちっちゃなものです。
715 名前:デフォルトの名無しさん [2014/02/13(木) 18:26:36.46 ] とりあえず今は手元にある Windows 7 64bits (チップセットはP55) に入れてお試しでCUDAプログラミングし、 C++で書いたシミュレータをCUDA対応に移植するところから始めたいとおもっています。 グラフィック出力を2本使いながらCUDAを使いたいのですが、私はどんなのを買ったらいいのでしょうか? 将来的には研究費をつぎ込むつもりですが、今は衝撃的に貧乏なので予算3万以内で。
716 名前:デフォルトの名無しさん mailto:sage [2014/02/13(木) 18:39:59.33 ] >グラフィック出力を2本使いながらCUDAを使いたいのですが ここんとこ詳しく。あとPCも。 まあビデオカードを1枚買うか2枚買うかくらいの違いでしかないとは思うけど。
717 名前:デフォルトの名無しさん mailto:sage [2014/02/13(木) 18:43:12.82 ] >>714 CUDA積んでるGPUのリスト https://developer.nvidia.com/cuda-gpus この中から予算の許す限り良い奴を買えばいい CUDA開発用のSDKは公開されてるから環境に合わせてインスコするんだ https://developer.nvidia.com/cuda-toolkit 導入方法やサンプルコードはググるか>>1-2 を参照 ……もっとも釣りじゃなければの話だが
718 名前:デフォルトの名無しさん [2014/02/13(木) 19:14:47.30 ] 早速レスが。有難うございます。 >>716 マザボは P7P55D ってので、PCI-Express 2.0 16X が2本あるのですが、 今はRadeon 2枚で4画面(1920x1200x2 と 1280x1024x2)出してます。 そのうちの片方を nvidia にしたいと考えてます。画面出力との併用って難しいでしょうか? >>717 色々ありますよね・・・。今は技術の練習として試そうと思うのですが、 GTX 660 を選択しようかとおもってるのですが、それはやめとけ、こっちがいいよ、とかありますか?
719 名前:デフォルトの名無しさん mailto:sage [2014/02/13(木) 19:26:36.61 ] >>714 その目的なら、PC買う前にまず本当にCUDA向きなのか、 実際に自分がそのプログラムを組むことはできるか、 などをクラウドで実験してみてはどうだろうか Amazon EC2 aws.amazon.com/jp/ec2/ 1時間CUDAマシンを借りるのに1ドルもかからない。
720 名前:716 mailto:sage [2014/02/13(木) 22:54:51.79 ] >>718 Radeon/Geforceの混在かぁ。ごめん俺はわからない。 1枚で画面出力とCUDA計算の併用自体はできる(長い間計算しっぱなしにせず、 ある程度の間隔で制御が戻るようにすれば。 計算しっぱなしでもタイムアウトしない範囲なら表示が完全に固まるわけではないし)。 GTX660でいいんじゃないかな。あるいはコスト抑えたいならもっと下でも。 ローエンドGPUでの実行時間がわかればCUDA Core数の比較で上位GPUにしたときの時間の見当もつくし。 あと2/18に28nmのMaxwell世代のGTX750が発売されるらしいという話もあるけど。
721 名前:デフォルトの名無しさん mailto:sage [2014/02/13(木) 22:58:14.70 ] >>719 と同じくEC2を推す 高いグラボ買って大して高速化できませんでしたじゃ目も当てられない
722 名前:デフォルトの名無しさん mailto:sage [2014/02/13(木) 23:16:31.67 ] >>719 有難うございます。実は仰る話は第二段階として既に計画していました。 ここがクリアしたら第三段階として本格的な予算を投じて 大量のGPUインスタンスを使って計算するかもしれないです。 でもその前に第一段階として、CUDAにあわせてソースを書き換えたり、 必要に応じてアルゴリズムも修正しなくてはならないと考えており、 そのトライ&エラーに例えば一ヶ月かかっちゃうなら安い奴を買ったほうがいいかな、 とりあえず3万円程度で使い倒してみようかな、と考えている次第です。
723 名前:デフォルトの名無しさん mailto:sage [2014/02/13(木) 23:23:35.47 ] >>721 そんな事情もあって3万円くらいで、なにがいいのかなと。
724 名前:デフォルトの名無しさん mailto:sage [2014/02/13(木) 23:50:08.03 ] しれっとアルゴリズムを発見したとか言うよねーw
725 名前:デフォルトの名無しさん mailto:sage [2014/02/13(木) 23:51:06.60 ] 連投すみません。 >>720 なるほどー。 最終的には256ビット幅でアルゴリズムを最適化するつもりだったので、 本当にメモリバスが256ビット幅なら、2/18まで待ってみようかな。 画面出力とCUDA計算の併用はそんなに心配しなくても良さそうなんですね。 Radeonとは、だめもとで混在させてみるつもりです。ありがとうございます。
726 名前:デフォルトの名無しさん mailto:sage [2014/02/13(木) 23:54:53.63 ] >>724 そんな言い方は野暮なんじゃね。何事も思いつきからでそ 失敗の責任を追うのは彼自身なんだしw
727 名前:デフォルトの名無しさん mailto:sage [2014/02/13(木) 23:57:20.97 ] >>724 結構たいへんでしたよ?要約すると、とあるトポロジーのハッシュを計算する問題なのですが、 条件分岐だらけのソースを、ごり押しで全パターン計算させたほうが結果的に高速らしい、ということが判ったのです。 ただ、現時点では「高速らしい」という段階でしかなく、本当に早いんだということは実際にやって見せるしかない状態です。
728 名前:716 mailto:sage [2014/02/14(金) 00:42:20.93 ] >>725 Maxwellについては2/18のは28nmのままだし本気出せるSDKもしばらく先だろうから たぶん結局Kepler買うことになり待つ意味は薄いと思います。 ただ一応知っておいたほうがいいと思って言ってみただけで。 あと整数演算オンリーならFermiが効率いいかもしれないけど今更Fermiに 合わせて作るというのもなんかロマンがないんですよね…
729 名前:デフォルトの名無しさん mailto:sage [2014/02/14(金) 01:15:43.13 ] スレチになっちゃうけどRadeon積んでるならOpenCLで試せばいいような気が。 Radeonのほうが速いだろうし。
730 名前:デフォルトの名無しさん mailto:sage [2014/02/14(金) 01:45:09.56 ] OpenCLは尚更ロマンがないな 旬ならkaveriでHSAか?といっても所詮ミドルレンジAPUだしなあ
731 名前:デフォルトの名無しさん mailto:sage [2014/02/14(金) 05:30:04.44 ] もっと野暮なんだけど、たった百万回×二百KBの演算ならCPUでごり押しした方が速い肝酢。
732 名前:デフォルトの名無しさん mailto:sage [2014/02/14(金) 07:03:57.86 ] 整数演算ならRadeonのほうが数倍速いよ。 CUDAもOpenCLも対して変わらん。
733 名前:デフォルトの名無しさん mailto:sage [2014/02/14(金) 07:06:11.60 ] >>722 本当に真面目にEC2の価格と比較・検討してみたのか疑問なんだが。 1日24時間ぶっ通しで使うわけでもあるまい。 あと、CUDA用のコードがエラーなく意図どおりの計算結果を出すかどうかは、 CUDA対応グラボを使わなくてもデバイスエミュレーションモードで確認できる。 それで動くか確認してからEC2で計算速度を実験すれば、かなり費用を抑えられると思うぞ。 まぁ、これも単なるひとつの手段でしかないし、 グラボを買う目的がCUDAだけでない場合もあるだろうから、 これ以上EC2を押したりはしない。
734 名前:デフォルトの名無しさん mailto:sage [2014/02/14(金) 07:33:44.94 ] CUDA 6.0RC公開きた https://developer.nvidia.com/rdp/cuda-60-rc-toolkit-download
735 名前:デフォルトの名無しさん mailto:sage [2014/02/14(金) 11:21:51.56 ] なんかCUDA excelとかあるな
736 名前:デフォルトの名無しさん mailto:sage [2014/02/14(金) 15:42:58.52 ] >>960 最近のなら www.nicovideo.jp/watch/sm22874825
737 名前:デフォルトの名無しさん mailto:sage [2014/02/14(金) 15:45:25.73 ] 誤爆orz
738 名前:デフォルトの名無しさん mailto:sage [2014/02/14(金) 19:32:15.90 ] GPUは1スレッドあたりの速度はCPUの何百倍も遅い 最低でも32ブロック×数百スレッド以上の並列計算できるような問題でないと力を発揮できない
739 名前:デフォルトの名無しさん mailto:sage [2014/02/14(金) 22:12:19.68 ] >>714 本当にCUDA初心者で、且つ上手く行った後に予算が着くなら 使い捨てにするつもりで1万5千円ぐらいの中古のGTX 580を 買ったほうがいいよ。下手なKeplerより速いし。
740 名前:デフォルトの名無しさん [2014/02/16(日) 17:30:13.79 ] MAXWELL世代はL2キャッシュが大幅に増えるらしい。これはどういう効果を生むの? videocardz.com/49557/exclusive-nvidia-maxwell-gm107-architecture-unveiled GM107 L2 cache has 2MB. GK107's cache has 256KB.
741 名前:デフォルトの名無しさん mailto:sage [2014/02/16(日) 19:06:36.41 ] というかキャッシュを増やさないと、 開き続けている演算能力とメモリ帯域のギャップがますます開いてしまう。
742 名前:デフォルトの名無しさん mailto:sage [2014/02/16(日) 20:26:20.43 ] メモリ周りをチューニングしなくても そこそこ性能がでる感じなのかな
743 名前:デフォルトの名無しさん mailto:sage [2014/02/17(月) 13:05:08.54 ] いろんな変更があるんだろうけど 研究や論文あさってみると階層化されたスケジューラとレジスタファイルの相乗効果と思われる ちょっと古いがこれがわかりやすいかも www.cs.utexas.edu/users/mgebhart/papers/MICRO_Slides.pdf www.cs.virginia.edu/~skadron/Papers/gebhart_tocs.pdf L2は大型warpのプリフェッチに使われる予感
744 名前:デフォルトの名無しさん mailto:sage [2014/02/22(土) 01:53:21.57 ] CUDAの本でてた。 https://gihyo.jp/dp/ebook/2014/978-4-7741-6361-1 買わなくてもサンプルコード落とせた。
745 名前:デフォルトの名無しさん mailto:sage [2014/02/22(土) 18:05:06.36 ] Maxwellはkelperより速くなってるの? Fermiのほうがいいという悲しいことになってない?
746 名前:デフォルトの名無しさん mailto:sage [2014/02/22(土) 19:11:00.32 ] CUDAって1PASSしか使えないけど、将来2PASS使えるようにならないのかね CUDAエンコが速いなら、その速さを活かして2PASSでエンコしたいんだが
747 名前:デフォルトの名無しさん mailto:sage [2014/02/22(土) 20:54:45.20 ] 何のソフトの話だよ。そういうのは作者に言え。
748 名前:デフォルトの名無しさん mailto:sage [2014/02/22(土) 21:20:39.06 ] >>746 お前がCUDAで2PASSでエンコ作ればOKだろ >>745 一般デスクトップ向けはGPGPUより素直にゲームに注力したほうが良いような気がするからな
749 名前:デフォルトの名無しさん mailto:sage [2014/02/22(土) 21:37:41.19 ] >>747 外人だから無理 >>748 ドライバがCUDAでの2PASSエンコに対応していない だからどんなソフトでもCUDAで2PASSエンコは出来ない
750 名前:デフォルトの名無しさん mailto:sage [2014/02/22(土) 21:45:32.45 ] ドライバが?なんか根本的な勘違いがありそうだな それPureVideoの話じゃないの? cudaに向いてるかどうかを別にすれば、マルチパスってのはストリームに対して複数回の処理を走らせる事を指す言葉に過ぎない。なのでドライバは関係無い
751 名前:デフォルトの名無しさん mailto:sage [2014/02/24(月) 13:03:48.10 ] Starting LU Decomposition (CUDA Dynamic Parallelism) GPU Device 0: "GeForce GTX 750 Ti" with compute capability 5.0 GPU device GeForce GTX 750 Ti has compute capabilities (SM 5.0) Compute LU decomposition of a random 1024x1024 matrix using CUDA Dynamic Paralle lism Launching single task from device... GPU perf(dgetrf)= 4.607 Gflops Checking results... done Tests suceeded ------------------------------------------------------------------------------ starting hyperQ... GPU Device 0: "GeForce GTX 750 Ti" with compute capability 5.0 > Detected Compute SM 5.0 hardware with 5 multi-processors Expected time for serial execution of 32 sets of kernels is between approx. 0.330s and 0.640s Expected time for fully concurrent execution of 32 sets of kernels is approx. 0.020s Measured time for sample = 0.053s C:\ProgramData\NVIDIA Corporation\CUDA Samples\v6.0\Bin\win64\Release>
752 名前:デフォルトの名無しさん mailto:sage [2014/02/26(水) 05:37:37.16 ] これって750TiでもhyperQ,Dynamic Parallelism使えるってこと?
753 名前:デフォルトの名無しさん mailto:sage [2014/02/26(水) 05:44:00.90 ] yes
754 名前:デフォルトの名無しさん mailto:sage [2014/02/26(水) 06:09:03.50 ] そうか。すごいな。使うだけならGT640も使えるんだっけ?
755 名前:デフォルトの名無しさん mailto:sage [2014/02/26(水) 09:50:04.54 ] うん
756 名前:デフォルトの名無しさん mailto:sage [2014/02/26(水) 11:18:45.60 ] 公式みると750tiはCompute Capability3.0になってるけど使えるの? 640は3.5だから合ってるけど
757 名前:デフォルトの名無しさん mailto:sage [2014/02/26(水) 12:14:54.40 ] 750TiはCC5.0です CUDA Device Query (Runtime API) version (CUDART static linking) Detected 1 CUDA Capable device(s) Device 0: "GeForce GTX 750 Ti" CUDA Driver Version / Runtime Version 6.0 / 6.0 CUDA Capability Major/Minor version number: 5.0 Total amount of global memory: 2048 MBytes (2147483648 bytes) ( 5) Multiprocessors, (128) CUDA Cores/MP: 640 CUDA Cores GPU Clock rate: 1163 MHz (1.16 GHz) Memory Clock rate: 2750 Mhz Memory Bus Width: 128-bit L2 Cache Size: 2097152 bytes Maximum Texture Dimension Size (x,y,z) 1D=(65536), 2D=(65536, 65536), 3D=(4096, 4096, 4096) Maximum Layered 1D Texture Size, (num) layers 1D=(16384), 2048 layers Maximum Layered 2D Texture Size, (num) layers 2D=(16384, 16384), 2048 layers Total amount of constant memory: 65536 bytes Total amount of shared memory per block: 49152 bytes Total number of registers available per block: 65536 Warp size: 32 Maximum number of threads per multiprocessor: 2048 Maximum number of threads per block: 1024 Max dimension size of a thread block (x,y,z): (1024, 1024, 64) Max dimension size of a grid size (x,y,z): (2147483647, 65535, 65535) Maximum memory pitch: 2147483647 bytes Texture alignment: 512 bytes Concurrent copy and kernel execution: Yes with 1 copy engine(s) Run time limit on kernels: Yes Integrated GPU sharing Host Memory: No
758 名前:デフォルトの名無しさん mailto:sage [2014/02/26(水) 12:35:53.56 ] 公式が間違ってるってことか 安いし買ってみるかな
759 名前:デフォルトの名無しさん mailto:sage [2014/02/26(水) 13:52:11.35 ] ついでにここでやってた32bit integer bit shiftのスループットやってみたら 68とか半端な数字になった けどkepler(GK110 tesla以外)の倍だねSMあたり blogs.yahoo.co.jp/natto_heaven/32775349.html Clock: 1163000 KHz, # of MPs: 5 Elapsed Time: 2774.579102 milliseconds # of Threads: 1024, # of SHLs : 1099511627776 Throughput: 68.147981
760 名前:デフォルトの名無しさん mailto:sage [2014/02/26(水) 14:39:08.09 ] https://devtalk.nvidia.com/default/topic/690631/cuda-programming-and-performance/so-whats-new-about-maxwell-/post/4127010/#4127010 https://devblogs.nvidia.com/parallelforall/5-things-you-should-know-about-new-maxwell-gpu-architecture/ shared が64KB使えるってか専用になったみたい
761 名前:デフォルトの名無しさん mailto:sage [2014/03/02(日) 02:29:33.00 ] >>748 OpenCL対応ソフトが地味に増えていってる状況でGPGPU性能削るとか自殺行為だろ
762 名前:デフォルトの名無しさん mailto:sage [2014/03/02(日) 03:47:16.71 ] GPGPU性能と言っても色々あるけどFermi→Keplerのときトータルでは削られたのでは。 2年前とはOpenCL/CUDAの対応状況が違うってことかね。 個人的にはMaxwellでGPGPU寄りに振ってくれると嬉しいんだけどね。
763 名前:デフォルトの名無しさん mailto:sage [2014/03/02(日) 05:18:44.03 ] >>759 見ればkepler(32)で削られてたbitshiftもtesla同様に64に増えてるし どこをみてgpgpu削られてたって言ってるわけ
764 名前:デフォルトの名無しさん mailto:sage [2014/03/02(日) 10:46:20.93 ] keplerん時にゲーム向けチップから倍精度が削られた事くらいか
765 名前:762 mailto:sage [2014/03/02(日) 11:22:55.08 ] >>763 は俺に言ってるの?
766 名前:デフォルトの名無しさん mailto:sage [2014/03/02(日) 14:07:18.50 ] IDがないから誰が誰だかわからねーよw
767 名前:デフォルトの名無しさん mailto:sage [2014/03/02(日) 15:31:43.89 ] keplerは倍精度削り過ぎだろ。 fermi以下とか酷過ぎる。
768 名前:デフォルトの名無しさん mailto:sage [2014/03/02(日) 15:52:36.88 ] 気持ちは分かるが倍精度は単にGPUとして使いたい層には無意味どころか足引っ張る要素だしな TITAN的な選択肢が今後も提供されるなら別に文句ないけどな 欲を言えば、もうちょい安い製品でも倍精度残したチップ用意してくれれば盤石かな
769 名前:デフォルトの名無しさん mailto:sage [2014/03/02(日) 19:03:02.96 ] FermiではTeslaの倍精度が単精度の1/2、GeForceではその1/4の1/8 KeplerではTeslaの倍精度が単精度の1/3、GeForceではその1/8の1/24 TeslaとGeForceの比率をKeplerでも1/4にするとか、 あるいはデフォでは1/8でもNVIDIAコントロールパネルで設定すると1/4 とか、そういう中間的な選択肢が欲しい。
770 名前:デフォルトの名無しさん mailto:sage [2014/03/02(日) 19:39:04.95 ] >>769 そういうの出したらtesla売れないから出さないよ。 俺は絶対性能は630や650クラスでいいから開発用にフルスペックのが欲しい ECC, GPUDirect, TCCDriver, Hyper-Q, DPなどが使えるやつ。 できればTeslaのflops/B比等が同じで、そのまま倍数かければ Teslaパフォーマンス特性がだいたい予測できるようなやつ。
771 名前:デフォルトの名無しさん mailto:sage [2014/03/02(日) 22:31:53.98 ] >>770 それいいね。 あとは同じ価格帯のボードが新世代で性能が 落ちるってのがなければいいんだけれど。
772 名前:デフォルトの名無しさん mailto:sage [2014/03/03(月) 19:45:26.63 ] >>768 ゲーム用VGAとGPGPU用ボードの2つを出せば良いんだよ (VGAでたとえるなら、ゲーム用VGAのGFとプロ用VGAのQuadroみたいに) ゲーム用にGPGPUを強化しても、ゲーマーにはあんまり価値がなく、そして爆熱になるだけし まぁ、GPGPU用の値段はゲーム用よりだいぶ高くなるだろうが
773 名前:デフォルトの名無しさん mailto:sage [2014/03/03(月) 20:12:15.14 ] >>772 それ現状のまま(GeForceとTesla)じゃね?
774 名前:デフォルトの名無しさん mailto:sage [2014/03/03(月) 21:21:49.23 ] ワラタ
775 名前:デフォルトの名無しさん mailto:sage [2014/03/03(月) 21:27:12.11 ] ロー・ミッドクラスVGA派生GPGPUはイラネで、今のteslaラインナップなんだろう
776 名前:デフォルトの名無しさん mailto:sage [2014/03/04(火) 10:00:41.96 ] 数年前までGPGPUは安いのが売りだったのに今じゃ高級品だもんなー。 そのうちIntelに負けるんじゃん。
777 名前:デフォルトの名無しさん mailto:sage [2014/03/04(火) 19:25:35.73 ] いまは、中GPGPU性能以下(低価格)でいいなら、AMDのHSA APUでって感じじゃない メモリーコピーいらんでお手軽にGPGPUできそうだし そして低GPGPU性能でいいならIntelのiGPUで良いやだろう
778 名前:デフォルトの名無しさん mailto:sage [2014/03/04(火) 20:38:12.99 ] いやぜんぜん
779 名前:デフォルトの名無しさん mailto:sage [2014/03/04(火) 20:40:03.72 ] 開発環境とトライバがだめだめで低シェアなAMDは論外
780 名前:デフォルトの名無しさん mailto:sage [2014/03/04(火) 21:09:46.72 ] Xeon Phiがどれだけ使い物になるかで状況が変わってくるだろうな。
781 名前:デフォルトの名無しさん mailto:sage [2014/03/04(火) 21:15:20.73 ] >>780 HPCシステムズ「ぶっちゃけ製品として尖り過ぎてて使いづらいんだよゴルァ!」 www.hpc.co.jp/benchmark20121113.html ↓ HPCシステムズ「まあ場合によってはコンパイルし直すだけで使える……かな?」 www.hpc.co.jp/benchmark20130409.html
782 名前:デフォルトの名無しさん mailto:sage [2014/03/04(火) 21:46:24.18 ] チューニングすればテレサと同じぐらいのレベルか
783 名前:デフォルトの名無しさん mailto:sage [2014/03/04(火) 23:35:20.92 ] つか、HSAでプログラミングしてみたいんだけど、 SDKか何か配布されてるわけ?
784 名前:デフォルトの名無しさん mailto:sage [2014/03/04(火) 23:53:05.20 ] hsa sdkでググるだけで公式ヒットしたけど、斜め読みではどうしろっていうのかはぶっちゃけよく分からんかったなあ シミュまで提供されてるようで、興味はあれどもkaveri機を組んでまで遊ぶ気力もなく
785 名前:デフォルトの名無しさん mailto:sage [2014/03/05(水) 00:35:12.84 ] 書籍出してほしいなぁ。
786 名前:デフォルトの名無しさん mailto:sage [2014/03/05(水) 14:00:02.54 ] >>696 昔「計算が遅いからメモリでなんとかしよう」 今「メモリが遅いから計算でなんとかしよう」 将来「何通りか計算結果をあらかじめ予想しとこう」
787 名前:デフォルトの名無しさん mailto:sage [2014/03/05(水) 20:12:09.63 ] amdのapp sdk使えばかってにやってくれるでしょ
788 名前:デフォルトの名無しさん mailto:sage [2014/03/05(水) 20:39:30.22 ] ha?
789 名前:デフォルトの名無しさん mailto:sage [2014/03/05(水) 22:09:53.32 ] へ?
790 名前:デフォルトの名無しさん mailto:sage [2014/03/05(水) 23:44:06.91 ] amdのコンパイラ使えばメインメモリとGPUのメモリをシームレスにみてくれるってことだよ そんなこともわからんのかバカちんが
791 名前:デフォルトの名無しさん mailto:sage [2014/03/05(水) 23:48:39.83 ] 予想の斜め下の返しきたなw そらな、もしHSAでシームレスじゃなかったら吹くがな
792 名前:デフォルトの名無しさん mailto:sage [2014/03/06(木) 06:39:04.19 ] dgpuじゃhsaは使えない ナライラナイ
793 名前:デフォルトの名無しさん mailto:sage [2014/03/08(土) 06:28:48.56 ] APUが、DDR3でもいいからクワッドアクセスすれば多少はましになる筈 DDR4や次世代になればさらにましになる
794 名前:デフォルトの名無しさん mailto:sage [2014/03/08(土) 06:39:39.30 ] 750tiの方が面白いわ
795 名前:デフォルトの名無しさん mailto:sage [2014/03/15(土) 13:06:52.12 ID:M4J1N6EC] devblogs.nvidia.com/parallelforall/cudacasts-episode-18-cuda-6-0-unified-memory/ CUDACasts Episode 18: CUDA 6.0 Unified Memory
796 名前:デフォルトの名無しさん mailto:sage [2014/03/19(水) 12:10:13.01 ID:6v1SjmcP] DELLとかHPのWSについているTeslaは他のPCでやっぱり使えないのかな
797 名前:デフォルトの名無しさん mailto:sage [2014/03/19(水) 20:14:26.71 ID:PnfH7c65] リモートデスクトップかsshdでいいじゃん(いいじゃん)
798 名前:デフォルトの名無しさん mailto:sage [2014/03/21(金) 17:28:21.60 ID:o3B7shKW] >>796 つかえるだろ。
799 名前:796 mailto:hanage [2014/03/22(土) 13:50:01.60 ID:a7qFo8sx] Quadroではロックされているとかで使えないと 聞いたので、teslaKもかな?と思った。 結局やめて違うの買った。 ドライバ当てただけでは機能はしないんだな、 これからCUDAインストールして勉強する。
800 名前:デフォルトの名無しさん mailto:sage [2014/03/23(日) 16:48:53.07 ID:cHhuAXVO] >>799 今時そんな金のかかることはしない。
801 名前:デフォルトの名無しさん mailto:sage [2014/03/27(木) 00:14:27.49 ID:tCKsDtJy] メモリ周りが多少良くなるみたいだね。 www.4gamer.net/games/251/G025177/20140326091/
802 名前:デフォルトの名無しさん mailto:sage [2014/03/27(木) 00:20:01.21 ID:wFte/jmz] ありゃ、Maxwellの次はVoltaじゃなかったっけと思ったけどよく読んだらPascalが割り込む形か。
803 名前:デフォルトの名無しさん mailto:sage [2014/03/27(木) 00:44:33.58 ID:Csq19V+D] PCでのGPGPUはHSAだからな Nvはスパコンなんかの非PCでがんばるしかないよな
804 名前:デフォルトの名無しさん mailto:sage [2014/03/27(木) 11:19:08.87 ID:2lNl+T/o] むしろHPCから組込みまで同じCUDAが使えますが すでに
805 名前:デフォルトの名無しさん mailto:sage [2014/03/27(木) 20:42:05.46 ID:Mbz0VA5O] PCはVGAレスが普通になってきているからね わざわざVGAをつける奴ってゲーマーや業務のCAD関係する奴とかで GPGPUのためにNvのVGAつける奴ってどれぐらいいるんだろ。
806 名前:デフォルトの名無しさん mailto:sage [2014/03/27(木) 20:58:44.66 ID:hvZ/Tb80] >>805 あとVGAつけるのはエンコード目的とかでIntelのLGA2011を使用している奴とかかな
807 名前:デフォルトの名無しさん mailto:こいつめちゃおもろいsage [2014/03/27(木) 21:35:27.25 ID:FegESPkl] 日本で一番Xeon Phiを無駄にしている系男子wとかもいるしな
808 名前:デフォルトの名無しさん mailto:sage [2014/03/28(金) 12:57:34.82 ID:TDRfQ9dS] >>805 ノシ
809 名前:デフォルトの名無しさん mailto:sage [2014/03/29(土) 16:39:32.17 ID:b7K/xNj3] 世の中には面白い事考えるもんがおるのう www.otb-japan.co.jp/dmpr/GPU/gpgpu-xeonPhi-E5-Hybrid.html
810 名前:デフォルトの名無しさん mailto:sage [2014/03/29(土) 16:44:00.72 ID:4Sjmsqlf] こりゃ面白いな。 試してみたいw ってかOpenACCってもうリリースされてるの??
811 名前:デフォルトの名無しさん mailto:sage [2014/03/29(土) 23:56:49.50 ID:lTQLq19d] >>809 やっべ、ちょっと欲しいなって思っちゃった 値段を見て熱は冷めた
812 名前:デフォルトの名無しさん mailto:sage [2014/03/30(日) 00:00:27.97 ID:stl8pn1Y] 値段ワロタw
813 名前:デフォルトの名無しさん mailto:sage [2014/03/30(日) 00:12:47.34 ID:tTPtoH97] >>812 むしろこの構成で高く付かない方がおかしいw 大雑把にはXeonで13万、Xeon Phiで37万、Taslaで39万掛かる感じ (更に8x 4GB DDR3 1600 ECC Reg.で10万掛かってるけどな!)
814 名前:デフォルトの名無しさん mailto:sage [2014/03/30(日) 00:31:08.95 ID:stl8pn1Y] ひゃぁ〜w
815 名前:デフォルトの名無しさん mailto:sage [2014/03/30(日) 02:22:44.74 ID:QNfBIqKl] >>813 eBayとかだと自作の方が安くなりそうだな。 といってもtesla K20だけで20万はCUDAらない。
816 名前:デフォルトの名無しさん mailto:sage [2014/03/30(日) 09:37:11.81 ID:stl8pn1Y] じつにCUDAらん。
817 名前:デフォルトの名無しさん mailto:sage [2014/03/30(日) 11:50:35.47 ID:tTPtoH97] >>815 まあ超ハイスペックだと自作のほうが安くつくしな ただ、Xeon Phiは一般販売されてなかった気がするんだが……
818 名前:デフォルトの名無しさん mailto:反エジソン陣営sage [2014/03/30(日) 15:11:23.08 ID:CQoCQWvX] そうだったのか。 まあ買って、単純に付けるだけでは動かんからな。 セッティングがめんどくさそう。 ttp://www.sekaimon.com/i221402553263
819 名前:デフォルトの名無しさん mailto:sage [2014/03/30(日) 18:21:35.60 ID:c0GW3Y/e] >>817 アメリカだと通販で売ってるよ。
820 名前:デフォルトの名無しさん mailto:sage [2014/03/30(日) 19:47:56.88 ID:/wxEDZQQ] >>819 アメリカの金持ちPC自作erはPhiを使ってPC自作するからか へたれ日本ではそんなことする奴いないだろからな
821 名前:デフォルトの名無しさん mailto:反エジソン陣営sage [2014/03/30(日) 21:20:53.45 ID:wRKHa685] それPCやない、WSかHPCや!! TITANの3wayに比べれば敷居低いと思う。
822 名前:デフォルトの名無しさん mailto:sage [2014/03/30(日) 23:10:13.83 ID:MguGVXnW] 自作ショップで普通はおいてないだけで 注文だせは取り寄せできるんじゃない? Teslaに比べてもニッチだろうから。
823 名前:デフォルトの名無しさん mailto:sage [2014/03/31(月) 00:35:49.63 ID:icX96Vv2] インテルがばら売りを許していない少なくとも日本では
824 名前:デフォルトの名無しさん mailto:sage [2014/03/31(月) 01:26:50.30 ID:r13zbHe8] へーそうだったのか。 まあコンパイラも追加で買わないといけないし、 そこそこの資金は要求されるから、そういうもんかもね。
825 名前:デフォルトの名無しさん mailto:sage [2014/03/31(月) 02:07:17.65 ID:oafVDW+W] >>820 よりどりみどりだよ。 https://www.google.com/webhp#q=xeon+phi&tbm=shop
826 名前:デフォルトの名無しさん mailto:sage [2014/03/31(月) 02:42:33.40 ID:oafVDW+W] >>825 のリンク、勝手にwww.google.co.jpにリダイレクトされちゃうな。
827 名前:デフォルトの名無しさん mailto:sage [2014/03/31(月) 20:18:51.35 ID:6pmx6lUa] 気が向いたらXeon Phiも付けてみたいよね。 これって現行の普通のi7だと厳しいのかね。
828 名前:デフォルトの名無しさん [2014/04/16(水) 08:25:39.12 ID:gbx/TG/2] ブラボ4枚差しを今週には購入予定だが、cuda並列はむずかしいのかね? おススメのHPか本はあるのかね
829 名前:デフォルトの名無しさん mailto:sage [2014/04/16(水) 08:31:11.50 ID:1cVHCjMz] きみらって皆採掘が目的なの?10年やって元がとれるくらい?
830 名前:デフォルトの名無しさん mailto:sage [2014/04/16(水) 09:10:52.39 ID:Im3wLyAK] >>829 今から採掘に回るとかただの情弱じゃねーかw 単純にゲームとプログラミングのためだよ
831 名前:デフォルトの名無しさん mailto:sage [2014/04/16(水) 11:53:39.91 ID:7M5hcjK/] 6の正式版が出た
832 名前:デフォルトの名無しさん mailto:sage [2014/04/16(水) 16:32:35.94 ID:oaV0pcFq] >>828 The CUDA Hanbook に書いてあったかも。 CUDA by Exampleにもちょっとだけ書いてあったかも。 gtc-ondemandにもなんかあったような。 曖昧で申し訳ない。
833 名前:デフォルトの名無しさん mailto:sage [2014/04/16(水) 19:54:40.11 ID:GTugkpJK] 5000円以下ぐらいの適当な安いビデオカードって CUDA使うぐらいならCPUで全部計算した方が速い感じなの? 例えば AUS 210-SL-TC1GD3-L とか
834 名前:デフォルトの名無しさん mailto:sage [2014/04/16(水) 20:08:41.61 ID:3pKdyFk1] 多分誰も答えてくれないと思います。
835 名前:デフォルトの名無しさん mailto:sage [2014/04/16(水) 20:22:53.52 ID:Wy9kGQMK] まずGPGPUの基本から勉強しましょう。
836 名前:デフォルトの名無しさん mailto:sage [2014/04/16(水) 20:38:13.05 ID:Im3wLyAK] >>833 一応CUDAコアは付いているようだな……>AUS 210-SL-TC1GD3-L(GeForce 210) www.nvidia.co.jp/object/product_geforce_210_jp.html 同じ「計算」でも、GPGPUが効く分野とそうじゃない分野があるけどな
837 名前:デフォルトの名無しさん mailto:sage [2014/04/16(水) 21:36:19.29 ID:1cVHCjMz] 同じ値段のCPUとGPUを比較した場合 500並列くらいまではCPUの方が勝ってるが10000並列だとGPUの方が圧倒的になる感じ。
838 名前:デフォルトの名無しさん mailto:sage [2014/04/16(水) 21:43:51.00 ID:NtFft60O] IntelがきっちりCPU向けに並列化したコードならxeonも速いよ、GPUメーカーの圧倒的数字は幻想だよせいぜい2倍だよ、みたいな主張をphiの宣伝時にやってて首を傾げたな その理屈ならxeon倍積みましょう、って宣伝すりゃええやろと 結局、極端なオーダーではやっぱりGPUやphiの方が有利なケースがあるんだろうなと理解したけど
839 名前:デフォルトの名無しさん mailto:sage [2014/04/16(水) 22:00:14.31 ID:Im3wLyAK] 例の奴貼っておきますね。どうしても問題の性質や書き方やコンパイラに 依存する部分が大きいからな…… www.hpc.co.jp/benchmark20121113.html www.hpc.co.jp/benchmark20130201.html www.hpc.co.jp/benchmark20130329.html
840 名前:デフォルトの名無しさん mailto:sage [2014/04/16(水) 22:46:25.66 ID:3pKdyFk1] やっぱこの人たちすごいな。 久々に見て回ったらK40の新機能のGPU Boostが他で使えるとか見つけてしまった。 水冷化してないけど試してみるか。
841 名前:デフォルトの名無しさん mailto:sage [2014/04/27(日) 14:27:08.90 ID:KCo2Uyz6] cuda 6.0正式リリースきた https://developer.nvidia.com/cuda-downloads
842 名前:デフォルトの名無しさん mailto:sage [2014/04/28(月) 03:07:08.53 ID:eSg57KtL] もしかしなくても、 また面倒くさいことこの上ない初期設定をしないといけないのか。 インストールとVS2012で拡張子変えて保存するだけで動くようになって欲しいよ。 エラーがでると、どの設定がミスったのかバカには分からんのですよ。
843 名前:デフォルトの名無しさん mailto:sage [2014/04/28(月) 11:41:14.51 ID:eFT4bAhD] >>842 自分が使ってるVSは2008 Pro SP1, 2010 Pro, 2012 Pro UP4 だけど、 CUDA Toolkit 入れた後、 新規プロジェクトなら「NVIDIA -> CUDA X.X」だけで、 既存のプロジェクトなら古いCUDA Tookkitと新しいのを両方入れて プロジェクトを右クリックで「ビルドのカスタマイズ(B)...」すれば動くよ。 少なくとも CUDA Toolkit 4.x -> 5.0 -> 5.5 RC -> 5.5 -> 6.0RC -> 6.0 はこの方法でできた。
844 名前:デフォルトの名無しさん mailto:sage [2014/04/29(火) 07:51:06.18 ID:xH63q4tk] >>843 VS ExpressだとNsight入らないんじゃない?
845 名前:デフォルトの名無しさん mailto:sage [2014/04/29(火) 21:03:18.01 ID:AVMxK0NV] 大して変わってないくせに開発環境変えるなよな
846 名前:デフォルトの名無しさん mailto:sage [2014/05/03(土) 04:22:53.23 ID:qVaKcd2l] これまで開発したプログラムをmaxwellアーキテクチャーで動作させるには 5.5までのtoolkitでptxを吐かせるのか、6.0に移行するしかない模様。
847 名前:デフォルトの名無しさん mailto:sage [2014/05/04(日) 16:05:09.46 ID:/x2IsFFD] >>846 3月末に、カーネルを15種類連続実行するプログラムを、 CUDA Toolkit 5.5でFermi(CC=2.0/2.1)用コンパイルした物、 Kepler(CC=3.0/3.5)用にコンパイルした物、 CUDA Tooklit 6.0でMaxwell(CC=5.0)用にコンパイルした物の3つで、 GeForce 750 + NSIGHT Visual Studio Editonで「All」でプロファイル採ってみた。 いずれの場合も、ほとんど速度が変わらなかったよ。 だから、無理にMaxwell(CC=5.0)用にする必要は無いかも。
848 名前:デフォルトの名無しさん mailto:sage [2014/05/06(火) 20:16:14.86 ID:OXY1qxhv] >>847 これ docs.nvidia.com/cuda/maxwell-compatibility-guide/#axzz30vrhsMg6 は釣りってこと? Σ(-д -;)
849 名前:デフォルトの名無しさん mailto:sage [2014/05/07(水) 05:55:14.20 ID:OEkku2Ok] >>848 >>846 で合ってると思う。 Gxx→FermiやFermi→Keplerのときも、 新アーキテクチャ非対応な古いToolkitで作ったcubinは使えなかったはず。
850 名前:デフォルトの名無しさん mailto:sage [2014/05/10(土) 00:08:49.78 ID:YhiaKf7O] Jetson買った人いる?
851 名前:デフォルトの名無しさん mailto:sage [2014/05/10(土) 01:02:28.38 ID:sYRhNUSv] Jetsonってなんだと思ってぐるぐるしたら、NvidiaのRasPiか RasPiより性能大分良いんだろうが、でも、2万超えは高いな
852 名前:デフォルトの名無しさん mailto:sage [2014/05/10(土) 16:02:02.59 ID:p0Sddlo6] 自動車用じゃん。スレチだろ
853 名前:デフォルトの名無しさん mailto:sage [2014/05/10(土) 20:33:48.55 ID:/nRhPCsz] べつに限定されてはいない 組み込み用といだけ
854 名前:デフォルトの名無しさん mailto:sage [2014/05/12(月) 23:57:23.84 ID:LAs79Y1U] この手の奴にBTデフォでついてんのみたことない 今後の組み込みの方向性的に必須なのに
855 名前:デフォルトの名無しさん mailto:sage [2014/05/13(火) 00:15:45.25 ID:CSl2SJJR] CUDA Tooklit を6.0にしたらGPU稼働率が下がったんだけど気のせい?
856 名前:デフォルトの名無しさん mailto:sage [2014/05/13(火) 00:29:27.71 ID:Iv7eBFJt] >>855 Ver変えたら能率が大きく違ったりするのはよくあることだからなあ……
857 名前:デフォルトの名無しさん mailto:sage [2014/05/13(火) 01:44:04.16 ID:CSl2SJJR] CUDA Tooklit を5.0から6.0にしたら 数値計算プログラムの挙動がおかしくなったorz おんなじような事になった人いますか?
858 名前:デフォルトの名無しさん mailto:sage [2014/05/13(火) 04:13:33.08 ID:CSl2SJJR] 連投すんません。数値計算上の安定化を入れたら解決しました。 浮動少数演算の癖がこれまでと違うのかも・・・。
859 名前:デフォルトの名無しさん mailto:sage [2014/05/13(火) 12:29:21.30 ID:pJVewP3A] 安定化って何したんですか?
860 名前:デフォルトの名無しさん mailto:sage [2014/05/13(火) 12:54:09.60 ID:X1Xq41se] ja.wikipedia.org/wiki/%E6%95%B0%E5%80%A4%E7%9A%84%E5%AE%89%E5%AE%9A%E6%80%A7
861 名前:デフォルトの名無しさん mailto:sage [2014/05/13(火) 14:08:06.19 ID:CSl2SJJR] >>859 非線形最小二乗法のプログラムで、 一回の反復で更新する解の量を少し減らしたら安定しました。 CPUプログラムよりもGPUプログラムの場合に、 初期値からとんでもなく離れていってしまう場合が多いように感じます。
862 名前:デフォルトの名無しさん mailto:sage [2014/05/13(火) 14:15:03.11 ID:CSl2SJJR] >>859 en.wikipedia.org/wiki/Nelder%E2%80%93Mead_method 4. Expansion のパラメータγを通常2とするところ、 1.9-2.0の間で初期値に応じて変化させるようにしました。
863 名前:デフォルトの名無しさん mailto:sage [2014/05/13(火) 15:22:39.27 ID:pJVewP3A] >>860-862 サンクス 誤差の拡大を抑えるってことなんですね でもCUDAバージョンの違いで問題が出るってなんだろ? へんな最適化がされてしまってるのかなあ
864 名前:デフォルトの名無しさん mailto:sage [2014/05/13(火) 20:36:41.53 ID:ckwx0yCj] 演算の挙動が論理的に変わるような変更ってあったっけ?
865 名前:デフォルトの名無しさん mailto:sage [2014/05/17(土) 07:57:45.29 ID:jeRfV2R/] developer.download.nvidia.com/compute/cuda/6_0/rel/docs/CUDA_6_Performance_Report.pdf CUDA 6 Report