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

このスレッドは、他のスレッドでは書き込めない超低レベル、
もしくは質問者自身何が何だが分からない質問を勇気を持って書き込むスレッドです。
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】
http://toro.2ch.net/test/read.cgi/tech/1314104886/

関連サイト
CUDA
http://www.nvidia.co.jp/object/cuda_home_new_jp.html

CUDAに触れてみる
http://chihara.naist.jp/people/STAFF/imura/computer/OpenGL/cuda1/disp_content

CUDA のインストール
http://blog.goo.ne.jp/sdpaninf/e/9533f75438b670a174af345f4a33bd51

NVIDIAの「GeForce 8800 GT(G92)」と次に控える64-bit GPUアーキテクチャ
http://pc.watch.impress.co.jp/docs/2007/1031/kaigai398.htm

CUDAを使う
http://tech.ckme.co.jp/cuda.shtml

NVIDIA CUDAを弄ってみた その2
http://dvd-r.sblo.jp/article/10422960.html

CUDAベンチ
http://wataco.air-nifty.com/syacho/2008/02/cuda_2044.html

KNOPPIX for CUDA
http://www.yasuoka.mech.keio.ac.jp/cuda/

テンプレここまでです。変更点はcudaさわってみた(http://gpgpu.jp/article/61432191.html)を
リンク切れのため外したことのみです。

>>1乙

>>1乙

CUDAでがんばっていきやしょう！

マンデルブロ集合を描いてフルHDサイズの画像を保存するプログラム書てワクワクしてたら、
CUDAよりCPUのシングルスレッドの方が速かった (´・ω・`)

この程度の計算だとメモリ転送がかなり足を引っ張るんだね

たくさんのお仕事がないとGPUさんはがんばれない^^；

メモリ転送イヤポだから、AMD、IntelのiGPUはCPUの管理のメモリをGPUからでもアクセス
出来るようにする方向なんだろ

ユニファイドメモリってなんか先祖返りな不思議な気分ね
あとからGPUだけ交換できない時代にもなって寂しくなりそうだ

実現したらCPUがそのまま浮動小数点モンスターになって、ドライバもDirect3DコマンドをCPU処理で完結して、
CPUが直でディスプレイコントローラ叩くことになるんだろかね

メモリ空間統合したほうが絶対イイね。
データのコピー転送オーバーヘッドはあほらいい。

コピー転送オーバーヘッドを調べる方法ないのか

16EiB の壁はいつ来るだろう

>>11
イベントで計測できるんじゃない？

書籍「CUDA BY EXAMPLE」の第7章テクスチャメモリを読んでるんだが、
意味が分からない。

テクスチャメモリは読み取り専用と言っておきながら、
普通に書き込んでもいるような気がする。

デバイス側に確保したメモリ data.dev_inSrc を texIn に、
デバイス側に確保したメモリ data.dev_outSrc を texOut に
それぞれテクスチャとしてバインドしている。

で、熱伝導を計算するカーネル関数の「引数」に、
1 フレーム毎に data.dev_inSrc と data.dev_outSrc を切り替えて渡している。
このカーネル関数の中ではそれらに値を書き込んでいる。
（もちろん、もう一方はテクスチャとして tex1Dfetch、あるいは tex2D で読み取ってる）

これって読み取り専用というよりは、たとえテクスチャとしてバインドしようが、
依然グローバルメモリとして使うこともでき、かつ tex1Dfetch などで読み取れば、
特別なキャッシュが働いて近傍への読み取りは速くなる、という事？

CUDAプログラミングを体験したいのですが，CUDAのできる格安ノートPCを教えてください．

CUDAのプログラミングを体験したいのならエミュレーションで十分

>>14
そういうことだと思う。
グラフィックスやってると普通の感覚なんだけど、
テクスチャ読み出ししてテクスチャにレンダリングするのは常套手段。
汎用コンピューティング時にテクスチャとしてデータを読むときの利点は
テクスチャ用の高帯域バスやキャッシュ、そしてフィルタリング用の固定機能ハードウェアを利用でき、
よりGPUを効率的に扱えることにあると思う。

>>17
すいません、ちょっと質問です。

テクスチャ用の高帯域バスを活用するには、
それのバス幅などが分からないといけない（他と比較できない）と思いますが、
deviceQuery.exe で調べても nVidia のサイトでスペック表を調べても、
どこにも載っていないような気がします。

普通のメモリバスやメモリクロック数などは分かるのですが、
テクスチャ用の高帯域バスについてはどこで調べればいいのでしょうか。

フィルタリング用の固定機能ハードウェアについても、
自分が使用しているグラフィックチップにどのような物が搭載されているかも、
分かりません。

そもそも、CUDAカーネルからテクスチャメモリの値を読み取る場合、
フィルタリングってされるのですか？

テクスチャ用の特別なキャッシュ機構がある事については、
「CUDA BY EXAMPLE」に載っていましたから分かりました。

>>18
自分は後藤さんの記事を参考にしているよ。
たくさんあるけど、いくつか紹介。

NVIDIAが次世代GPUアーキテクチャ「Kepler」のベールを剥いだ
http://pc.watch.impress.co.jp/docs/column/kaigai/20120322_520640.html

NVIDIA Fermiのマルチスレッディングアーキテクチャ
http://pc.watch.impress.co.jp/docs/column/kaigai/20091105_326442.html

NVIDIAの1TFLOPS GPU
「GeForce GTX 280」がついに登場
http://pc.watch.impress.co.jp/docs/2008/0617/kaigai446.htm

固定機能関係についてはDirectXなどのグラフィックスAPIと同調しているから
そちらの知識が必要になるね。

固定機能にはフィルタリングやアドレッシングがあるけど、
CUDAではこれらをバインド時に設定するみたいだね。

CUDA テクニカルトレーニング
Vol I:CUDA プログラミング入門
http://www.nvidia.co.jp/docs/IO/59373/VolumeI.pdf
（103ページ目のスライド）

>>19
ありがとうございます。
参考にさせていただきます。


すいません、私が使用しているのは Quadro K2000M なんですけど、
テクスチャのバスの帯域幅とかって、どこで分かりますか？
これはフィールレートと呼ばれるものとは別物？

いま nVidia のサイトで調べているのですが、なかなか見当たらないです。
公式には出てない情報なのでしょうか。

テクスチャフィルレートが相当すると言えなくもないけど、フォーマットによって変わるから参考にしかならない
概算はTex数×コアクロック（シェーダクロックではない）で見積もれる

テクスチャメモリを使うとテクスチャキャッシュを通るから、汎用のキャッシュがない上にコアレスアクセスの条件が厳しいG80〜GT21xでは有効だった
Fermiはテクスチャユニット数の比率が減らされた上に、テクスチャキャッシュより汎用キャッシュの方が大容量になったので、むしろ遅くなることもあった
完全に予想だが、Keplerは（線形補間やテクスチャ端の丸め処理を手動でやる必要がなければ）おそらくテクスチャメモリを使っても使わなくてもそんなに変わらない

>>20
そのGPUは最新世代のKeplerアーキテクチャだね。

Keplerの前のFermi世代からはメモリ階層が大きく改変されて、
テクスチャ転送に最適化された上りのパスがなくなった。

http://pc.watch.impress.co.jp/img/pcw/docs/359/423/06.jpg

依然としてテクスチャL1キャッシュを利用できるメリットはあるけどね。

いずれにせよ、内部バスがどれくらいの帯域であるかは公開されていないと思うよ。
クロスバスイッチ接続で調停しながらでもあるから、ちゃんとした数字も出せないだろうし。
Fermiからはキャッシュが噛むようになったから、なおさら。

仮定と実測の両輪でうまく最適化して詰めていくことが醍醐味だろうね。
面倒だけど･･･ｗ
まぁ、ハード屋やってると、こういうのは楽しい。

フィルレートはグラフィックスにおいて、画像の画素を埋めていく（フィル）する速さのことだから、
グラフィックス処理用のROPユニットの能力が影響してくると思うし、あまり参考にはならないかもね。


http://www.nvidia.com/content/PDF/product-comparison/Product-Comparison-Quadro-mobile-series.pdf

しかし、このGPU、CUDAコア数に対してメモリ帯域が残念すぎないか？
Keplerアーキ自体が以前と比べてそういう傾向あるけど、それにしてもヒドイｗ
キャッシュがあるから大丈夫なんかな？
どうであれ、演算/ロード比が相当大きくないと性能出すの難しいかもね。

>>21
> 概算はTex数×コアクロック（シェーダクロックではない）で見積もれる

ありがとうございます。

Tex数というのはテクスチャユニットの数ですかね。
今自分のチップにどれくらいの数が乗ってるか調べてます。


>>22
> しかし、このGPU、CUDAコア数に対してメモリ帯域が残念すぎないか？
> Keplerアーキ自体が以前と比べてそういう傾向あるけど、それにしてもヒドイｗ

そうなんですか。
ThinkPad で CUDA 使えるイイ奴といったらこれしかなかったもので。

> どうであれ、演算/ロード比が相当大きくないと性能出すの難しいかもね。

がんばります。

PTXでブロックまたがってすべてのスレッドでグローバルメモリの同期やりたい時ってmembar.glでいいんだよね多分。

ホスト側のコードだけを書いた cu ファイルと、
デバイス側のコードだけを書いた ptx ファイルとをリンクして
ひとつの exe ファイルを作る方法はあるでしょうか。

もしあれば、やり方を教えてください。

>>25
ホスト側が面倒になるけどDriver APIとか

>>26
やはりそれしかないですか・・・

なんでそんなことをしたいのかが気になる。

>>28
ptx のアセンブラコードを理解する必要がでてきました。
「PTX: Parallel Thread Execution ISA Version 2.3」のPDFは読んでますが、
やはり実際にアセンブラコードを書いたりして実験しないと難しいです。

そこで、nvcc が cu ファイルに書かれたカーネル関数を、
どのようなアセンブラコードにコンパイルするのか、
そのアセンブラコードに変更を施したら結果はどのように変わるか、
などの実験をいろいろやっています。

今はカーネル関数が書かれた cu ファイルを nvcc で ptx ファイルにコンパイルし、
ホスト側で Driver API を使ってそれをロードして実行しています。

ptx ファイルを多少いじるだけなら再コンパイルの必要は無く、
また cu ファイルを多少いじるだけでも、こちらの再コンパイルだけで済みます。
しかしカーネル関数の引数を変えたり、使うデータを変えたりするなら、
ホスト側のコードも供に再コンパイルする必要があり、手間がかかります。
実験が数回くらいならいいですが、何回もやってるとけっこう面倒です。

nvcc host.cpp dev.ptx などと一気にコンパイルできたらさぞ快適だろうなと思い、
質問した次第です。

Makefile

>>30
あぁ、そっちでコンパイルするファイルやコンパイル方法を制御するわけですね。

挑戦してみます。

PTXのコードをインラインアセンブラを使って直接cuファイルの
中にかけばいいじゃん。

>>32
知りませんでした。
「NVIDIA CUDA C Programming Guide Version 4.2」を「inline」で検索してみましたが、
__noinline__ や __forceinline__ の記述しかなかったです。
どこに詳細が載っているのでしょうか。

他にも、ptx のコードを cu ファイル内に書くのでしたら、
文字列として書いた ptx コードの先頭アドレスを適当な変数に入れて、
cuModuleLoadData 関数でロードすることでも実現できますね。

ただ問題は、それだと C 言語で書いたカーネル関数が、
nvcc によってどのような ctx コードにコンパイルされるか、
という部分が調べられない事です。

>>33
つ http://developer.download.nvidia.com/compute/DevZone/docs/html/C/doc/Using_Inline_PTX_Assembly_In_CUDA.pdf

試してないけど、nvccが出力したPTXのコードをインラインアセンブラの形式で
書き換えることも出来るんじゃない?

>>34
ありがとうございます。
なんというか、少々独特のインラインアセンブラ構文ですね。

今の環境より実験がやりやすくなるか調べてみます。

>>35
GCCのインラインアセンブリ構文がこういうのだよ

>>36
http://d.hatena.ne.jp/wocota/20090628/1246188338
これ見ると、たしかに同じですね。

インラインアセンブラはその昔 VC++ でしか使ったことがなかったもので

インラインアセンブラは今回の目的には合いませんでした。

インラインアセンブラ自体は問題なく使えて、なかなか面白いのですが、
nvcc で出力した ptx のコードをそのままインラインにしたのでは使えず、
けっこうな修正を余儀なくされます。

なかなか慣れないこともあって作業量はむしろ増えてしまうので、
今回は make を使ってやる方向でがんばってみます。
（こちらだと、今までの延長線上の考え方で何とかいけるので）

みなさん、ありがとうございました。

CUDA5で美味しい事あるの？

>>39
新機能を使わないんだったら全然美味しくない。
CUDA5でビルドしたらかなり遅くなった。

CUDA 5 Production Release Now Available
CUDA Downloads | NVIDIA Developer Zone
http://developer.nvidia.com/cuda/cuda-downloads

早く５の報告しやがれ

4Gamer.net ― NVIDIA，「CUDA 5」を正式発表。第2世代Kepler「GK110」に向けた準備が整う
http://www.4gamer.net/games/076/G007660/20121016013/

>>42
普通に動いてるよ。

>>43
Nsightはプロファイラーも付いてるのか。
こりゃいい。

>>43
読んでみたけど
これはGeＦｏｒｃｅ切り捨てってこと？

今まで十分遊んだろ
これからはまともにGPGPUしたかったら、金出してTesla買えや
ていう風に読める

それは、Kepler発表の時からゲーム用のKepler1と
GPGPU向けのKepler2があるってことになってた。
Kepler1があまりにGPGPUに向いてなくてGeforce680あたりを
買った人はがっかりしてたよ。

>>43
Eclipse用のNsightも出てLinuxやMacでも開発しやすくなるのは大きいかも。

>>46
Dynamic ParallelismはGK110以降での対応でMaxwell世代ではコンシューマ向けでも対応するのでは？
GPUDirectはクラスタ向けの機能で差別化されても仕方ない気がする。

>>47
ゲーマーにGPGPUっていらんだろうからな
要らないのを付けて高い値段･高消費電力になって売れないものになるなら削れだろ
CUDAする奴はとんがったことする奴だろ。そんな奴ならKepler2のTeslaぐらい買うだろうからな
買えない貧乏人はAMDのradeonに移行しろだな

Mac はどうか知らんが、Linux は Windows 版に比べて、
どうしてもドライバのチューニングが徹底されていない感じがする。

SDK 内のサンプルを動かしてみても、
Windows 上で動かしたときより明らかにフレーム数が落ちる。

そりゃ、いくらDRIが実装されたからといってX11なんていう
太古のグラフィックAPI使ってるんだから、そんなもんじゃないの？
本気を出させたければWayland待ち。

>>51
遅いのはCUDAじゃなくて、その結果を表示する
グラフィックスライブラリの方ってこと？

確かに、フレーム数が低いと分かって時点で Linux パーティション消したから、
グラフィックスを伴わない純粋な計算で比較したことはないなぁ

今度ためしてみよ

フレーム数計る時点で、グラフィックカードに描画させてるんだよね？
その描画をグラフィックス・ライブラリが足引っぱってるんじゃないかってこと
CUDA自身はプログラムをGPGPU用のアセンブリ言語に変えて
GPGPUに実行させるだけだから、あまり差が出るとは考えにくい。

そう言えば Yellow Dog Linux for CUDA 使ってる人いる？
どんな感じなの？

Linuxなら、GUI止めないとカーネルによっては処理速度ががた落ちする。
使えるGVRAMも激減する。

CUDA + GUIつっても、サンプルでXが関わるところなんてウィンドウの枠だけじゃないか？
あとはOpenGLで描画されていると思うが

>>55
Windowsのほうがガタ落ちだし、使えるメモリも少ない。
グラフィックスを使うと遅くなるのはXの問題だから。

dynamic parallelism は GeForce じゃ無理なんですか

うん。

調べたなかではＧＤＲＡＭのみのように見えるんだけど、
テスラだとＬ１、Ｌ２、シェアードメモリもＥＣＣ保護されてるの？
それともＬ１、Ｌ２くらいの容量なら気にしなくてもいいのかな？

レジスタも。

>>58
今のところTesla K20のみだったはず

一般人向けは2014年まで待てとか遅すぎる

GK110はいつになったら一般向けで出てくるのやら…

>>64
ないと思うのは俺だけか

>>65
gtx780とかじゃないか？
来年の春だった気がする。

GTX 780はKepler1の改良版だって聞いたぞ。

一般人向けでダイナミックなんちゃらが使えるのはMaxwellからとか

AMDが2013年中に簡単にOpenCL対応アプリをかけるようにしてきたらどうするんだろ

NVIDIA Visual Profiler v4.2をCentos6.2で使おうとしてるんだけど、

No Timeline
Application timeline is required for the analysis.

と出て解析できない。
調べたらLD_LIBRARY_PATHに/usr/local/cuda/lib64やらを追加せよとあったんでやってみたけど状況変わらず。
どなたか同様な症状に出くわした方はいらっしゃいませんか？

>>69
CUDAプログラミングはまだまだ敷居が高いね･･･

nvcc ***.cu -O2 -Xcompiler -O2
のようにO2を重ねるのは無意味ですか？
前者のO2はGPU用，後者のO2はCPU用と勝手に思っていたんですが，
同じ事を繰り返しているような気がしてきました．

>>71
意味があるのか、どのような意味があるのかまでは分からんが、
とりあえず、「同じ事を繰り返しているのかどうか」については、
出力されたファイルを比較すれば直ぐに分かると思うぞ。

バイナリで比較してもいいし、アセンブラコードで比較してもいい。

ptxで２つある場合，前者のみ，後者のみ，両方無い場合を比較しましたが，
冒頭の***.iファイルの名前が微かに違うのみで差はありませんでした．
両方消しても差が出ないのは？ですが，
重ねても意味は無さそうであることが分かりました．

>>71
ありがとうございました．

>>73
今のgccのディフォルトが-O2相当なんで、書かなくても変わらないのはその所為。
試しに、-O3とか-O1との組み合わせを試してみたら?

登録ユーザーサイトが復旧したよ

k20はやっぱり高いな。
38万だそうだ。
20万切ってくれないと買えない。

dynamic parallelism対応のGeforce（GTX8XX?)が出たら
画像とか動画を扱うソフトは瞬く間にCUDA完全対応になるのかね?

んなわけない

dynamic parallelismができるからCUDAが劇的に簡単になるわけじゃないから。
Reductionとかで効果はあるけど。

そもそもReductionはマルチパスにしないで
2パスで済ませた方がいいのは、
CUDAのreductionトレーニングでも明らか

CUDAのプログラム作って動かしたいです
自分のMacbookは、グラフィックのチップがIntel GMA X3100なんですけど、
NVIDIAじゃないとCUDAは使えないんですか？

ここで聞いて良いのか分からないので、不適切なら誘導お願いします。

GeForceの省電力の状態(P0〜P12)をGetLastInputInfo-GetTickCountに
応じて切り替えるようなソフトを作りたいのですが、
P0〜P12を切り替えるAPI関数はありませんか?

NVAPIをhackすればできる

CUDAカーネルの中で呼び出す関数に特定の処理を入れるとカーネル自体が読み込まれなくなります
具体的にはプロファイラで実行時間見てみるとカーネル自体が表示されず、一瞬で動作終了する状況です
一応、その特定の処理の部分をコメントアウトするときちんと実行されます（当然正しい結果は出ませんが）
こういったことはどういう状況で起こり得るのでしょうか？

>>84
カーネル呼び出した時にエラーが起きてるんでしょ。
エラーチェックしていないんじゃないの?

>>84
cudaGetLastError()は何と言っている?

>>85
即レスありがとうございます
正にその通りでした。単にメモリの要求量がおかしかっただけみたいです
初歩的すぎるミスに自己嫌悪…

NVIDIAR Nsight? Visual Studio Edition 3.0 CUDA Preview
Nsight Visual Studio Edition Early Access | NVIDIA Developer Zone
https://developer.nvidia.com/rdp/nsight-visual-studio-edition-early-access

Nvidia Geforce forum is back from the dead
http://www.fudzilla.com/home/item/29337-nvidia-geforce-forum-is-back-from-the-dead

CRS形式の行列格納サンプルコードってどこかにない？

いくらでもあるだろ
圧縮方法を理解できたらサンプルもいらんな

1 2 3 4
2 5 6 7
3 6 8 9
4 7 9 10

>>91
圧縮方法はわかったんですがコードに上手く起こすことができなくて困っていたんです。何かいいサンプルがあれば教えていただけると助かります。

馬鹿には無理。

CUDA5は既存のGPUに入れると遅くなるの？

研究室でCUDA用にGTX680搭載PCの導入が決定してしまったんだが評判悪いとはいえ流石に今使ってる560Tiよりは性能いいよね？

Tesla K20きたぞ

>>96
予算処理上の都合だったらしい。

最近プログラム入門した
CUDAとか聞くとワクワクするけど物理の知識も科学の知識も特にないので
数百万スレッド並列で処理するネタが思いつけなくて悲しい思いになる
もっとちゃんと勉強しておけば良かった

京が3位に

東工大の学生たちはもうGK110貰ってるの？

Intelがついに来るぞ
http://pc.watch.impress.co.jp/docs/news/20121113_572526.html

ソースの改変が少しでパラレル計算ができるとのことだが、実際の所どうなんだろうね。

>>101
nVIDIAが押されて、もうちっと貧乏客を引き込むマーケティングをやってくれんかな。

一般のビデオカードで定格の80%までクロックを公式に落とせかつその速度なら
GPGPUの動作を保証。
これを是非やってほしい。仲がよいベンダーがいくつかあるし。

開発環境やソフトウェアの安定性とか含めて、XeonPhi強そうだなぁ

XeonPhiは高いぞ
安いGPUは安い
Tesla買うならXeonPhiのほうがよさそうだが

半年ぐらいしたら、$500くらいのローエンドXeon Phiが出るだろうから、純粋にアクセラレータとしてのteslaは厳しいかもなあ。

Phi触ってみてぇ。
OpenMPで簡単マルチコアプログラミング♪


スレッドオーバーヘッドが小さいことを願う･･･

SSEとかAVXみたいなのをちゃんと使える人じゃないと
TESLAのような性能はでないよ。
512bit演算命令が命だから。

ただのロジックを複数スレッド回したい人なら、
TESLAより速いかもね。かなりの無駄だが(笑)

512bit演算命令ってのがあるのか？
AVXでも256bitだが・・・

VPUてので512ビット命令を処理するようだな

ま、経験上はベクトル命令はCUDAよりは扱いやすいよ

うん、イントリンシックでベクトリ処理書くの楽♪
条件分岐がめんどいけど、LNIはマスクレジスタをサポートしてたからだいぶ楽に書けそう。
しかも512bitもあるなんて最高すぐる。

あー、Phi触りてぇ〜。

OEM向け1000個ロットでXeon Phi 5110Pが2650ドル
らすぃ
なんか価格でもTeslaやばそうだな

Intel，スーパーコンピュータ向けアクセラレータ「Xeon Phi 5110P」発表。60基のx86コアを1チップ上に集積
ttp://www.4gamer.net/games/049/G004963/20121111001/

むしろ値下げ合戦になればよい。

合戦になるほど数競争起きる市場でもないべ

CUDAの強み：先行者利益、CUDAが一応動く環境が多め、設計製造がGPUと共用なので低コスト
Xeon Phiの強み：たぶん使いやすさ

って感じだと思う。HPCを本気でやる人たちはXeon Phiのほうに目がいくんじゃないかな。
Xeon Phiはそれはそれで制約があるんだろうけど、CUDAよりは融通が利きそうだから。

Geforce持ってるしCUDAで遊ぶのはいいけどXeon Phi買うとかありえんわっていう一般人としては、
KeplerはあきらめるとしてMaxwellで再びFermi並にGeforceにもGPGPUの機能を盛り込んでほしいと思う。

しかしFermiのときにNVIDIAはCUDA使いの増殖とCUDAアプリの誕生の期待をこめて
Fermiにもそれなりに機能を持たせたんだと思う。しかし今後CUDAをうまく活用するアプリが
HPC以外で出てくるかというと、結構諦めモードなんじゃないかと。
つまりMaxwellもGeforce製品はGPGPU捨ててくるんちゃうかと。
つまりCUDA使いのおまいらがんばってくださいおながいします

長文の上に間違えてーらorz
Fermiにもそれなりに機能を持たせたんだと思う→Geforceにもそれなりに機能を持たせたんだと思う

>>115
nVIDIAの株を空売りすれば儲かるということか。

phiはCPUに内蔵させるGPUコアと共通化させて
コスト落としつつマーケットシエア取る作戦かな？

そしたら、本気でnVidia終わるな

>>112
Phi、扱い易そうだな。
ベクタ演算器処理の記述法が気になるとこだし、
nVidiaがアセンブラのように複雑ってディスってたけど、
イントリンシック記述だったら簡単だし、
条件分岐のマスクまでサポートしてくれたら文句なしだ。

これ、マジで触ってみたいな。

CUDAは開発環境タダだけどXeon PhiはIntel Compiler必須だよね

ものいりだねえ、 Phi

>>120
GCCが対応するって書いてあったぞ

そもそもintelコンパイラもLinux版はフリーだよな

ここのスレ見てたら、phiを月2、3万位でレンタルする商売が出来そうな気がしてきた…

GPUだってAmazonのクラウドで借りたりできるし
Phiもそういうの出るだろう

>>114
まあ、ATI Streamが出てきたからといってTeslaが安くなったというわけじゃないからな。
しかし、今回の場合Phiの場合はコードの書きやすさからすると、CUDAの比じゃないから、
>>115にあるお互いの強みを生かして、切磋琢磨して値下げ合戦してほしいわ。
両方のコードを書いている身としては安くなってくれればどっちでもいいんだが。

Phi、4、5万で買えるようにならないかなぁ〜。

Phいらないだろ
一晩中PC動かせばいいだけだろ

動画エンコ用途でもあるまいにｗ
すでに一年中計算回してるような人に、これなら一ヶ月で済むよ、って訴求するのが筋の製品だろ

数時間、数日動かして、後から些細なバグに気づいた時の何とも言えない気持ち

これを何とか少しでも解消してくれるシステムが欲しいな

バグを直したら、その部分だけ再計算すればいいような仕組み

>>130
とりあえず賽の河原症候群と名付けておくよ

phiは本体CPUもXeon使った時の協調性とかで
パワー増すんだろうなぁ
Teslaやばいなぁ…

…投げ売りになってくれると嬉しいなぁ

しかし投げ売りの後に待っているのが撤退だとしたら…？
Xeon Phiには縁がなさそうだから気軽に触れるCUDAにがんばってほしいなぁ

自作ゲームにCUDA利用してる奴っている？

いるなら、何に使ってる？

使えもしないのに欲しくなってとりあえずダウンロードしているんだけど
クソミソに通信速度遅い。
25〜60KB/sをウロウロしてるけどそんなもん？
ttp://developer.download.nvidia.com/compute/cuda/5_0/rel-update-1/installers/cuda_5.0.35_winvista_win7_win8_general_64-1.msi

>>135
ウチじゃ2.0MB/s位出ているぞ。

まじすか
一端回線切ってIPアドレス変更とかしても速度出ない・・・

OCN保土ヶ谷

>135
遅すぎ
間に無線とか入れてない？

人柱はまだか
http://www.amazon.co.jp/dp/B00A31Q6QI/

>>139
高すぎ。

今だったら34万くらいで買えるだろ。

CPU内蔵のiGPUをPCの表示用に、dGPUをCUDA　GPGPU専用にする場合
やdGPUを2つ使って片方をPCの表示用に、もう片方のdGPUをCUDA　GPGPU専用にする場合
ってそれらが出来る(出来ない)マザー、CPU・APU、dGPUってある?
出来るのなら、これをPC表示用、これはGPGPU用って設定とかするの
する場合どうするんですか?

>>141
GTX 580/590とASUS Maximus V GeneとCore i7-3770Kの組み合わせなら出来た。
BIOSでどちらを表示用に使うか設定できる。

>>142
劇速れすありがとう
マザーのBIOSに設定があれば、iGPUとdGPUの場合は出来ると
思っていいのかな

>>143
CUDAのデバイス指定はアプリケーション次第だよ。BIOSは関係ない。
ちゃんとどのデバイスを使うか指定できるようになっていれば問題ないよ。

>>144
CPU内蔵の（CUDAが使えない）iGPUとグラボ側の（CUDAが使える）dGPUがあったとして、
今dGPUを表示用に使用して、iGUの方は眠らせるようにBIOSが設定されているのなら、
CUDAを使うとひとつのdGPUで表示もCUDAも使うことになると思う。

この場合はBIOSでiGPUを表示用に設定させないとダメなんじゃないか？


あと、ついでに俺も聞きたいんだが、そうやってiGPUで表示してdGPUでCUDAする場合、
cudaGLSetGLDevice関数などを使ったCUDAとOpenGLドライバとの相互運用はできるの？
（DirectXとの相互運用でもいいけど）

>>145
だから、BIOSで設定するのは画面の表示だけってことなんだよ。
CUDAでの利用はそれとは全く別に行えるよ。
BIOSでiGPUを表示用に選んでマザボにディスプレイをつないでから、
CUDA対応アプリでdGPUを選べばいいだけの話。

OpenGLは使ったことないけど、CUDAを使った限りでは相互運用は全く問題ない
ように見える。

>>146
すまん、言い方が悪かった。
その「BIOSでiGPUを表示用に選んでマザボにディスプレイをつないでから」が、
CUDAを使ったプログラム側からは操作できないから、
BIOSをいじる必要があるよねという（当たり前と言えば当たり前の）確認だけだったんだ。

> OpenGLは使ったことないけど、CUDAを使った限りでは相互運用は全く問題ない
> ように見える。

ん？
cudaGLSetGLDevice関数を使った相互運用は、例えばCUDAの結果がVRAMに入ってて、
それを直接OpenGLのテクスチャとして使える（CPUやメインメモリを介さず）、
という事だと俺は認識してるんだが、表示用とCUDA用で分かれててもできるのか？

もしかしたら、俺の認識を根底から改めねばならんかも・・・

誰かこの辺り分かる人いる？

>>145
グラフィックライブラリと相互運用する場合は
出力用GPUとCUDA用GPUは同じな必要があるんじゃない？
俺はそうしてる．
確かめたことがあるわけじゃないから無責任な言い方になるけど．
GTX 580(1)から二画面，GTX 580(2)から一画面のトリプルディスプレイやった時に
SDKのスモークパーティクルとか起動しなかった記憶がある．

表示用デバイスでなくてもOpenGLは動かせるから、cudaと連携できると思う。

SDKのマーチングキューブのサンプルで、defines.hの中の #define SAMPLE_VOLUME ってところが
0だとあらかじめ用意された関数が、1（デフォ）だとファイルが読み込まれるんだけど、ここを0にしてもなにも表示されない
コードはそこ以外いじってないんだけどほかにも変更しなきゃいけない部分とかあるのかしら

>>150
ごめん自己解決した
VSでコードいじってたんだけど、すべてリビルドしたら表示されるようになった

Fermiでスレッドブッロクを512以上を指定すると、カーネルが起動しない。
Fermiはブロックごとに1024スレッド対応しているはずなので、
までレジスタが足りないからなのか、
シンプルなカーネルだと1024スレッドまでいける。
動かないならエラーで落ちて欲しいんだが。

質問です

CUDA+VisualStudio2012Deskという環境でプログラミングしているのですが、
Intellisenceがうまく動かないんです。

__global__ void kernel(){...}
main() {... kernel<<<1,1>>>(); ...}
の、<<< >>>だけうまく動かないんですよ。

ビルドは一応できて実行もできるんですが、気持ち悪いので何とかならないでしょうか

>>153
いんてりせんすがC++以外の表記に対応していないんだろ。
自分でぷらぐいんを書けばなんとかなるんじゃね?
尤も、いんてりせんすに頼るような奴に書けるかどうかは知らんが。

2012じゃCUDAの環境がまだまだなんじゃね？
2010だとようやくこなれてきた感じがあるが。

>>154 >>155
ありがとうございます
無理そうなので諦めます

こんな感じのカーネルがあって、

kernelfunc<<<1,16>>()
{
　　int ix = threadIdx.x;
　　
if(ix < 14)
　　{
　　 何らかの処理
　　　__syncthreads();
　　}
　　何らかの処理
}

MPIとかSMPなどではこのような処理は帰ってこなくなるけど
CUDAのカーネルでは問題なく動く。
__syncthreads()っていうのは、分岐があってもWarp単位では分岐から外れたスレッドは単に何もしないだけで、
__syncthreads()がどっかで呼ばれたらとりあえず足並みを揃えることはする。
という理解でいいのかな？

大体そんな感じ。分岐から外れたスレッドが生きているか死んでいるかは兎も角も。

>>158
ありがとう。

CC3.5のGeForceまだ？

K20が売れなくなるから当分無し

CUDAの性能でいったらGTX580>GTX680なんですか？

>>162
http://blog.accelereyes.com/blog/2012/04/26/benchmarking-kepler-gtx-680/

倍精度を使うならGTX580が圧倒的。
単精度ならモノによるが、基本680は足回りが遅い感じ。

>>163
ありがとです。

ブロック数がSM数以上の場合、ブロックでの動作が終了したら次のブロックにいくんですか？

次のブロックっつーか、残りのブロックだな。

CG法をCUDAで実装してMatrixMarketの疎行列を求解しようとしているんですが
連立一次方程式のb要素はどのように設定したらいいのでしょうか？
MatrixMarketで与えられているのはAの疎行列だけなのでbをどのように設定したらいいのかわかりません

テクスチャメモリって使っててあまり早くならないんだけど，実際の効果ってどれくらいなの？

質問です。

今までCソースをCUDAに置き換える作業をしていてうまくいっていたのですが、
C++ソースをCUDAに置き換えようとしたときに
error C2059: 構文エラー : '<'
というエラーが出ます。
Cを変える時と違い、C++では何か特別なことをしないといけないのですか？
考えられる原因などありましたら教えていただきたいのですが...

ちなみにVisualStudio2008でCUDA4.2という環境です。

>>169
ソースの該当箇所の引用ぐらいしろよ。

>>169
nvccでコンパイルする対象(*.cu)ではテンプレートは使えなかったと思う。
C++からCのライブラリを呼ぶ要領で分割コンパイルすればホスト側で使う分にはOK。

>>170
申し訳ないです。カーネルの部分で以下のように記述しています。

// execute the kernel
kernel <<< block, thread >>> (d_Org, d_Cur, d_SubShift, d_StrideCur, d_StrideOrg, d_Comp, d_Hor, d_Ver,
uiSad, iSrchRngHorLeft, iSrchRngHorRight, iSrchRngVerTop, iSrchRngVerBottom, i_StrideOrg, i_Cols, i_Rows, piRefY, iRefStride, d_SAD, m_uiCost, m_iHor, m_iVer );

>>172
extern Cをしてないだけじゃないのか？
それにしても引数多いな。

関数のテンプレートは >>171
template <typename Float, bool checkBottom>
__device__ void gpu_impl_sub(Float const * __restrict__ f, Float * __restrict__ fn, dim3 const &fDim, 以下略
みたいな感じで使ってるけど問題ないよ。

.cu単体ではコンパイルできるんですけどビルドしようとするとCUDA特有の記述の部分で
エラーを吐きます。
でよく見てみるとその前に
nvcc fatal : A single input file is required for a non-link phase when an outputfile is specified
みたいなエラー吐いてるんでリンク関連で間違えてるんですかね…

ただSDKのソースは動くんでここら辺は間違っていないと思うんですけど…

>>168
Fermiより前のL2キャッシュが無いようなGPUだと効果あるよ。
あとハードウェア補間を活かせる分野

>>176
fermiでも場合によっては効果はある。
でも一番いいのは線形補完で使う時だな。

>>175
リンク前の個々の .cuファイルのコンパイルの時点でのエラーのように見える。
どこかで nvcc -c -o a.o a.cu b.cu みたいな呼び出ししてるんじゃ。

SMが10個あって処理の数が合計100個ある場合SMはそれぞれ10回動く？
それとも重い処理があって時間がかかってるSMがあれば空いてる別のSMが担当する？

なんとも答えにくい初心者の質問だな

質問
TESLA K20買おうと思ってるんだけど、プログラミングとは別にエンコでも使おうかなと思ってる。
やっぱりGTX680とかGTX580と較べて全然性能違う？（体感）
開発にはVS2012Pro持ってるからそれに追加しようと思ってる

宝の持ち腐れ臭

>>182
一部のスーパープログラマー以外皆大なり小なり宝の持ち腐れだろう

>>180
色々見たけど書いてないから自明なのかな？

>>181
GTX680でええやん、いくらなんでも値段が凄まじすぎる＞K20
それにGPUエンコは品質が……フィルタはともかく、AVCに落とすにはQSV使う方が速いし

>>185
まーなー、BOINCもやってるからTESLAちゃんにBOINCさせたら
どうなるのかなって気持ちもあったりするんだよな

まぁ一応言っておくと、性能以前に
Dynamic Parallelismの対応有無の差が大きい >GTX680とK20

どうせ数年待てば下位にも降りてくるんだから
待ちなされ

そして降りてくる時にはさらに凄い何かに目移りする訳ですね

質問です。

CUDA4.0に対応したCUDA C Best Practices Guideを探しているのですが、どこかに公開いないでしょうか

>>190
ググればでてくるだろうに。
ttp://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html

所用で確認できず、返信遅れました。すいません

>>191
冒頭部分に

This Best Practices Guide is a manual to help developers obtain the best performance from the NVIDIA® CUDA™ architecture using version 5.0 of the CUDA Toolkit.

とあるので、これはCUDA5.0対応ではないでしょうか。
おそらくCUDA5.0対応のものでも、CUDA4.0の内容は満たすとは思いますが、どこが5.0のものかが判別できないと思うので、できればCUDA4.0対応のものがほしいのですが･･･

DPでのQソートのソース見ちゃうと
無しでやるのがアホらしく感じてくるな

CC3.5がコンシューマーに降りてくるかも♪

GK110を搭載するGeForce GTX 780 Titan　2013年2月末？
http://northwood.blog60.fc2.com/blog-entry-6531.html

winrarをcudaから展開するdllってなんか無いかな？
pw付きでもこちらからの操作で出来るやつ

>>195
GPU使ってパス解読できるソフトなら売られてるが、
CUDA使ってRAR解凍するソフトってあったっけ……

それってそんなに意味あるの？

>>196
パスワード解析なんかはGPUに向いてないんじゃないか？

>>198
zip用→http://www.internal.co.jp/products/util/passgetter/about/
RAR用→http://www.golubev.com/rargpu.htm
それなりに速度は出るっぽい

>>198
ブルートフォースはGPUの独壇場じゃない？
FPGAの専用プロセッサを除いて。

ブルートフォース
プルートフォース

>>200
>>198は、GPUは整数演算がそれ程得意では無い事を言いたいのでは？

VirtexクラスのFPGAでチューニングして200MHzくらいで動かしたらハイエンドGPUより速い？

>>202
http://http.developer.nvidia.com/GPUGems3/gpugems3_ch36.html

ごめん、くっそ遅かった

>>203
20KLUTsで20GBpsでる

GeForce 8800の整数は遅いけど、Fermiはそれほど遅くない

GeForceGTX Titanでは倍精度演算性能が削られている。単精度演算性能は4.70TFlops、倍精度演算性能はこれまでの“Kepler”と同様に単精度の1/24となり、196GFlopsにとどまる。
Pixel Fillrateは49GPixel/s、Texture Fillrateは196GTexel/sである。メモリ帯域は384-bitインターフェースもあいまって288.4GB/sに達する。

http://northwood.blog60.fc2.com/blog-entry-6558.html

GTX580 197.6GFLOPS
GTX590 311GFLOPS
あれ？

Fermiは大サービスだったんだよ。
もう倍精度欲しいならTesla買うしか無い。
K20でも30万くらいだろ？

自分は単精度しか使わないからTitan買うけど。

本当Fermiは大サービスだったよなぁ。整数やビット演算もまともだったし。
CUDA的にkepler refreshは大きく期待できることもないのでmaxwellが頼りなのだが

>CUDA的にkepler refreshは大きく期待できることもない
いやいやDPがあるだろ

Xeon Phiの発売で貧乏人以外には完全にオワコンでしょ

GPGPUはAMDですよ

>>209
Dynamic Parallelism（綴り覚えた）ってそんなに大きいの？
なんていうか使えるハードウェア資源が増えるわけじゃないし…って思って。

一方でXeon Phiがそれほど脅威って感じもしないんだよなぁ。
1. プログラムが組みやすい
2. 実効性能を引き出しやすい
あれを使ったプログラムの組み方を知らないので当てずっぽうなんだけど、
1.は大差ないんじゃないかなって。並列処理の組み方がそんなに素晴らしく変わることは想像できない。
2.でCUDAにコストや電力で勝利なんてことにでもならない限りは棲み分けるかと。
まあ貧乏なんでそうあってほしいっていう願望も入ってるんだけど。

インテル? Xeon Phi　5110PとNVIDIA Tesla K20の行列積における実効性能比較
http://www.hpc.co.jp/benchmark20130201.html

単精度計算ならK20、倍精度計算ならXeon Phi 5110P　という結論らしい

>>212
OpenMPだからPC用のプログラムがそのままで動くよ

http://devgurus.amd.com/thread/159457

7970はDGEMM 665GFLOPSらしい。

>>214
動くけど、専用ベクトル命令使わないと速度でないよね？

>>213
このK20Mと5110Pっておいくら万円するんだろう。

プログラムが簡単に書けることと、速度が出し切れることは大抵両立できない

>>216
インテルコンパイラを使えばいい。
それに、インテルにはMKLがあるから、
線形問題であれば、既存のコードにディレクティブを挿入するだけで速くなる。

>>217
どちらも30万くらいだろ。

Tesla買うならXeon Phiになるだろうが
GTXなんとかでもCUDAは使えるからな

◇GeForce GTX Titan 演算性能
　 SP：4.5TFlops
　 DP：1.3TFlops

単精度浮動小数点演算性能が4.5TFlops、倍精度浮動小数点演算性能が
1.3TFlopsと伝えられており、倍精度浮動小数点演算性能は
単精度浮動小数点演算性能の2/7程度となっています。
2月13日の情報では倍精度は単精度の1/32に制限されるという話も
出ていましたが、今回のスペックの通りならばGK110の
倍精度浮動小数点演算性能に制限がかかっている様子はなさそうです。

http://northwood.blog60.fc2.com/blog-entry-6565.html

>>221
そんなうまい話は無いと思うけどね

Titanたっけー。
誰か研究室とかで買って報告よろ。

初日か初週でゲーマーが争奪戦して在庫切れで終わりだろ
悠長に評価して予算付けてって頃にはもう終わってる

年度末で10万くらい予算余ってるから何か買っていいよ。



ってのが、今年は来ない。

>>225
10万で買えるの？

>>260
118000円で予約したからさすがに10万ジャストは無理

K20の約半額か

>>228
安いと言えば安いが、個人で遊ぶのはきつい価格だった。

>>229
でも580や480と違って倍精度もtesla相当に設定できるようだし、
ECCと演算保証がないことに目をつぶればお得といえばお得じゃない？
後々の中古売却を考えるなら値崩れはTeslaのほうが少ないだろうけど。

「法人さんはTeslaを買ってくださいよ」（2/21）
http://www.gdm.or.jp/voices/2013/0222/21236

「GeForce GTX TITAN」を「Tesla K20X」の代替え品として考えている
法人ユーザーも多いようで、早くも在庫の問い合わせが数件あるという。


お前らか…

まー自分の金じゃないしとりあえず問い合わせぐらいはするでしょ

実際の性能向上率をみたり、書き換えが必要な場合それの専攻開発用とかに買うんだろ。
普通本番機ではサーバーと一緒にTESLA導入するよ。

しかしTITANはDynamic Parallelismが無いのが不便だね
枝分かれが動的に決まる深さ優先探索を分散処理させたいのだけど

ちょｗ
Dynamic Parallelismないって致命的じゃんｗ

オワタ
しかし、GTX TitanではKepler GPUの新機能である
Dynamic ParallelismとHyper-Qという機能が省かれている。

カ○タムBIOSでTesla化出来ないかな？

Dynamic ParallelismつけたらTesla売れなくなるだろｗｗ

最も格安なCUDAできるノートPCはどれ？

倍精度付けてくれたから、Dynamic Parallelismもてっきり対応してるものかと。。
CUDAのページにはGTX TITAN CC 3.5って書いてるしね。

それに以前、nVIDIAはパフォーマンスは別としても、
同じプログラムが動くことが重要だと訴えてたはずだ。

納得できん。

最近はグラ用とコンピュート用で分化させる方針みたいだよ。
残念ながらね･･･。

Hyper-Q Dynamic palallerlismはついているようだが？
http://pc.watch.impress.co.jp/docs/news/20130219_588387.html

>>241
>なお、既報の通り、HPC向けにGK110では「Hyper-Q」および「Dynamic Parallelism」という機能も追加されている。

http://news.mynavi.jp/articles/2013/02/20/geforce_gtx_titan/index.html
>しかし、GTX TitanではKepler GPUの新機能であるDynamic ParallelismとHyper-Qという機能が省かれている。

>>240
Fermiで懲りたんだろな
VGAにGPGPUで使う機能をてんこ盛りして爆熱にしてもしょうがないし
高く売れるGPGPU機を安く売っているVGAで代替できるじゃ金儲けできないし

TitanはXeon Phiが出てなかったら出なかったかもな

Fremiで動いていたコードでKeplerで動かすとカーネルが起動しないらしくて、
計算できないんだが、Keplerだとなにか気をつけることってあるのかな

Fermiで動いてたのはたまたまであって、
注意深く解析すると、そのコードはスレッドの起動順序などによって
アクセス違反を起こす、とかね。

もしかしてスレタイって下らないのクダとCUDAを掛けてるの？

Xeon Phiって結局まだ個人では単品で入手できないね

>>247
すげーーーーーーーーーーーーーーーー
良く気付いたな

CUDAはCUDAらない。

今夜は良く冷えますね

おいおい、TitanでDP使えるらしいじゃないか！
訂正記事書かれてるぞ
http://headlines.yahoo.co.jp/hl?a=20130220-00000025-mycomj-sci
2/28追記:記事掲載当初、「GTX TitanではKepler GPUの新機能であるDynamic ParallelismとHyper-Qという機能が省かれている」と記載させていただいておりまし
たが、GTX Titanでもこれらの機能をサポートしていることが明らかとなりましたので、当該部分ならびにそこに付随する文章
を修正させていただきました。

>>252
キター
DPって、倍精度と紛らわしいよね

マジかよｗ
この訂正記事に気づいてない人の分だけ売上落ちるぞこれ

このスレ見てて良かったと思う瞬間
↓

│　 ↑
└─┘
おらっしゃあぁぁ！！！
　∩∧ ∧
　ヽ（ ﾟДﾟ）
　　　＼⊂＼
　　　 Ｏ-､　）〜
　　　　　　∪

計算ミスしないなら欲しい

計算エラーって、確率的に起こるもんなの？
それとも同じところで再現？

ハードウェア的には同じところだとは思うが、プログラム上の同じところで
再現してくれるとは限らない。

計算ミスするコアは必ずミスしてくれるかっていうと…詳しい人教えてください

メモリエラーの場合は確率的じゃない？

Dynamic Parallelismあり＋Geforce並のエラー精度
Dynamic Parallelismなし＋Tesla並のエラー精度

処理内容にもよるけどこの2択なら前者のほうが欲しい。
なんていうかGPGPUで何かさせるときは計算ミスは織り込んだ上で作る感じだし。
そういう意味でもTitanはおもしろいと思うけど、なぜにDynParal不可という報道を放置したし・・

その記者がIntelから金渡されてたりして

品薄もんだからよかったものの、
大量に捌きたいもんだったらNVidiaから訴えられてもおかしくないレベルｗ

nVIDIA Japan公式Twitterだと使えるってアナウンスが
随分前にあったけど、使えるで確定なのねｗ

くそ、DPない情報に踊らされて買いそびれたじゃねぇか･･･
まだ店で売ってるのかな？

どさくさに紛れてGPUDirectも対応してないかな？

>>261
ECC無してTesla C1070、C2070、C2075、GTX580を使ってきた経験上、
計算エラーが起きたことは無い。
TSUBAMEのように大量に組めばECCも重要だろうけど、
デスクトップレベルなら気にすることは無い。
エラーがでたらやり直せばいい。
それが嫌ならK20を買えという事だ。

>>267
＞エラーがでたらやり直せばいい。

ほんと、ただそれだけのことだよね。
処理結果を一定間隔で保存しとけば、
万一エラーが出ても、少し前からやり直すだけでいい。

ECCなんかにダイやサイクルを浪費するより、
よっぽど効率的。

そもそもエラーだと気付ければいいけどな
いつの間にか数値の一部が化けてたら気付かないぞ

>>269
あ、そういえばｗ

まぁ、エラー「訂正」までせずに、「検知」に留めておいて
ロジックを減らすのはいいんじゃないかなぁ。

シミュレーションに使っているけど、計算ミスが問題になったことはないなぁ。
あ、そもそも精度が要るところはCPUで計算してたわ。

>>269
大事な計算が必要なら信頼性の高い
システムを使えばいいだけ。
指数部が化けたら気づきやすいだろうし、
仮数部が化けても気づかないならそもそも
それほど感度がいる解析ではないので重要ではないよ。
そもそも研究なりで回す場合はいくつも計算するだろうから、
明らかにおかしい値は人間のフィルタがかかるでしょ。

そんなに心配なら、同じ計算を別のボードでもやって比較すれば良い。
同じボードでやっても、あまり意味は無いぞ？

計算分野は狭いようで果てしなく広い分野だからな
エラー検知がどの程度重要かどうかは分野ごとに異なるから
自分の身の回りだけの話をしてもしょうがないだろ

これでTeslaと同じくらいの価格だな
http://akiba-pc.watch.impress.co.jp/img/ah/docs/590/125/html/s3201.jpg.html

CUDAとOpenCV併用するとき、.cuファイル内でOpenCV使うことってできないの？
コンパイル時にインクルードファイル周りでエラー吐いたからとりあえずOpenCV周りだけ.cppファイルに隔離したんだけど。
後々不便そうだから対策があれば教えてほしい。

>>276
CUDAのカーネル関係はlib化して、
ビルド時にリンクする方法を取ればいいんじゃね？
NVCCってそんなに賢くないから、CPU周りで遅くなるし。

ふと、CUDAで自分で書いたカーネルを実行しつつ、OpenCV の cv::gpu:: の機能も同時に使うと、どうなっちゃうんだろう、と思った。

>>278
どちらも実行されるよ。
普通のマルチスレッドのプログラミングと
考え方は変わらん。

http://news.mynavi.jp/series/nvidia_kepler_gpu/001/
【連載】
GPUの実効性能を上げるDynamic ParallelismとHyper-Q

>>280
安藤先生の記事や。
ありがたや〜。

http://www.eevblog.com/forum/projects/hacking-nvidia-cards-into-their-professional-counterparts/

おいお前ら
GTX690がレジスタ変えるだけでTesraになるらしいぞ

半田班ちょっと人柱頼む

>>282
どうやって発見したんだよ！ｗ

英語読める割に綴り間違うんだな

>>282
レジスタってGPUのレジスタじゃないぞ。
抵抗のレジスタだ。

しかし、興味深いなあ。

>>285
半田班ってんだからわかってくれよ
>>284
携帯だったんだよ興奮してたんだよ

チップ自体は同じだから理論上可能というのは分かるんだけど、
よもやこんな方法でイネーブルにできるとはｗ

>>283
俺が知りたいよ

早めに買っとかないと対策されそうな気がするけどこのチップ抵抗は怖いよな。。

>>287
しかも690のがクロック高いから性能いいらしいぞ(;´Д`)
買ってみてやってみて

>>289
マジかよｗ

あー、久々おもしろい話題でわくわくするｗ
やっぱハードウェアはこうでないとなｗ

マジレスすると、quadroやteslaでの構成と同等なものがないから、
ドライバ側が柔軟な作りに立っていないと不具合が出るだろうな。
例えばハードウェアのモニタリング用の石が載ってなかったりするだろうから。

nvidia的には、geforceがteslaになったところで、購買層が違うから放置だろうが、quadro化はマーケティング的にやばいので塞ぎにくるだろうな。

>>291
ただそれはドライバを更新しなければいいだけの話だよね。

>>292
しばらくは放置なんじゃなかな。
さすがにドライバだけでは対応しづらい気がする。
基板のリビジョンが変更になった時から使えなくなると思う。

K10化より、TITANをK20Xにしたい。

おっ
また発見されたのかよ
むかしクワドロ化の方法が発見されたのに、まだまともな対策してなかったのか

これ、カーネルドライバがボードに直接プロダクトIDを聞きにいくからカーネルフックじゃ駄目って書いてあるけど
ドライバ改造できればﾜﾝﾁｬﾝ有るってことじゃね。

GeForceをQuadro/Tesla化出来るからって、Quadro/Teslaの売り上げが
減るとは思えない。

tesla化するって言っても、gk110じゃないしな。gpgpuやってる人からすれば、
そんなにありがたみがないんじゃね

NVIDIA CUDA Gets Python Support | techPowerUp
http://www.techpowerup.com/181585/NVIDIA-CUDA-Gets-Python-Support.html

これネイティブ対応ってことなのかね。
パイソンってそんなに人気あったんだ・・・
ruby人気は国内限定？

うん

Visual StudioもPython使えるようになるパッチが出てるし
OpenCVもPython対応(一応)してるしLibreOfficeもマクロでPython使えるし
そしてCUDAまでも…ということで俺の身の周りではPython株が右肩上がり。

ネット上のサンプルが2.x用のソースだと動かなかったりして躓いたこともあったり
今でも躓いたりしてるけど、これから3.x用の情報がもっと増えてくれると嬉しいと思う。

SL#（えすえるしゃーぷ）とは、GPUで実行されるプログラマブルシェーダーを、超高級言語である
C#で書けてしまうという夢のようなオープンソースのフレームワークである。
http://monobook.org/wiki/SL_Sharp

C#ｲﾗﾈ

DXエラーだったのでdirectXかと思いここで聞きます
色々ググったのですが解決に至らず

http://i.imgur.com/XFoHQke.jpg
まず黒画面突入後にビープ音1回
待機マークでEsc押すと�@が発生、続いて�Aが発生

某ゲームのランチャー起動後にこのエラーが出ます
誰かご教授下さい

スレチだし
ゲームの事は開発元に聞け
あとDX8のレンダリングにすら対応してないクソPCを捨てろ

なんでDirectXでこのスレなんだ？

本出るの久しぶりじゃない？
http://www.amazon.co.jp/dp/4061538209

>>308
おお〜、CUDA5本か！
注目だな〜。

って、誰が書いてるかと思いきや、
伊藤智義さんって、「栄光なき天才たち」シリーズの原作だった人じゃまいかｗ
「宇宙を夢みた男たち」は感動した！
まさか、こんなところでまた名前を見ることになるとは感慨深い。

http://next.rikunabi.com/tech/docs/ct_s03600.jsp?p=001216
つか、この人、すごい人なんだな！
原作やってたときは東大生のときだったのか･･･。

宣伝乙

もしかして、東大の初代grapeに関わった人？

そだよ！

もしかして、TitanってP2P使えない？
cudaDeviceCanAccessPeerで0が返ってくるのは想定の範囲内だったけど、
cudaDeviceEnablePeerAccessするとエラーになってしまう。


OS:CentOS 6.4 64bit
Driver:313.30
MB:P9X79
Titan 2枚をそれぞれPCIex gen3でx16接続

>>157
書籍「CUDA BY EXAMPLE」の「5.3.2 ドット積の(誤った)最適化」の項には、
真逆な事が書いてあって悩んでいる。

この項の概要は、
「__syncthreads()」をif文の中にいれたときに、
スレッドの一部が「__syncthreads()」に到達しないものが発生して、
プロセッサが事実上ハングするというもの。

しかし、サンプルプログラムを作成して実験しても、上記事象は発生せず、
正常終了する。
（「cudaGetLastError()」の結果も「cudaSuccess」が返る。）

書籍の内容は正しいのだろうか？
知っている人教えてください。

>>314
多分「CUDA BY EXAMPLE」が間違っているか、CUDAのバージョンが新しくなったことによって変わったのかもしれないが、基本的に__syncthreads()はスレッド間のフェンスみたいなものだから。SIMDの概念で考えればわかると思う。

PTX ISAのドキュメントには以下のように書いてある。
Barriers are executed on a per-warp basis as if all the threads in a warp are active. Thus, if any
thread in a warp executes a bar instruction, it is as if all the threads in the warp have executed
the bar instruction. All threads in the warp are stalled until the barrier completes, and the
arrival count for the barrier is incremented by the warp size (not the number of active threads
in the warp). In conditionally executed code, a bar instruction should only be used if it is known
that all threads evaluate the condition identically (the warp does not diverge). Since barriers
are executed on a per-warp basis, the optional thread count must be a multiple of the warp
size.

要は、_syncthreadsはwarp内の一つでも実行されたら
そのwarpの全てのthreadが実行したのと同じと看做される。
だから同一block内の全てのwarpそれぞれで、
いずれかのthreadがアクティブな状態で_syncthreadを実行すれば辻褄はあう。

逆にMPIのbarrierで辻褄を合わせるように、全threadでbarrierを実行しようと
分岐の両側に_syncthreadを置いて、実際にwarp内で両パスにthreadが別れてしまうと
困ることになるだろう。

>>315 >>316 >>317
PTXレベルでの解説ありがとうございます。
なるほど、「CUDA BY EXAMPLE」を読んだだけの頃は
「MPIのbarrier」のような仕様をイメージしていましたが、
そうではないんですね。
「__syncthreads()」の分岐に入らなかったスレッド用に、
別の「__syncthreads()」を書く必要は無いということか。。。
SIMDの概念ならその仕様でないとダメなんですね。

ものすごく初心者丸出しな質問で申し訳無いのですが、質問させてください。
CUDA5.0をVS2008で使うにはどういった手順が必要になるのでしょうか？
ネット上だと2010や2012についてばかり見つかって2008について見つけられませんでした。
どうかよろしくお願いします。

マルチgpu環境での使い方が分からず困っています。
何か良い資料は有りませんでしょうか。

ハードウェアとCUDAのバージョン依存があるんでね？

>>320
つ http://www.meriken2ch.com/programming/merikens-tripcode-finder
マルチGPUをサポートしてて、ソースも付いてるよ。

compute capability 3.0 cuda 5.0　です。

>>320
1GPU毎に1CPUスレッドを割り当てて、それぞれのスレッドで
cudaSetDevice()を実行すればいいって話じゃないの?

>>320
デバイスドライバとの関係でメインスレッドでcudaSetDevice()すると巧くないので、
>324のようにcreate_pthread()した先でcudaSetDevice()するよろし。

cudaSetDevice()を呼び出した後はシングルgpuの場合と同様なんですよね。

gpuが二つある場合に単一スレッドで cudaSetDevice(0)を呼び出して作業した後に
cudaSetDevice(1)を呼び出した場合はどうなるのでしょうか。

二度目以降は無視されるよ。

各スレッドでcudaSetDevice()を呼び出して各スレッド毎に各gpuを割り当てますが、
複数のスレッドでcudaSetDevice()を呼び出さずにカーネルを実行したらどうなるのでしょうか。
また、複数のスレッドでcudaSetDevice()使って同一のデバイスを指定はできるのでしょうか。

>>328
何でそんなことしたいのかさっぱりわからん。

自分で試して確認しようとしていないようだけど、そういう環境がないのか？
何の為に学ぼうとしてるのかしらんが、そうやって無限に問い続けるつもりか？

環境がなくて試せないんです。すみません。

数千円でボードかえるのに環境がないとか甘えだろ。

>>326
CUDA C Programming Guide 5.0
3.2.6.2 Device Selection

環境がなくて試せないのに、なんでそんなに重箱の隅みたいなこと聞くのか

環境の構築(購入)に踏み切ろうか悩んでいるとか

マルチGPU環境とかこだわらなければ1万円以下でできるのにな。

現在CUDAで画像処理を行っているのですが，GPUを長時間使用するような
プログラムを走らせると，「ディスプレイドライバの応答が停止しました」または
PCが落ちる，という事が頻発します．
現在は，ソースコード上で使用するブロック数とスレッド数の数を減らすことで工夫して対応していますが，実装上不便です．
上記に対して，ソースコード上で使用するスレッド数とブロック数を制限せずに，解決する方法は何かございますか？
環境はWindows7,CUDA5.0，VC2008，GT610です．
宜しくお願いします．

タイムアウトの設定をレジストリで変えた？
たしかwin7だとデフォで1秒か2秒くらいになってて、それ以上カーネルが動きっぱなしに
なると「ディスプレイドライバの応答が〜」になるので変更する必要がある。

可能ならディスプレイ表示用（ショボいので十分）とコンピューティング用GPUと2つ接続したいところだね。

最近はCPUオンチップGPU+CUDA用GPUという構成を組みやすい。

>>338-340
ありがとうございます．
タイムアウトの設定は最後の手段としたいと思います．デフォルトでは２秒だそうです．
表示用と処理用のGPU二つ用意するのはいいですね．
残念ながら，このPCのBIOSではオンボードとGPUを共存できないうえ，
GPU二枚刺しも出来ないので，新しく機材を調達しようと思います．

タイムアウトが最初の手段だと思うけど。。
それで不便だったらお金のかかる対策に移ればいいわけだし。

GT610なんて使っている時点でCUDAの意味がないがな。
スレッドブロック数のチューニングにすらならん。

>>340
今はiGPU搭載CPUが普通で、それがCUDAお手軽構成って感じだよね

>>341
お薦めはAMDのAPUだよ。
値段が安いのにCPU部はIntelの高性能のi7並みの性能で、GPUはGT610より性能が良い

linuxでやれよハゲ

>>344
iGPUでCUDAができるかよ。やるならOpenCLだ。
AMDのCPUがi7並であるわけ無いだろ。
スレ違い。

GPGPUはHSAでAMDが主流になるって感じになってきたよな
一部のスパーコンピューティングでCUDAがつかわれ、それ以外はHSAのAMDになるのかな

>>347
性能が桁違いだから仕方がない。自作のアプリではTITANは
7970の6割ぐらいの性能だし。CUDAとOpenCLって
やってることはほとんど同じだから、移行するのもそんなに
難しくなかった。

あとSouthern Islands系のAMDのGPUは、CUDAに比べると分岐処理のペナルティが
はるかに小さいのでプログラミングしやすいというのもある。
(EvergreenやNorthern Islandsは共有メモリが遅かったり
ベクトル型を使わないと性能が出なかったりするのでこの限りでない)

()

>>349
お、GCNはだいぶいいんだね。

teslaでも使える温度計測ソフトはないものか

>>352
何の温度計測なんだ？
ボードの温度なのか、何かの測定器に連動するものなのかで答えが違うと思うが。

チクショウ、GTX780は倍精度切られているらしい。
Titanに逝ってくる。

またどっかの抵抗取り替えたりするとTITAN化するような仕掛けが潜んでたりして

TITANをK20xにしたいんだけど

TITANにECCつけたもの（SMXを一つくらい品質保証のため削るかもしれないが）をK20の後継として出るんじゃないかな？
今年のSCあたりで発表されそうだ

ゲフォがFermiが倍精度1/4でKepler1/8で結局性能悪化って酷くない?
ラデへの乗り換えを検討したほうがいいかな?

CUDA使ったフリーウェアを作ってるのに、
げふぉに将来性がなくなったら詰む。

>>359
CUDA で計算している部分だけ OpenCL に変えれば済む話では？
どうせモジュール化してるんでしょ

>>360
そうかも。
OpenCLってデバイスごとにコンパイルが必要なんでしたっけ?
実行時コンパイルとかよくわからんです。
理想的にはバイナリを一個作っとけばどこでも使えるといいのですけれど。
入門書読めばその辺の話は載ってますか?

スレ違いスマソ。

>>361
その辺の話は全て入門書、あるいは入門ページに載っている

>>362
さんクス

>>360
カーネルだけなら簡単だけど、メモリ周りとか考えると簡単にはいかないよ。

>>364
CUDAで書く時はグローバルメモリへのコアレスアクセスとか
シェアードメモリ使ってチューニングしてるけど、
OpenCLはそのへん配慮できるのかな。

>>365
普通にできるよ。

>>364
それは「げふぉに将来性がなくなったら詰む」ほどの事か？

結局、アーキテクチャ毎に最適なコードを書かなければならない。
そのように書かず、可搬性を意識しすぎるとGPU(などの並列プロセッサ)を使う旨みがない。
なんというジレンマ。

つまり、コンパイラがまだまだ馬鹿ということか

teslak20って一番安い所ではいくらで買えるんだろ？
ヤフオクで三年間分有効の保証が残ってる奴が30万なんだが安いのかな？

http://news.nvidia.com/t/213824/48003696/16930/10/
nvからkayla(armようcuda開発機)のおすすめメールが来た

tegra3とGPU別ｗ

>>370
titanでいいやん。
ヤフオクで買う程度なら個人利用なんだろ？
会えてteslaを買う理由がないわ。

マルチGPUするなら、GPUDirectの有無は大きいよ>TeslaとTitan

titanでもGPUDirectはついてるよ。ないのはRDMAの機能の方だ。

CUDA5.5はどうですか

いいえ、どうではありません。

MPIプログラムのプロファイルを取れるようになったのはでかいな。
https://developer.nvidia.com/cuda-toolkit

cuda-gdbのステップ実行で、
コンソールが返ってこなくなる事がありますが、
原因を確認する方法はありますか？
（普通のgdbも使った経験が無いのですが。。。）

海外amazonでGK208なGT640かったら
やっぱSM35だった

TOP500でxeon-phiがトップなたぞ

正直微妙だ

　　　　　　Efficiency (%) Mflops/Watt
Tianhe-2　 61.68　　　　　　1901.54
Titan　　　　64.88　　　　　　2142.77

Starting LU Decomposition (CUDA Dynamic Parallelism)
GPU Device 0: "GeForce GT 640" with compute capability 3.5

GPU device GeForce GT 640 has compute capabilities (SM 3.5)
Compute LU decomposition of a random 1024x1024 matrix using CUDA Dynamic Parallelism
Launching single task from device...
GPU perf(dgetrf)= 3.358 Gflops
Checking results... done
Tests suceeded
------------------------------------------------------------------------------
starting hyperQ...
GPU Device 0: "GeForce GT 640" with compute capability 3.5

> Detected Compute SM 3.5 hardware with 2 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.050s

>Detected Compute SM 3.5 hardware with 2 multi-processors

SMX2器ってこと?

そう
deviceQueryDrv.exe Starting...

CUDA Device Query (Driver API) statically linked version
Detected 1 CUDA Capable device(s)

Device 0: "GeForce GT 640"
CUDA Driver Version: 5.5
CUDA Capability Major/Minor version number: 3.5
Total amount of global memory: 1024 MBytes (1073414144 bytes)
( 2) Multiprocessors x (192) CUDA Cores/MP: 384 CUDA Cores
GPU Clock rate: 1046 MHz (1.05 GHz)
Memory Clock rate: 2505 Mhz
Memory Bus Width: 64-bit
L2 Cache Size: 524288 bytes

Dynamic Parallelismも使えるの？

バス幅とL2キャッシュ典

>385
hyperQとDynamic Parallelismは使えるようだね>382

GK110でL2が1536KBだから512KBって結構でかいな

GK208きた

cudaGetDevicePropertiesで
maxThreadsPerBlockが1024とでたので、
kernel<<<32, 1024>>>()
とやったらKeplerはおKでFermiではだめだった。
基本的にはKeplerでも1ブロックあたり512スレッドが上限だったっけ?

使えるようにするまで2時間かかった

CUDAは導入がメンドイね。
HSAはよ。

>>392
Linuxだとホントしんどいよね。
何が悲しゅうてドライバインストールのために
カーネルを再コンパイルせにゃならんのか。

Linuxでの導入はなおさら OpenCL の方が楽だ

プログラムはめんどいが

Linuxのカーネルのコンパイルといってもモジュールを
コンパイルするだけなのだけれど、それってそんなに面倒?
リブートも必要ないし。

GUIを使わないモードでリブートして、
カーネルコンパイルして、
もう一回リブートしないとならないと思ったけど。

GUIのサービス止めてモジュールをrmmodして
nVidiaが提供するスクリプトを
使ってモジュールだけをコンパイルして
insmodしてからGUIのサービス再開するだけだよ。

Bumblebee はどうなの？

普段は CPU 内蔵のグラフィック機能を使って、
CUDA やる時、正確に言えば CUDA の結果を OpenGL でレンダリングするアプリの場合だけ
Bumblebee で nVIDIA カードの方を使うってできるの？

もしできるのなら Linux に乗り換えてみようかな。
他のことで Windows がちょっと使いづらくなってて、
でも踏ん切りつかなくて迷ってる。

>>395
適切なOSとGCCのバージョンをそろえるのが面倒。
ちょっとでもNVIDIA推奨環境と違うと
Getting Startedだけ読んででインストールするのは絶対無理
だと思う。

>>399
別に面倒でも何でもない。
CUDAならメジャーなディストリならほぼ動く。
面倒なのは単にLinuxの知識がないだけ。
例外は組み込み系で使う場合。
そもそもCUDAの開発やるのに、

sudo bash NVIDIA-Linux-x86_64-xxx.xx.run

できない奴なんていないだろ。

>>400
その後 nouveau が邪魔だって言われて
「cuda nouveau install」でググるんですよね

多倍長整数の計算におすすめのライブラリとかある？

>>400
X79の最新のチップセット使うとUbuntu 12以上じゃないと
動かなくてですね、そいつのデフォのバージョンのgccだと
CUDAが対応しないんですわ。

>>403
単にnvccのベースがgcc4.4までだからだろ。
Ubuntuならソフトウェアセンターでインストールしてalternativeで切り替えればいいだけ。
これはCUDAに限らず、インテルコンパイラでも必要。

>>404
はいはい情弱ですみませんねえ。
みんながみんなGetting Startedだけ読んでインストールできたら
「Ubuntu 12.**でCUDA 5.0入れてみた」系のブログを書く人も読む人いないですよ。
ふーんだ。

CUDA 5.0がだめなら、5.5 RCを試せばいいじゃない。

cuda5.0のgccは4.6だろ？
それよりnVidiaはFedora16のサポートが切れてることについてどう思ってるんだろう。

Ubuntuにせよ、FedoraにせよNVIDIAは最近Linuxに対してあんまりやる気ないな。
リーナスに中指立てられて批判されたからかな？

windowsが最高の開発環境だし

GPGPUはAMDになってしまったから、Nvはやる気でないだろ

はい？

適切なスレが分からなかったので、ここで質問します。
今のCUDAはCUDA CとOpenCLでバックエンドが共通になっていると聞きましたが、
今もしくは将来のCUDAで、HSA(Heterogeneous System Architecture)を
共通のバックエンドで動かすことは技術的に可能ですか?

nvidiaに聞け
公開資料にない事の予定問われても
スレの住人はnvidia関係者な訳じゃないし
関係者が居たとしても、2chで非公開の予定情報の可否なんか答える訳ないだろ

メジャーになればそれなりの対応もあるだろうが
影も形もないものを・・

技術的に可能かどうかと言われれば可能でしょ。
メモリ空間が共通化されれば、GPUの演算器がCPUのSIMD演算器のように扱えるわけだし。
ただCUDAである必要があるかどうかはNVIDIAが判断するんじゃないか？

N×1行列とM×N行列を計算して結果をテクスチャに書き込むという単純な処理で
これを合計512スレッド(Mに関して並列化)で実行しているんだけど(N=3000 M=512)
各ブロックを16×16スレッドの2ブロックよりも
各ブロックは16×1スレッドの32ブロックのほうが2〜3%速度が速いという不可解な結果が出てしまっている
何でこんなことが起こるんだろう

適度に粒度下がってスケジューラの効率上がったとか？

local memoryにレジスタが溢れているとか

こういう予測しにくい挙動こそがGPGPUのクソなところ。

そういえばNVIDIAはPGIを買収したらしいね
ここに書くことでもないかも知れないが

>>416
Nを並列化せずに本当に512スレッドしか使ってないんだったら、16warp*32threadより
32warp*16threadの方がわずかに効率が良いというだけの話じゃないのか？分岐やなんかで。

>>416
メモリの配置次第で不可解でも何でもないと思うが。

【AMD涙目】デファクトスタンダードへの道を突き進むCUDA・・IBMやGoogleも陣営に加わる
http://engawa.2ch.net/test/read.cgi/poverty/1375883942/

デファクトスタンダードへの道を突き進むCUDA　IBMやGoogleも陣営に加わる
http://kohada.2ch.net/test/read.cgi/pcnews/1375893221/

性能の低さを政治力でカバーするのか？

まぁニュースサイトへの宣伝広告費でカバーしてるところより良いと思うよ

カーネル関数内で，乱数が生成された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で配列の先頭を示してるから，
実質同じ事だと考えてたんだけど，違うの？

>>427
今手元にないから記憶で書くけど、D_cの型がdouble []じゃないからだと思う。
D_cの定義をdouble D_c[is_110 * is_110 * is_110]にしてみたらどうなる?

>>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__でやるのは無理があるのかなぁ？

>>429
うーん、よく判らんな。
少なくとも、cudaMemcpyFromSymbol()はCのコードとしては特殊(トリッキー)な仕様だから
使い難いんだよね。まるでプリプロセッサマクロのように。
だから私のところではリファクタリングの結果、使うのをやめてしまっている。
変に悩むくらいなら、代替手段を考えたほうがいいかもよ。

>>430
特殊なのかー．
アドバイス通り，普通にホストでMallocして，カーネル関数にポインタ渡す方式にします．
ありがとう！

ドライバをアップデートすると演算性能上がりますか?落ちますか?

__global__の関数内でグローバルメモリの内容を一気にコピーしたいんだけどそういう方法ってある？
一要素ずつやった方が無難？

>>433

コピーの手前でカーネル切って、新たに
__global__ copy(double *a, double *b)
{
int tid = blockIdx * blockDim.x + threadIdx.x;
b[tid] = a[tid]
}
を実行するのじゃだめ?

あるいはダイナミックパラレリズムで何とかなるのかな。

カーネル分けんでもindexの生成を工夫すれば同じことでしょ。
それがDtoDなら。DtoHやHtoDなら設計を見直すべき。

ツールキット5.5プロダクションリリースってどういう意味ですか?
ベータ版とかじゃなくて正式に使えるってこと?

curandGenerateUniformって1.0を含むけど1.0を含まない乱数生成ってないの？

それ使う所の式変形で対処できないの？

>>435
ブロックまたいで同期する必要がある場合は?

.cubin ファイルをエディターで・・・と
公式スライドにあるのですが、本当ですか？

cudaErrorInvalidValueってカーネルがエラーはいたんだけどどういう状況でなる？

>>441
詳細はAPIマニュアル見てねだけど、
ホントにカーネルで出てる?
APIの引数を間違ってるんじゃない?

実行はできるし結果もそれっぽいけど画面が不安定になるなあ

>>443
カーネル実行時間が長すぎると画面が真っ暗になる場合がある。
制限時間はレジストリをいじって変えられたはず。

>>444
レジストリはいじったあとなんだけどそうなる

>>445
メモリアクセスが間違ってても落ちることあるですよ。

Geforce GT 530 で CUDAインストーラー(5.5)が「対応デバイスが無い」とかで失敗するので
古いドライバに変えてみたり再起動繰り返したりしたのですが、
developer.nvidia.com/cuda-gpus
に、GT530が載ってませんでした・・・。
マシンに「nVidia GEFORCE with CUDA」のシールあるし、GPU対応のソフトも動いているのですが
開発用としては使えないってことでしょうか？

セットアップがミスってんでしょ。

セットアップのミスの原因ってなにかありますか？
展開後は高速インストールかカスタムか選ぶだけだし、
どちらを選んでも失敗します・・・。
プログレスバー見てる感じ、Toolkitのインストール中１０％くらいでエラーが出ます。

c:\NVIDIA\CUDAへの展開でかなり時間かかりますが、
インストール失敗すると、この下のインストーラー各種全部削除されて
最初からやり直す仕様なんですね・・・。（成功しても消えるのかもしれませんが）

ちゃんとリリースノートやインストールノート読んでやってないところ
あとはほんとに５３０のせいかどうかを他のグラボ（nvidiaね）に変えて切り分けして原因しぼっていくしかないだろ。

可能性だけなら
グラボ、PCパーツの不良
DLしたソフトウエアの不良
システム不良
などたくさんあんだからさ

俺も今5.5のToolkit のインストールでずっこけてる。
Windows XP service pack 3
Quadro FX 580

Toolkit と sample以外はカスタムでインストールできたんだけど、
Toolkitが8割ぐらい行ったところで失敗しました、てなる。

インスコディレクトリにいくらかコピーできてるみたいなんだけど、
環境変数なんかは設定されてない。

リリースノート見てもGUI使うかmsiをシェルで実行しろとしか書いてなくて
困ったぽよ

もう朝だお。。。
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よく分かってないのか・・・。
スレチなら他行きますので・・・。

ここは俺の日記かお・・・。
VS2012でプロジェクト→プロパティ→構成プロパティ→デバッグ→環境
を選んで編集モードへ。
ここで「マクロ>>」をクリックすると設定されたCuda用環境変数もちらほらありますが
$(CudaToolkitDir)の値がからっぽでした。
普通に動かせてる方、この辺の弄り方教えてください。
設定するパスはこれですかね？→C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.5

>>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\

>>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の貧弱環境ですが
ひと通り習得できてきちんと開発できるようになったら
新しいグラフィクスカード買わないとな・・・。

追加質問です。

ビルド中に警告が大量に出ますが、手動インストールによるものなのか判断できません。
> \include\math_functions.h :
> warning C4819: ファイルは、現在のコード ページ (932) で表示できない文字を
> 含んでいます。データの損失を防ぐために、ファイルを Unicode 形式で保存してください。」
> というのが、たくさん出ます。

普通にインストーラーが成功した方も同じようになりますか？

エディタで開いたところ、
math_functions.hはUTF-8N、
cuda.hはSJISと表示されました。

>>456
> > warning C4819: ファイルは、現在のコード ページ (932) で表示できない文字を
> > 含んでいます。データの損失を防ぐために、ファイルを Unicode 形式で保存してください。」
> > というのが、たくさん出ます。
>
> 普通にインストーラーが成功した方も同じようになりますか？

CUDAにかぎらず、海外のコードは、よくその警告がでる。
自分は、プロパティマネージャで、C4819を無効にしちゃった。
あと、自分が書いたコードはUTF-8(BOM無し)にしてる。

すみません。
まだまだ初めたばかりの初心者なのですが、
質問させてください。

Tesla C2075を使っていて、DeviceQuery.exeでスペックを見たところ
Total amount of global memoryが1.2GB 程度しかありませんでした。
仕様では5GB程度あるはずなのですが・・・。
OSが32bit(win7)であることが関係してたりしますか？

>>458
そうかも。
GPUのメモリはPCのメモリマップ上に割り付けられるけど、32bitでは4GBしかないので。
メモリマップは、Windowsのデバイスマネージャで「表示」→「リソース (種類別)」の「メモリ」で判る。

64bitでもVRAM全部を割り付けたりしないよ
多分、互換性の為だろうけど

Windows 7 service pack 1　64bit
Geforce TITAN
の構成でも試してみたが5.5インストールできなかった。

Windowsでインストーラ使って5.5入れられた人いますか?

>>454さんはするっとインストールできちゃった感じなんでしょうか。

連投すまん。
Visual Studioが2008 Express しか入っていないのが原因かな?

>>461-463
何の問題もなくインストールできたよ

>>462
>>463
toolkit 5.0, 5.5 はVS StdかProが必要だよ。
4.0, 3.0 は忘れた。

おかげさまでどうにか自作ソースを5.5でコンパイルするところまでこぎつけました。
インストーラでToolkitとSample 以外を先にインストールして、
Toolkitを単独でインストール、ファイルのコピーが終わったところでこけるので、
>>454の手順で環境変数を設定して、
Sampleはインストーラを立ち上げて展開されたものを所定の位置にこぴぺしました。

CUDAでホストのmain関数（グローバル関数の呼び出し込）で時間計測にclock()ってつかって問題あったっけ？

>>467
OpenMPと一緒に使うとCPUタイムがでちゃって正確な経過時間を計れない場合があります。
精度が1msecというのも場合によっては足りないかもしれません。

カーネル内でmallocとかnewとかってできる？
できるとしたらどのメモリが確保されるん？
グローバル？ローカル？

グローカル

ﾜﾛﾀｗ

CUDA C Programming Guide の
B.18. Dynamic Global Memory Allocation and
Operations
にcompute capability 2.0以上でカーネル内のmallocが使える
みたいな事が書いてあるね。GPU側のグローバルメモリじゃないかな。

>>472
サンクス！
GPUのグローバルメモリのヒープ領域(デフォルト6MB)に確保されるのね

>>468
CPUの方でも処理（マルチスレッドではなく）してるからそういう方式では問題ないってことでいいのかな？

>>474
だったらたぶん。大丈夫

カーネル関数でグローバル変数を使う方法ありますか？
渡したい値がたくさんあるのでカーネル関数の引数がかなり多くなってしまい面倒なのですが

__device__付けたらいいんですね自己解決しました

グローバル化するより構造体作って引数減らしたら駄目なん？

>>476
構造体作ってポインタで渡す。

構造体渡しはつまり全メンバーコピーだからね

cudaするのに、１１インチくらいのディスプレイのノートPCってあるの？
kakaku.comで調べても、14インチくらいのしかないけど、だれか教えてください。

なぜそこでノート限定なのか

ちっこいノートはインテルHDグラフィックスなんでねーの?
と適当に書いてみる。

いい加減、グラボを換装できるノートが欲しいよな

普通にあるけど割高だよ。

以前は、IONの載ったちいさいノートPCでCUDAできたのに・・・

どこかが展示してた外付けグラボが欲しいんだが、発売する気配なしｗ

外付けグラボは過去に何製品か出たけど
あまりにもニッチすぎる、発熱や性能の問題でどれもヒットしなかった

>>488
そうなのかぁ･･･。
残念。

くだらないしつもんですが、
1デバイス内のグローバルメモリ内でデータをコピーするのに、
cudaMemcpyDeviceToDeviceを使うのと、
カーネルを書くのではどちらが速いですか?

簡単なんだから、実測してみたら?
コピー専用のカーネル作ってたらロードで不利だけど、メインの処理に抱き合わせるならきっと速いよ。

>>481
一台GPU積んだlinux機作ってsshできるようにすると快適だよ

最近cudaを触りはじめた者なのですが質問があります

多次元配列（2とか3次元）の中で指定した値に一番近い値を抽出する方法をcudaで処理するにはなにが一番適しているのですか？

色々ggってみて、とりあえずソートして絶対値をとって比較すれば良いのではと思い
バイトニックソートなどで試そうとしましたがソースコード付の物はどれも一次元配列の物ばかりでどうすればいいのかわかりません

それにその方法をとっても総当たりより速度が余計遅くなるのではないかとも考えましたがどうなのでしょうか？

抽象的な質問になってすみませんが誰か回答をお願いします

>>493
kd木構築してNNでいいんでないの
どう並列化するかは細かい状況による

総当りを（並列で）やるしかないと思うよ。あれをreductionって言うのかわからないけど。
たとえばMaxを求めるのにソートは必要なくてO(n)でできるのと同様に、
一番近い値を求める（だけ）ならばソートする必要はないはず。

ああ、多次元配列は変化せずに指定する値をコロコロ変えるなら最初にソートしておけば
その後はO(logn)でいけるね。条件次第。

>>494-496
返信ありがとうございます

とりあえず自分の知識が全く足りてない事が分かったのでもう少し色々調べてみます

あと、できれば496あたりの詳しい解説を知りたいのですがO(long)とはなんでしょうか？

>>497
O(long)じゃなく、O(logn)な
ロガリズムnのオーダー

>>497
O(n)とかについては「ランダウの記号」でググるかウィキペディると説明がみつかるはずです。
（ウィキペディアの記事は「一般的なオーダー」よりも前のところは正直俺はよくわからないけど…）

>>499
了解です
回答ありがとうございました

また分からなくなったら来ます

その多次元配列は、一次元配列に投影できないのだろうか。

>>501
(i, j, k)の三次元配列だったら
index = k * N^2 + j * N + i
で一次元になるね。

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]
ってスピルしちゃうんですが．
どこを確認すればいいんでしょうか

>>503
2.xと3.0は最大レジスタ63個までだから
処理少なくするかコード見直して不要冗長な部分を削るとか
あるいはレジスタ制限がゆるい1.xか3.5にする

>>504
63個までだったなんて知らなかった…．
18K Shared使ってるのと，三次元グリッドなんでCP3.5にしたら，ちゃんとコンパイル通って動作しました．
ありがとうございます！

カーネル呼び出しの行で実行がストップしてしまうのですが（次の行に進まない）、
原因がわかりません。
しかし、nvccの「-G」オプションを付けると、この事象が発生せず、
普通に動作します。
原因のわかる方、あるいは原因の特定方法がわかる方、
教えて下さい。

>>506
カーネル内で不正なメモリアドレスにアクセスすると
問答無用で落ちる場合はありますね。

-Gをつけた上で、カーネル実行の後ろに
printf("%s\n", cudaGetErrorString(cudaGetLastError());
と入れてみてはどうでしょうか。

>>506

>>507さんが言ってるように，まずはエラー原因を拾ってみると良いのでは？
30： Unknownが出そうな気はするけど．
デバッグ情報付与するとメモリ確保位置が異なるので，
デバッグ版では偶然に計算に影響を与えない領域にアクセスしていたという可能性も．

まず思いつくのは，ある配列の要素Nに対してthreadを生成しているならば，
そのthread数がNを超えて生成されるので（BlockDims*GridDims>N），
N以上のIdxのthreadが配列外アクセスしてるとか．
あとは，カーネル関数内でPrintf()して確認していくとか．

>>506
追記
CudaGetLastError()は，カーネル関数がエラー起こしてても拾わない事があるので，
そのまえのカーネルが原因という可能性もありますよ．
前に起動したカーネルがあるなら，その後に主要変数をmemcpyして調べるといいのでは？

回答ありがとうございます。

>>507
なるほど、メモリアクセスが怪しいのですね。

確認してみましたが、エラーとななっていない事がわかりました。
「cudaGetLastError()」の代わりに、「cudaPeekAtLastError()」でも
試したのですが、エラーとなりませんでした。

>>508
メモリ確保位置が異なるということで、「-G」を付けずに試してみましたが、
やはりカーネル呼び出しの行で実行がストップしてしまい、
エラー取得の行まで進みません。
（ログを入れて確認しました。）

配列のインデックスのバグの可能性があるのですね。
細かく確認してみます。

>>509
問題の起きているカーネルが、最初に実行されるカーネルになります。
カーネル実行の前では、「cudaMemsetAsync」やら「cudaMemcpyAsync」等を
実行していたりしますが、戻り値を確認しましたが、正常値でした。
（ストリームを使っています。）
非同期関数なので、すぐにはエラーがでないそうですが。
この辺も怪しいですね。確認します。

>>510
＞配列のインデックスのバグの可能性
それはもちろんあるんだけど，ThreadIdxが計算したいNを超えていないか？という事を言いたかったわけで．
例えば8,000個の計算をしたかったときに，1,024Threads/block，8Blockでカーネル起動したとき，
192個余分にThreadを起動しているので，その処理は大丈夫？って意味だったんだけど

＞カーネル呼び出しの行で実行がストップしてしまい、エラー取得の行まで進みません
ホストからカーネル起動したらすぐに次の行にステップして，
カーネルの動作は投げっぱなしだったと思うんだけど…違ったっけ？
なんか，メモリアクセスエラーよりもカーネル起動すら出来ていない様な気がするので，
簡単なテストカーネル作って起動させて確認するといいのでは？

意外と凡ミスでＨＷ制限越えた値で起動させようとしてるとか
パラメータの指定のような気がする

c/パラメータの指定/パラメータの指定ミス

おかしな動作をしている箇所を特定しました。
（「cuda-gdb」で再現しないと、デバッグに時間がかかりますね。。。）

ワープ内で３２回ループ（ワープ内スレッド数分のループ）している箇所があり、
その中で１つ１つのスレッドが順番に「とあるシェアードメモリ変数」を
書き換えているのですが、別のスレッドが書き換えたはずの値が、
正しく取れていない動作をしていました。
追跡はこれからですが、その後の処理でその変数を使っているところで、
何かしらカーネルが止まる動作になっているものと思います。

とりあえず、その変数に「volatile」を付けたところ、
カーネルが動く様になりました。

まだまだ、これで直ったかどうかは、これからじっくり試験しないとダメですが、
変な箇所が発見できたので、大きく前進しました。
ありがとうございました。

511>>
threadIdx/blockIdxについて確認してみました。
大丈夫でした。
カーネル呼び出しは、通常であればご指摘の通りの動作をしますが、
今回のバグでは、カーネル呼び出しで処理がストップしていました。

>>514
別のスレッドが書き換えたシェアードメモリの内容を参照しょうとしてるの？
排他処理とかどうなってんの？

>>515

排他処理は、ワープ内の1番目のワープが代表してロックを取っています。
その後、514で書いた32回ループの処理を行い、
再び1番目のワープが代表してロックを解放しています。

すいません、
ワープ内の「1番目のワープ」とか書いてますね。
誤記です。
「1番目のスレッド」です。

あるプログラムでビジュアルプロファイラー使ってみたら、
カーネルの実行時間が数マイクロ秒で、
カーネルの起動とcudaThreadSynchronizeの
オーバーヘッドが数百マイクロ秒だた。

FermiとKeplerを比べると、
Keplerの方がカーネル実行時間は短くなっているのに、
オーバーヘッドがでかくなって、トータルで遅くなっている。

カーネルの実行時間が数十から数百ミリ秒のプログラムなら
問題にならないんだけど、カーネルちっこいと効率悪い。

エヌヴィディアさん何とかしてください。

>>518
それはもうGPUの宿命だね。
カーネル起動やデータ転送のオーバーヘッドがゴミみたいなレベルの
ど〜んとデカい並列処理をやらすのがGPU。
一時期、AMDがHSAで細粒度並列処理に振ることも検討していたが、
結局、粗粒度での効率を追求するアーキテクチャに絞ったようだ。

将来の解決策の一つがハイブリッドアーキテクチャなんだろうなあ
レイテンシコアがオーバヘッドのでかい処理を担当するという

以前(>>493)の質問の続きというか似たような質問なのですが誰か教えてください

あれから少し調べて配列から二分探索で値を探索し、バイトニックでソートしようと思いcでプログラムを組みました(>>494に関してはよくわからなったので止めておきました)

そこで思ったのですが1*nの配列ならその方法で出来ても目標であるn*nの配列から複数回の入力に近似した値を抽出するという作業を行うにはどうすればいいのかわかりません

自分が考えた方法としては無理矢理配列を1*nに置き換えて探索するという方法ですがそれは処理時間的に有用な方法なのでしょうか？
また根本的に探索法の選出に誤りがあるなどの点があれば教えてください

最後にc言語からcudaを適用する際に現状で理解できたのはソートの配列入れ替えを並列化するという点なのですがそこをプログラムとして書くときに注意すべき点があれば教えてください

理解力の乏しい素人の質問ですみませんが誰か回答をお願いします

ゲームだと処理された画像がすぐに表示されるのに、
GPGPUだとオーバーヘッドが大きいのは何でなん？

JITしてるから

>>522
ゲームの場合は、

初期化時にモデルやテクスチャなどの容量大きめのデータをCPUからGPUのVRAMに転送。
ランタイム時はGPU⇔VRAMの広帯域を利用してでモデルやテクスチャなどの容量大きめのデータをやりとりし、
CPU⇔GPUの間（比較的帯域の狭いPCIE）では行列や描画コマンドといった小容量のデータ転送しか行わない。
また、処理結果はCPUに戻さず、GPUからビデオ出力するだけ。

大して、GPGPUでは、

ランタイム時にもCPUからGPUへ処理対象となるデータを転送するため、
帯域の狭いPCIEに頻繁に大量のデータを流すことになる上、
処理された結果をCPUに戻すために往復分の帯域を必要とする。

結局ゲームなどのグラフィック処理がGPU自体はもちろん、
周辺システムに関しても相性の良いアプリケーションとなっている。

>>521
何次元配列で各次元の長さはどのぐらいあるん?

>>525
二次元配列で大きさとしては512*512ぐらいを条件と考えています

>>521
まずさ、1次元配列で近似値抽出するコード書きなよ。
それで2次元（というか多次元）配列を1次元として処理するイディオムがあるからそれ使って処理する。
そのイディオムを知らないなら入門書買って読みな。

>>526
CUDAのスレッドの生成はブロックもグリッドも
三次元構造までサポートしている。
「CUDA by Example 汎用GPUプログラミング入門」
に画像(2次元配列)処理の例が載ってる。

今扱ってる問題が二次元だったら、
2次元構造にスレッドを起動すればいいんじゃないかな。
配列サイズもその規模なら比較的容易にプログラムできそう。

配列中の一番近い値の検索って要するに一番近い値を
持っているインデックス(i,j)の検索だから、
まず2次元のグリッドで入力値と配列値の差の絶対値を
全要素に対して計算して、
次に、j方向に差の絶対値とインデックス(i,j)をセットでリダクションして、
各iの中の最適なjを見つける。ここの段階でリダクションされたデータは一次元になる。
つぎにiに関してリダクションすれば入力値に一番近い値を持っている
インデックスが見つかるはず。

あと、多次元配列と言えどメモリ上では一次元で管理されているから、
コアレスアクセスになるように気をつける必要がある。

まあこれがパフォーマンス的にいいアルゴリズムなのかは責任持てんけどw

ふと思ったんだが、C++で
(value, i, j)の構造体作って、
そいつの512*512長のvector作って配列の値とインデックス放り込んで
valueでsortしてlower_boundとかupper_bound使えば
CUDAいらなくね?

ふと思ったのなら検証してくれ

速度は知らんけどできると思うよ。検証は知りたい人がやればいいと思う。

>あれから少し調べて配列から二分探索で値を探索し、
>バイトニックでソートしようと思いcでプログラムを組みました
とあるけどc(非CUDA)で組んでちゃんと結果はでたのかな？

さすがに、実は課題は各要素がスカラーではなくベクトルで、指定したベクトルに
（ノルム的に）一番近いものを抽出したかったなんてオチはないと思うけど。
そしたら問題の根本が変わるし。>>493で「絶対値をとって比較」の絶対値がふと気になりました。

>>527
回答ありがとうございます

お答えいただいた通りまずは一次元でcudaを適用しその後そのイディオムとやらのやり方を調べてみようと思います

>>528-531

回答ありがとうございます

皆さんの意見を参考にもう少し自分で考えてみます

ありがとうございます！

そう、自分で考えることが大事だ。

がんばれよ。

c(非CUDA)で組めてなかったのか

非cudaで組まないで、どうやってデバッグするつもりだったんだろう……

アーキテクチャーが変わるたびに
「既存のプログラムが遅くなった」という悲鳴を聞くのだが、
どうよ?

どうもこうもない。
各アーキテクチャごとに最適な記述が変わるというだけだ。

そのままのコードで速くなるx86が特殊だと思ったほうがイイ。

同じハードウェア・同じソースでSDKのverを上げてビルドすると速度が落ちたというケースも…

新しいアーキテクチャーが出てきた時にそれに対応する
コードを注ぎ足していければいいのだが、
マクロCUDA_ARCHって使えるの?
定義されてないとコンパイラに怒られるんだけど、
nvccが勝手に定義してくれるもんじゃないの?

ダイナミックパラレリズムのカーネル起動コストは
ホストからの起動よりも軽いですか?

どうして自分で検証して確認しようとしないのですか？

既に誰かが試したことを再度試すのは時間の無駄

回答を待つ時間は無駄ではないと

そいつの試験方法が正しいということを証明する方法を述べよ

くだすれやんw

重いとは何度か聞いているが自分で試したことはないな
今のところD-Pを使う機会がないというのもあるが

D-P試してみたよ。
とりま、孫カーネルまでネストしちゃうとオーバーヘッドが
かなりきつそう。

子世代ぐらいまでにして、うまくグリッドのサイズを調整すれば、
普通にストリーム使う以上に大量のカーネルを同時実行できるから
いいかも。

今のところD-Pを使える機械を持っていません＞＜

CUDAは一般用とだとどういう分野が得意なの？
フォトショみたいに単純な画像処理だけなん？

レイテンシだとかオーバーヘッドみたいな文章を読んでると、
CPUで処理したほうがいいからGPGPUが普及しないんじゃないかと勘ぐってしまうんだが・・

>>550
GPUのアーキテクチャにうま〜くハマる処理ならイイんだよ。
んで、結局グラフィック処理が一番イイというオチになるｗ

TSUBAME2は一般用ではないのか？

CUDAで囲碁を強くすることはできんのかな

浮動小数点演算をたくさん使うならGPGPUが優位になる可能性はある。単精度で済むならなおよし。
整数演算はKeplerからぱっとしなくなってるんだよなぁ。おそらくMaxwellも。

>>553
囲碁は評価関数がネックになってるからなあ
着手可能手の膨大さに関しては上手くはまれば期待できるけど、盤状態の評価はifだらけのCPU向きなコードにならざるを得ないから無理目に感じる
積み手もないから枝刈り判定が出来なくてDPの出番もなさげ

モンテカルロ木探索とかどうよ

木探索とか一番苦手なんじゃないのか
あとメモリ大量に使う処理も苦手

上界/下界と探索済/未探索空間の共有以外は依存性がないから
並列計算に向かないわけじゃないと思うがなぁ。

並列計算にむかないんじゃなくてCUDAにむかないということだろう。
ワープ内のスレッドがあるifーelse文でそれぞれちがうステートメントを実行するような場合、
両方実行されるためにペナルティが発生するから。

完全な単純全手探索の速度ならCPUぶっちぎるだろうけど、10の360乗ともなるとたとえ千倍速かろうが焼け石に水だなw

test

780tiってCuda5.5フルに使える？
公式いってもCuda対応としかかいてなかったんだけど

>>562
Tesla向け機能以外は使えるよ

>>557
モンテカルロ木は普通の木探索とは違うよ

>>563
ダイナミックなんちゃらは不可？

>>565
可

現役CUDA使いの方々，教えて下さい。
久々に趣味で学生時代のコードを触ろうと思って調べ始めたら，
Kepler世代になって色々と変わってて，たった3年で浦島状態に…

そこでお聞きしたいんですが，
チューニングのコツとかでFermi迄と大幅に変わってたりします？
特に，リダクションとか，キャッシュ回りとか，
メモリパディングとかのアドバイスを頂きたいです。

ワープ内のリダクションは共有メモリを介さずシャッフル関数で
できるようになりましたね。

>>557
そこそこメモリバンド幅あるんでねがったか?

http://www.anandtech.com/show/7515/nvidia-announces-cuda-6-unified-memory-for-cuda

NVIDIA Announces CUDA 6: Unified Memory for CUDA

Unified Memoryはホストデバイス間の通信コストがなくなるのかね。

何その魔法の技術は？

>>571
http://news.mynavi.jp/news/2013/11/15/539/
CUDA6で出るみたい。
ソースコードレベルでMemcpyを消せるだけなのか、
実際にPCI Expressの転送が無くなるのかは分かんないね。

有るとしたら
・仮想アドレスレベルでのポインタ共有
・CPUのメモリ内容がGPUから直接参照可能(今でも可能)なだけでなく、
　GPUのメモリ内容がCPUから直接参照可能になる。
・GPUから他のGPUのメモリへの直接参照(GPU-DirectはCPUメモリを介した間接コピー)
てなところじゃね。

明示的コピーをしなくてもオンデマンドでデータ転送が
行われるようになるだけで、PCIeの転送は無くならないだろう。

アクセスパターンによっては、明示的コピーをした方が
速いとかは十分にあり得る。

てかどう考えたらPCIeの転送が無くなるという発想に至るんだ？
APUみたいにVRAMとメインメモリが完全に共有されない限り無理だろ。

Intelがスパコン用かなんかで似たようなことしてなかったっけ？

x86なAPUを持ってないNVIDIAがPCIe転送を回避する方法はないよな
HSA構想を足蹴にできる現実的な新手法でも閃いたんだろうか

>>576
そんなんあった？今時IA64でもないだろし、
Phi上で独立したLinux走らせるという荒技っぷりしか記憶にないのだけど
あ、IrisのSRAMがCPU共有のキャッシュになってるとか読んだような。あれは何か野心感じたな、よく知らんけど

次世代GPUではarm組み込むんでGPU側だけで完結とか

>>578
ますますソースコードの可搬性が無くなりそうな悪寒。

>>578
それはそれで、組み込み用に特化しそうだｗ

それなんてcell？

phiに近くなるだけでは
http://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で同等の処理を実行する。

シリコンだけでその構造とったら意味が変わるでしょ

？

http://devblogs.nvidia.com/parallelforall/unified-memory-in-cuda-6/
サンプル
簡単な説明あり

CufftでMKLであるDftiSetValeに該当すものってある？

DftiSetValeなんてもんはない。

単なるミス。
status = DftiSetValue(desc_handle,config_param, config_val)
だよ。

だから該当すものなんてないってば。

なんだないのか。

またよくわからない状態になったので再度質問に来ました。
よろしくお願いします。(過去の質問>>493 >>521)

以前の状態からいろいろやってみて二次元配列においてCUDAを用いて
数値探索をできるようになったのですがどうも領域の確保というか
ブロックやスレッドの使い方が理解できずスレッドの最大数(?)である
512個以上のデータを扱おうとした場合に正しくない結果が出てしまう状態になってしました。

そもそも使っているのがthreadIdx.xだけでそこに代用として
blockDim.x*blockIdx.x + threadIdx.xなどを入れてみてやったところ
一定周期でソートされていて全体ではソートされていないという出力になってしました。

まだよくわかったないことが多いですが512以上のデータを扱う場合に
どのようブロックなどを扱えばよいのでしょうか？

スレッドの最大数が512？？、

>>592
すみません言い方が悪かったです。

www.gdep.jp/page/view/253　を見た感じ1ブロックの最大スレッド数が512らしいので
たぶん1ブロック分しかとれていないのだろうと思ったので512だと書きました。

スレッドが１ブロックでのスレッドが512を超えるとスルーされるな。
GTXだからかもしれんが。

>>593
いや、プログラミングガイド見ろよ。
CCによってその制限はちがうよ

もでなぜか512を超えるカーネルが起動しないケースがある。
レジスタが足りんのかわからんが。

>>591
まず
blockDim.x*blockIdx.x + threadIdx.x
の意味は分かってる？
左の項は自分がどのblockにいるかを計算していて、右は自分がそのblockの何番目のthreadなのかを示してる

つまりこれは、CUDAを一次元的に使うなら、それぞれのスレッドに固有の値になるってわけ

Tesla k40。。。NVIDIAは日本で売る気無いんだな。たかすぎw。

最近CUDAに興味持ちました。
識者の方教えてください。
CUDAやる時はGPUは画面表示用と計算に使うGPUの2つが必要になるのですか？

>>599
ひとつでいいけど、GPGPUの負荷が高いと画面表示がかなり重くなるよ。

一つでもできるけど、画面表示が重いと計算速度が落ちるよ。
3dゲームなんか動かしてたら、メモリもろくに使えなくなるし。

表示のリアルタイム性が重要ならおなじGPU。
そうでない処理ならば表示のコストなんて屁みたいなもん。
複数GPU使うのは計算自体を分散したくなってからでいい。

>>601
CUDAで計算しながら3Dゲームでもするのかと思ったｗ

何かのゲームでDirect3Dで通常のゲーム描画しつつ、
CUDAを演算に使ってたぞ。

CUDA使う3Dゲームなら一つのGPUでCUDAとDirectXのバッファ共有だろうな。

SDKにCUDAとDirectX同時に使うサンプルいっぱいあるよ

皆様ありがとうございました。
早速買いに行ってきます！

開発中にOSごとクラッシュしたりディスプレイが落ちたりすることがあるから
面倒なんだけどGPU2付けるのも面倒なんだよな

>>608
クラウドを使え

教えて頂きたいのですが、エンコード等でCUDAを使う場合、
グラフィックカードに搭載されているCUDAコアプロセッサの
搭載数の多い少ないだけで、性能を判断してよいのでしょうか？

例えばGTX Titanのように、倍精度に制限がない場合、
CUDAでの処理性能に違いが出たりしますか？

いいえ
GeForceでは3Dグラフィックスに使われる浮動小数点演算が最も重視され
エンコード等に使われる整数演算は重視されません
GPUのアーキテクチャにより整数演算性能は大きく左右されるので単純にCUDAコアの多少で決まるわけではありません

>>611
わかり易く教えて頂いて、ありがとう御座います。

バイトニックソートに関して質問があります。
カーネル関数にバイトニックソートを実装してみたのですが、複数ブロックを立てた時のソート結果が正しくなりません。
ちなみにソースコードは「CUDA高速GPUプログラミング入門」に掲載されているものを使用しています。
(こちらのHPにも同じものがあります→www.shuwasystem.co.jp/support/7980html/2578.html)
過去ログを読んでみたところ、>>591さんと似たような症状みたいです。
bitonicSort<<<2, 512>>>(...)みたいな感じで関数呼び出ししているのですが、
何かプログラムに手を加える必要があるのでしょうか？

>>>613
#define BLOCK_SIZE 256

以外にいじったとこある?

>>614
レスありがとうございます！
底以外は特に何も手を加えていない状態なんですが、
実行したところ、512スレッド毎にバイトニックソートがかかっているみたいで、全体のソートが出来ていません。
<<<1, 512>>>のように1ブロックしか立てなかった場合はしっかりとソートがかかるのですが、
複数ブロックを立てるとやはり1ブロック中のスレッド毎にしかソートがかからないようです。
何かハードの制約でもあるのでしょうか？

それは恐らくハードの制約ではなくブロック内でソートすると言うサンプルの制限なのではないだろうか。
つーか、サンプルを理解できないなら使うなよ。

>>616
レスありがとうございます。
自分なりにいろいろと勉強はしているのですが、サンプルが掲載されている書籍がもうひとつ説明が少なくて苦労してました。
ブロック内でソートするとかいう説明もなく、ひたすらデバッグやトレースを繰り返してはいるものの、
解決策が見つからなかったのでこちらで質問させていただきました。
CUDAのバイトニックソートのプログラムはほとんどが1ブロックしか立てていないものばかりなので、
参考になる情報も少ないので・・・

>>615
書き換えない状態では256スレッドのブロックが四つという設定の筈ですが、
サンプルコードがそのままでも動かないということですか?

>>615
というか、
#define SIZE 1024

に相当する部分をいじってないですか?
カーネルにも SIZEが使われてるので、
下手に数字をべた書きすると整合性が
取れなくなりそうな希ガスるですよ。

>>618
迅速なレスありがとうございます。
今サンプルコードのままで実行して結果を調べてみたところ、
ちゃんとソートされたり、>>615のようになってソートされなかったりします。
ああ、余計に訳が分からなくなってきた・・・

>>619
レスありがとうございます。
defineの部分はいじってないです。
SIZEはサンプル通りの1024のままですね。

その本読むよりも、他の資料、例えば CUDA by Example を読んだほうが、基礎が掴めると思うよ。

>>622
レスありがとうございます。
そうですね、自分の勉強不足もあると思います。
もう少し勉強して頑張ってみます。

勉強不足もあると思うけど、理解してから次に進むっていう意識が欠如してるよな
blockDim.x*blockIdx.x + threadIdx.x
の意味すら理解せずにアルゴリズムがどうこう言ったって理解できるわけないだろ。

まあくだすれだし。
またーりしようよ。

>>620
まずサンプルコードをいじらずにそのまま使って
何度も実行して、問題が再現するかどうかちぇっくしませう。

http://on-demand.gputechconf.com/supercomputing/2013/video/SC3108-New-Features-CUDA%206%20-GPU-Acceleration.mp4

関数の呼び出しについて質問があります　回答をお願いします

ｃ言語でつくられたプログラムの一部をCUDAに適用しそれを
本体(cppで作られたプログラム)とは別にCUDA用プログラムとして(拡張子cuとして別に作っておいて)
作り本体から呼び出そうとしたところ未解決の外部シンボルとしてエラーが吐かれてしまうのですが
どのように処置したらC言語のプログラムからCUDAのプログラムを呼び出すことができるのでしょうか？

CUDA SDKのサンプルはビルド＆実行できた？

>>627
Cはc++またはg++とかでコンパイルしている？
ccで生成した中間生製オブジェクトとはリンクできなかったような

>>628
回答ありがとうございます

サンプルとは「〜\NVIDIA Corporation\CUDA Samples\v5.0」の中にあるサンプルのことでしょうか？
もしそのサンプルなら全てではありませんが一部をやってみたところ問題なく動きました。

また外部シンボル等でいろいろ調べてみたところライブラリのリンクが等々と書いていたので
構成プロパティでいくつか追加してみても特にエラーは直りませんでした。

>>629
回答ありがとうございます

コンパイルについては問題ないと思います。
ネットに書いてあったような設定をしてやったところ.cuのプログラム一つだけの場合
問題なくコンパイルされ期待通りの実行結果が出てくれたのでCのプログラム内で
CUDAのプログラムの関数を呼び出す際に問題が発生したのではないかと思います。

というかそのやり方がわかりません・・・

.cuファイルを複数使ってるのかな？…と思ったらCUDA5からそれでもOKっぽいんだね

NVIDIA，「CUDA 5」を正式発表。第2世代Kepler「GK110」に向けた準備が整う ニコニコニュース
http://news.nicovideo.jp/watch/nw401167

cuファイルが1つだけならビルドできるサンプルにコピペすれば
どうにかなる（問題の切り分けを始められる）と思ったけど

1つ抜けてた、サンプルというのは>>630の言うそれのことで合ってますです。

>>632-634
回答ありがとうございます

cuファイルは一つだけなので該当サンプルと内容をすり替えれば良いということでしょうか？
よろしければそのサンプルはどのサンプルなのかお教えください。
よろしくお願いします

そうです＞内容をすり替えればいい

ごめん、今使ってるPCにCUDA SDKを入れてない（非Geforce機）ので何があるか
わからんけど初級の短めのやつにすればいいと思う。

あと（少なくともcuda sdk3か4の頃だと）プロジェクトに相対pathか絶対pathかが使われてて、
フォルダの場所を移動させるとそれに合わせて設定変更しない限りビルドか実行かが
うまくいかなかった覚えがあるので注意してくださーい

コペンバローナ「ンーｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗ」

ごめん、問題のポイントを誤解してたかも。
（CとC++とcuの入ったサンプルは見た覚えがない）

とりあえず俺の発言は忘れてください。ごめんなさい。

ペッコンバローナーｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗ

>>631
未解決の外部シンボルのエラーって、リンク（コンパイル）の時じゃなくて、実行時に出るの？

>>628
ペッコンペッコンペッコンバローナーｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗwwwwwwwwwwwwwwwwwwwwwwwwwwww

ロ・・・ロバｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗ

コペンハーゲ「ンーｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗ」

コペンバロー「ナｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗ」

バコナロ「バコーンｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗwｗｗｗｗｗｗｗｗｗｗｗｗｗ」

、　　′　　　　　、　’、　　′　　　　　’　　　；.
　｀',．　’　　　　　　’、　　　′　’　　　．　・
　、′・．　’　　　；　　　’、　’、′‘　．・”
　　　　　　　　　　’、′・　　’、．・”；　　” 　’
　　　　　　’、　　(;;ﾉ;;　(′‘ ・．　’、′”；
　　　　 ’、′・　 (　（´;^｀⌒)∴⌒｀．・　 　”　；　′・　　　　　　　　　　 　, '´｀ヽ.-──-,'´｀ヽ.
　　　　､　’、　’・　､´⌒,;y'⌒(（´;;;;;ﾉ､"'人　　　　　　　　　　　　　　 　　/　　　　ﾞi::::::::::::::ﾞi　 　ﾞ:
　　　　　_､(⌒ ;;;:;´'从 ;'　　 ;:;;) ;⌒ ;; :） )､　　　　　　　　　　　　ミ　　　/　　　　　;ﾞ:;::::::;:::::::! 　　!
　　　　　（　´;`ヾ,;⌒)´　 从⌒ ;) ｀⌒ ）⌒:｀．・　　　　　　　　　　　　,.;ﾞ　　r'^ｰ､;ﾞ:;ｨ:::ﾊ::λi,ｒ'ヽ, i
　‘:　；゜+° ′、:::::．　:::（ ::;; ﾉ ´⌒（,ゞ､⌒） ;;:::)::ﾉ　　　　　　　　　 i::i 　　| 　 iﾞﾉiﾉﾚ　ﾚ' ﾞ!!i　 | !
　....................｀:::､　ﾉ　 ...;:;_)　　...::ノ　 ソ　...::ノ　　　　　　　　　ミ　　 ﾊ::ﾞ､ 　i 　 ! ＞　　　＜!| 　!;
　　　　　:::::::::｀- ´::::::::::::::::::::::::::::::::::::::::::::::::::　　　　　　　　　　　 　 ﾉ:::λﾞｰ| 　 |〃 　＿_ 〃i　 l'
　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　 　 〈ﾉ::ｲ::ﾉ::::!　　!　 　'､　ﾉ ノ'.　;
　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　 　 〈/!::;ｲ::/,ｒ'´｀ﾞﾞi＞‐_-_ｔ.´ｒﾞ´ヾ
　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　ミ　 　V＾ﾚ´i　　　,.〉｀'ﾞ'i／ﾞ｀i,_,..ﾉ
┏━━━━━━━━━━━━━━━━━━━┓　　　　　　　　　　｀ｰiﾞ´　　/>　 　i
┃　「キャー♪」　　　　　　　　　　　　　　. 　 . .......　┃　　　　　　　　　　　 |　　　´　　　|
┃　　てゐちゃんは　きょうも　たのしそうだ！！ ┃ミ　　 　,.へ.　　　　,ﾉ　 　 ,　,.　　!
┗━━━━━━━━━━━━━━━━━━━┛ 　　　 /_＿_ﾞﾆ=-‐´ ,　　 /　 〉　 ゝ.
　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　 /´￣　　　　＿ ノ　／　 /　　ﾉ　〉

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

お、Transmetaの人が関わってるのか。

プログラム実行中のGPUの温度をモニターしたいのでNVMLを試してみようと思っています。
ここで、nvmlDevice_tとCUdeviceの対応はどのようにとればいいんでしょう？
CUDAのデバイスindexとNVMLのindexは必ずしも一致しないという記述はあったのですが、
じゃあどうすればいいのか、というところを見つけられませんでした。

GPU2枚差して、CPU介さずにデータ共有ってできる？

Teslaならできるらしい。持ってないんで試したことないが。
https://developer.nvidia.com/gpudirect

試してみたら、GeForceでもできました。
ありがとうございます

>>652
詳細お願いします。

>>649
普通にインデックスでいいよ
汎用に作るならデバイス数を取得して、それぞれのnvmlDeviceをインデックスで取得して、いろんな情報とればいい

GeForce,Quadroはメインメモリ→ボードのDMACしか持ってないよね？

なんでそう思ったのかが気になる。

NNみたいなモロにメモリ律速な計算だとろくに速度出ないな
帯域80GB/s使って160GTlopsとかになる

結局どういう問題なら高速化できるんだ

メモリへのアクセスが少ない、扱うデータサイズが小さい、分岐がない
最低数万スレッド以上で並列計算可能な問題であること

メモリ量と計算量が比例する問題しか普段扱ってないんだよなあ
暗号解読とか？

Geforce GT520(VRAM: DDR3 1GB)でもCore2Duo E4300に比べたらFFTを高速化できるかな？

マンデルブロが超得意

データ量Nに対して計算量がN^1より大きいオーダーで
増えていくような処理

巨大な元データが必要でも、それ自体は変えずに
少量のパラメータを与えて再計算を繰り返すような処理
しかも結果をグラフィックス表示すればOKな用途

>>662
＞しかも結果をグラフィックス表示すればOKな用途
GPU⇒CPUが入ると途端にスループット落ちることになるもんね･･･。

人工ニューラルネットワークなんかは、
データ量N、i段目のニューロン数n_iとすると、
計算量=NΠ_i n_i だから実はあんまり適してねえのか

GPUもCPUも足回りが全然ついていかないんだよな
NvidiaもAMDもFlops値ばかり競ってるけどメモリ帯域はこの数年で1割程度しか増えていない
完全に頭打ちの傾向

そして効率的な演算とデータアクセスの比率は高まるばかり･･･

石の性能が良くなっても仕方ないよな。

プロセッサの性能が無駄になってる。

まぁ、VoltaでスタックドDRAM使うみたいだから、いくらか改善されるかもね。

　　　　　　☆ ﾁﾝ　　　　　ﾏﾁｸﾀﾋﾞﾚﾀ〜
　　　　　　　　　　　　　　　　　　　　　　　　ﾏﾁｸﾀﾋﾞﾚﾀ〜
　　 　 　 ☆　ﾁﾝ　　〃　 ∧＿∧　　　／￣￣￣￣￣￣￣￣￣￣￣￣￣
　 　 　 　　ヽ　＿＿_＼（＼・∀・）　＜　データまだ〜？
　 　 　 　 　 　 ＼＿／⊂　⊂＿ )　　 ＼＿＿＿＿＿＿＿＿＿＿＿＿＿
　 　 　 　 　 ／￣￣￣￣￣￣ ／|
　　　　　　　|￣￣￣￣￣￣￣|　 |
　　　　　　　|　 CPU・GPU　 　|／

GPUはバス幅を狭くすることでコストダウンを図ってるんだから仕方ないな。

それこそ、バス幅求めるならベクトル計算機でも使えと。全レジスタに対して本当の同時操作が出来るぞ。

>>669
現状を表す最適なAA乙ｗ

このアンバランスな状態を解消できるのはプロセスルールが物理的限界に到達した後だろうな。

しかしその頃には光コンピュータが実用化されていたのだった……

速さが足りない！！

俺が遅い・・・ 俺がスロウリィ?!

HOLY隊員のクーダーです

FFTぐらいしか応用が思いつかねぇ。

>>676
FFTに向いているなら自動的に円周率計算もバリバリなはずだが、ググっても
「円周率の小数点以下8000兆桁めをGeForceで求める方法」
(http://www.4gamer.net/games/120/G012093/20130323002/)
といった話ぐらいしか出てこねぇ……

音声処理におけるFIRフィルタを想定してるぜ・・・。

世間が持てはやすのがFLOPS値ばかりだから一向に帯域増える方向にいかんな

帯域はコストが高く付くからな。

バランス取ろうと思ったら、途端に価格が跳ね上がる。
一般人じゃ手の届かない価格になるよ。

>>680
別にHPC用なら値段高くても買う奴いるじゃんか……

普及してて値段が安いからGPGPUがもてはやされてるわけでさ。

値段が高くなればベクトル計算機のプロセッサをPC向けに販売して使ったほうが良いって。

>>669
わらった。
GPGPUの一般用途での最大の問題点はCPU<=>GPU間データ転送。一般用途ではそれを解消したAMDのAPUでHSAする方が良いからな
いくらGPUがすごくても、メモリ転送に時間掛かってはお手軽に使えないからな

kaveri出たらHSA酷使した絶賛ベンチが次々と出てnvidia叩きレスで溢れかえると思ったら思いのほか静かで不思議

>>683
データ転送せず極力内部で計算するようにしても結局GPU側の帯域で足引っ張られる
780Tiで単精度5.76Tflopsに対して330GB/sだから足回りが70倍も遅い

>>684
言い出しっぺの法則
>>685
だが待ってほしい
70倍遅いなら70倍転送せずに計算すればトントンではないだろうか

HSA使ってみたいんだけど、具体的にどうすればいいの？
VisualStudioで始められる？？

CPUGPU間の転送が足を引っ張ってるってイメージはないな
シェアードメモリやキャッシュ以外のVRAM・GPU間がただただ遅いのだ

レイテンシ？

基本的にI/Oが遅いんだよ。
これが何とかなったらいいけど、何とかするとコストがかさむから一般向けでは無理。

一般向=>一般向CUDA用途==スパコン

なぜそうなる。数十万でも買うのかよ。

重ーい超越関数をバリバリ使う計算ならメモリ転送はさほど器にしなくて良いのでは。

三角関数がそこそこ速いから最初に三角関数テーブルを作っておいて纏めて計算するんだけど、
キャッシュに乗らないとべらぼうに遅くなるw。

今や、テーブルにしてメモリから読み出すよりも、
手前で計算で作ったほうが速いからなｗ

昔「計算が遅いからメモリでなんとかしよう」
今「メモリが遅いから計算でなんとかしよう」
将来「？？？」

PS3もちょうどその技術トレンドを読んで企画されたけど、ちょっと早漏すぎたな。

>>689
VRAMのレイテンシは数百クロックもある上にピーク速度でも計算速度より何百倍も遅い

>>697
早漏てw
そこは先駆者として評価してやっていいんじゃないの。十分出回ったしハード的にもソフト的にも注目されて、長めのゲーム機サイクルの中で研究されたんだからアーキテクチャとしては幸せな方でしょ

ソニーさんのビジネス的にどうだったのかは知らんけど

>>698
え、マジで？
>>699
さすがに逆ザヤはNG

サブプロセッサの性能は兎も角、メイン側が遅過ぎ。
メインとサブの間のメモリ空間も狭いし。
あれでよくゲームに活かせたと思うよ。

>>694
三角関数テーブルって精度的にはどうなん?
多項式補間とかするの?

用途によるだろう

多項式補間といっても奥が深くてだな……
単なるテイラー展開(途中打ち切り)とよく練られた多項式との差はダンチ

例：
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(テイラー展開)。しかし、

途中送信してしまった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
(出典：http://www.amazon.co.jp/dp/456301382X)

>>702
私(>694)のところで使うのは周波数空間像の畳み込みだから、三角関数の引き数は格子上の点の距離。
なので、補間の必要もないの。ついでに、cufft相当も自前で実装した。

テイラー展開とか教科書に載ってるだけで、
関数近似の方法としては、ほぼ実用されてねえよ

>>705
URLが見つかりません

最大誤差が小さくても、cos(0)が0.99999981155になる関数なんて使いたくないな。
0みたいな重要点でおかしな値が出ると致命傷になることが多い。

>>707
テイラー展開の誤差範囲の理論値が明確であるメリットは結構大きい

>>708
URLの最後の)がいらない。

スレ違いかもしらんが、
gccとかのソースを見れば超越関数の実装が分かったりするのかな。

>>708
URLの最後の)がいらない。

>>712
FPUがサポートしている超越関数はソースがないかもね。
iccならSSE版の並列演算用の超越関数が実装されているんだけど。

質問です。CUDAを初めて使おうと思うのですが何を買っていいのかわかりません。
当方プログラマです。整数演算主体の力学シミュレータを自作しています。
その中にある絶望的に激重な評価関数が高速化できたらなぁと夢見ています。
その関数は同じデータセット（200キロバイトくらい）を、さまざまな初期値で評価するのですが、
条件分岐が殆ど発生しないアルゴリズムを発見しました。CUDA 向けなんじゃないかと使ったこともないのに妄想しております。
１回の評価計算そのものがめちゃくちゃ重い（単純に１００万回くらいループさせているだけ）で、
ループさせるプログラムそのものは数キロバイトも無いちっちゃなものです。

とりあえず今は手元にある Windows 7 64bits (チップセットはP55) に入れてお試しでCUDAプログラミングし、
C++で書いたシミュレータをCUDA対応に移植するところから始めたいとおもっています。
グラフィック出力を２本使いながらCUDAを使いたいのですが、私はどんなのを買ったらいいのでしょうか？
将来的には研究費をつぎ込むつもりですが、今は衝撃的に貧乏なので予算３万以内で。

＞グラフィック出力を２本使いながらCUDAを使いたいのですが
ここんとこ詳しく。あとPCも。
まあビデオカードを1枚買うか2枚買うかくらいの違いでしかないとは思うけど。

>>714
CUDA積んでるGPUのリスト
https://developer.nvidia.com/cuda-gpus
この中から予算の許す限り良い奴を買えばいい

CUDA開発用のSDKは公開されてるから環境に合わせてインスコするんだ
https://developer.nvidia.com/cuda-toolkit

導入方法やサンプルコードはググるか>>1-2を参照


……もっとも釣りじゃなければの話だが

早速レスが。有難うございます。

>>716 マザボは P7P55D ってので、PCI-Express 2.0 16X が２本あるのですが、
今はRadeon ２枚で４画面（1920x1200x2 と 1280x1024x2）出してます。
そのうちの片方を nvidia にしたいと考えてます。画面出力との併用って難しいでしょうか？

>>717 色々ありますよね・・・。今は技術の練習として試そうと思うのですが、
GTX 660 を選択しようかとおもってるのですが、それはやめとけ、こっちがいいよ、とかありますか？

>>714
その目的なら、PC買う前にまず本当にCUDA向きなのか、
実際に自分がそのプログラムを組むことはできるか、
などをクラウドで実験してみてはどうだろうか

Amazon EC2
http://aws.amazon.com/jp/ec2/

1時間CUDAマシンを借りるのに1ドルもかからない。

>>718
Radeon/Geforceの混在かぁ。ごめん俺はわからない。
1枚で画面出力とCUDA計算の併用自体はできる（長い間計算しっぱなしにせず、
ある程度の間隔で制御が戻るようにすれば。
計算しっぱなしでもタイムアウトしない範囲なら表示が完全に固まるわけではないし）。

GTX660でいいんじゃないかな。あるいはコスト抑えたいならもっと下でも。
ローエンドGPUでの実行時間がわかればCUDA Core数の比較で上位GPUにしたときの時間の見当もつくし。
あと2/18に28nmのMaxwell世代のGTX750が発売されるらしいという話もあるけど。

>>719と同じくEC2を推す
高いグラボ買って大して高速化できませんでしたじゃ目も当てられない

>>719 有難うございます。実は仰る話は第二段階として既に計画していました。
ここがクリアしたら第三段階として本格的な予算を投じて
大量のGPUインスタンスを使って計算するかもしれないです。

でもその前に第一段階として、CUDAにあわせてソースを書き換えたり、
必要に応じてアルゴリズムも修正しなくてはならないと考えており、
そのトライ＆エラーに例えば一ヶ月かかっちゃうなら安い奴を買ったほうがいいかな、
とりあえず３万円程度で使い倒してみようかな、と考えている次第です。

>>721
そんな事情もあって3万円くらいで、なにがいいのかなと。

しれっとアルゴリズムを発見したとか言うよねーｗ

連投すみません。 >>720 なるほどー。
最終的には256ビット幅でアルゴリズムを最適化するつもりだったので、
本当にメモリバスが256ビット幅なら、2/18まで待ってみようかな。

画面出力とCUDA計算の併用はそんなに心配しなくても良さそうなんですね。
Radeonとは、だめもとで混在させてみるつもりです。ありがとうございます。

>>724
そんな言い方は野暮なんじゃね。何事も思いつきからでそ
失敗の責任を追うのは彼自身なんだしw

>>724
結構たいへんでしたよ？要約すると、とあるトポロジーのハッシュを計算する問題なのですが、
条件分岐だらけのソースを、ごり押しで全パターン計算させたほうが結果的に高速らしい、ということが判ったのです。
ただ、現時点では「高速らしい」という段階でしかなく、本当に早いんだということは実際にやって見せるしかない状態です。

>>725
Maxwellについては2/18のは28nmのままだし本気出せるSDKもしばらく先だろうから
たぶん結局Kepler買うことになり待つ意味は薄いと思います。
ただ一応知っておいたほうがいいと思って言ってみただけで。

あと整数演算オンリーならFermiが効率いいかもしれないけど今更Fermiに
合わせて作るというのもなんかロマンがないんですよね…

スレチになっちゃうけどRadeon積んでるならOpenCLで試せばいいような気が。
Radeonのほうが速いだろうし。

OpenCLは尚更ロマンがないな
旬ならkaveriでHSAか？といっても所詮ミドルレンジAPUだしなあ

もっと野暮なんだけど、たった百万回×二百KBの演算ならCPUでごり押しした方が速い肝酢。

整数演算ならRadeonのほうが数倍速いよ。
CUDAもOpenCLも対して変わらん。

>>722
本当に真面目にEC2の価格と比較・検討してみたのか疑問なんだが。
1日24時間ぶっ通しで使うわけでもあるまい。

あと、CUDA用のコードがエラーなく意図どおりの計算結果を出すかどうかは、
CUDA対応グラボを使わなくてもデバイスエミュレーションモードで確認できる。
それで動くか確認してからEC2で計算速度を実験すれば、かなり費用を抑えられると思うぞ。

まぁ、これも単なるひとつの手段でしかないし、
グラボを買う目的がCUDAだけでない場合もあるだろうから、
これ以上EC2を押したりはしない。

CUDA 6.0RC公開きた
https://developer.nvidia.com/rdp/cuda-60-rc-toolkit-download

なんかCUDA excelとかあるな

>>960
最近のなら
http://www.nicovideo.jp/watch/sm22874825

誤爆orz

GPUは1スレッドあたりの速度はCPUの何百倍も遅い
最低でも32ブロック×数百スレッド以上の並列計算できるような問題でないと力を発揮できない

>>714
本当にCUDA初心者で、且つ上手く行った後に予算が着くなら
使い捨てにするつもりで1万5千円ぐらいの中古のGTX 580を
買ったほうがいいよ。下手なKeplerより速いし。

MAXWELL世代はL2キャッシュが大幅に増えるらしい。これはどういう効果を生むの？
http://videocardz.com/49557/exclusive-nvidia-maxwell-gm107-architecture-unveiled

GM107 L2 cache has 2MB. GK107's cache has 256KB.

というかキャッシュを増やさないと、
開き続けている演算能力とメモリ帯域のギャップがますます開いてしまう。

メモリ周りをチューニングしなくても
そこそこ性能がでる感じなのかな

いろんな変更があるんだろうけど
研究や論文あさってみると階層化されたスケジューラとレジスタファイルの相乗効果と思われる

ちょっと古いがこれがわかりやすいかも
http://www.cs.utexas.edu/users/mgebhart/papers/MICRO_Slides.pdf

http://www.cs.virginia.edu/~skadron/Papers/gebhart_tocs.pdf

L2は大型warpのプリフェッチに使われる予感

CUDAの本でてた。
https://gihyo.jp/dp/ebook/2014/978-4-7741-6361-1
買わなくてもサンプルコード落とせた。

Maxwellはkelperより速くなってるの？
Fermiのほうがいいという悲しいことになってない？

CUDAって1PASSしか使えないけど、将来2PASS使えるようにならないのかね
CUDAエンコが速いなら、その速さを活かして2PASSでエンコしたいんだが

何のソフトの話だよ。そういうのは作者に言え。

>>746
お前がCUDAで2PASSでエンコ作ればOKだろ

>>745
一般デスクトップ向けはGPGPUより素直にゲームに注力したほうが良いような気がするからな

>>747
外人だから無理

>>748
ドライバがCUDAでの2PASSエンコに対応していない
だからどんなソフトでもCUDAで2PASSエンコは出来ない

ドライバが？なんか根本的な勘違いがありそうだな
それPureVideoの話じゃないの？

cudaに向いてるかどうかを別にすれば、マルチパスってのはストリームに対して複数回の処理を走らせる事を指す言葉に過ぎない。なのでドライバは関係無い

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>

これって750TiでもhyperQ,Dynamic Parallelism使えるってこと？

yes

そうか。すごいな。使うだけならGT640も使えるんだっけ？

うん

公式みると750tiはCompute Capability3.0になってるけど使えるの？
640は3.5だから合ってるけど

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

公式が間違ってるってことか
安いし買ってみるかな

ついでにここでやってた32bit integer bit shiftのスループットやってみたら
68とか半端な数字になった
けどkepler(GK110 tesla以外)の倍だねSMあたり
http://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

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使えるってか専用になったみたい

>>748
OpenCL対応ソフトが地味に増えていってる状況でGPGPU性能削るとか自殺行為だろ

GPGPU性能と言っても色々あるけどFermi→Keplerのときトータルでは削られたのでは。
2年前とはOpenCL/CUDAの対応状況が違うってことかね。
個人的にはMaxwellでGPGPU寄りに振ってくれると嬉しいんだけどね。

>>759見ればkepler（３２）で削られてたbitshiftもtesla同様に64に増えてるし
どこをみてgpgpu削られてたって言ってるわけ

keplerん時にゲーム向けチップから倍精度が削られた事くらいか

>>763は俺に言ってるの？

IDがないから誰が誰だかわからねーよw

keplerは倍精度削り過ぎだろ。
fermi以下とか酷過ぎる。

気持ちは分かるが倍精度は単にGPUとして使いたい層には無意味どころか足引っ張る要素だしな
TITAN的な選択肢が今後も提供されるなら別に文句ないけどな
欲を言えば、もうちょい安い製品でも倍精度残したチップ用意してくれれば盤石かな

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
とか、そういう中間的な選択肢が欲しい。

>>769
そういうの出したらtesla売れないから出さないよ。

俺は絶対性能は630や650クラスでいいから開発用にフルスペックのが欲しい
ECC, GPUDirect, TCCDriver, Hyper-Q, DPなどが使えるやつ。
できればTeslaのflops/B比等が同じで、そのまま倍数かければ
Teslaパフォーマンス特性がだいたい予測できるようなやつ。

>>770
それいいね。

あとは同じ価格帯のボードが新世代で性能が
落ちるってのがなければいいんだけれど。

>>768
ゲーム用VGAとGPGPU用ボードの２つを出せば良いんだよ
（VGAでたとえるなら、ゲーム用VGAのGFとプロ用VGAのQuadroみたいに)
ゲーム用にGPGPUを強化しても、ゲーマーにはあんまり価値がなく、そして爆熱になるだけし
まぁ、GPGPU用の値段はゲーム用よりだいぶ高くなるだろうが

>>772
それ現状のまま(GeForceとTesla)じゃね？

ワラタ

ロー･ミッドクラスVGA派生GPGPUはイラネで、今のteslaラインナップなんだろう

数年前までGPGPUは安いのが売りだったのに今じゃ高級品だもんなー。
そのうちIntelに負けるんじゃん。

いまは、中GPGPU性能以下(低価格)でいいなら、AMDのHSA APUでって感じじゃない
メモリーコピーいらんでお手軽にGPGPUできそうだし
そして低GPGPU性能でいいならIntelのiGPUで良いやだろう

いやぜんぜん

開発環境とトライバがだめだめで低シェアなAMDは論外

Xeon Phiがどれだけ使い物になるかで状況が変わってくるだろうな。

>>780
HPCシステムズ「ぶっちゃけ製品として尖り過ぎてて使いづらいんだよゴルァ！」
http://www.hpc.co.jp/benchmark20121113.html
↓
HPCシステムズ「まあ場合によってはコンパイルし直すだけで使える……かな？」
http://www.hpc.co.jp/benchmark20130409.html

チューニングすればテレサと同じぐらいのレベルか

つか、HSAでプログラミングしてみたいんだけど、
SDKか何か配布されてるわけ？

hsa sdkでググるだけで公式ヒットしたけど、斜め読みではどうしろっていうのかはぶっちゃけよく分からんかったなあ
シミュまで提供されてるようで、興味はあれどもkaveri機を組んでまで遊ぶ気力もなく

書籍出してほしいなぁ。

>>696

昔「計算が遅いからメモリでなんとかしよう」
今「メモリが遅いから計算でなんとかしよう」
将来「何通りか計算結果をあらかじめ予想しとこう」

amdのapp sdk使えばかってにやってくれるでしょ

ha?

へ？

amdのコンパイラ使えばメインメモリとGPUのメモリをシームレスにみてくれるってことだよ
そんなこともわからんのかバカちんが

予想の斜め下の返しきたなw
そらな、もしHSAでシームレスじゃなかったら吹くがな

dgpuじゃhsaは使えない

ナライラナイ

APUが、DDR3でもいいからクワッドアクセスすれば多少はましになる筈

DDR4や次世代になればさらにましになる

750tiの方が面白いわ

http://devblogs.nvidia.com/parallelforall/cudacasts-episode-18-cuda-6-0-unified-memory/
CUDACasts Episode 18: CUDA 6.0 Unified Memory

DELLとかHPのWSについているTeslaは他のPCでやっぱり使えないのかな

リモートデスクトップかsshdでいいじゃん（いいじゃん）

>>796
つかえるだろ。

Quadroではロックされているとかで使えないと
聞いたので、teslaKもかな？と思った。
結局やめて違うの買った。

ドライバ当てただけでは機能はしないんだな、
これからCUDAインストールして勉強する。

>>799
今時そんな金のかかることはしない。

メモリ周りが多少良くなるみたいだね。
http://www.4gamer.net/games/251/G025177/20140326091/

ありゃ、Maxwellの次はVoltaじゃなかったっけと思ったけどよく読んだらPascalが割り込む形か。

PCでのGPGPUはHSAだからな
Nvはスパコンなんかの非PCでがんばるしかないよな

むしろHPCから組込みまで同じCUDAが使えますが
すでに

PCはVGAレスが普通になってきているからね
わざわざVGAをつける奴ってゲーマーや業務のCAD関係する奴とかで
GPGPUのためにNvのVGAつける奴ってどれぐらいいるんだろ。

>>805
あとVGAつけるのはエンコード目的とかでIntelのLGA2011を使用している奴とかかな

日本で一番Xeon Phiを無駄にしている系男子wとかもいるしな

>>805
ﾉｼ

世の中には面白い事考えるもんがおるのう

http://www.otb-japan.co.jp/dmpr/GPU/gpgpu-xeonPhi-E5-Hybrid.html

こりゃ面白いな。
試してみたいｗ

ってかOpenACCってもうリリースされてるの？？

>>809
やっべ、ちょっと欲しいなって思っちゃった
値段を見て熱は冷めた

値段ﾜﾛﾀｗ

>>812
むしろこの構成で高く付かない方がおかしいｗ
大雑把にはXeonで13万、Xeon Phiで37万、Taslaで39万掛かる感じ
（更に8x 4GB DDR3 1600 ECC Reg.で10万掛かってるけどな！）

ひゃぁ〜ｗ

>>813
eBayとかだと自作の方が安くなりそうだな。
といってもtesla K20だけで20万はCUDAらない。

じつにCUDAらん。

>>815
まあ超ハイスペックだと自作のほうが安くつくしな
ただ、Xeon Phiは一般販売されてなかった気がするんだが……

そうだったのか。
まあ買って、単純に付けるだけでは動かんからな。
セッティングがめんどくさそう。
ttp://www.sekaimon.com/i221402553263

>>817
アメリカだと通販で売ってるよ。

>>819
アメリカの金持ちPC自作erはPhiを使ってPC自作するからか
へたれ日本ではそんなことする奴いないだろからな

それPCやない、WSかHPCや!!
TITANの3wayに比べれば敷居低いと思う。

自作ショップで普通はおいてないだけで
注文だせは取り寄せできるんじゃない？
Teslaに比べてもニッチだろうから。

インテルがばら売りを許していない少なくとも日本では

へーそうだったのか。
まあコンパイラも追加で買わないといけないし、
そこそこの資金は要求されるから、そういうもんかもね。

>>820
よりどりみどりだよ。
https://www.google.com/webhp#q=xeon+phi&tbm=shop

>>825のリンク、勝手にwww.google.co.jpにリダイレクトされちゃうな。

気が向いたらXeon Phiも付けてみたいよね。
これって現行の普通のi7だと厳しいのかね。

ブラボ4枚差しを今週には購入予定だが、cuda並列はむずかしいのかね？
おススメのＨＰか本はあるのかね

きみらって皆採掘が目的なの？１０年やって元がとれるくらい？

>>829
今から採掘に回るとかただの情弱じゃねーかｗ
単純にゲームとプログラミングのためだよ

6の正式版が出た

>>828
The CUDA Hanbook に書いてあったかも。
CUDA by Exampleにもちょっとだけ書いてあったかも。
gtc-ondemandにもなんかあったような。
曖昧で申し訳ない。

5000円以下ぐらいの適当な安いビデオカードって
CUDA使うぐらいならCPUで全部計算した方が速い感じなの？
例えば
AUS 210-SL-TC1GD3-L
とか

多分誰も答えてくれないと思います。

まずGPGPUの基本から勉強しましょう。

>>833
一応CUDAコアは付いているようだな……＞AUS 210-SL-TC1GD3-L(GeForce 210)
http://www.nvidia.co.jp/object/product_geforce_210_jp.html
同じ「計算」でも、GPGPUが効く分野とそうじゃない分野があるけどな

同じ値段のCPUとGPUを比較した場合
500並列くらいまではCPUの方が勝ってるが10000並列だとGPUの方が圧倒的になる感じ。

IntelがきっちりCPU向けに並列化したコードならxeonも速いよ、GPUメーカーの圧倒的数字は幻想だよせいぜい2倍だよ、みたいな主張をphiの宣伝時にやってて首を傾げたな
その理屈ならxeon倍積みましょう、って宣伝すりゃええやろと

結局、極端なオーダーではやっぱりGPUやphiの方が有利なケースがあるんだろうなと理解したけど

例の奴貼っておきますね。どうしても問題の性質や書き方やコンパイラに
依存する部分が大きいからな……
http://www.hpc.co.jp/benchmark20121113.html
http://www.hpc.co.jp/benchmark20130201.html
http://www.hpc.co.jp/benchmark20130329.html

やっぱこの人たちすごいな。
久々に見て回ったらK40の新機能のGPU Boostが他で使えるとか見つけてしまった。
水冷化してないけど試してみるか。

cuda 6.0正式リリースきた
https://developer.nvidia.com/cuda-downloads

もしかしなくても、
また面倒くさいことこの上ない初期設定をしないといけないのか。
インストールとVS2012で拡張子変えて保存するだけで動くようになって欲しいよ。
エラーがでると、どの設定がミスったのかバカには分からんのですよ。

>>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 はこの方法でできた。

>>843
VS ExpressだとNsight入らないんじゃない?

大して変わってないくせに開発環境変えるなよな

これまで開発したプログラムをmaxwellアーキテクチャーで動作させるには
5.5までのtoolkitでptxを吐かせるのか、6.0に移行するしかない模様。

>>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)用にする必要は無いかも。

>>847
これ
http://docs.nvidia.com/cuda/maxwell-compatibility-guide/#axzz30vrhsMg6
は釣りってこと? Σ(-д -;)

>>848
>>846 で合ってると思う。
Gxx→FermiやFermi→Keplerのときも、
新アーキテクチャ非対応な古いToolkitで作ったcubinは使えなかったはず。

Jetson買った人いる？

Jetsonってなんだと思ってぐるぐるしたら、NvidiaのRasPiか
RasPiより性能大分良いんだろうが、でも、2万超えは高いな

自動車用じゃん。スレチだろ

べつに限定されてはいない

組み込み用といだけ

この手の奴にBTデフォでついてんのみたことない
今後の組み込みの方向性的に必須なのに

CUDA Tooklit を6.0にしたらGPU稼働率が下がったんだけど気のせい?

>>855
Ver変えたら能率が大きく違ったりするのはよくあることだからなあ……

CUDA Tooklit を5.0から6.0にしたら
数値計算プログラムの挙動がおかしくなったorz
おんなじような事になった人いますか?

連投すんません。数値計算上の安定化を入れたら解決しました。
浮動少数演算の癖がこれまでと違うのかも・・・。

安定化って何したんですか？

http://ja.wikipedia.org/wiki/%E6%95%B0%E5%80%A4%E7%9A%84%E5%AE%89%E5%AE%9A%E6%80%A7

>>859
非線形最小二乗法のプログラムで、
一回の反復で更新する解の量を少し減らしたら安定しました。
CPUプログラムよりもGPUプログラムの場合に、
初期値からとんでもなく離れていってしまう場合が多いように感じます。

>>859
http://en.wikipedia.org/wiki/Nelder%E2%80%93Mead_method
4. Expansion
のパラメータγを通常2とするところ、
1.9-2.0の間で初期値に応じて変化させるようにしました。

>>860-862
サンクス
誤差の拡大を抑えるってことなんですね
でもCUDAバージョンの違いで問題が出るってなんだろ？
へんな最適化がされてしまってるのかなあ

演算の挙動が論理的に変わるような変更ってあったっけ？

http://developer.download.nvidia.com/compute/cuda/6_0/rel/docs/CUDA_6_Performance_Report.pdf

CUDA 6 Report

はじめて CUDA いじってるんだけど、
cu のコンパイルって こんなに時間がかかるものなの？
ホスト側の修正しかしてないときでも
一分近くかかってる。そういうもんなのかな。。。

nvcc にも /MP スイッチみたいなのがあるのですか？

ちなみにソースは正味100行足らずの試験的なもので、ビルドに40秒。そんなもん？
コンパイル環境は [email protected], メモリ16GB
Visual Studio 2012 Express
ソースもコンパイラもSSDにおいてある。

ホスト側とデバイス側と極力分離してみては如何でしょう。
VisualStudioでどう設定するのか知らんけど。

コンパイラにverbose出力ないのん

とりま、使うデバイス以外のcompute capablityはオフにしとけば。

GPGPU良いな。一部のプログラムには革命的な変化じゃないか？

そう、"一部の"プログラムにはな･･･

…それは分かるけど、面白味が薄くなるような…

ピーキーすぎて俺には無理だよ

もうCUDA6.5の話？
http://www.alexstjohn.com/WP/2014/06/01/cuda-6-5/

Windows 上で CUDA と MPI を組み合わせる場合、
お勧めの実装は？

IntelMPI は大学でライセンス持っているので自由に使えるけど
他でも使いたいならMSのほうがいいの？

CUDA SDK 付属の simpleMPI は、HPC Pack SDK 2008
入れろとコメントにあるけど。

>>876
openmpi
intelやmpich2系はlocalonlyオプションつけないと動かない。

>>877

thanks! でも Windows サポートが不安だなぁ。

binary support for a Microsoft Windows Open MPI build has been discontinued

MS の MPI じゃダメなのかかな？とりあえず HPC Pack SDK 2008
入れて CUDA SDK 付属の simpleMPI を動かしてみます。

cuda fortranについて初歩的な質問なのですが、
PGI workstationで、-cublasとオプションを付けてコンパイルしようとすると、
『fatal error LNK1104: cannnot open file 'libcublas.lib'』とエラーが出てしまいます。
この場合、どうしたら良いですかね?
どなたか教えてくださると助かります。

すいません。スレチでした。
超初心者用のスレで質問します。

>>875
driverはもう6.5なんだね
第二世代maxwellが近いのか？

CUDA Device Query (Driver API) statically linked version
Detected 1 CUDA Capable device(s)

Device 0: "GeForce GTX 750 Ti"
CUDA Driver Version: 6.5
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

http://devblogs.nvidia.com/parallelforall/low-power-sensing-autonomy-nvidia-jetson-tk1/#more-3339
単精度だがなるほど

>>882
マーケットが広がるのはよい事だね。
CUDAに還元されると折りがたい。

6.5 RCきたぞ。ようやくVS2013対応か。

vs2008はまだサポートされるの?

なくなった

nForceのビデオカードが３つささるマザーボードを手に入れたのですが
OpenCVのCUDAを使った超解像度計算を高速にするのに
適当なビデオカードを３つ刺してみようかと思っています
しかしビデオカードが複数あるとかえって動かすのが難しくなってしまうのでしょうか
OpenCVなら、その辺も自動でやってくれたりしないのでしょうか

http://bighow.net/4371742-Error_OpenCV_with_CUDA_using_TBB_for_multiple_GPUs.html