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

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


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

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/
2 名前：デフォルトの名無しさん mailto:sage [2012/09/23(日) 23:18:20.77 ]: 関連サイト
CUDA
www.nvidia.co.jp/object/cuda_home_new_jp.html

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

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

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

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

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

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

KNOPPIX for CUDA
www.yasuoka.mech.keio.ac.jp/cuda/
3 名前：1 忍法帖【Lv=38,xxxPT】(2+0：5) mailto:sage [2012/09/23(日) 23:20:02.29 ]: テンプレここまでです。変更点はcudaさわってみた(gpgpu.jp/article/61432191.html)を
リンク切れのため外したことのみです。
4 名前：デフォルトの名無しさん mailto:sage [2012/09/24(月) 00:35:21.83 ]: >>1乙
5 名前：デフォルトの名無しさん mailto:sage [2012/09/24(月) 00:53:05.44 ]: >>1乙

CUDAでがんばっていきやしょう！
6 名前：デフォルトの名無しさん mailto:sage [2012/09/25(火) 20:45:00.45 ]: マンデルブロ集合を描いてフルHDサイズの画像を保存するプログラム書てワクワクしてたら、
CUDAよりCPUのシングルスレッドの方が速かった (´・ω・`)

この程度の計算だとメモリ転送がかなり足を引っ張るんだね
7 名前：デフォルトの名無しさん mailto:sage [2012/09/25(火) 20:47:36.80 ]: たくさんのお仕事がないとGPUさんはがんばれない^^；
8 名前：デフォルトの名無しさん mailto:sage [2012/09/25(火) 22:01:29.76 ]: メモリ転送イヤポだから、AMD、IntelのiGPUはCPUの管理のメモリをGPUからでもアクセス
出来るようにする方向なんだろ
9 名前：デフォルトの名無しさん mailto:sage [2012/09/25(火) 23:05:27.26 ]: ユニファイドメモリってなんか先祖返りな不思議な気分ね
あとからGPUだけ交換できない時代にもなって寂しくなりそうだ

実現したらCPUがそのまま浮動小数点モンスターになって、ドライバもDirect3DコマンドをCPU処理で完結して、
CPUが直でディスプレイコントローラ叩くことになるんだろかね
10 名前：デフォルトの名無しさん mailto:sage [2012/09/25(火) 23:34:41.10 ]: メモリ空間統合したほうが絶対イイね。
データのコピー転送オーバーヘッドはあほらいい。
11 名前：デフォルトの名無しさん mailto:sage [2012/09/26(水) 17:04:17.48 ]: コピー転送オーバーヘッドを調べる方法ないのか
12 名前：デフォルトの名無しさん [2012/09/26(水) 18:23:45.59 ]: 16EiB の壁はいつ来るだろう
13 名前：デフォルトの名無しさん mailto:sage [2012/09/26(水) 20:30:47.96 ]: >>11
イベントで計測できるんじゃない？
14 名前：デフォルトの名無しさん mailto:sage [2012/10/08(月) 16:16:36.47 ]: 書籍「CUDA BY EXAMPLE」の第7章テクスチャメモリを読んでるんだが、
意味が分からない。

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

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

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

これって読み取り専用というよりは、たとえテクスチャとしてバインドしようが、
依然グローバルメモリとして使うこともでき、かつ tex1Dfetch などで読み取れば、
特別なキャッシュが働いて近傍への読み取りは速くなる、という事？
15 名前：hoge [2012/10/08(月) 16:59:48.04 ]: CUDAプログラミングを体験したいのですが，CUDAのできる格安ノートPCを教えてください．
16 名前：デフォルトの名無しさん mailto:sage [2012/10/08(月) 17:45:27.76 ]: CUDAのプログラミングを体験したいのならエミュレーションで十分
17 名前：デフォルトの名無しさん mailto:sage [2012/10/08(月) 18:24:26.17 ]: >>14
そういうことだと思う。
グラフィックスやってると普通の感覚なんだけど、
テクスチャ読み出ししてテクスチャにレンダリングするのは常套手段。
汎用コンピューティング時にテクスチャとしてデータを読むときの利点は
テクスチャ用の高帯域バスやキャッシュ、そしてフィルタリング用の固定機能ハードウェアを利用でき、
よりGPUを効率的に扱えることにあると思う。
18 名前：デフォルトの名無しさん mailto:sage [2012/10/08(月) 19:06:04.44 ]: >>17
すいません、ちょっと質問です。

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

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

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

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

テクスチャ用の特別なキャッシュ機構がある事については、
「CUDA BY EXAMPLE」に載っていましたから分かりました。
19 名前：17 mailto:sage [2012/10/08(月) 21:21:05.42 ]: >>18
自分は後藤さんの記事を参考にしているよ。
たくさんあるけど、いくつか紹介。

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

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

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

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

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

CUDA テクニカルトレーニング
Vol I:CUDA プログラミング入門
www.nvidia.co.jp/docs/IO/59373/VolumeI.pdf
（103ページ目のスライド）
20 名前：デフォルトの名無しさん mailto:sage [2012/10/08(月) 22:47:17.84 ]: >>19
ありがとうございます。
参考にさせていただきます。

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

いま nVidia のサイトで調べているのですが、なかなか見当たらないです。
公式には出てない情報なのでしょうか。
21 名前：デフォルトの名無しさん mailto:sage [2012/10/09(火) 00:35:20.42 ]: テクスチャフィルレートが相当すると言えなくもないけど、フォーマットによって変わるから参考にしかならない
概算はTex数×コアクロック（シェーダクロックではない）で見積もれる

テクスチャメモリを使うとテクスチャキャッシュを通るから、汎用のキャッシュがない上にコアレスアクセスの条件が厳しいG80～GT21xでは有効だった
Fermiはテクスチャユニット数の比率が減らされた上に、テクスチャキャッシュより汎用キャッシュの方が大容量になったので、むしろ遅くなることもあった
完全に予想だが、Keplerは（線形補間やテクスチャ端の丸め処理を手動でやる必要がなければ）おそらくテクスチャメモリを使っても使わなくてもそんなに変わらない
22 名前：17 mailto:sage [2012/10/09(火) 00:40:24.46 ]: >>20
そのGPUは最新世代のKeplerアーキテクチャだね。

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

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

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

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

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

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

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

しかし、このGPU、CUDAコア数に対してメモリ帯域が残念すぎないか？
Keplerアーキ自体が以前と比べてそういう傾向あるけど、それにしてもヒドイｗ
キャッシュがあるから大丈夫なんかな？
どうであれ、演算/ロード比が相当大きくないと性能出すの難しいかもね。
23 名前：デフォルトの名無しさん mailto:sage [2012/10/09(火) 19:04:03.29 ]: >>21
> 概算はTex数×コアクロック（シェーダクロックではない）で見積もれる

ありがとうございます。

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

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

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

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

がんばります。
24 名前：デフォルトの名無しさん mailto:sage [2012/10/12(金) 16:08:53.39 ]: PTXでブロックまたがってすべてのスレッドでグローバルメモリの同期やりたい時ってmembar.glでいいんだよね多分。
25 名前：デフォルトの名無しさん mailto:sage [2012/10/13(土) 23:43:36.29 ]: ホスト側のコードだけを書いた cu ファイルと、
デバイス側のコードだけを書いた ptx ファイルとをリンクして
ひとつの exe ファイルを作る方法はあるでしょうか。

もしあれば、やり方を教えてください。
26 名前：デフォルトの名無しさん [2012/10/14(日) 00:10:02.48 ]: >>25
ホスト側が面倒になるけどDriver APIとか
27 名前：デフォルトの名無しさん mailto:sage [2012/10/14(日) 01:00:51.58 ]: >>26
やはりそれしかないですか・・・
28 名前：デフォルトの名無しさん mailto:sage [2012/10/14(日) 01:04:08.72 ]: なんでそんなことをしたいのかが気になる。
29 名前：デフォルトの名無しさん mailto:sage [2012/10/14(日) 01:52:04.34 ]: >>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 などと一気にコンパイルできたらさぞ快適だろうなと思い、
質問した次第です。
30 名前：デフォルトの名無しさん [2012/10/14(日) 01:55:27.84 ]: Makefile
31 名前：デフォルトの名無しさん mailto:sage [2012/10/14(日) 02:04:07.01 ]: >>30
あぁ、そっちでコンパイルするファイルやコンパイル方法を制御するわけですね。

挑戦してみます。
32 名前：デフォルトの名無しさん mailto:sage [2012/10/14(日) 04:55:38.24 ]: PTXのコードをインラインアセンブラを使って直接cuファイルの
中にかけばいいじゃん。
33 名前：デフォルトの名無しさん mailto:sage [2012/10/14(日) 09:11:24.52 ]: >>32
知りませんでした。
「NVIDIA CUDA C Programming Guide Version 4.2」を「inline」で検索してみましたが、
__noinline__ や __forceinline__ の記述しかなかったです。
どこに詳細が載っているのでしょうか。

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

ただ問題は、それだと C 言語で書いたカーネル関数が、
nvcc によってどのような ctx コードにコンパイルされるか、
という部分が調べられない事です。
34 名前：デフォルトの名無しさん mailto:sage [2012/10/14(日) 09:16:31.85 ]: >>33
つ developer.download.nvidia.com/compute/DevZone/docs/html/C/doc/Using_Inline_PTX_Assembly_In_CUDA.pdf

試してないけど、nvccが出力したPTXのコードをインラインアセンブラの形式で
書き換えることも出来るんじゃない?
35 名前：デフォルトの名無しさん mailto:sage [2012/10/14(日) 09:35:03.51 ]: >>34
ありがとうございます。
なんというか、少々独特のインラインアセンブラ構文ですね。

今の環境より実験がやりやすくなるか調べてみます。
36 名前：デフォルトの名無しさん mailto:sage [2012/10/14(日) 14:53:23.16 ]: >>35
GCCのインラインアセンブリ構文がこういうのだよ
37 名前：デフォルトの名無しさん mailto:sage [2012/10/14(日) 15:28:15.65 ]: >>36
d.hatena.ne.jp/wocota/20090628/1246188338
これ見ると、たしかに同じですね。

インラインアセンブラはその昔 VC++ でしか使ったことがなかったもので
38 名前：デフォルトの名無しさん mailto:sage [2012/10/14(日) 20:42:43.14 ]: インラインアセンブラは今回の目的には合いませんでした。

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

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

みなさん、ありがとうございました。
39 名前：デフォルトの名無しさん mailto:sage [2012/10/16(火) 07:24:30.85 ]: CUDA5で美味しい事あるの？
40 名前：デフォルトの名無しさん mailto:sage [2012/10/16(火) 07:37:48.25 ]: >>39
新機能を使わないんだったら全然美味しくない。
CUDA5でビルドしたらかなり遅くなった。
41 名前：デフォルトの名無しさん [2012/10/16(火) 09:54:43.52 ]: CUDA 5 Production Release Now Available
CUDA Downloads | NVIDIA Developer Zone
developer.nvidia.com/cuda/cuda-downloads
42 名前：デフォルトの名無しさん mailto:sage [2012/10/16(火) 22:33:51.56 ]: 早く５の報告しやがれ
43 名前：デフォルトの名無しさん mailto:sage [2012/10/17(水) 03:16:50.02 ]: 4Gamer.net ― NVIDIA，「CUDA 5」を正式発表。第2世代Kepler「GK110」に向けた準備が整う
www.4gamer.net/games/076/G007660/20121016013/
44 名前：デフォルトの名無しさん mailto:sage [2012/10/17(水) 08:26:25.67 ]: >>42
普通に動いてるよ。
45 名前：デフォルトの名無しさん mailto:sage [2012/10/17(水) 23:16:16.73 ]: >>43
Nsightはプロファイラーも付いてるのか。
こりゃいい。
46 名前：デフォルトの名無しさん mailto:sage [2012/10/18(木) 10:47:37.44 ]: >>43
読んでみたけど
これはGeＦｏｒｃｅ切り捨てってこと？

今まで十分遊んだろ
これからはまともにGPGPUしたかったら、金出してTesla買えや
ていう風に読める
47 名前：やんやん ◆yanyan72E. mailto:sage [2012/10/18(木) 12:03:55.32 ]: それは、Kepler発表の時からゲーム用のKepler1と
GPGPU向けのKepler2があるってことになってた。
Kepler1があまりにGPGPUに向いてなくてGeforce680あたりを
買った人はがっかりしてたよ。
48 名前：デフォルトの名無しさん mailto:sage [2012/10/18(木) 12:23:14.57 ]: >>43
Eclipse用のNsightも出てLinuxやMacでも開発しやすくなるのは大きいかも。

>>46
Dynamic ParallelismはGK110以降での対応でMaxwell世代ではコンシューマ向けでも対応するのでは？
GPUDirectはクラスタ向けの機能で差別化されても仕方ない気がする。
49 名前：デフォルトの名無しさん mailto:sage [2012/10/18(木) 15:18:50.07 ]: >>47
ゲーマーにGPGPUっていらんだろうからな
要らないのを付けて高い値段･高消費電力になって売れないものになるなら削れだろ
CUDAする奴はとんがったことする奴だろ。そんな奴ならKepler2のTeslaぐらい買うだろうからな
買えない貧乏人はAMDのradeonに移行しろだな
50 名前：デフォルトの名無しさん mailto:sage [2012/10/18(木) 19:01:56.13 ]: Mac はどうか知らんが、Linux は Windows 版に比べて、
どうしてもドライバのチューニングが徹底されていない感じがする。

SDK 内のサンプルを動かしてみても、
Windows 上で動かしたときより明らかにフレーム数が落ちる。
51 名前：やんやん ◆yanyan72E. mailto:sage [2012/10/18(木) 19:34:39.18 ]: そりゃ、いくらDRIが実装されたからといってX11なんていう
太古のグラフィックAPI使ってるんだから、そんなもんじゃないの？
本気を出させたければWayland待ち。
52 名前：デフォルトの名無しさん mailto:sage [2012/10/18(木) 19:55:50.80 ]: >>51
遅いのはCUDAじゃなくて、その結果を表示する
グラフィックスライブラリの方ってこと？

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

今度ためしてみよ
53 名前：やんやん ◆yanyan72E. mailto:sage [2012/10/18(木) 22:37:06.54 ]: フレーム数計る時点で、グラフィックカードに描画させてるんだよね？
その描画をグラフィックス・ライブラリが足引っぱってるんじゃないかってこと
CUDA自身はプログラムをGPGPU用のアセンブリ言語に変えて
GPGPUに実行させるだけだから、あまり差が出るとは考えにくい。
54 名前：デフォルトの名無しさん mailto:sage [2012/10/19(金) 00:22:17.74 ]: そう言えば Yellow Dog Linux for CUDA 使ってる人いる？
どんな感じなの？
55 名前：デフォルトの名無しさん mailto:sage [2012/10/19(金) 12:39:47.00 ]: Linuxなら、GUI止めないとカーネルによっては処理速度ががた落ちする。
使えるGVRAMも激減する。
56 名前：デフォルトの名無しさん mailto:sage [2012/10/19(金) 19:37:02.47 ]: CUDA + GUIつっても、サンプルでXが関わるところなんてウィンドウの枠だけじゃないか？
あとはOpenGLで描画されていると思うが
57 名前：デフォルトの名無しさん mailto:sage [2012/10/20(土) 00:42:13.70 ]: >>55
Windowsのほうがガタ落ちだし、使えるメモリも少ない。
グラフィックスを使うと遅くなるのはXの問題だから。
58 名前：デフォルトの名無しさん mailto:sage [2012/10/20(土) 16:27:41.31 ]: dynamic parallelism は GeForce じゃ無理なんですか
59 名前：デフォルトの名無しさん mailto:sage [2012/10/20(土) 17:35:11.22 ]: うん。
60 名前：デフォルトの名無しさん mailto:sage [2012/10/20(土) 17:41:18.83 ]: 調べたなかではＧＤＲＡＭのみのように見えるんだけど、
テスラだとＬ１、Ｌ２、シェアードメモリもＥＣＣ保護されてるの？
それともＬ１、Ｌ２くらいの容量なら気にしなくてもいいのかな？
61 名前：デフォルトの名無しさん mailto:sage [2012/10/20(土) 20:24:55.79 ]: レジスタも。
62 名前：デフォルトの名無しさん mailto:sage [2012/10/20(土) 20:47:40.75 ]: >>58
今のところTesla K20のみだったはず
63 名前：デフォルトの名無しさん mailto:sage [2012/10/22(月) 20:36:15.81 ]: 一般人向けは2014年まで待てとか遅すぎる
64 名前：デフォルトの名無しさん mailto:sage [2012/10/22(月) 21:06:37.69 ]: GK110はいつになったら一般向けで出てくるのやら…
65 名前：デフォルトの名無しさん mailto:sage [2012/10/22(月) 23:32:49.15 ]: >>64
ないと思うのは俺だけか
66 名前：デフォルトの名無しさん mailto:sage [2012/10/22(月) 23:35:31.31 ]: >>65
gtx780とかじゃないか？
来年の春だった気がする。
67 名前：デフォルトの名無しさん mailto:sage [2012/10/23(火) 04:53:17.04 ]: GTX 780はKepler1の改良版だって聞いたぞ。
68 名前：デフォルトの名無しさん mailto:sage [2012/10/23(火) 08:47:20.31 ]: 一般人向けでダイナミックなんちゃらが使えるのはMaxwellからとか

AMDが2013年中に簡単にOpenCL対応アプリをかけるようにしてきたらどうするんだろ
69 名前：デフォルトの名無しさん mailto:sage [2012/10/23(火) 16:30:26.65 ]: NVIDIA Visual Profiler v4.2をCentos6.2で使おうとしてるんだけど、

No Timeline
Application timeline is required for the analysis.

と出て解析できない。
調べたらLD_LIBRARY_PATHに/usr/local/cuda/lib64やらを追加せよとあったんでやってみたけど状況変わらず。
どなたか同様な症状に出くわした方はいらっしゃいませんか？
70 名前：デフォルトの名無しさん mailto:sage [2012/10/23(火) 22:26:20.49 ]: >>69
CUDAプログラミングはまだまだ敷居が高いね･･･
71 名前：デフォルトの名無しさん mailto:sage [2012/10/24(水) 10:52:47.74 ]: nvcc ***.cu -O2 -Xcompiler -O2
のようにO2を重ねるのは無意味ですか？
前者のO2はGPU用，後者のO2はCPU用と勝手に思っていたんですが，
同じ事を繰り返しているような気がしてきました．
72 名前：デフォルトの名無しさん mailto:sage [2012/10/24(水) 12:32:18.86 ]: >>71
意味があるのか、どのような意味があるのかまでは分からんが、
とりあえず、「同じ事を繰り返しているのかどうか」については、
出力されたファイルを比較すれば直ぐに分かると思うぞ。

バイナリで比較してもいいし、アセンブラコードで比較してもいい。
73 名前：71 mailto:sage [2012/10/24(水) 13:25:43.83 ]: ptxで２つある場合，前者のみ，後者のみ，両方無い場合を比較しましたが，
冒頭の***.iファイルの名前が微かに違うのみで差はありませんでした．
両方消しても差が出ないのは？ですが，
重ねても意味は無さそうであることが分かりました．

>>71
ありがとうございました．
74 名前：デフォルトの名無しさん mailto:sage [2012/10/24(水) 14:43:07.65 ]: >>73
今のgccのディフォルトが-O2相当なんで、書かなくても変わらないのはその所為。
試しに、-O3とか-O1との組み合わせを試してみたら?
75 名前：デフォルトの名無しさん [2012/10/25(木) 04:28:58.35 ]: 登録ユーザーサイトが復旧したよ
76 名前：デフォルトの名無しさん mailto:sage [2012/10/25(木) 21:42:51.21 ]: k20はやっぱり高いな。
38万だそうだ。
20万切ってくれないと買えない。
77 名前：デフォルトの名無しさん mailto:sage [2012/10/27(土) 22:36:10.01 ]: dynamic parallelism対応のGeforce（GTX8XX?)が出たら
画像とか動画を扱うソフトは瞬く間にCUDA完全対応になるのかね?
78 名前：デフォルトの名無しさん mailto:sage [2012/10/28(日) 00:23:20.16 ]: んなわけない
79 名前：デフォルトの名無しさん mailto:sage [2012/10/28(日) 00:40:32.58 ]: dynamic parallelismができるからCUDAが劇的に簡単になるわけじゃないから。
Reductionとかで効果はあるけど。
80 名前：デフォルトの名無しさん mailto:sage [2012/10/28(日) 03:58:28.19 ]: そもそもReductionはマルチパスにしないで
2パスで済ませた方がいいのは、
CUDAのreductionトレーニングでも明らか
81 名前：デフォルトの名無しさん mailto:sage [2012/10/29(月) 13:40:34.37 ]: CUDAのプログラム作って動かしたいです
自分のMacbookは、グラフィックのチップがIntel GMA X3100なんですけど、
NVIDIAじゃないとCUDAは使えないんですか？
82 名前：デフォルトの名無しさん mailto:sage [2012/10/29(月) 15:35:40.18 ]: ここで聞いて良いのか分からないので、不適切なら誘導お願いします。

GeForceの省電力の状態(P0～P12)をGetLastInputInfo-GetTickCountに
応じて切り替えるようなソフトを作りたいのですが、
P0～P12を切り替えるAPI関数はありませんか?
83 名前：デフォルトの名無しさん mailto:sage [2012/10/29(月) 18:51:19.74 ]: NVAPIをhackすればできる
84 名前：デフォルトの名無しさん mailto:sage [2012/10/31(水) 14:40:39.17 ]: CUDAカーネルの中で呼び出す関数に特定の処理を入れるとカーネル自体が読み込まれなくなります
具体的にはプロファイラで実行時間見てみるとカーネル自体が表示されず、一瞬で動作終了する状況です
一応、その特定の処理の部分をコメントアウトするときちんと実行されます（当然正しい結果は出ませんが）
こういったことはどういう状況で起こり得るのでしょうか？
85 名前：デフォルトの名無しさん mailto:sage [2012/10/31(水) 14:43:09.18 ]: >>84
カーネル呼び出した時にエラーが起きてるんでしょ。
エラーチェックしていないんじゃないの?
86 名前：デフォルトの名無しさん mailto:sage [2012/10/31(水) 14:54:52.27 ]: >>84
cudaGetLastError()は何と言っている?
87 名前：デフォルトの名無しさん mailto:sage [2012/10/31(水) 14:58:04.91 ]: >>85
即レスありがとうございます
正にその通りでした。単にメモリの要求量がおかしかっただけみたいです
初歩的すぎるミスに自己嫌悪…
88 名前：デフォルトの名無しさん [2012/10/31(水) 16:49:46.93 ]: 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
89 名前：デフォルトの名無しさん [2012/11/03(土) 02:08:23.38 ]: Nvidia Geforce forum is back from the dead
www.fudzilla.com/home/item/29337-nvidia-geforce-forum-is-back-from-the-dead
90 名前：デフォルトの名無しさん mailto:sage [2012/11/07(水) 15:17:37.33 ]: CRS形式の行列格納サンプルコードってどこかにない？
91 名前：デフォルトの名無しさん mailto:sage [2012/11/07(水) 15:59:46.17 ]: いくらでもあるだろ
圧縮方法を理解できたらサンプルもいらんな

1 2 3 4
2 5 6 7
3 6 8 9
4 7 9 10
92 名前：デフォルトの名無しさん mailto:sage [2012/11/08(木) 02:56:28.41 ]: >>91
圧縮方法はわかったんですがコードに上手く起こすことができなくて困っていたんです。何かいいサンプルがあれば教えていただけると助かります。
93 名前：デフォルトの名無しさん mailto:sage [2012/11/08(木) 10:51:49.85 ]: 馬鹿には無理。
94 名前：デフォルトの名無しさん [2012/11/12(月) 06:21:00.64 ]: CUDA5は既存のGPUに入れると遅くなるの？
95 名前：デフォルトの名無しさん mailto:sage [2012/11/12(月) 14:32:24.27 ]: 研究室でCUDA用にGTX680搭載PCの導入が決定してしまったんだが評判悪いとはいえ流石に今使ってる560Tiよりは性能いいよね？
96 名前：デフォルトの名無しさん mailto:sage [2012/11/12(月) 23:48:40.00 ]: Tesla K20きたぞ
97 名前：95 mailto:sage [2012/11/13(火) 01:28:48.03 ]: >>96
予算処理上の都合だったらしい。
98 名前：デフォルトの名無しさん mailto:sage [2012/11/13(火) 03:25:03.66 ]: 最近プログラム入門した
CUDAとか聞くとワクワクするけど物理の知識も科学の知識も特にないので
数百万スレッド並列で処理するネタが思いつけなくて悲しい思いになる
もっとちゃんと勉強しておけば良かった
99 名前：デフォルトの名無しさん mailto:sage [2012/11/13(火) 05:46:01.18 ]: 京が3位に
100 名前：デフォルトの名無しさん [2012/11/13(火) 06:14:28.20 ]: 東工大の学生たちはもうGK110貰ってるの？
101 名前：デフォルトの名無しさん mailto:sage [2012/11/13(火) 23:01:28.72 ]: Intelがついに来るぞ
pc.watch.impress.co.jp/docs/news/20121113_572526.html

ソースの改変が少しでパラレル計算ができるとのことだが、実際の所どうなんだろうね。
102 名前：デフォルトの名無しさん mailto:sage [2012/11/14(水) 00:25:44.33 ]: >>101
nVIDIAが押されて、もうちっと貧乏客を引き込むマーケティングをやってくれんかな。

一般のビデオカードで定格の80%までクロックを公式に落とせかつその速度なら
GPGPUの動作を保証。
これを是非やってほしい。仲がよいベンダーがいくつかあるし。
103 名前：デフォルトの名無しさん mailto:sage [2012/11/14(水) 16:56:20.91 ]: 開発環境やソフトウェアの安定性とか含めて、XeonPhi強そうだなぁ
104 名前：デフォルトの名無しさん mailto:sage [2012/11/14(水) 18:03:25.42 ]: XeonPhiは高いぞ
安いGPUは安い
Tesla買うならXeonPhiのほうがよさそうだが
105 名前：デフォルトの名無しさん mailto:sage [2012/11/14(水) 19:32:13.13 ]: 半年ぐらいしたら、$500くらいのローエンドXeon Phiが出るだろうから、純粋にアクセラレータとしてのteslaは厳しいかもなあ。
106 名前：デフォルトの名無しさん mailto:sage [2012/11/14(水) 20:08:20.86 ]: Phi触ってみてぇ。
OpenMPで簡単マルチコアプログラミング♪

スレッドオーバーヘッドが小さいことを願う･･･
107 名前：デフォルトの名無しさん [2012/11/14(水) 20:25:15.46 ]: SSEとかAVXみたいなのをちゃんと使える人じゃないと
TESLAのような性能はでないよ。
512bit演算命令が命だから。

ただのロジックを複数スレッド回したい人なら、
TESLAより速いかもね。かなりの無駄だが(笑)
108 名前：デフォルトの名無しさん mailto:sage [2012/11/14(水) 20:30:01.47 ]: 512bit演算命令ってのがあるのか？
AVXでも256bitだが・・・
109 名前：デフォルトの名無しさん mailto:sage [2012/11/14(水) 20:32:10.61 ]: VPUてので512ビット命令を処理するようだな
110 名前：デフォルトの名無しさん mailto:sage [2012/11/14(水) 20:33:10.12 ]: ま、経験上はベクトル命令はCUDAよりは扱いやすいよ
111 名前：デフォルトの名無しさん mailto:sage [2012/11/14(水) 20:37:54.55 ]: うん、イントリンシックでベクトリ処理書くの楽♪
条件分岐がめんどいけど、LNIはマスクレジスタをサポートしてたからだいぶ楽に書けそう。
しかも512bitもあるなんて最高すぐる。

あー、Phi触りてぇ～。
112 名前：デフォルトの名無しさん mailto:sage [2012/11/14(水) 21:32:04.73 ]: OEM向け1000個ロットでXeon Phi 5110Pが2650ドル
らすぃ
なんか価格でもTeslaやばそうだな

Intel，スーパーコンピュータ向けアクセラレータ「Xeon Phi 5110P」発表。60基のx86コアを1チップ上に集積
ttp://www.4gamer.net/games/049/G004963/20121111001/
113 名前：デフォルトの名無しさん [2012/11/14(水) 21:33:53.30 ]: むしろ値下げ合戦になればよい。
114 名前：デフォルトの名無しさん mailto:sage [2012/11/14(水) 21:47:31.12 ]: 合戦になるほど数競争起きる市場でもないべ
115 名前：デフォルトの名無しさん mailto:sage [2012/11/14(水) 21:53:57.69 ]: 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使いのおまいらがんばってくださいおながいします
116 名前：デフォルトの名無しさん mailto:sage [2012/11/14(水) 21:56:39.37 ]: 長文の上に間違えてーらorz
Fermiにもそれなりに機能を持たせたんだと思う→Geforceにもそれなりに機能を持たせたんだと思う
117 名前：デフォルトの名無しさん mailto:sage [2012/11/14(水) 22:59:58.87 ]: >>115
nVIDIAの株を空売りすれば儲かるということか。
118 名前：デフォルトの名無しさん mailto:sage [2012/11/14(水) 23:57:37.97 ]: phiはCPUに内蔵させるGPUコアと共通化させて
コスト落としつつマーケットシエア取る作戦かな？

そしたら、本気でnVidia終わるな
119 名前：デフォルトの名無しさん mailto:sage [2012/11/15(木) 00:11:23.74 ]: >>112
Phi、扱い易そうだな。
ベクタ演算器処理の記述法が気になるとこだし、
nVidiaがアセンブラのように複雑ってディスってたけど、
イントリンシック記述だったら簡単だし、
条件分岐のマスクまでサポートしてくれたら文句なしだ。

これ、マジで触ってみたいな。
120 名前：デフォルトの名無しさん mailto:sage [2012/11/15(木) 04:26:43.84 ]: CUDAは開発環境タダだけどXeon PhiはIntel Compiler必須だよね
121 名前：デフォルトの名無しさん mailto:sage [2012/11/15(木) 05:55:49.56 ]: ものいりだねえ、 Phi
122 名前：デフォルトの名無しさん mailto:sage [2012/11/15(木) 11:55:29.66 ]: >>120
GCCが対応するって書いてあったぞ
123 名前：デフォルトの名無しさん mailto:sage [2012/11/15(木) 15:34:52.68 ]: そもそもintelコンパイラもLinux版はフリーだよな
124 名前：デフォルトの名無しさん mailto:sage [2012/11/15(木) 18:21:21.39 ]: ここのスレ見てたら、phiを月2、3万位でレンタルする商売が出来そうな気がしてきた…
125 名前：デフォルトの名無しさん mailto:sage [2012/11/15(木) 20:33:05.20 ]: GPUだってAmazonのクラウドで借りたりできるし
Phiもそういうの出るだろう
126 名前：デフォルトの名無しさん mailto:sage [2012/11/15(木) 23:35:19.70 ]: >>114
まあ、ATI Streamが出てきたからといってTeslaが安くなったというわけじゃないからな。
しかし、今回の場合Phiの場合はコードの書きやすさからすると、CUDAの比じゃないから、
>>115にあるお互いの強みを生かして、切磋琢磨して値下げ合戦してほしいわ。
両方のコードを書いている身としては安くなってくれればどっちでもいいんだが。
127 名前：デフォルトの名無しさん mailto:sage [2012/11/15(木) 23:36:48.09 ]: Phi、4、5万で買えるようにならないかなぁ～。
128 名前：デフォルトの名無しさん mailto:sage [2012/11/17(土) 12:12:13.46 ]: Phいらないだろ
一晩中PC動かせばいいだけだろ
129 名前：デフォルトの名無しさん mailto:sage [2012/11/17(土) 12:35:26.08 ]: 動画エンコ用途でもあるまいにｗ
すでに一年中計算回してるような人に、これなら一ヶ月で済むよ、って訴求するのが筋の製品だろ
130 名前：デフォルトの名無しさん mailto:sage [2012/11/17(土) 13:00:33.16 ]: 数時間、数日動かして、後から些細なバグに気づいた時の何とも言えない気持ち

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

バグを直したら、その部分だけ再計算すればいいような仕組み
131 名前：デフォルトの名無しさん mailto:sage [2012/11/17(土) 14:05:25.11 ]: >>130
とりあえず賽の河原症候群と名付けておくよ
132 名前：デフォルトの名無しさん mailto:sage [2012/11/19(月) 21:17:02.44 ]: phiは本体CPUもXeon使った時の協調性とかで
パワー増すんだろうなぁ
Teslaやばいなぁ…

…投げ売りになってくれると嬉しいなぁ
133 名前：デフォルトの名無しさん mailto:sage [2012/11/19(月) 22:16:34.19 ]: しかし投げ売りの後に待っているのが撤退だとしたら…？
Xeon Phiには縁がなさそうだから気軽に触れるCUDAにがんばってほしいなぁ
134 名前：デフォルトの名無しさん mailto:sage [2012/11/19(月) 22:19:19.97 ]: 自作ゲームにCUDA利用してる奴っている？

いるなら、何に使ってる？
135 名前：デフォルトの名無しさん mailto:sage [2012/11/22(木) 18:48:11.46 ]: 使えもしないのに欲しくなってとりあえずダウンロードしているんだけど
クソミソに通信速度遅い。
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
136 名前：デフォルトの名無しさん mailto:sage [2012/11/22(木) 19:03:52.66 ]: >>135
ウチじゃ2.0MB/s位出ているぞ。
137 名前：デフォルトの名無しさん mailto:sage [2012/11/22(木) 20:53:34.08 ]: まじすか
一端回線切ってIPアドレス変更とかしても速度出ない・・・

OCN保土ヶ谷
138 名前：デフォルトの名無しさん mailto:sage [2012/11/23(金) 01:07:04.36 ]: >135
遅すぎ
間に無線とか入れてない？
139 名前：デフォルトの名無しさん mailto:sage [2012/11/30(金) 22:59:09.24 ]: 人柱はまだか
www.amazon.co.jp/dp/B00A31Q6QI/
140 名前：デフォルトの名無しさん mailto:sage [2012/12/01(土) 19:38:58.57 ]: >>139
高すぎ。

今だったら34万くらいで買えるだろ。
141 名前：デフォルトの名無しさん mailto:sage [2012/12/01(土) 23:01:57.20 ]: CPU内蔵のiGPUをPCの表示用に、dGPUをCUDA　GPGPU専用にする場合
やdGPUを2つ使って片方をPCの表示用に、もう片方のdGPUをCUDA　GPGPU専用にする場合
ってそれらが出来る(出来ない)マザー、CPU・APU、dGPUってある?
出来るのなら、これをPC表示用、これはGPGPU用って設定とかするの
する場合どうするんですか?
142 名前：デフォルトの名無しさん mailto:sage [2012/12/01(土) 23:12:12.90 ]: >>141
GTX 580/590とASUS Maximus V GeneとCore i7-3770Kの組み合わせなら出来た。
BIOSでどちらを表示用に使うか設定できる。
143 名前：デフォルトの名無しさん mailto:sage [2012/12/01(土) 23:59:21.54 ]: >>142
劇速れすありがとう
マザーのBIOSに設定があれば、iGPUとdGPUの場合は出来ると
思っていいのかな
144 名前：デフォルトの名無しさん mailto:sage [2012/12/02(日) 00:02:38.36 ]: >>143
CUDAのデバイス指定はアプリケーション次第だよ。BIOSは関係ない。
ちゃんとどのデバイスを使うか指定できるようになっていれば問題ないよ。
145 名前：デフォルトの名無しさん mailto:sage [2012/12/02(日) 10:03:50.91 ]: >>144
CPU内蔵の（CUDAが使えない）iGPUとグラボ側の（CUDAが使える）dGPUがあったとして、
今dGPUを表示用に使用して、iGUの方は眠らせるようにBIOSが設定されているのなら、
CUDAを使うとひとつのdGPUで表示もCUDAも使うことになると思う。

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

あと、ついでに俺も聞きたいんだが、そうやってiGPUで表示してdGPUでCUDAする場合、
cudaGLSetGLDevice関数などを使ったCUDAとOpenGLドライバとの相互運用はできるの？
（DirectXとの相互運用でもいいけど）
146 名前：デフォルトの名無しさん mailto:sage [2012/12/02(日) 10:25:11.74 ]: >>145
だから、BIOSで設定するのは画面の表示だけってことなんだよ。
CUDAでの利用はそれとは全く別に行えるよ。
BIOSでiGPUを表示用に選んでマザボにディスプレイをつないでから、
CUDA対応アプリでdGPUを選べばいいだけの話。

OpenGLは使ったことないけど、CUDAを使った限りでは相互運用は全く問題ない
ように見える。
147 名前：デフォルトの名無しさん mailto:sage [2012/12/02(日) 10:35:51.19 ]: >>146
すまん、言い方が悪かった。
その「BIOSでiGPUを表示用に選んでマザボにディスプレイをつないでから」が、
CUDAを使ったプログラム側からは操作できないから、
BIOSをいじる必要があるよねという（当たり前と言えば当たり前の）確認だけだったんだ。

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

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

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

誰かこの辺り分かる人いる？
148 名前：デフォルトの名無しさん mailto:sage [2012/12/02(日) 10:45:06.90 ]: >>145
グラフィックライブラリと相互運用する場合は
出力用GPUとCUDA用GPUは同じな必要があるんじゃない？
俺はそうしてる．
確かめたことがあるわけじゃないから無責任な言い方になるけど．
GTX 580(1)から二画面，GTX 580(2)から一画面のトリプルディスプレイやった時に
SDKのスモークパーティクルとか起動しなかった記憶がある．
149 名前：デフォルトの名無しさん mailto:sage [2012/12/02(日) 16:46:04.44 ]: 表示用デバイスでなくてもOpenGLは動かせるから、cudaと連携できると思う。
150 名前：デフォルトの名無しさん mailto:sage [2012/12/21(金) 19:28:45.07 ]: SDKのマーチングキューブのサンプルで、defines.hの中の #define SAMPLE_VOLUME ってところが
0だとあらかじめ用意された関数が、1（デフォ）だとファイルが読み込まれるんだけど、ここを0にしてもなにも表示されない
コードはそこ以外いじってないんだけどほかにも変更しなきゃいけない部分とかあるのかしら
151 名前：デフォルトの名無しさん mailto:sage [2012/12/22(土) 19:54:35.28 ]: >>150
ごめん自己解決した
VSでコードいじってたんだけど、すべてリビルドしたら表示されるようになった
152 名前：デフォルトの名無しさん mailto:sage [2012/12/26(水) 08:21:21.01 ]: Fermiでスレッドブッロクを512以上を指定すると、カーネルが起動しない。
Fermiはブロックごとに1024スレッド対応しているはずなので、
までレジスタが足りないからなのか、
シンプルなカーネルだと1024スレッドまでいける。
動かないならエラーで落ちて欲しいんだが。
153 名前： ◆4hloUmTGPY [2012/12/27(木) 10:24:29.87 ]: 質問です

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

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

ビルドは一応できて実行もできるんですが、気持ち悪いので何とかならないでしょうか
154 名前：デフォルトの名無しさん mailto:sage [2012/12/27(木) 11:09:28.44 ]: >>153
いんてりせんすがC++以外の表記に対応していないんだろ。
自分でぷらぐいんを書けばなんとかなるんじゃね?
尤も、いんてりせんすに頼るような奴に書けるかどうかは知らんが。
155 名前：デフォルトの名無しさん mailto:sage [2012/12/27(木) 22:41:04.67 ]: 2012じゃCUDAの環境がまだまだなんじゃね？
2010だとようやくこなれてきた感じがあるが。
156 名前： ◆4hloUmTGPY [2012/12/28(金) 11:54:50.91 ]: >>154 >>155
ありがとうございます
無理そうなので諦めます
157 名前：デフォルトの名無しさん mailto:sage [2012/12/29(土) 11:26:38.86 ]: こんな感じのカーネルがあって、

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

MPIとかSMPなどではこのような処理は帰ってこなくなるけど
CUDAのカーネルでは問題なく動く。
__syncthreads()っていうのは、分岐があってもWarp単位では分岐から外れたスレッドは単に何もしないだけで、
__syncthreads()がどっかで呼ばれたらとりあえず足並みを揃えることはする。
という理解でいいのかな？
158 名前：デフォルトの名無しさん mailto:sage [2012/12/29(土) 22:44:34.71 ]: 大体そんな感じ。分岐から外れたスレッドが生きているか死んでいるかは兎も角も。
159 名前：デフォルトの名無しさん mailto:sage [2012/12/30(日) 13:32:23.86 ]: >>158
ありがとう。
160 名前：デフォルトの名無しさん [2012/12/30(日) 15:06:29.69 ]: CC3.5のGeForceまだ？
161 名前：デフォルトの名無しさん mailto:sage [2013/01/03(木) 11:10:10.18 ]: K20が売れなくなるから当分無し
162 名前：デフォルトの名無しさん [2013/01/05(土) 10:06:43.42 ]: CUDAの性能でいったらGTX580>GTX680なんですか？
163 名前：デフォルトの名無しさん mailto:sage [2013/01/05(土) 15:19:19.49 ]: >>162
blog.accelereyes.com/blog/2012/04/26/benchmarking-kepler-gtx-680/

倍精度を使うならGTX580が圧倒的。
単精度ならモノによるが、基本680は足回りが遅い感じ。
164 名前：デフォルトの名無しさん [2013/01/05(土) 21:13:33.60 ]: >>163
ありがとです。
165 名前：デフォルトの名無しさん mailto:sage [2013/01/10(木) 18:01:12.19 ]: ブロック数がSM数以上の場合、ブロックでの動作が終了したら次のブロックにいくんですか？
166 名前：デフォルトの名無しさん mailto:sage [2013/01/10(木) 18:27:54.28 ]: 次のブロックっつーか、残りのブロックだな。
167 名前：デフォルトの名無しさん mailto:sage [2013/01/14(月) 00:28:51.24 ]: CG法をCUDAで実装してMatrixMarketの疎行列を求解しようとしているんですが
連立一次方程式のb要素はどのように設定したらいいのでしょうか？
MatrixMarketで与えられているのはAの疎行列だけなのでbをどのように設定したらいいのかわかりません
168 名前：デフォルトの名無しさん mailto:sage [2013/01/18(金) 16:14:19.91 ]: テクスチャメモリって使っててあまり早くならないんだけど，実際の効果ってどれくらいなの？
169 名前：デフォルトの名無しさん mailto:sage [2013/01/18(金) 18:10:26.59 ]: 質問です。

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

ちなみにVisualStudio2008でCUDA4.2という環境です。
170 名前：デフォルトの名無しさん mailto:sage [2013/01/18(金) 18:12:36.87 ]: >>169
ソースの該当箇所の引用ぐらいしろよ。
171 名前：デフォルトの名無しさん mailto:sage [2013/01/18(金) 18:17:57.52 ]: >>169
nvccでコンパイルする対象(*.cu)ではテンプレートは使えなかったと思う。
C++からCのライブラリを呼ぶ要領で分割コンパイルすればホスト側で使う分にはOK。
172 名前：デフォルトの名無しさん mailto:sage [2013/01/18(金) 18:30:58.43 ]: >>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 );
173 名前：デフォルトの名無しさん mailto:sage [2013/01/19(土) 20:26:41.74 ]: >>172
extern Cをしてないだけじゃないのか？
それにしても引数多いな。
174 名前：デフォルトの名無しさん mailto:sage [2013/01/19(土) 20:55:07.35 ]: 関数のテンプレートは >>171
template <typename Float, bool checkBottom>
__device__ void gpu_impl_sub(Float const * __restrict__ f, Float * __restrict__ fn, dim3 const &fDim, 以下略
みたいな感じで使ってるけど問題ないよ。
175 名前：169 mailto:sage [2013/01/20(日) 00:40:40.98 ]: .cu単体ではコンパイルできるんですけどビルドしようとするとCUDA特有の記述の部分で
エラーを吐きます。
でよく見てみるとその前に
nvcc fatal : A single input file is required for a non-link phase when an outputfile is specified
みたいなエラー吐いてるんでリンク関連で間違えてるんですかね…

ただSDKのソースは動くんでここら辺は間違っていないと思うんですけど…
176 名前：デフォルトの名無しさん [2013/01/20(日) 02:11:59.82 ]: >>168
Fermiより前のL2キャッシュが無いようなGPUだと効果あるよ。
あとハードウェア補間を活かせる分野
177 名前：デフォルトの名無しさん mailto:sage [2013/01/20(日) 02:20:35.01 ]: >>176
fermiでも場合によっては効果はある。
でも一番いいのは線形補完で使う時だな。
178 名前：デフォルトの名無しさん mailto:sage [2013/01/21(月) 13:30:50.30 ]: >>175
リンク前の個々の .cuファイルのコンパイルの時点でのエラーのように見える。
どこかで nvcc -c -o a.o a.cu b.cu みたいな呼び出ししてるんじゃ。
179 名前：デフォルトの名無しさん mailto:sage [2013/01/23(水) 14:40:54.94 ]: SMが10個あって処理の数が合計100個ある場合SMはそれぞれ10回動く？
それとも重い処理があって時間がかかってるSMがあれば空いてる別のSMが担当する？
180 名前：デフォルトの名無しさん mailto:sage [2013/01/25(金) 23:00:07.64 ]: なんとも答えにくい初心者の質問だな
181 名前：デフォルトの名無しさん [2013/01/26(土) 02:29:08.92 ]: 質問
TESLA K20買おうと思ってるんだけど、プログラミングとは別にエンコでも使おうかなと思ってる。
やっぱりGTX680とかGTX580と較べて全然性能違う？（体感）
開発にはVS2012Pro持ってるからそれに追加しようと思ってる
182 名前：デフォルトの名無しさん [2013/01/26(土) 02:39:37.68 ]: 宝の持ち腐れ臭
183 名前：デフォルトの名無しさん [2013/01/26(土) 02:44:45.43 ]: >>182
一部のスーパープログラマー以外皆大なり小なり宝の持ち腐れだろう
184 名前：デフォルトの名無しさん mailto:sage [2013/01/26(土) 04:14:43.40 ]: >>180
色々見たけど書いてないから自明なのかな？
185 名前：デフォルトの名無しさん mailto:sage [2013/01/26(土) 09:11:31.41 ]: >>181
GTX680でええやん、いくらなんでも値段が凄まじすぎる＞K20
それにGPUエンコは品質が……フィルタはともかく、AVCに落とすにはQSV使う方が速いし
186 名前：デフォルトの名無しさん [2013/01/26(土) 12:22:12.43 ]: >>185
まーなー、BOINCもやってるからTESLAちゃんにBOINCさせたら
どうなるのかなって気持ちもあったりするんだよな
187 名前：デフォルトの名無しさん [2013/01/26(土) 13:45:33.61 ]: まぁ一応言っておくと、性能以前に
Dynamic Parallelismの対応有無の差が大きい >GTX680とK20
188 名前：デフォルトの名無しさん [2013/01/28(月) 04:32:35.73 ]: どうせ数年待てば下位にも降りてくるんだから
待ちなされ
189 名前：デフォルトの名無しさん mailto:sage [2013/01/28(月) 13:15:40.00 ]: そして降りてくる時にはさらに凄い何かに目移りする訳ですね
190 名前：デフォルトの名無しさん mailto:sage [2013/01/28(月) 19:07:10.24 ]: 質問です。

CUDA4.0に対応したCUDA C Best Practices Guideを探しているのですが、どこかに公開いないでしょうか
191 名前：デフォルトの名無しさん mailto:sage [2013/01/29(火) 00:35:26.56 ]: >>190
ググればでてくるだろうに。
ttp://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html
192 名前：190 mailto:sage [2013/02/01(金) 11:59:03.69 ]: 所用で確認できず、返信遅れました。すいません

>>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対応のものがほしいのですが･･･
193 名前：デフォルトの名無しさん mailto:sage [2013/02/02(土) 10:55:34.50 ]: DPでのQソートのソース見ちゃうと
無しでやるのがアホらしく感じてくるな
194 名前：デフォルトの名無しさん [2013/02/07(木) 00:52:32.17 ]: CC3.5がコンシューマーに降りてくるかも♪

GK110を搭載するGeForce GTX 780 Titan　2013年2月末？
northwood.blog60.fc2.com/blog-entry-6531.html
195 名前：デフォルトの名無しさん [2013/02/11(月) 00:01:58.59 ]: winrarをcudaから展開するdllってなんか無いかな？
pw付きでもこちらからの操作で出来るやつ
196 名前：デフォルトの名無しさん mailto:sage [2013/02/11(月) 02:10:48.39 ]: >>195
GPU使ってパス解読できるソフトなら売られてるが、
CUDA使ってRAR解凍するソフトってあったっけ……
197 名前：デフォルトの名無しさん mailto:sage [2013/02/11(月) 07:57:27.35 ]: それってそんなに意味あるの？
198 名前：デフォルトの名無しさん mailto:sage [2013/02/11(月) 12:16:16.19 ]: >>196
パスワード解析なんかはGPUに向いてないんじゃないか？
199 名前：デフォルトの名無しさん mailto:sage [2013/02/11(月) 16:18:44.80 ]: >>198
zip用→www.internal.co.jp/products/util/passgetter/about/
RAR用→www.golubev.com/rargpu.htm
それなりに速度は出るっぽい
200 名前：デフォルトの名無しさん mailto:sage [2013/02/12(火) 11:46:35.13 ]: >>198
ブルートフォースはGPUの独壇場じゃない？
FPGAの専用プロセッサを除いて。
201 名前：デフォルトの名無しさん mailto:sage [2013/02/12(火) 12:50:32.08 ]: ブルートフォース
プルートフォース
202 名前：デフォルトの名無しさん mailto:sage [2013/02/12(火) 22:26:05.34 ]: >>200
>>198は、GPUは整数演算がそれ程得意では無い事を言いたいのでは？
203 名前：デフォルトの名無しさん mailto:sage [2013/02/12(火) 22:26:14.43 ]: VirtexクラスのFPGAでチューニングして200MHzくらいで動かしたらハイエンドGPUより速い？
204 名前：デフォルトの名無しさん mailto:sage [2013/02/13(水) 00:38:36.16 ]: >>202
http.developer.nvidia.com/GPUGems3/gpugems3_ch36.html

ごめん、くっそ遅かった

>>203
20KLUTsで20GBpsでる
205 名前：デフォルトの名無しさん mailto:sage [2013/02/14(木) 00:30:55.49 ]: GeForce 8800の整数は遅いけど、Fermiはそれほど遅くない
206 名前：デフォルトの名無しさん mailto:sage [2013/02/14(木) 02:55:25.27 ]: GeForceGTX Titanでは倍精度演算性能が削られている。単精度演算性能は4.70TFlops、倍精度演算性能はこれまでの“Kepler”と同様に単精度の1/24となり、196GFlopsにとどまる。
Pixel Fillrateは49GPixel/s、Texture Fillrateは196GTexel/sである。メモリ帯域は384-bitインターフェースもあいまって288.4GB/sに達する。

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

GTX580 197.6GFLOPS
GTX590 311GFLOPS
あれ？
207 名前：デフォルトの名無しさん [2013/02/14(木) 03:13:03.14 ]: Fermiは大サービスだったんだよ。
もう倍精度欲しいならTesla買うしか無い。
K20でも30万くらいだろ？

自分は単精度しか使わないからTitan買うけど。
208 名前：デフォルトの名無しさん mailto:sage [2013/02/14(木) 03:30:00.81 ]: 本当Fermiは大サービスだったよなぁ。整数やビット演算もまともだったし。
CUDA的にkepler refreshは大きく期待できることもないのでmaxwellが頼りなのだが
209 名前：デフォルトの名無しさん [2013/02/14(木) 12:17:13.77 ]: >CUDA的にkepler refreshは大きく期待できることもない
いやいやDPがあるだろ
210 名前：デフォルトの名無しさん mailto:sage [2013/02/14(木) 17:07:23.36 ]: Xeon Phiの発売で貧乏人以外には完全にオワコンでしょ
211 名前：デフォルトの名無しさん mailto:sage [2013/02/14(木) 18:52:59.63 ]: GPGPUはAMDですよ
212 名前：デフォルトの名無しさん mailto:sage [2013/02/14(木) 22:50:06.45 ]: >>209
Dynamic Parallelism（綴り覚えた）ってそんなに大きいの？
なんていうか使えるハードウェア資源が増えるわけじゃないし…って思って。

一方でXeon Phiがそれほど脅威って感じもしないんだよなぁ。
1. プログラムが組みやすい
2. 実効性能を引き出しやすい
あれを使ったプログラムの組み方を知らないので当てずっぽうなんだけど、
1.は大差ないんじゃないかなって。並列処理の組み方がそんなに素晴らしく変わることは想像できない。
2.でCUDAにコストや電力で勝利なんてことにでもならない限りは棲み分けるかと。
まあ貧乏なんでそうあってほしいっていう願望も入ってるんだけど。
213 名前：デフォルトの名無しさん mailto:sage [2013/02/15(金) 00:32:37.20 ]: インテル? Xeon Phi　5110PとNVIDIA Tesla K20の行列積における実効性能比較
www.hpc.co.jp/benchmark20130201.html

単精度計算ならK20、倍精度計算ならXeon Phi 5110P　という結論らしい
214 名前：デフォルトの名無しさん mailto:sage [2013/02/15(金) 01:40:16.84 ]: >>212
OpenMPだからPC用のプログラムがそのままで動くよ
215 名前：デフォルトの名無しさん mailto:sage [2013/02/15(金) 08:15:58.27 ]: devgurus.amd.com/thread/159457

7970はDGEMM 665GFLOPSらしい。
216 名前：デフォルトの名無しさん [2013/02/15(金) 08:55:15.74 ]: >>214
動くけど、専用ベクトル命令使わないと速度でないよね？
217 名前：デフォルトの名無しさん mailto:sage [2013/02/15(金) 18:51:43.68 ]: >>213
このK20Mと5110Pっておいくら万円するんだろう。
218 名前：デフォルトの名無しさん mailto:sage [2013/02/15(金) 22:17:58.89 ]: プログラムが簡単に書けることと、速度が出し切れることは大抵両立できない
219 名前：デフォルトの名無しさん mailto:sage [2013/02/16(土) 00:03:14.05 ]: >>216
インテルコンパイラを使えばいい。
それに、インテルにはMKLがあるから、
線形問題であれば、既存のコードにディレクティブを挿入するだけで速くなる。

>>217
どちらも30万くらいだろ。
220 名前：デフォルトの名無しさん mailto:sage [2013/02/17(日) 11:06:58.52 ]: Tesla買うならXeon Phiになるだろうが
GTXなんとかでもCUDAは使えるからな
221 名前：デフォルトの名無しさん [2013/02/17(日) 18:30:01.37 ]: ◇GeForce GTX Titan 演算性能
　 SP：4.5TFlops
　 DP：1.3TFlops

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

northwood.blog60.fc2.com/blog-entry-6565.html
222 名前：デフォルトの名無しさん mailto:sage [2013/02/18(月) 00:12:16.06 ]: >>221
そんなうまい話は無いと思うけどね
223 名前：デフォルトの名無しさん mailto:sage [2013/02/18(月) 20:17:49.22 ]: Titanたっけー。
誰か研究室とかで買って報告よろ。
224 名前：デフォルトの名無しさん mailto:sage [2013/02/19(火) 00:57:26.08 ]: 初日か初週でゲーマーが争奪戦して在庫切れで終わりだろ
悠長に評価して予算付けてって頃にはもう終わってる
225 名前：デフォルトの名無しさん [2013/02/19(火) 01:11:14.40 ]: 年度末で10万くらい予算余ってるから何か買っていいよ。

ってのが、今年は来ない。
226 名前：デフォルトの名無しさん mailto:sage [2013/02/19(火) 04:58:05.78 ]: >>225
10万で買えるの？
227 名前：デフォルトの名無しさん mailto:sage [2013/02/20(水) 00:31:37.18 ]: >>260
118000円で予約したからさすがに10万ジャストは無理
228 名前：デフォルトの名無しさん mailto:sage [2013/02/21(木) 19:46:59.76 ]: K20の約半額か
229 名前：デフォルトの名無しさん mailto:sage [2013/02/21(木) 19:59:41.22 ]: >>228
安いと言えば安いが、個人で遊ぶのはきつい価格だった。
230 名前：デフォルトの名無しさん mailto:sage [2013/02/21(木) 20:25:32.67 ]: >>229
でも580や480と違って倍精度もtesla相当に設定できるようだし、
ECCと演算保証がないことに目をつぶればお得といえばお得じゃない？
後々の中古売却を考えるなら値崩れはTeslaのほうが少ないだろうけど。
231 名前：デフォルトの名無しさん mailto:sage [2013/02/23(土) 01:24:32.93 ]: 「法人さんはTeslaを買ってくださいよ」（2/21）
www.gdm.or.jp/voices/2013/0222/21236

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

お前らか…
232 名前：デフォルトの名無しさん mailto:sage [2013/02/23(土) 07:48:27.16 ]: まー自分の金じゃないしとりあえず問い合わせぐらいはするでしょ
233 名前：デフォルトの名無しさん mailto:sage [2013/02/23(土) 22:46:08.24 ]: 実際の性能向上率をみたり、書き換えが必要な場合それの専攻開発用とかに買うんだろ。
普通本番機ではサーバーと一緒にTESLA導入するよ。
234 名前：デフォルトの名無しさん mailto:sage [2013/02/24(日) 01:05:09.57 ]: しかしTITANはDynamic Parallelismが無いのが不便だね
枝分かれが動的に決まる深さ優先探索を分散処理させたいのだけど
235 名前：デフォルトの名無しさん mailto:sage [2013/02/24(日) 01:08:24.14 ]: ちょｗ
Dynamic Parallelismないって致命的じゃんｗ
236 名前：デフォルトの名無しさん [2013/02/24(日) 01:15:58.14 ]: オワタ
しかし、GTX TitanではKepler GPUの新機能である
Dynamic ParallelismとHyper-Qという機能が省かれている。

カ○タムBIOSでTesla化出来ないかな？
237 名前：デフォルトの名無しさん mailto:sage [2013/02/24(日) 12:09:11.09 ]: Dynamic ParallelismつけたらTesla売れなくなるだろｗｗ
238 名前：デフォルトの名無しさん [2013/02/24(日) 14:45:08.55 ]: 最も格安なCUDAできるノートPCはどれ？
239 名前：デフォルトの名無しさん [2013/02/24(日) 14:45:50.04 ]: 倍精度付けてくれたから、Dynamic Parallelismもてっきり対応してるものかと。。
CUDAのページにはGTX TITAN CC 3.5って書いてるしね。

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

納得できん。
240 名前：デフォルトの名無しさん mailto:sage [2013/02/24(日) 14:50:45.05 ]: 最近はグラ用とコンピュート用で分化させる方針みたいだよ。
残念ながらね･･･。
241 名前：デフォルトの名無しさん [2013/02/25(月) 00:41:40.15 ]: Hyper-Q Dynamic palallerlismはついているようだが？
pc.watch.impress.co.jp/docs/news/20130219_588387.html
242 名前：デフォルトの名無しさん [2013/02/25(月) 01:03:16.82 ]: >>241
>なお、既報の通り、HPC向けにGK110では「Hyper-Q」および「Dynamic Parallelism」という機能も追加されている。

news.mynavi.jp/articles/2013/02/20/geforce_gtx_titan/index.html
>しかし、GTX TitanではKepler GPUの新機能であるDynamic ParallelismとHyper-Qという機能が省かれている。
243 名前：デフォルトの名無しさん mailto:sage [2013/02/25(月) 23:22:58.37 ]: >>240
Fermiで懲りたんだろな
VGAにGPGPUで使う機能をてんこ盛りして爆熱にしてもしょうがないし
高く売れるGPGPU機を安く売っているVGAで代替できるじゃ金儲けできないし
244 名前：デフォルトの名無しさん mailto:sage [2013/02/26(火) 04:32:52.99 ]: TitanはXeon Phiが出てなかったら出なかったかもな
245 名前：デフォルトの名無しさん mailto:sage [2013/02/26(火) 09:55:19.15 ]: Fremiで動いていたコードでKeplerで動かすとカーネルが起動しないらしくて、
計算できないんだが、Keplerだとなにか気をつけることってあるのかな
246 名前：デフォルトの名無しさん [2013/02/26(火) 11:17:27.91 ]: Fermiで動いてたのはたまたまであって、
注意深く解析すると、そのコードはスレッドの起動順序などによって
アクセス違反を起こす、とかね。
247 名前：デフォルトの名無しさん [2013/02/26(火) 11:32:45.34 ]: もしかしてスレタイって下らないのクダとCUDAを掛けてるの？
248 名前：デフォルトの名無しさん mailto:sage [2013/02/26(火) 11:49:25.57 ]: Xeon Phiって結局まだ個人では単品で入手できないね
249 名前：デフォルトの名無しさん [2013/02/26(火) 13:09:38.68 ]: >>247
すげーーーーーーーーーーーーーーーー
良く気付いたな
250 名前：デフォルトの名無しさん mailto:sage [2013/02/26(火) 21:09:38.74 ]: CUDAはCUDAらない。
251 名前：デフォルトの名無しさん mailto:sage [2013/02/27(水) 01:58:36.28 ]: 今夜は良く冷えますね
252 名前：デフォルトの名無しさん mailto:sage [2013/03/01(金) 08:07:53.55 ]: おいおい、TitanでDP使えるらしいじゃないか！
訂正記事書かれてるぞ
headlines.yahoo.co.jp/hl?a=20130220-00000025-mycomj-sci
2/28追記:記事掲載当初、「GTX TitanではKepler GPUの新機能であるDynamic ParallelismとHyper-Qという機能が省かれている」と記載させていただいておりまし
たが、GTX Titanでもこれらの機能をサポートしていることが明らかとなりましたので、当該部分ならびにそこに付随する文章
を修正させていただきました。
253 名前：デフォルトの名無しさん [2013/03/01(金) 09:42:16.44 ]: >>252
キター
DPって、倍精度と紛らわしいよね
254 名前：デフォルトの名無しさん mailto:sage [2013/03/01(金) 11:19:34.68 ]: マジかよｗ
この訂正記事に気づいてない人の分だけ売上落ちるぞこれ
255 名前：デフォルトの名無しさん mailto:sage [2013/03/01(金) 12:55:23.44 ]: このスレ見てて良かったと思う瞬間
↓
256 名前：デフォルトの名無しさん mailto:sage [2013/03/01(金) 16:59:12.85 ]: │　 ↑
└─┘
おらっしゃあぁぁ！！！
　∩∧ ∧
　ヽ（ﾟДﾟ）
　　　＼⊂＼
　　　Ｏ-､　）～
　　　　　　∪
257 名前：デフォルトの名無しさん mailto:sage [2013/03/01(金) 17:17:18.74 ]: 計算ミスしないなら欲しい
258 名前：デフォルトの名無しさん mailto:sage [2013/03/01(金) 18:55:26.41 ]: 計算エラーって、確率的に起こるもんなの？
それとも同じところで再現？
259 名前：デフォルトの名無しさん mailto:sage [2013/03/01(金) 19:58:01.79 ]: ハードウェア的には同じところだとは思うが、プログラム上の同じところで
再現してくれるとは限らない。

計算ミスするコアは必ずミスしてくれるかっていうと…詳しい人教えてください
260 名前：デフォルトの名無しさん mailto:sage [2013/03/01(金) 20:12:57.40 ]: メモリエラーの場合は確率的じゃない？
261 名前：デフォルトの名無しさん mailto:sage [2013/03/01(金) 20:40:09.11 ]: Dynamic Parallelismあり＋Geforce並のエラー精度
Dynamic Parallelismなし＋Tesla並のエラー精度

処理内容にもよるけどこの2択なら前者のほうが欲しい。
なんていうかGPGPUで何かさせるときは計算ミスは織り込んだ上で作る感じだし。
そういう意味でもTitanはおもしろいと思うけど、なぜにDynParal不可という報道を放置したし・・
262 名前：デフォルトの名無しさん mailto:sage [2013/03/01(金) 21:02:25.16 ]: その記者がIntelから金渡されてたりして
263 名前：デフォルトの名無しさん mailto:sage [2013/03/01(金) 23:24:13.26 ]: 品薄もんだからよかったものの、
大量に捌きたいもんだったらNVidiaから訴えられてもおかしくないレベルｗ
264 名前：デフォルトの名無しさん mailto:sage [2013/03/01(金) 23:26:48.44 ]: nVIDIA Japan公式Twitterだと使えるってアナウンスが
随分前にあったけど、使えるで確定なのねｗ
265 名前：デフォルトの名無しさん mailto:sage [2013/03/01(金) 23:31:11.08 ]: くそ、DPない情報に踊らされて買いそびれたじゃねぇか･･･
まだ店で売ってるのかな？
266 名前：デフォルトの名無しさん [2013/03/01(金) 23:42:31.95 ]: どさくさに紛れてGPUDirectも対応してないかな？
267 名前：デフォルトの名無しさん mailto:sage [2013/03/02(土) 09:08:59.28 ]: >>261
ECC無してTesla C1070、C2070、C2075、GTX580を使ってきた経験上、
計算エラーが起きたことは無い。
TSUBAMEのように大量に組めばECCも重要だろうけど、
デスクトップレベルなら気にすることは無い。
エラーがでたらやり直せばいい。
それが嫌ならK20を買えという事だ。
268 名前：デフォルトの名無しさん mailto:sage [2013/03/02(土) 09:43:58.97 ]: >>267
＞エラーがでたらやり直せばいい。

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

ECCなんかにダイやサイクルを浪費するより、
よっぽど効率的。
269 名前：デフォルトの名無しさん mailto:sage [2013/03/02(土) 11:20:49.48 ]: そもそもエラーだと気付ければいいけどな
いつの間にか数値の一部が化けてたら気付かないぞ
270 名前：268 mailto:sage [2013/03/02(土) 11:37:11.88 ]: >>269
あ、そういえばｗ

まぁ、エラー「訂正」までせずに、「検知」に留めておいて
ロジックを減らすのはいいんじゃないかなぁ。
271 名前：デフォルトの名無しさん mailto:sage [2013/03/02(土) 11:53:59.26 ]: シミュレーションに使っているけど、計算ミスが問題になったことはないなぁ。
あ、そもそも精度が要るところはCPUで計算してたわ。
272 名前：デフォルトの名無しさん mailto:sage [2013/03/02(土) 16:30:38.24 ]: >>269
大事な計算が必要なら信頼性の高い
システムを使えばいいだけ。
指数部が化けたら気づきやすいだろうし、
仮数部が化けても気づかないならそもそも
それほど感度がいる解析ではないので重要ではないよ。
そもそも研究なりで回す場合はいくつも計算するだろうから、
明らかにおかしい値は人間のフィルタがかかるでしょ。
273 名前：デフォルトの名無しさん [2013/03/02(土) 17:05:17.55 ]: そんなに心配なら、同じ計算を別のボードでもやって比較すれば良い。
同じボードでやっても、あまり意味は無いぞ？
274 名前：デフォルトの名無しさん mailto:sage [2013/03/03(日) 22:00:13.88 ]: 計算分野は狭いようで果てしなく広い分野だからな
エラー検知がどの程度重要かどうかは分野ごとに異なるから
自分の身の回りだけの話をしてもしょうがないだろ
275 名前：デフォルトの名無しさん mailto:sage [2013/03/03(日) 23:45:41.01 ]: これでTeslaと同じくらいの価格だな
akiba-pc.watch.impress.co.jp/img/ah/docs/590/125/html/s3201.jpg.html
276 名前：デフォルトの名無しさん mailto:sage [2013/03/06(水) 17:18:26.89 ]: CUDAとOpenCV併用するとき、.cuファイル内でOpenCV使うことってできないの？
コンパイル時にインクルードファイル周りでエラー吐いたからとりあえずOpenCV周りだけ.cppファイルに隔離したんだけど。
後々不便そうだから対策があれば教えてほしい。
277 名前：デフォルトの名無しさん mailto:sage [2013/03/06(水) 23:06:53.58 ]: >>276
CUDAのカーネル関係はlib化して、
ビルド時にリンクする方法を取ればいいんじゃね？
NVCCってそんなに賢くないから、CPU周りで遅くなるし。
278 名前：デフォルトの名無しさん mailto:sage [2013/03/06(水) 23:12:33.76 ]: ふと、CUDAで自分で書いたカーネルを実行しつつ、OpenCV の cv::gpu:: の機能も同時に使うと、どうなっちゃうんだろう、と思った。
279 名前：デフォルトの名無しさん mailto:sage [2013/03/07(木) 16:52:32.98 ]: >>278
どちらも実行されるよ。
普通のマルチスレッドのプログラミングと
考え方は変わらん。
280 名前：デフォルトの名無しさん [2013/03/19(火) 17:09:17.32 ]: news.mynavi.jp/series/nvidia_kepler_gpu/001/
【連載】
GPUの実効性能を上げるDynamic ParallelismとHyper-Q
281 名前：デフォルトの名無しさん mailto:sage [2013/03/19(火) 19:27:04.30 ]: >>280
安藤先生の記事や。
ありがたや～。
282 名前：デフォルトの名無しさん mailto:sage [2013/03/20(水) 00:32:25.98 ]: www.eevblog.com/forum/projects/hacking-nvidia-cards-into-their-professional-counterparts/

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

半田班ちょっと人柱頼む
283 名前：デフォルトの名無しさん mailto:sage [2013/03/20(水) 07:16:49.30 ]: >>282
どうやって発見したんだよ！ｗ
284 名前：デフォルトの名無しさん mailto:sage [2013/03/20(水) 08:04:17.97 ]: 英語読める割に綴り間違うんだな
285 名前：デフォルトの名無しさん mailto:sage [2013/03/20(水) 09:36:22.71 ]: >>282
レジスタってGPUのレジスタじゃないぞ。
抵抗のレジスタだ。

しかし、興味深いなあ。
286 名前：デフォルトの名無しさん mailto:sage [2013/03/20(水) 09:59:06.90 ]: >>285
半田班ってんだからわかってくれよ
>>284
携帯だったんだよ興奮してたんだよ
287 名前：デフォルトの名無しさん mailto:sage [2013/03/20(水) 10:07:36.15 ]: チップ自体は同じだから理論上可能というのは分かるんだけど、
よもやこんな方法でイネーブルにできるとはｗ
288 名前：デフォルトの名無しさん mailto:sage [2013/03/20(水) 10:07:47.70 ]: >>283
俺が知りたいよ

早めに買っとかないと対策されそうな気がするけどこのチップ抵抗は怖いよな。。
289 名前：デフォルトの名無しさん mailto:sage [2013/03/20(水) 10:08:50.58 ]: >>287
しかも690のがクロック高いから性能いいらしいぞ(;´Д`)
買ってみてやってみて
290 名前：デフォルトの名無しさん mailto:sage [2013/03/20(水) 10:13:47.22 ]: >>289
マジかよｗ

あー、久々おもしろい話題でわくわくするｗ
やっぱハードウェアはこうでないとなｗ
291 名前：デフォルトの名無しさん [2013/03/20(水) 10:35:15.50 ]: マジレスすると、quadroやteslaでの構成と同等なものがないから、
ドライバ側が柔軟な作りに立っていないと不具合が出るだろうな。
例えばハードウェアのモニタリング用の石が載ってなかったりするだろうから。

nvidia的には、geforceがteslaになったところで、購買層が違うから放置だろうが、quadro化はマーケティング的にやばいので塞ぎにくるだろうな。
292 名前：デフォルトの名無しさん mailto:sage [2013/03/20(水) 12:36:37.99 ]: >>291
ただそれはドライバを更新しなければいいだけの話だよね。
293 名前：デフォルトの名無しさん mailto:sage [2013/03/20(水) 13:02:33.89 ]: >>292
しばらくは放置なんじゃなかな。
さすがにドライバだけでは対応しづらい気がする。
基板のリビジョンが変更になった時から使えなくなると思う。
294 名前：デフォルトの名無しさん [2013/03/20(水) 15:20:27.90 ]: K10化より、TITANをK20Xにしたい。
295 名前：デフォルトの名無しさん mailto:sage [2013/03/20(水) 16:16:11.67 ]: おっ
また発見されたのかよ
むかしクワドロ化の方法が発見されたのに、まだまともな対策してなかったのか
296 名前：デフォルトの名無しさん mailto:sage [2013/03/20(水) 18:48:30.94 ]: これ、カーネルドライバがボードに直接プロダクトIDを聞きにいくからカーネルフックじゃ駄目って書いてあるけど
ドライバ改造できればﾜﾝﾁｬﾝ有るってことじゃね。
297 名前：デフォルトの名無しさん [2013/03/20(水) 20:54:10.54 ]: GeForceをQuadro/Tesla化出来るからって、Quadro/Teslaの売り上げが
減るとは思えない。
298 名前：デフォルトの名無しさん mailto:sage [2013/03/20(水) 22:40:19.86 ]: tesla化するって言っても、gk110じゃないしな。gpgpuやってる人からすれば、
そんなにありがたみがないんじゃね
299 名前：デフォルトの名無しさん [2013/03/24(日) 16:59:47.30 ]: NVIDIA CUDA Gets Python Support | techPowerUp
www.techpowerup.com/181585/NVIDIA-CUDA-Gets-Python-Support.html
300 名前：デフォルトの名無しさん mailto:sage [2013/03/24(日) 19:08:27.70 ]: これネイティブ対応ってことなのかね。
パイソンってそんなに人気あったんだ・・・
ruby人気は国内限定？
301 名前：デフォルトの名無しさん mailto:sage [2013/03/24(日) 19:23:07.90 ]: うん
302 名前：デフォルトの名無しさん mailto:sage [2013/03/24(日) 19:23:52.69 ]: Visual StudioもPython使えるようになるパッチが出てるし
OpenCVもPython対応(一応)してるしLibreOfficeもマクロでPython使えるし
そしてCUDAまでも…ということで俺の身の周りではPython株が右肩上がり。

ネット上のサンプルが2.x用のソースだと動かなかったりして躓いたこともあったり
今でも躓いたりしてるけど、これから3.x用の情報がもっと増えてくれると嬉しいと思う。
303 名前：デフォルトの名無しさん mailto:sage [2013/03/27(水) 13:42:12.44 ]: SL#（えすえるしゃーぷ）とは、GPUで実行されるプログラマブルシェーダーを、超高級言語である
C#で書けてしまうという夢のようなオープンソースのフレームワークである。
monobook.org/wiki/SL_Sharp
304 名前：デフォルトの名無しさん mailto:sage [2013/03/27(水) 19:32:08.19 ]: C#ｲﾗﾈ
305 名前：デフォルトの名無しさん [2013/03/30(土) 17:24:17.17 ]: DXエラーだったのでdirectXかと思いここで聞きます
色々ググったのですが解決に至らず

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

某ゲームのランチャー起動後にこのエラーが出ます
誰かご教授下さい
306 名前：デフォルトの名無しさん mailto:sage [2013/03/30(土) 23:26:35.66 ]: スレチだし
ゲームの事は開発元に聞け
あとDX8のレンダリングにすら対応してないクソPCを捨てろ
307 名前：デフォルトの名無しさん mailto:sage [2013/04/02(火) 08:57:50.31 ]: なんでDirectXでこのスレなんだ？
308 名前：デフォルトの名無しさん mailto:sage [2013/04/05(金) 14:24:31.80 ]: 本出るの久しぶりじゃない？
www.amazon.co.jp/dp/4061538209
309 名前：デフォルトの名無しさん mailto:sage [2013/04/05(金) 20:17:57.08 ]: >>308
おお～、CUDA5本か！
注目だな～。

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

next.rikunabi.com/tech/docs/ct_s03600.jsp?p=001216
つか、この人、すごい人なんだな！
原作やってたときは東大生のときだったのか･･･。
310 名前：デフォルトの名無しさん [2013/04/05(金) 22:20:39.03 ]: 宣伝乙
311 名前：デフォルトの名無しさん [2013/04/06(土) 08:59:19.78 ]: もしかして、東大の初代grapeに関わった人？
312 名前：デフォルトの名無しさん mailto:sage [2013/04/06(土) 09:40:17.13 ]: そだよ！
313 名前：デフォルトの名無しさん [2013/04/12(金) 02:37:27.22 ]: もしかして、TitanってP2P使えない？
cudaDeviceCanAccessPeerで0が返ってくるのは想定の範囲内だったけど、
cudaDeviceEnablePeerAccessするとエラーになってしまう。

OS:CentOS 6.4 64bit
Driver:313.30
MB:P9X79
Titan 2枚をそれぞれPCIex gen3でx16接続
314 名前：デフォルトの名無しさん [2013/04/12(金) 17:39:21.97 ]: >>157
書籍「CUDA BY EXAMPLE」の「5.3.2 ドット積の(誤った)最適化」の項には、
真逆な事が書いてあって悩んでいる。

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

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

書籍の内容は正しいのだろうか？
知っている人教えてください。
315 名前：デフォルトの名無しさん mailto:sage [2013/04/14(日) 00:10:01.63 ]: >>314
多分「CUDA BY EXAMPLE」が間違っているか、CUDAのバージョンが新しくなったことによって変わったのかもしれないが、基本的に__syncthreads()はスレッド間のフェンスみたいなものだから。SIMDの概念で考えればわかると思う。
316 名前：デフォルトの名無しさん mailto:sage [2013/04/14(日) 02:24:20.82 ]: 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.
317 名前：デフォルトの名無しさん mailto:sage [2013/04/14(日) 02:44:35.83 ]: 要は、_syncthreadsはwarp内の一つでも実行されたら
そのwarpの全てのthreadが実行したのと同じと看做される。
だから同一block内の全てのwarpそれぞれで、
いずれかのthreadがアクティブな状態で_syncthreadを実行すれば辻褄はあう。

逆にMPIのbarrierで辻褄を合わせるように、全threadでbarrierを実行しようと
分岐の両側に_syncthreadを置いて、実際にwarp内で両パスにthreadが別れてしまうと
困ることになるだろう。
318 名前：314 [2013/04/14(日) 05:11:45.30 ]: >>315 >>316 >>317
PTXレベルでの解説ありがとうございます。
なるほど、「CUDA BY EXAMPLE」を読んだだけの頃は
「MPIのbarrier」のような仕様をイメージしていましたが、
そうではないんですね。
「__syncthreads()」の分岐に入らなかったスレッド用に、
別の「__syncthreads()」を書く必要は無いということか。。。
SIMDの概念ならその仕様でないとダメなんですね。
319 名前：デフォルトの名無しさん mailto:sage [2013/04/23(火) 12:25:54.57 ]: ものすごく初心者丸出しな質問で申し訳無いのですが、質問させてください。
CUDA5.0をVS2008で使うにはどういった手順が必要になるのでしょうか？
ネット上だと2010や2012についてばかり見つかって2008について見つけられませんでした。
どうかよろしくお願いします。
320 名前：デフォルトの名無しさん [2013/04/24(水) 11:40:18.49 ]: マルチgpu環境での使い方が分からず困っています。
何か良い資料は有りませんでしょうか。
321 名前：デフォルトの名無しさん mailto:sage [2013/04/25(木) 22:41:10.57 ]: ハードウェアとCUDAのバージョン依存があるんでね？
322 名前：デフォルトの名無しさん mailto:sage [2013/04/26(金) 06:03:02.19 ]: >>320
つ www.meriken2ch.com/programming/merikens-tripcode-finder
マルチGPUをサポートしてて、ソースも付いてるよ。
323 名前：デフォルトの名無しさん [2013/04/26(金) 08:14:11.43 ]: compute capability 3.0 cuda 5.0　です。
324 名前：デフォルトの名無しさん mailto:sage [2013/04/26(金) 09:28:57.17 ]: >>320
1GPU毎に1CPUスレッドを割り当てて、それぞれのスレッドで
cudaSetDevice()を実行すればいいって話じゃないの?
325 名前：デフォルトの名無しさん mailto:sage [2013/04/26(金) 14:32:31.09 ]: >>320
デバイスドライバとの関係でメインスレッドでcudaSetDevice()すると巧くないので、
>324のようにcreate_pthread()した先でcudaSetDevice()するよろし。
326 名前：デフォルトの名無しさん [2013/04/26(金) 16:55:32.79 ]: cudaSetDevice()を呼び出した後はシングルgpuの場合と同様なんですよね。

gpuが二つある場合に単一スレッドで cudaSetDevice(0)を呼び出して作業した後に
cudaSetDevice(1)を呼び出した場合はどうなるのでしょうか。
327 名前：デフォルトの名無しさん mailto:sage [2013/04/26(金) 18:34:49.73 ]: 二度目以降は無視されるよ。
328 名前：デフォルトの名無しさん [2013/04/26(金) 21:44:04.64 ]: 各スレッドでcudaSetDevice()を呼び出して各スレッド毎に各gpuを割り当てますが、
複数のスレッドでcudaSetDevice()を呼び出さずにカーネルを実行したらどうなるのでしょうか。
また、複数のスレッドでcudaSetDevice()使って同一のデバイスを指定はできるのでしょうか。
329 名前：デフォルトの名無しさん mailto:sage [2013/04/26(金) 22:55:17.00 ]: >>328
何でそんなことしたいのかさっぱりわからん。
330 名前：デフォルトの名無しさん mailto:sage [2013/04/26(金) 22:59:08.02 ]: 自分で試して確認しようとしていないようだけど、そういう環境がないのか？
何の為に学ぼうとしてるのかしらんが、そうやって無限に問い続けるつもりか？
331 名前：デフォルトの名無しさん [2013/04/27(土) 03:09:44.84 ]: 環境がなくて試せないんです。すみません。
332 名前：デフォルトの名無しさん [2013/04/27(土) 04:12:10.93 ]: 数千円でボードかえるのに環境がないとか甘えだろ。
333 名前：デフォルトの名無しさん mailto:sage [2013/04/27(土) 07:55:25.89 ]: >>326
CUDA C Programming Guide 5.0
3.2.6.2 Device Selection
334 名前：デフォルトの名無しさん mailto:sage [2013/04/28(日) 17:13:33.38 ]: 環境がなくて試せないのに、なんでそんなに重箱の隅みたいなこと聞くのか
335 名前：デフォルトの名無しさん mailto:sage [2013/04/28(日) 17:49:38.30 ]: 環境の構築(購入)に踏み切ろうか悩んでいるとか
336 名前：デフォルトの名無しさん mailto:sage [2013/04/29(月) 18:00:10.69 ]: マルチGPU環境とかこだわらなければ1万円以下でできるのにな。
337 名前：デフォルトの名無しさん mailto:sage [2013/05/05(日) 13:11:10.63 ]: 現在CUDAで画像処理を行っているのですが，GPUを長時間使用するような
プログラムを走らせると，「ディスプレイドライバの応答が停止しました」または
PCが落ちる，という事が頻発します．
現在は，ソースコード上で使用するブロック数とスレッド数の数を減らすことで工夫して対応していますが，実装上不便です．
上記に対して，ソースコード上で使用するスレッド数とブロック数を制限せずに，解決する方法は何かございますか？
環境はWindows7,CUDA5.0，VC2008，GT610です．
宜しくお願いします．
338 名前：デフォルトの名無しさん mailto:sage [2013/05/05(日) 13:37:43.82 ]: タイムアウトの設定をレジストリで変えた？
たしかwin7だとデフォで1秒か2秒くらいになってて、それ以上カーネルが動きっぱなしに
なると「ディスプレイドライバの応答が～」になるので変更する必要がある。
339 名前：デフォルトの名無しさん mailto:sage [2013/05/05(日) 13:49:16.71 ]: 可能ならディスプレイ表示用（ショボいので十分）とコンピューティング用GPUと2つ接続したいところだね。
340 名前：デフォルトの名無しさん [2013/05/05(日) 14:00:03.24 ]: 最近はCPUオンチップGPU+CUDA用GPUという構成を組みやすい。
341 名前：デフォルトの名無しさん mailto:sage [2013/05/05(日) 16:00:33.60 ]: >>338-340
ありがとうございます．
タイムアウトの設定は最後の手段としたいと思います．デフォルトでは２秒だそうです．
表示用と処理用のGPU二つ用意するのはいいですね．
残念ながら，このPCのBIOSではオンボードとGPUを共存できないうえ，
GPU二枚刺しも出来ないので，新しく機材を調達しようと思います．
342 名前：デフォルトの名無しさん mailto:sage [2013/05/05(日) 16:15:31.63 ]: タイムアウトが最初の手段だと思うけど。。
それで不便だったらお金のかかる対策に移ればいいわけだし。
343 名前：デフォルトの名無しさん mailto:sage [2013/05/05(日) 16:26:07.35 ]: GT610なんて使っている時点でCUDAの意味がないがな。
スレッドブロック数のチューニングにすらならん。
344 名前：デフォルトの名無しさん mailto:sage [2013/05/05(日) 23:29:26.34 ]: >>340
今はiGPU搭載CPUが普通で、それがCUDAお手軽構成って感じだよね

>>341
お薦めはAMDのAPUだよ。
値段が安いのにCPU部はIntelの高性能のi7並みの性能で、GPUはGT610より性能が良い
345 名前：デフォルトの名無しさん mailto:sage [2013/05/06(月) 20:16:35.75 ]: linuxでやれよハゲ
346 名前：デフォルトの名無しさん mailto:sage [2013/05/07(火) 07:37:55.74 ]: >>344
iGPUでCUDAができるかよ。やるならOpenCLだ。
AMDのCPUがi7並であるわけ無いだろ。
スレ違い。
347 名前：デフォルトの名無しさん mailto:sage [2013/05/07(火) 20:27:42.82 ]: GPGPUはHSAでAMDが主流になるって感じになってきたよな
一部のスパーコンピューティングでCUDAがつかわれ、それ以外はHSAのAMDになるのかな
348 名前：デフォルトの名無しさん mailto:sage [2013/05/08(水) 22:17:31.73 ]: >>347
性能が桁違いだから仕方がない。自作のアプリではTITANは
7970の6割ぐらいの性能だし。CUDAとOpenCLって
やってることはほとんど同じだから、移行するのもそんなに
難しくなかった。
349 名前：デフォルトの名無しさん mailto:sage [2013/05/08(水) 22:54:07.61 ]: あとSouthern Islands系のAMDのGPUは、CUDAに比べると分岐処理のペナルティが
はるかに小さいのでプログラミングしやすいというのもある。
(EvergreenやNorthern Islandsは共有メモリが遅かったり
ベクトル型を使わないと性能が出なかったりするのでこの限りでない)
350 名前：デフォルトの名無しさん [2013/05/09(木) 07:45:11.56 ]: ()
351 名前：デフォルトの名無しさん mailto:sage [2013/05/09(木) 21:33:31.18 ]: >>349
お、GCNはだいぶいいんだね。
352 名前：デフォルトの名無しさん mailto:sage [2013/05/16(木) 00:41:04.04 ]: teslaでも使える温度計測ソフトはないものか
353 名前：デフォルトの名無しさん [2013/05/19(日) 14:04:36.33 ]: >>352
何の温度計測なんだ？
ボードの温度なのか、何かの測定器に連動するものなのかで答えが違うと思うが。
354 名前：デフォルトの名無しさん mailto:sage [2013/05/26(日) 13:54:06.19 ]: チクショウ、GTX780は倍精度切られているらしい。
Titanに逝ってくる。
355 名前：デフォルトの名無しさん mailto:sage [2013/05/30(木) 23:33:46.73 ]: またどっかの抵抗取り替えたりするとTITAN化するような仕掛けが潜んでたりして
356 名前：デフォルトの名無しさん mailto:sage [2013/05/31(金) 00:42:40.83 ]: TITANをK20xにしたいんだけど
357 名前：デフォルトの名無しさん mailto:sage [2013/06/01(土) 17:17:27.06 ]: TITANにECCつけたもの（SMXを一つくらい品質保証のため削るかもしれないが）をK20の後継として出るんじゃないかな？
今年のSCあたりで発表されそうだ
358 名前：デフォルトの名無しさん mailto:sage [2013/06/01(土) 18:12:25.50 ]: ゲフォがFermiが倍精度1/4でKepler1/8で結局性能悪化って酷くない?
ラデへの乗り換えを検討したほうがいいかな?
359 名前：デフォルトの名無しさん mailto:sage [2013/06/01(土) 18:38:27.12 ]: CUDA使ったフリーウェアを作ってるのに、
げふぉに将来性がなくなったら詰む。
360 名前：デフォルトの名無しさん mailto:sage [2013/06/01(土) 19:05:28.96 ]: >>359
CUDA で計算している部分だけ OpenCL に変えれば済む話では？
どうせモジュール化してるんでしょ
361 名前：359 mailto:sage [2013/06/01(土) 19:33:33.76 ]: >>360
そうかも。
OpenCLってデバイスごとにコンパイルが必要なんでしたっけ?
実行時コンパイルとかよくわからんです。
理想的にはバイナリを一個作っとけばどこでも使えるといいのですけれど。
入門書読めばその辺の話は載ってますか?

スレ違いスマソ。
362 名前：デフォルトの名無しさん mailto:sage [2013/06/01(土) 21:06:59.29 ]: >>361
その辺の話は全て入門書、あるいは入門ページに載っている
363 名前：359 mailto:sage [2013/06/01(土) 22:07:59.04 ]: >>362
さんクス
364 名前：デフォルトの名無しさん mailto:sage [2013/06/02(日) 23:02:35.27 ]: >>360
カーネルだけなら簡単だけど、メモリ周りとか考えると簡単にはいかないよ。
365 名前：デフォルトの名無しさん mailto:sage [2013/06/03(月) 01:40:28.77 ]: >>364
CUDAで書く時はグローバルメモリへのコアレスアクセスとか
シェアードメモリ使ってチューニングしてるけど、
OpenCLはそのへん配慮できるのかな。
366 名前：デフォルトの名無しさん mailto:sage [2013/06/03(月) 04:03:55.61 ]: >>365
普通にできるよ。
367 名前：デフォルトの名無しさん mailto:sage [2013/06/03(月) 07:03:45.92 ]: >>364
それは「げふぉに将来性がなくなったら詰む」ほどの事か？
368 名前：デフォルトの名無しさん mailto:sage [2013/06/03(月) 20:38:07.29 ]: 結局、アーキテクチャ毎に最適なコードを書かなければならない。
そのように書かず、可搬性を意識しすぎるとGPU(などの並列プロセッサ)を使う旨みがない。
なんというジレンマ。
369 名前：デフォルトの名無しさん mailto:sage [2013/06/03(月) 21:37:12.36 ]: つまり、コンパイラがまだまだ馬鹿ということか
370 名前：デフォルトの名無しさん mailto:sage [2013/06/05(水) 06:41:20.54 ]: teslak20って一番安い所ではいくらで買えるんだろ？
ヤフオクで三年間分有効の保証が残ってる奴が30万なんだが安いのかな？
371 名前：デフォルトの名無しさん mailto:sage [2013/06/05(水) 08:49:33.35 ]: news.nvidia.com/t/213824/48003696/16930/10/
nvからkayla(armようcuda開発機)のおすすめメールが来た

tegra3とGPU別ｗ
372 名前：デフォルトの名無しさん [2013/06/07(金) 08:47:44.96 ]: >>370
titanでいいやん。
ヤフオクで買う程度なら個人利用なんだろ？
会えてteslaを買う理由がないわ。
373 名前：デフォルトの名無しさん mailto:sage [2013/06/07(金) 09:27:17.14 ]: マルチGPUするなら、GPUDirectの有無は大きいよ>TeslaとTitan
374 名前：デフォルトの名無しさん [2013/06/08(土) 09:03:33.73 ]: titanでもGPUDirectはついてるよ。ないのはRDMAの機能の方だ。
375 名前：デフォルトの名無しさん mailto:sage [2013/06/10(月) 13:22:19.65 ]: CUDA5.5はどうですか
376 名前：デフォルトの名無しさん mailto:sage [2013/06/10(月) 14:29:57.78 ]: いいえ、どうではありません。
377 名前：デフォルトの名無しさん mailto:sage [2013/06/11(火) 09:31:22.59 ]: MPIプログラムのプロファイルを取れるようになったのはでかいな。
https://developer.nvidia.com/cuda-toolkit
378 名前：デフォルトの名無しさん [2013/06/15(土) 23:05:03.06 ]: cuda-gdbのステップ実行で、
コンソールが返ってこなくなる事がありますが、
原因を確認する方法はありますか？
（普通のgdbも使った経験が無いのですが。。。）
379 名前：デフォルトの名無しさん mailto:sage [2013/06/17(月) 15:41:11.18 ]: 海外amazonでGK208なGT640かったら
やっぱSM35だった
380 名前：デフォルトの名無しさん mailto:sage [2013/06/18(火) 05:37:29.42 ]: TOP500でxeon-phiがトップなたぞ
381 名前：デフォルトの名無しさん mailto:sage [2013/06/18(火) 12:06:25.31 ]: 正直微妙だ

　　　　　　Efficiency (%) Mflops/Watt
Tianhe-2　 61.68　　　　　　1901.54
Titan　　　　64.88　　　　　　2142.77
382 名前：デフォルトの名無しさん mailto:sage [2013/06/18(火) 12:15:47.25 ]: 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
383 名前：デフォルトの名無しさん mailto:sage [2013/06/19(水) 09:00:43.77 ]: >Detected Compute SM 3.5 hardware with 2 multi-processors

SMX2器ってこと?
384 名前：デフォルトの名無しさん mailto:sage [2013/06/19(水) 13:12:02.16 ]: そう
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
385 名前：デフォルトの名無しさん mailto:sage [2013/06/19(水) 22:01:10.70 ]: Dynamic Parallelismも使えるの？
386 名前：デフォルトの名無しさん mailto:sage [2013/06/19(水) 22:38:31.81 ]: バス幅とL2キャッシュ典
387 名前：デフォルトの名無しさん mailto:sage [2013/06/20(木) 10:08:10.18 ]: >385
hyperQとDynamic Parallelismは使えるようだね>382
388 名前：デフォルトの名無しさん mailto:sage [2013/06/20(木) 11:17:52.11 ]: GK110でL2が1536KBだから512KBって結構でかいな
389 名前：デフォルトの名無しさん mailto:sage [2013/06/22(土) 09:56:35.61 ]: GK208きた
390 名前：デフォルトの名無しさん mailto:sage [2013/06/24(月) 14:06:27.76 ]: cudaGetDevicePropertiesで
maxThreadsPerBlockが1024とでたので、
kernel<<<32, 1024>>>()
とやったらKeplerはおKでFermiではだめだった。
基本的にはKeplerでも1ブロックあたり512スレッドが上限だったっけ?
391 名前：デフォルトの名無しさん [2013/06/30(日) 20:50:09.84 ]: 使えるようにするまで2時間かかった
392 名前：デフォルトの名無しさん mailto:sage [2013/06/30(日) 20:58:07.76 ]: CUDAは導入がメンドイね。
HSAはよ。
393 名前：デフォルトの名無しさん mailto:sage [2013/07/01(月) NY:AN:NY.AN ]: >>392
Linuxだとホントしんどいよね。
何が悲しゅうてドライバインストールのために
カーネルを再コンパイルせにゃならんのか。
394 名前：デフォルトの名無しさん mailto:sage [2013/07/01(月) NY:AN:NY.AN ]: Linuxでの導入はなおさら OpenCL の方が楽だ

プログラムはめんどいが
395 名前：やんやん ◆yanyan72E. mailto:sage [2013/07/01(月) NY:AN:NY.AN ]: Linuxのカーネルのコンパイルといってもモジュールを
コンパイルするだけなのだけれど、それってそんなに面倒?
リブートも必要ないし。
396 名前：デフォルトの名無しさん mailto:sage [2013/07/02(火) NY:AN:NY.AN ]: GUIを使わないモードでリブートして、
カーネルコンパイルして、
もう一回リブートしないとならないと思ったけど。
397 名前：やんやん ◆yanyan5.Xudd mailto:sage [2013/07/02(火) NY:AN:NY.AN ]: GUIのサービス止めてモジュールをrmmodして
nVidiaが提供するスクリプトを
使ってモジュールだけをコンパイルして
insmodしてからGUIのサービス再開するだけだよ。
398 名前：デフォルトの名無しさん mailto:sage [2013/07/02(火) NY:AN:NY.AN ]: Bumblebee はどうなの？

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

もしできるのなら Linux に乗り換えてみようかな。
他のことで Windows がちょっと使いづらくなってて、
でも踏ん切りつかなくて迷ってる。
399 名前：デフォルトの名無しさん mailto:sage [2013/07/03(水) NY:AN:NY.AN ]: >>395
適切なOSとGCCのバージョンをそろえるのが面倒。
ちょっとでもNVIDIA推奨環境と違うと
Getting Startedだけ読んででインストールするのは絶対無理
だと思う。
400 名前：デフォルトの名無しさん mailto:sage [2013/07/03(水) NY:AN:NY.AN ]: >>399
別に面倒でも何でもない。
CUDAならメジャーなディストリならほぼ動く。
面倒なのは単にLinuxの知識がないだけ。
例外は組み込み系で使う場合。
そもそもCUDAの開発やるのに、

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

できない奴なんていないだろ。
401 名前：デフォルトの名無しさん mailto:sage [2013/07/04(木) NY:AN:NY.AN ]: >>400
その後 nouveau が邪魔だって言われて
「cuda nouveau install」でググるんですよね
402 名前：デフォルトの名無しさん mailto:sage [2013/07/04(木) NY:AN:NY.AN ]: 多倍長整数の計算におすすめのライブラリとかある？
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のインストール中１０％くらいでエラーが出ます。

c:\NVIDIA\CUDAへの展開でかなり時間かかりますが、
インストール失敗すると、この下のインストーラー各種全部削除されて
最初からやり直す仕様なんですね・・・。（成功しても消えるのかもしれませんが）
450 名前：デフォルトの名無しさん mailto:sage [2013/09/11(水) 00:52:29.92 ]: ちゃんとリリースノートやインストールノート読んでやってないところ
あとはほんとに５３０のせいかどうかを他のグラボ（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 ]: ﾜﾛﾀｗ
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するのに、１１インチくらいのディスプレイのノート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 ]: どこかが展示してた外付けグラボが欲しいんだが、発売する気配なしｗ
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 ]: 意外と凡ミスでＨＷ制限越えた値で起動させようとしてるとか
パラメータの指定のような気がする
513 名前：デフォルトの名無しさん mailto:sage [2013/10/14(月) 18:15:28.26 ]: c/パラメータの指定/パラメータの指定ミス
514 名前：506 [2013/10/15(火) 01:01:28.01 ]: おかしな動作をしている箇所を特定しました。
（「cuda-gdb」で再現しないと、デバッグに時間がかかりますね。。。）

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

とりあえず、その変数に「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のアーキテクチャにうま～くハマる処理ならイイんだよ。
んで、結局グラフィック処理が一番イイというオチになるｗ
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
それはそれで、組み込み用に特化しそうだｗ
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 ]: スレッドが１ブロックでのスレッドが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ゲームでもするのかと思ったｗ
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 ]: 関数の呼び出しについて質問があります　回答をお願いします

ｃ言語でつくられたプログラムの一部を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 ]: コペンバローナ「ンーｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗ」
638 名前：628 mailto:sage [2014/01/16(木) 15:46:41.43 ]: ごめん、問題のポイントを誤解してたかも。
（CとC++とcuの入ったサンプルは見た覚えがない）

とりあえず俺の発言は忘れてください。ごめんなさい。
639 名前：デフォルトの名無しさん [2014/01/16(木) 15:47:58.43 ]: ペッコンバローナーｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗ
640 名前：デフォルトの名無しさん mailto:sage [2014/01/16(木) 17:03:38.03 ]: >>631
未解決の外部シンボルのエラーって、リンク（コンパイル）の時じゃなくて、実行時に出るの？
641 名前：デフォルトの名無しさん [2014/01/16(木) 18:25:52.58 ]: >>628
ペッコンペッコンペッコンバローナーｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗwwwwwwwwwwwwwwwwwwwwwwwwwwww
642 名前：デフォルトの名無しさん [2014/01/16(木) 18:29:18.81 ]: ロ・・・ロバｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗ
643 名前：デフォルトの名無しさん [2014/01/16(木) 18:32:40.24 ]: コペンハーゲ「ンーｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗ」
644 名前：デフォルトの名無しさん [2014/01/16(木) 18:33:44.44 ]: コペンバロー「ナｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗ」
645 名前：デフォルトの名無しさん [2014/01/16(木) 18:34:43.10 ]: バコナロ「バコーンｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗｗwｗｗｗｗｗｗｗｗｗｗｗｗｗ」
646 名前：デフォルトの名無しさん [2014/01/16(木) 18:58:08.66 ]: 、　　′　　　　　、　’、　　′　　　　　’　　　；.
　｀',．　’　　　　　　’、　　　′　’　　　．　・
　、′・．　’　　　；　　　’、　’、′‘　．・”
　　　　　　　　　　’、′・　　’、．・”；　　” 　’
　　　　　　’、　　(;;ﾉ;;　(′‘ ・．　’、′”；
　　　　 ’、′・　 (　（´;^｀⌒)∴⌒｀．・　　”　；　′・　　　　　　　　　　　, '´｀ヽ.-──-,'´｀ヽ.
　　　　､　’、　’・　､´⌒,;y'⌒(（´;;;;;ﾉ､"'人　　　　　　　　　　　　　　　　/　　　　ﾞi::::::::::::::ﾞi　　ﾞ:
　　　　　_､(⌒ ;;;:;´'从 ;'　　 ;:;;) ;⌒ ;; :） )､　　　　　　　　　　　　ミ　　　/　　　　　;ﾞ:;::::::;:::::::! 　　!
　　　　　（　´;`ヾ,;⌒)´　从⌒ ;) ｀⌒ ）⌒:｀．・　　　　　　　　　　　　,.;ﾞ　　r'^ｰ､;ﾞ:;ｨ:::ﾊ::λi,ｒ'ヽ, i
　‘:　；゜+° ′、:::::．　:::（ ::;; ﾉ ´⌒（,ゞ､⌒） ;;:::)::ﾉ　　　　　　　　　 i::i 　　| 　 iﾞﾉiﾉﾚ　ﾚ' ﾞ!!i　 | !
　....................｀:::､　ﾉ　 ...;:;_)　　...::ノ　ソ　...::ノ　　　　　　　　　ミ　　ﾊ::ﾞ､　i 　 ! ＞　　　＜!| 　!;
　　　　　:::::::::｀- ´::::::::::::::::::::::::::::::::::::::::::::::::::　　　　　　　　　　　　ﾉ:::λﾞｰ| 　 |〃　＿_ 〃i　 l'
　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　〈ﾉ::ｲ::ﾉ::::!　　!　　'､　ﾉノ'.　;
　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　〈/!::;ｲ::/,ｒ'´｀ﾞﾞi＞‐_-_ｔ.´ｒﾞ´ヾ
　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　ミ　　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乙ｗ
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 ]: 今や、テーブルにしてメモリから読み出すよりも、
手前で計算で作ったほうが速いからなｗ
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 向けなんじゃないかと使ったこともないのに妄想しております。
１回の評価計算そのものがめちゃくちゃ重い（単純に１００万回くらいループさせているだけ）で、
ループさせるプログラムそのものは数キロバイトも無いちっちゃなものです。
715 名前：デフォルトの名無しさん [2014/02/13(木) 18:26:36.46 ]: とりあえず今は手元にある Windows 7 64bits (チップセットはP55) に入れてお試しでCUDAプログラミングし、
C++で書いたシミュレータをCUDA対応に移植するところから始めたいとおもっています。
グラフィック出力を２本使いながらCUDAを使いたいのですが、私はどんなのを買ったらいいのでしょうか？
将来的には研究費をつぎ込むつもりですが、今は衝撃的に貧乏なので予算３万以内で。
716 名前：デフォルトの名無しさん mailto:sage [2014/02/13(木) 18:39:59.33 ]: ＞グラフィック出力を２本使いながら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 が２本あるのですが、
今はRadeon ２枚で４画面（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にあわせてソースを書き換えたり、
必要に応じてアルゴリズムも修正しなくてはならないと考えており、
そのトライ＆エラーに例えば一ヶ月かかっちゃうなら安い奴を買ったほうがいいかな、
とりあえず３万円程度で使い倒してみようかな、と考えている次第です。
723 名前：デフォルトの名無しさん mailto:sage [2014/02/13(木) 23:23:35.47 ]: >>721
そんな事情もあって3万円くらいで、なにがいいのかなと。
724 名前：デフォルトの名無しさん mailto:sage [2014/02/13(木) 23:50:08.03 ]: しれっとアルゴリズムを発見したとか言うよねーｗ
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（３２）で削られてた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用ボードの２つを出せば良いんだよ
（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]: こりゃ面白いな。
試してみたいｗ

ってか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]: 値段ﾜﾛﾀｗ
813 名前：デフォルトの名無しさん mailto:sage [2014/03/30(日) 00:12:47.34 ID:tTPtoH97]: >>812
むしろこの構成で高く付かない方がおかしいｗ
大雑把には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]: ひゃぁ～ｗ
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並列はむずかしいのかね？
おススメのＨＰか本はあるのかね
829 名前：デフォルトの名無しさん mailto:sage [2014/04/16(水) 08:31:11.50 ID:1cVHCjMz]: きみらって皆採掘が目的なの？１０年やって元がとれるくらい？
830 名前：デフォルトの名無しさん mailto:sage [2014/04/16(水) 09:10:52.39 ID:Im3wLyAK]: >>829
今から採掘に回るとかただの情弱じゃねーかｗ
単純にゲームとプログラミングのためだよ
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

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

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