【GPGPU】くだすれCUDAスレ part6【NVIDIA】at TECH
【GPGPU】くだすれCUDAスレ part6【NVIDIA】 - 暇つぶし2ch1:デフォルトの名無しさん
12/09/23 23:17:47.58
このスレッドは、他のスレッドでは書き込めない超低レベル、
もしくは質問者自身何が何だが分からない質問を勇気を持って書き込むスレッドです。
CUDA使いが優しくコメントを返しますが、
お礼はCUDAの布教と初心者の救済をお願いします。

CUDA・HomePage
URLリンク(developer.nvidia.com)

関連スレ
GPGPU#5
スレリンク(tech板)

前スレ
【GPGPU】くだすれCUDAスレ【NVIDIA】
スレリンク(tech板)
【GPGPU】くだすれCUDAスレ pert2【NVIDIA】
スレリンク(tech板)
【GPGPU】くだすれCUDAスレ pert3【NVIDIA】
スレリンク(tech板)
【GPGPU】くだすれCUDAスレ pert4【NVIDIA】
スレリンク(tech板)
【GPGPU】くだすれCUDAスレ part5【NVIDIA】
スレリンク(tech板)


2:デフォルトの名無しさん
12/09/23 23:18:20.77
関連サイト
CUDA
URLリンク(www.nvidia.co.jp)

CUDAに触れてみる
URLリンク(chihara.naist.jp)

CUDA のインストール
URLリンク(blog.goo.ne.jp)

NVIDIAの「GeForce 8800 GT(G92)」と次に控える64-bit GPUアーキテクチャ
URLリンク(pc.watch.impress.co.jp)

CUDAを使う
URLリンク(tech.ckme.co.jp)

NVIDIA CUDAを弄ってみた その2
URLリンク(dvd-r.sblo.jp)

CUDAベンチ
URLリンク(wataco.air-nifty.com)

KNOPPIX for CUDA
URLリンク(www.yasuoka.mech.keio.ac.jp)


3:1 忍法帖【Lv=38,xxxPT】(2+0:5)
12/09/23 23:20:02.29
テンプレここまでです。変更点はcudaさわってみた(URLリンク(gpgpu.jp))を
リンク切れのため外したことのみです。


4:デフォルトの名無しさん
12/09/24 00:35:21.83
>>1

5:デフォルトの名無しさん
12/09/24 00:53:05.44
>>1

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

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

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

7:デフォルトの名無しさん
12/09/25 20:47:36.80
たくさんのお仕事がないとGPUさんはがんばれない^^;

8:デフォルトの名無しさん
12/09/25 22:01:29.76
メモリ転送イヤポだから、AMD、IntelのiGPUはCPUの管理のメモリをGPUからでもアクセス
出来るようにする方向なんだろ

9:デフォルトの名無しさん
12/09/25 23:05:27.26
ユニファイドメモリってなんか先祖返りな不思議な気分ね
あとからGPUだけ交換できない時代にもなって寂しくなりそうだ

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

10:デフォルトの名無しさん
12/09/25 23:34:41.10
メモリ空間統合したほうが絶対イイね。
データのコピー転送オーバーヘッドはあほらいい。

11:デフォルトの名無しさん
12/09/26 17:04:17.48
コピー転送オーバーヘッドを調べる方法ないのか

12:デフォルトの名無しさん
12/09/26 18:23:45.59
16EiB の壁はいつ来るだろう

13:デフォルトの名無しさん
12/09/26 20:30:47.96
>>11
イベントで計測できるんじゃない?

14:デフォルトの名無しさん
12/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
12/10/08 16:59:48.04
CUDAプログラミングを体験したいのですが,CUDAのできる格安ノートPCを教えてください.


16:デフォルトの名無しさん
12/10/08 17:45:27.76
CUDAのプログラミングを体験したいのならエミュレーションで十分

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

18:デフォルトの名無しさん
12/10/08 19:06:04.44
>>17
すいません、ちょっと質問です。

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

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

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

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

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

19:17
12/10/08 21:21:05.42
>>18
自分は後藤さんの記事を参考にしているよ。
たくさんあるけど、いくつか紹介。

NVIDIAが次世代GPUアーキテクチャ「Kepler」のベールを剥いだ
URLリンク(pc.watch.impress.co.jp)

NVIDIA Fermiのマルチスレッディングアーキテクチャ
URLリンク(pc.watch.impress.co.jp)

NVIDIAの1TFLOPS GPU
「GeForce GTX 280」がついに登場
URLリンク(pc.watch.impress.co.jp)

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

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

CUDA テクニカルトレーニング
Vol I:CUDA プログラミング入門
URLリンク(www.nvidia.co.jp)
(103ページ目のスライド)

20:デフォルトの名無しさん
12/10/08 22:47:17.84
>>19
ありがとうございます。
参考にさせていただきます。


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

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

21:デフォルトの名無しさん
12/10/09 00:35:20.42
テクスチャフィルレートが相当すると言えなくもないけど、フォーマットによって変わるから参考にしかならない
概算はTex数×コアクロック(シェーダクロックではない)で見積もれる

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

22:17
12/10/09 00:40:24.46
>>20
そのGPUは最新世代のKeplerアーキテクチャだね。

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

URLリンク(pc.watch.impress.co.jp)

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

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

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

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


URLリンク(www.nvidia.com)

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

23:デフォルトの名無しさん
12/10/09 19:04:03.29
>>21
> 概算はTex数×コアクロック(シェーダクロックではない)で見積もれる

ありがとうございます。

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


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

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

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

がんばります。

24:デフォルトの名無しさん
12/10/12 16:08:53.39
PTXでブロックまたがってすべてのスレッドでグローバルメモリの同期やりたい時ってmembar.glでいいんだよね多分。

25:デフォルトの名無しさん
12/10/13 23:43:36.29
ホスト側のコードだけを書いた cu ファイルと、
デバイス側のコードだけを書いた ptx ファイルとをリンクして
ひとつの exe ファイルを作る方法はあるでしょうか。

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

26:デフォルトの名無しさん
12/10/14 00:10:02.48
>>25
ホスト側が面倒になるけどDriver APIとか


27:デフォルトの名無しさん
12/10/14 01:00:51.58
>>26
やはりそれしかないですか・・・

28:デフォルトの名無しさん
12/10/14 01:04:08.72
なんでそんなことをしたいのかが気になる。

29:デフォルトの名無しさん
12/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:デフォルトの名無しさん
12/10/14 01:55:27.84
Makefile

31:デフォルトの名無しさん
12/10/14 02:04:07.01
>>30
あぁ、そっちでコンパイルするファイルやコンパイル方法を制御するわけですね。

挑戦してみます。

32:デフォルトの名無しさん
12/10/14 04:55:38.24
PTXのコードをインラインアセンブラを使って直接cuファイルの
中にかけばいいじゃん。

33:デフォルトの名無しさん
12/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:デフォルトの名無しさん
12/10/14 09:16:31.85
>>33
URLリンク(developer.download.nvidia.com)

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

35:デフォルトの名無しさん
12/10/14 09:35:03.51
>>34
ありがとうございます。
なんというか、少々独特のインラインアセンブラ構文ですね。

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

36:デフォルトの名無しさん
12/10/14 14:53:23.16
>>35
GCCのインラインアセンブリ構文がこういうのだよ




37:デフォルトの名無しさん
12/10/14 15:28:15.65
>>36
URLリンク(d.hatena.ne.jp)
これ見ると、たしかに同じですね。

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

38:デフォルトの名無しさん
12/10/14 20:42:43.14
インラインアセンブラは今回の目的には合いませんでした。

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

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

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

39:デフォルトの名無しさん
12/10/16 07:24:30.85
CUDA5で美味しい事あるの?

40:デフォルトの名無しさん
12/10/16 07:37:48.25
>>39
新機能を使わないんだったら全然美味しくない。
CUDA5でビルドしたらかなり遅くなった。

41:デフォルトの名無しさん
12/10/16 09:54:43.52
CUDA 5 Production Release Now Available
CUDA Downloads | NVIDIA Developer Zone
URLリンク(developer.nvidia.com)


42:デフォルトの名無しさん
12/10/16 22:33:51.56
早く5の報告しやがれ

43:デフォルトの名無しさん
12/10/17 03:16:50.02
4Gamer.net ― NVIDIA,「CUDA 5」を正式発表。第2世代Kepler「GK110」に向けた準備が整う
URLリンク(www.4gamer.net)


44:デフォルトの名無しさん
12/10/17 08:26:25.67
>>42
普通に動いてるよ。

45:デフォルトの名無しさん
12/10/17 23:16:16.73
>>43
Nsightはプロファイラーも付いてるのか。
こりゃいい。

46:デフォルトの名無しさん
12/10/18 10:47:37.44
>>43
読んでみたけど
これはGeForce切り捨てってこと?

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

47:やんやん ◆yanyan72E.
12/10/18 12:03:55.32
それは、Kepler発表の時からゲーム用のKepler1と
GPGPU向けのKepler2があるってことになってた。
Kepler1があまりにGPGPUに向いてなくてGeforce680あたりを
買った人はがっかりしてたよ。

48:デフォルトの名無しさん
12/10/18 12:23:14.57
>>43
Eclipse用のNsightも出てLinuxやMacでも開発しやすくなるのは大きいかも。

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

49:デフォルトの名無しさん
12/10/18 15:18:50.07
>>47
ゲーマーにGPGPUっていらんだろうからな
要らないのを付けて高い値段・高消費電力になって売れないものになるなら削れだろ
CUDAする奴はとんがったことする奴だろ。そんな奴ならKepler2のTeslaぐらい買うだろうからな
買えない貧乏人はAMDのradeonに移行しろだな

50:デフォルトの名無しさん
12/10/18 19:01:56.13
Mac はどうか知らんが、Linux は Windows 版に比べて、
どうしてもドライバのチューニングが徹底されていない感じがする。

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

51:やんやん ◆yanyan72E.
12/10/18 19:34:39.18
そりゃ、いくらDRIが実装されたからといってX11なんていう
太古のグラフィックAPI使ってるんだから、そんなもんじゃないの?
本気を出させたければWayland待ち。

52:デフォルトの名無しさん
12/10/18 19:55:50.80
>>51
遅いのはCUDAじゃなくて、その結果を表示する
グラフィックスライブラリの方ってこと?

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

今度ためしてみよ

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


54:デフォルトの名無しさん
12/10/19 00:22:17.74
そう言えば Yellow Dog Linux for CUDA 使ってる人いる?
どんな感じなの?

55:デフォルトの名無しさん
12/10/19 12:39:47.00
Linuxなら、GUI止めないとカーネルによっては処理速度ががた落ちする。
使えるGVRAMも激減する。

56:デフォルトの名無しさん
12/10/19 19:37:02.47
CUDA + GUIつっても、サンプルでXが関わるところなんてウィンドウの枠だけじゃないか?
あとはOpenGLで描画されていると思うが

57:デフォルトの名無しさん
12/10/20 00:42:13.70
>>55
Windowsのほうがガタ落ちだし、使えるメモリも少ない。
グラフィックスを使うと遅くなるのはXの問題だから。

58:デフォルトの名無しさん
12/10/20 16:27:41.31
dynamic parallelism は GeForce じゃ無理なんですか

59:デフォルトの名無しさん
12/10/20 17:35:11.22
うん。

60:デフォルトの名無しさん
12/10/20 17:41:18.83
調べたなかではGDRAMのみのように見えるんだけど、
テスラだとL1、L2、シェアードメモリもECC保護されてるの?
それともL1、L2くらいの容量なら気にしなくてもいいのかな?

61:デフォルトの名無しさん
12/10/20 20:24:55.79
レジスタも。

62:デフォルトの名無しさん
12/10/20 20:47:40.75
>>58
今のところTesla K20のみだったはず

63:デフォルトの名無しさん
12/10/22 20:36:15.81
一般人向けは2014年まで待てとか遅すぎる

64:デフォルトの名無しさん
12/10/22 21:06:37.69
GK110はいつになったら一般向けで出てくるのやら…

65:デフォルトの名無しさん
12/10/22 23:32:49.15
>>64
ないと思うのは俺だけか

66:デフォルトの名無しさん
12/10/22 23:35:31.31
>>65
gtx780とかじゃないか?
来年の春だった気がする。

67:デフォルトの名無しさん
12/10/23 04:53:17.04
GTX 780はKepler1の改良版だって聞いたぞ。

68:デフォルトの名無しさん
12/10/23 08:47:20.31
一般人向けでダイナミックなんちゃらが使えるのはMaxwellからとか

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

69:デフォルトの名無しさん
12/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:デフォルトの名無しさん
12/10/23 22:26:20.49
>>69
CUDAプログラミングはまだまだ敷居が高いね・・・

71:デフォルトの名無しさん
12/10/24 10:52:47.74
nvcc ***.cu -O2 -Xcompiler -O2
のようにO2を重ねるのは無意味ですか?
前者のO2はGPU用,後者のO2はCPU用と勝手に思っていたんですが,
同じ事を繰り返しているような気がしてきました.

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

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

73:71
12/10/24 13:25:43.83
ptxで2つある場合,前者のみ,後者のみ,両方無い場合を比較しましたが,
冒頭の***.iファイルの名前が微かに違うのみで差はありませんでした.
両方消しても差が出ないのは?ですが,
重ねても意味は無さそうであることが分かりました.

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

74:デフォルトの名無しさん
12/10/24 14:43:07.65
>>73
今のgccのディフォルトが-O2相当なんで、書かなくても変わらないのはその所為。
試しに、-O3とか-O1との組み合わせを試してみたら?

75:デフォルトの名無しさん
12/10/25 04:28:58.35
登録ユーザーサイトが復旧したよ

76:デフォルトの名無しさん
12/10/25 21:42:51.21
k20はやっぱり高いな。
38万だそうだ。
20万切ってくれないと買えない。

77:デフォルトの名無しさん
12/10/27 22:36:10.01
dynamic parallelism対応のGeforce(GTX8XX?)が出たら
画像とか動画を扱うソフトは瞬く間にCUDA完全対応になるのかね?

78:デフォルトの名無しさん
12/10/28 00:23:20.16
んなわけない

79:デフォルトの名無しさん
12/10/28 00:40:32.58
dynamic parallelismができるからCUDAが劇的に簡単になるわけじゃないから。
Reductionとかで効果はあるけど。

80:デフォルトの名無しさん
12/10/28 03:58:28.19
そもそもReductionはマルチパスにしないで
2パスで済ませた方がいいのは、
CUDAのreductionトレーニングでも明らか


81:デフォルトの名無しさん
12/10/29 13:40:34.37
CUDAのプログラム作って動かしたいです
自分のMacbookは、グラフィックのチップがIntel GMA X3100なんですけど、
NVIDIAじゃないとCUDAは使えないんですか?

82:デフォルトの名無しさん
12/10/29 15:35:40.18
ここで聞いて良いのか分からないので、不適切なら誘導お願いします。

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

83:デフォルトの名無しさん
12/10/29 18:51:19.74
NVAPIをhackすればできる

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

85:デフォルトの名無しさん
12/10/31 14:43:09.18
>>84
カーネル呼び出した時にエラーが起きてるんでしょ。
エラーチェックしていないんじゃないの?

86:デフォルトの名無しさん
12/10/31 14:54:52.27
>>84
cudaGetLastError()は何と言っている?

87:デフォルトの名無しさん
12/10/31 14:58:04.91
>>85
即レスありがとうございます
正にその通りでした。単にメモリの要求量がおかしかっただけみたいです
初歩的すぎるミスに自己嫌悪…

88:デフォルトの名無しさん
12/10/31 16:49:46.93
NVIDIAR Nsight? Visual Studio Edition 3.0 CUDA Preview
Nsight Visual Studio Edition Early Access | NVIDIA Developer Zone
URLリンク(developer.nvidia.com)


89:デフォルトの名無しさん
12/11/03 02:08:23.38
Nvidia Geforce forum is back from the dead
URLリンク(www.fudzilla.com)


90:デフォルトの名無しさん
12/11/07 15:17:37.33
CRS形式の行列格納サンプルコードってどこかにない?

91:デフォルトの名無しさん
12/11/07 15:59:46.17
いくらでもあるだろ
圧縮方法を理解できたらサンプルもいらんな

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

92:デフォルトの名無しさん
12/11/08 02:56:28.41
>>91
圧縮方法はわかったんですがコードに上手く起こすことができなくて困っていたんです。何かいいサンプルがあれば教えていただけると助かります。

93:デフォルトの名無しさん
12/11/08 10:51:49.85
馬鹿には無理。

94:デフォルトの名無しさん
12/11/12 06:21:00.64
CUDA5は既存のGPUに入れると遅くなるの?

95:デフォルトの名無しさん
12/11/12 14:32:24.27
研究室でCUDA用にGTX680搭載PCの導入が決定してしまったんだが評判悪いとはいえ流石に今使ってる560Tiよりは性能いいよね?

96:デフォルトの名無しさん
12/11/12 23:48:40.00
Tesla K20きたぞ

97:95
12/11/13 01:28:48.03
>>96
予算処理上の都合だったらしい。

98:デフォルトの名無しさん
12/11/13 03:25:03.66
最近プログラム入門した
CUDAとか聞くとワクワクするけど物理の知識も科学の知識も特にないので
数百万スレッド並列で処理するネタが思いつけなくて悲しい思いになる
もっとちゃんと勉強しておけば良かった

99:デフォルトの名無しさん
12/11/13 05:46:01.18
京が3位に

100:デフォルトの名無しさん
12/11/13 06:14:28.20
東工大の学生たちはもうGK110貰ってるの?

101:デフォルトの名無しさん
12/11/13 23:01:28.72
Intelがついに来るぞ
URLリンク(pc.watch.impress.co.jp)

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

102:デフォルトの名無しさん
12/11/14 00:25:44.33
>>101
nVIDIAが押されて、もうちっと貧乏客を引き込むマーケティングをやってくれんかな。

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

103:デフォルトの名無しさん
12/11/14 16:56:20.91
開発環境やソフトウェアの安定性とか含めて、XeonPhi強そうだなぁ

104:デフォルトの名無しさん
12/11/14 18:03:25.42
XeonPhiは高いぞ
安いGPUは安い
Tesla買うならXeonPhiのほうがよさそうだが

105:デフォルトの名無しさん
12/11/14 19:32:13.13
半年ぐらいしたら、$500くらいのローエンドXeon Phiが出るだろうから、純粋にアクセラレータとしてのteslaは厳しいかもなあ。

106:デフォルトの名無しさん
12/11/14 20:08:20.86
Phi触ってみてぇ。
OpenMPで簡単マルチコアプログラミング♪


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

107:デフォルトの名無しさん
12/11/14 20:25:15.46
SSEとかAVXみたいなのをちゃんと使える人じゃないと
TESLAのような性能はでないよ。
512bit演算命令が命だから。

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

108:デフォルトの名無しさん
12/11/14 20:30:01.47
512bit演算命令ってのがあるのか?
AVXでも256bitだが・・・

109:デフォルトの名無しさん
12/11/14 20:32:10.61
VPUてので512ビット命令を処理するようだな

110:デフォルトの名無しさん
12/11/14 20:33:10.12
ま、経験上はベクトル命令はCUDAよりは扱いやすいよ

111:デフォルトの名無しさん
12/11/14 20:37:54.55
うん、イントリンシックでベクトリ処理書くの楽♪
条件分岐がめんどいけど、LNIはマスクレジスタをサポートしてたからだいぶ楽に書けそう。
しかも512bitもあるなんて最高すぐる。

あー、Phi触りてぇ~。

112:デフォルトの名無しさん
12/11/14 21:32:04.73
OEM向け1000個ロットでXeon Phi 5110Pが2650ドル
らすぃ
なんか価格でもTeslaやばそうだな

Intel,スーパーコンピュータ向けアクセラレータ「Xeon Phi 5110P」発表。60基のx86コアを1チップ上に集積
URLリンク(www.4gamer.net)

113:デフォルトの名無しさん
12/11/14 21:33:53.30
むしろ値下げ合戦になればよい。

114:デフォルトの名無しさん
12/11/14 21:47:31.12
合戦になるほど数競争起きる市場でもないべ

115:デフォルトの名無しさん
12/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:デフォルトの名無しさん
12/11/14 21:56:39.37
長文の上に間違えてーらorz
Fermiにもそれなりに機能を持たせたんだと思う→Geforceにもそれなりに機能を持たせたんだと思う

117:デフォルトの名無しさん
12/11/14 22:59:58.87
>>115
nVIDIAの株を空売りすれば儲かるということか。

118:デフォルトの名無しさん
12/11/14 23:57:37.97
phiはCPUに内蔵させるGPUコアと共通化させて
コスト落としつつマーケットシエア取る作戦かな?

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

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

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

120:デフォルトの名無しさん
12/11/15 04:26:43.84
CUDAは開発環境タダだけどXeon PhiはIntel Compiler必須だよね

121:デフォルトの名無しさん
12/11/15 05:55:49.56
ものいりだねえ、 Phi

122:デフォルトの名無しさん
12/11/15 11:55:29.66
>>120
GCCが対応するって書いてあったぞ

123:デフォルトの名無しさん
12/11/15 15:34:52.68
そもそもintelコンパイラもLinux版はフリーだよな

124:デフォルトの名無しさん
12/11/15 18:21:21.39
ここのスレ見てたら、phiを月2、3万位でレンタルする商売が出来そうな気がしてきた…

125:デフォルトの名無しさん
12/11/15 20:33:05.20
GPUだってAmazonのクラウドで借りたりできるし
Phiもそういうの出るだろう

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

127:デフォルトの名無しさん
12/11/15 23:36:48.09
Phi、4、5万で買えるようにならないかなぁ~。

128:デフォルトの名無しさん
12/11/17 12:12:13.46
Phいらないだろ
一晩中PC動かせばいいだけだろ

129:デフォルトの名無しさん
12/11/17 12:35:26.08
動画エンコ用途でもあるまいにw
すでに一年中計算回してるような人に、これなら一ヶ月で済むよ、って訴求するのが筋の製品だろ

130:デフォルトの名無しさん
12/11/17 13:00:33.16
数時間、数日動かして、後から些細なバグに気づいた時の何とも言えない気持ち

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

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

131:デフォルトの名無しさん
12/11/17 14:05:25.11
>>130
とりあえず賽の河原症候群と名付けておくよ

132:デフォルトの名無しさん
12/11/19 21:17:02.44
phiは本体CPUもXeon使った時の協調性とかで
パワー増すんだろうなぁ
Teslaやばいなぁ…

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

133:デフォルトの名無しさん
12/11/19 22:16:34.19
しかし投げ売りの後に待っているのが撤退だとしたら…?
Xeon Phiには縁がなさそうだから気軽に触れるCUDAにがんばってほしいなぁ

134:デフォルトの名無しさん
12/11/19 22:19:19.97
自作ゲームにCUDA利用してる奴っている?

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

135:デフォルトの名無しさん
12/11/22 18:48:11.46
使えもしないのに欲しくなってとりあえずダウンロードしているんだけど
クソミソに通信速度遅い。
25~60KB/sをウロウロしてるけどそんなもん?
URLリンク(developer.download.nvidia.com)

136:デフォルトの名無しさん
12/11/22 19:03:52.66
>>135
ウチじゃ2.0MB/s位出ているぞ。

137:デフォルトの名無しさん
12/11/22 20:53:34.08
まじすか
一端回線切ってIPアドレス変更とかしても速度出ない・・・

OCN保土ヶ谷

138:デフォルトの名無しさん
12/11/23 01:07:04.36
>135
遅すぎ
間に無線とか入れてない?

139:デフォルトの名無しさん
12/11/30 22:59:09.24
人柱はまだか
URLリンク(www.amazon.co.jp)

140:デフォルトの名無しさん
12/12/01 19:38:58.57
>>139
高すぎ。

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

141:デフォルトの名無しさん
12/12/01 23:01:57.20
CPU内蔵のiGPUをPCの表示用に、dGPUをCUDA GPGPU専用にする場合
やdGPUを2つ使って片方をPCの表示用に、もう片方のdGPUをCUDA GPGPU専用にする場合
ってそれらが出来る(出来ない)マザー、CPU・APU、dGPUってある?
出来るのなら、これをPC表示用、これはGPGPU用って設定とかするの
する場合どうするんですか?

142:デフォルトの名無しさん
12/12/01 23:12:12.90
>>141
GTX 580/590とASUS Maximus V GeneとCore i7-3770Kの組み合わせなら出来た。
BIOSでどちらを表示用に使うか設定できる。

143:デフォルトの名無しさん
12/12/01 23:59:21.54
>>142
劇速れすありがとう
マザーのBIOSに設定があれば、iGPUとdGPUの場合は出来ると
思っていいのかな

144:デフォルトの名無しさん
12/12/02 00:02:38.36
>>143
CUDAのデバイス指定はアプリケーション次第だよ。BIOSは関係ない。
ちゃんとどのデバイスを使うか指定できるようになっていれば問題ないよ。

145:デフォルトの名無しさん
12/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:デフォルトの名無しさん
12/12/02 10:25:11.74
>>145
だから、BIOSで設定するのは画面の表示だけってことなんだよ。
CUDAでの利用はそれとは全く別に行えるよ。
BIOSでiGPUを表示用に選んでマザボにディスプレイをつないでから、
CUDA対応アプリでdGPUを選べばいいだけの話。

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

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

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

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

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

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

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

149:デフォルトの名無しさん
12/12/02 16:46:04.44
表示用デバイスでなくてもOpenGLは動かせるから、cudaと連携できると思う。

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

151:デフォルトの名無しさん
12/12/22 19:54:35.28
>>150
ごめん自己解決した
VSでコードいじってたんだけど、すべてリビルドしたら表示されるようになった

152:デフォルトの名無しさん
12/12/26 08:21:21.01
Fermiでスレッドブッロクを512以上を指定すると、カーネルが起動しない。
Fermiはブロックごとに1024スレッド対応しているはずなので、
までレジスタが足りないからなのか、
シンプルなカーネルだと1024スレッドまでいける。
動かないならエラーで落ちて欲しいんだが。

153: ◆4hloUmTGPY
12/12/27 10:24:29.87
質問です

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

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

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

154:デフォルトの名無しさん
12/12/27 11:09:28.44
>>153
いんてりせんすがC++以外の表記に対応していないんだろ。
自分でぷらぐいんを書けばなんとかなるんじゃね?
尤も、いんてりせんすに頼るような奴に書けるかどうかは知らんが。

155:デフォルトの名無しさん
12/12/27 22:41:04.67
2012じゃCUDAの環境がまだまだなんじゃね?
2010だとようやくこなれてきた感じがあるが。

156: ◆4hloUmTGPY
12/12/28 11:54:50.91
>>154 >>155
ありがとうございます
無理そうなので諦めます

157:デフォルトの名無しさん
12/12/29 11:26:38.86
こんな感じのカーネルがあって、

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

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

158:デフォルトの名無しさん
12/12/29 22:44:34.71
大体そんな感じ。分岐から外れたスレッドが生きているか死んでいるかは兎も角も。

159:デフォルトの名無しさん
12/12/30 13:32:23.86
>>158
ありがとう。

160:デフォルトの名無しさん
12/12/30 15:06:29.69
CC3.5のGeForceまだ?

161:デフォルトの名無しさん
13/01/03 11:10:10.18
K20が売れなくなるから当分無し

162:デフォルトの名無しさん
13/01/05 10:06:43.42
CUDAの性能でいったらGTX580>GTX680なんですか?

163:デフォルトの名無しさん
13/01/05 15:19:19.49
>>162
URLリンク(blog.accelereyes.com)

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

164:デフォルトの名無しさん
13/01/05 21:13:33.60
>>163
ありがとです。

165:デフォルトの名無しさん
13/01/10 18:01:12.19
ブロック数がSM数以上の場合、ブロックでの動作が終了したら次のブロックにいくんですか?

166:デフォルトの名無しさん
13/01/10 18:27:54.28
次のブロックっつーか、残りのブロックだな。

167:デフォルトの名無しさん
13/01/14 00:28:51.24
CG法をCUDAで実装してMatrixMarketの疎行列を求解しようとしているんですが
連立一次方程式のb要素はどのように設定したらいいのでしょうか?
MatrixMarketで与えられているのはAの疎行列だけなのでbをどのように設定したらいいのかわかりません

168:デフォルトの名無しさん
13/01/18 16:14:19.91
テクスチャメモリって使っててあまり早くならないんだけど,実際の効果ってどれくらいなの?

169:デフォルトの名無しさん
13/01/18 18:10:26.59
質問です。

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

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

170:デフォルトの名無しさん
13/01/18 18:12:36.87
>>169
ソースの該当箇所の引用ぐらいしろよ。

171:デフォルトの名無しさん
13/01/18 18:17:57.52
>>169
nvccでコンパイルする対象(*.cu)ではテンプレートは使えなかったと思う。
C++からCのライブラリを呼ぶ要領で分割コンパイルすればホスト側で使う分にはOK。

172:デフォルトの名無しさん
13/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:デフォルトの名無しさん
13/01/19 20:26:41.74
>>172
extern Cをしてないだけじゃないのか?
それにしても引数多いな。

174:デフォルトの名無しさん
13/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
13/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:デフォルトの名無しさん
13/01/20 02:11:59.82
>>168
Fermiより前のL2キャッシュが無いようなGPUだと効果あるよ。
あとハードウェア補間を活かせる分野

177:デフォルトの名無しさん
13/01/20 02:20:35.01
>>176
fermiでも場合によっては効果はある。
でも一番いいのは線形補完で使う時だな。

178:デフォルトの名無しさん
13/01/21 13:30:50.30
>>175
リンク前の個々の .cuファイルのコンパイルの時点でのエラーのように見える。
どこかで nvcc -c -o a.o a.cu b.cu みたいな呼び出ししてるんじゃ。

179:デフォルトの名無しさん
13/01/23 14:40:54.94
SMが10個あって処理の数が合計100個ある場合SMはそれぞれ10回動く?
それとも重い処理があって時間がかかってるSMがあれば空いてる別のSMが担当する?

180:デフォルトの名無しさん
13/01/25 23:00:07.64
なんとも答えにくい初心者の質問だな

181:デフォルトの名無しさん
13/01/26 02:29:08.92
質問
TESLA K20買おうと思ってるんだけど、プログラミングとは別にエンコでも使おうかなと思ってる。
やっぱりGTX680とかGTX580と較べて全然性能違う?(体感)
開発にはVS2012Pro持ってるからそれに追加しようと思ってる

182:デフォルトの名無しさん
13/01/26 02:39:37.68
宝の持ち腐れ臭

183:デフォルトの名無しさん
13/01/26 02:44:45.43
>>182
一部のスーパープログラマー以外皆大なり小なり宝の持ち腐れだろう

184:デフォルトの名無しさん
13/01/26 04:14:43.40
>>180
色々見たけど書いてないから自明なのかな?

185:デフォルトの名無しさん
13/01/26 09:11:31.41
>>181
GTX680でええやん、いくらなんでも値段が凄まじすぎる>K20
それにGPUエンコは品質が……フィルタはともかく、AVCに落とすにはQSV使う方が速いし

186:デフォルトの名無しさん
13/01/26 12:22:12.43
>>185
まーなー、BOINCもやってるからTESLAちゃんにBOINCさせたら
どうなるのかなって気持ちもあったりするんだよな

187:デフォルトの名無しさん
13/01/26 13:45:33.61
まぁ一応言っておくと、性能以前に
Dynamic Parallelismの対応有無の差が大きい >GTX680とK20

188:デフォルトの名無しさん
13/01/28 04:32:35.73
どうせ数年待てば下位にも降りてくるんだから
待ちなされ

189:デフォルトの名無しさん
13/01/28 13:15:40.00
そして降りてくる時にはさらに凄い何かに目移りする訳ですね

190:デフォルトの名無しさん
13/01/28 19:07:10.24
質問です。

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

191:デフォルトの名無しさん
13/01/29 00:35:26.56
>>190
ググればでてくるだろうに。
URLリンク(docs.nvidia.com)

192:190
13/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:デフォルトの名無しさん
13/02/02 10:55:34.50
DPでのQソートのソース見ちゃうと
無しでやるのがアホらしく感じてくるな

194:デフォルトの名無しさん
13/02/07 00:52:32.17
CC3.5がコンシューマーに降りてくるかも♪

GK110を搭載するGeForce GTX 780 Titan 2013年2月末?
URLリンク(northwood.blog60.fc2.com)

195:デフォルトの名無しさん
13/02/11 00:01:58.59
winrarをcudaから展開するdllってなんか無いかな?
pw付きでもこちらからの操作で出来るやつ

196:デフォルトの名無しさん
13/02/11 02:10:48.39
>>195
GPU使ってパス解読できるソフトなら売られてるが、
CUDA使ってRAR解凍するソフトってあったっけ……

197:デフォルトの名無しさん
13/02/11 07:57:27.35
それってそんなに意味あるの?

198:デフォルトの名無しさん
13/02/11 12:16:16.19
>>196
パスワード解析なんかはGPUに向いてないんじゃないか?

199:デフォルトの名無しさん
13/02/11 16:18:44.80
>>198
zip用→URLリンク(www.internal.co.jp)
RAR用→URLリンク(www.golubev.com)
それなりに速度は出るっぽい

200:デフォルトの名無しさん
13/02/12 11:46:35.13
>>198
ブルートフォースはGPUの独壇場じゃない?
FPGAの専用プロセッサを除いて。

201:デフォルトの名無しさん
13/02/12 12:50:32.08
ブルートフォース
プルートフォース

202:デフォルトの名無しさん
13/02/12 22:26:05.34
>>200
>>198は、GPUは整数演算がそれ程得意では無い事を言いたいのでは?

203:デフォルトの名無しさん
13/02/12 22:26:14.43
VirtexクラスのFPGAでチューニングして200MHzくらいで動かしたらハイエンドGPUより速い?

204:デフォルトの名無しさん
13/02/13 00:38:36.16
>>202
URLリンク(http.developer.nvidia.com)

ごめん、くっそ遅かった

>>203
20KLUTsで20GBpsでる

205:デフォルトの名無しさん
13/02/14 00:30:55.49
GeForce 8800の整数は遅いけど、Fermiはそれほど遅くない

206:デフォルトの名無しさん
13/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に達する。

URLリンク(northwood.blog60.fc2.com)

GTX580 197.6GFLOPS
GTX590 311GFLOPS
あれ?

207:デフォルトの名無しさん
13/02/14 03:13:03.14
Fermiは大サービスだったんだよ。
もう倍精度欲しいならTesla買うしか無い。
K20でも30万くらいだろ?

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

208:デフォルトの名無しさん
13/02/14 03:30:00.81
本当Fermiは大サービスだったよなぁ。整数やビット演算もまともだったし。
CUDA的にkepler refreshは大きく期待できることもないのでmaxwellが頼りなのだが

209:デフォルトの名無しさん
13/02/14 12:17:13.77
>CUDA的にkepler refreshは大きく期待できることもない
いやいやDPがあるだろ

210:デフォルトの名無しさん
13/02/14 17:07:23.36
Xeon Phiの発売で貧乏人以外には完全にオワコンでしょ

211:デフォルトの名無しさん
13/02/14 18:52:59.63
GPGPUはAMDですよ

212:デフォルトの名無しさん
13/02/14 22:50:06.45
>>209
Dynamic Parallelism(綴り覚えた)ってそんなに大きいの?
なんていうか使えるハードウェア資源が増えるわけじゃないし…って思って。

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

213:デフォルトの名無しさん
13/02/15 00:32:37.20
インテル? Xeon Phi 5110PとNVIDIA Tesla K20の行列積における実効性能比較
URLリンク(www.hpc.co.jp)

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

214:デフォルトの名無しさん
13/02/15 01:40:16.84
>>212
OpenMPだからPC用のプログラムがそのままで動くよ

215:デフォルトの名無しさん
13/02/15 08:15:58.27
URLリンク(devgurus.amd.com)

7970はDGEMM 665GFLOPSらしい。

216:デフォルトの名無しさん
13/02/15 08:55:15.74
>>214
動くけど、専用ベクトル命令使わないと速度でないよね?

217:デフォルトの名無しさん
13/02/15 18:51:43.68
>>213
このK20Mと5110Pっておいくら万円するんだろう。

218:デフォルトの名無しさん
13/02/15 22:17:58.89
プログラムが簡単に書けることと、速度が出し切れることは大抵両立できない

219:デフォルトの名無しさん
13/02/16 00:03:14.05
>>216
インテルコンパイラを使えばいい。
それに、インテルにはMKLがあるから、
線形問題であれば、既存のコードにディレクティブを挿入するだけで速くなる。

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

220:デフォルトの名無しさん
13/02/17 11:06:58.52
Tesla買うならXeon Phiになるだろうが
GTXなんとかでもCUDAは使えるからな

221:デフォルトの名無しさん
13/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の
倍精度浮動小数点演算性能に制限がかかっている様子はなさそうです。

URLリンク(northwood.blog60.fc2.com)

222:デフォルトの名無しさん
13/02/18 00:12:16.06
>>221
そんなうまい話は無いと思うけどね

223:デフォルトの名無しさん
13/02/18 20:17:49.22
Titanたっけー。
誰か研究室とかで買って報告よろ。

224:デフォルトの名無しさん
13/02/19 00:57:26.08
初日か初週でゲーマーが争奪戦して在庫切れで終わりだろ
悠長に評価して予算付けてって頃にはもう終わってる

225:デフォルトの名無しさん
13/02/19 01:11:14.40
年度末で10万くらい予算余ってるから何か買っていいよ。



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

226:デフォルトの名無しさん
13/02/19 04:58:05.78
>>225
10万で買えるの?

227:デフォルトの名無しさん
13/02/20 00:31:37.18
>>260
118000円で予約したからさすがに10万ジャストは無理

228:デフォルトの名無しさん
13/02/21 19:46:59.76
K20の約半額か

229:デフォルトの名無しさん
13/02/21 19:59:41.22
>>228
安いと言えば安いが、個人で遊ぶのはきつい価格だった。

230:デフォルトの名無しさん
13/02/21 20:25:32.67
>>229
でも580や480と違って倍精度もtesla相当に設定できるようだし、
ECCと演算保証がないことに目をつぶればお得といえばお得じゃない?
後々の中古売却を考えるなら値崩れはTeslaのほうが少ないだろうけど。

231:デフォルトの名無しさん
13/02/23 01:24:32.93
「法人さんはTeslaを買ってくださいよ」(2/21)
URLリンク(www.gdm.or.jp)

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


お前らか…

232:デフォルトの名無しさん
13/02/23 07:48:27.16
まー自分の金じゃないしとりあえず問い合わせぐらいはするでしょ

233:デフォルトの名無しさん
13/02/23 22:46:08.24
実際の性能向上率をみたり、書き換えが必要な場合それの専攻開発用とかに買うんだろ。
普通本番機ではサーバーと一緒にTESLA導入するよ。

234:デフォルトの名無しさん
13/02/24 01:05:09.57
しかしTITANはDynamic Parallelismが無いのが不便だね
枝分かれが動的に決まる深さ優先探索を分散処理させたいのだけど

235:デフォルトの名無しさん
13/02/24 01:08:24.14
ちょw
Dynamic Parallelismないって致命的じゃんw

236:デフォルトの名無しさん
13/02/24 01:15:58.14
オワタ
しかし、GTX TitanではKepler GPUの新機能である
Dynamic ParallelismとHyper-Qという機能が省かれている。

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

237:デフォルトの名無しさん
13/02/24 12:09:11.09
Dynamic ParallelismつけたらTesla売れなくなるだろww

238:デフォルトの名無しさん
13/02/24 14:45:08.55
最も格安なCUDAできるノートPCはどれ?

239:デフォルトの名無しさん
13/02/24 14:45:50.04
倍精度付けてくれたから、Dynamic Parallelismもてっきり対応してるものかと。。
CUDAのページにはGTX TITAN CC 3.5って書いてるしね。

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

納得できん。

240:デフォルトの名無しさん
13/02/24 14:50:45.05
最近はグラ用とコンピュート用で分化させる方針みたいだよ。
残念ながらね・・・。

241:デフォルトの名無しさん
13/02/25 00:41:40.15
Hyper-Q Dynamic palallerlismはついているようだが?
URLリンク(pc.watch.impress.co.jp)

242:デフォルトの名無しさん
13/02/25 01:03:16.82
>>241
>なお、既報の通り、HPC向けにGK110では「Hyper-Q」および「Dynamic Parallelism」という機能も追加されている。

URLリンク(news.mynavi.jp)
>しかし、GTX TitanではKepler GPUの新機能であるDynamic ParallelismとHyper-Qという機能が省かれている。

243:デフォルトの名無しさん
13/02/25 23:22:58.37
>>240
Fermiで懲りたんだろな
VGAにGPGPUで使う機能をてんこ盛りして爆熱にしてもしょうがないし
高く売れるGPGPU機を安く売っているVGAで代替できるじゃ金儲けできないし

244:デフォルトの名無しさん
13/02/26 04:32:52.99
TitanはXeon Phiが出てなかったら出なかったかもな

245:デフォルトの名無しさん
13/02/26 09:55:19.15
Fremiで動いていたコードでKeplerで動かすとカーネルが起動しないらしくて、
計算できないんだが、Keplerだとなにか気をつけることってあるのかな

246:デフォルトの名無しさん
13/02/26 11:17:27.91
Fermiで動いてたのはたまたまであって、
注意深く解析すると、そのコードはスレッドの起動順序などによって
アクセス違反を起こす、とかね。

247:デフォルトの名無しさん
13/02/26 11:32:45.34
もしかしてスレタイって下らないのクダとCUDAを掛けてるの?

248:デフォルトの名無しさん
13/02/26 11:49:25.57
Xeon Phiって結局まだ個人では単品で入手できないね

249:デフォルトの名無しさん
13/02/26 13:09:38.68
>>247
すげーーーーーーーーーーーーーーーー
良く気付いたな

250:デフォルトの名無しさん
13/02/26 21:09:38.74
CUDAはCUDAらない。

251:デフォルトの名無しさん
13/02/27 01:58:36.28
今夜は良く冷えますね

252:デフォルトの名無しさん
13/03/01 08:07:53.55
おいおい、TitanでDP使えるらしいじゃないか!
訂正記事書かれてるぞ
URLリンク(headlines.yahoo.co.jp)
2/28追記:記事掲載当初、「GTX TitanではKepler GPUの新機能であるDynamic ParallelismとHyper-Qという機能が省かれている」と記載させていただいておりまし
たが、GTX Titanでもこれらの機能をサポートしていることが明らかとなりましたので、当該部分ならびにそこに付随する文章
を修正させていただきました。

253:デフォルトの名無しさん
13/03/01 09:42:16.44
>>252
キター
DPって、倍精度と紛らわしいよね

254:デフォルトの名無しさん
13/03/01 11:19:34.68
マジかよw
この訂正記事に気づいてない人の分だけ売上落ちるぞこれ

255:デフォルトの名無しさん
13/03/01 12:55:23.44
このスレ見てて良かったと思う瞬間


256:デフォルトの名無しさん
13/03/01 16:59:12.85
│  ↑
└─┘
おらっしゃあぁぁ!!!
 ∩∧ ∧
 ヽ( ゚Д゚)
   \⊂\
    O-、 )~
      ∪

257:デフォルトの名無しさん
13/03/01 17:17:18.74
計算ミスしないなら欲しい

258:デフォルトの名無しさん
13/03/01 18:55:26.41
計算エラーって、確率的に起こるもんなの?
それとも同じところで再現?

259:デフォルトの名無しさん
13/03/01 19:58:01.79
ハードウェア的には同じところだとは思うが、プログラム上の同じところで
再現してくれるとは限らない。

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

260:デフォルトの名無しさん
13/03/01 20:12:57.40
メモリエラーの場合は確率的じゃない?

261:デフォルトの名無しさん
13/03/01 20:40:09.11
Dynamic Parallelismあり+Geforce並のエラー精度
Dynamic Parallelismなし+Tesla並のエラー精度

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

262:デフォルトの名無しさん
13/03/01 21:02:25.16
その記者がIntelから金渡されてたりして

263:デフォルトの名無しさん
13/03/01 23:24:13.26
品薄もんだからよかったものの、
大量に捌きたいもんだったらNVidiaから訴えられてもおかしくないレベルw

264:デフォルトの名無しさん
13/03/01 23:26:48.44
nVIDIA Japan公式Twitterだと使えるってアナウンスが
随分前にあったけど、使えるで確定なのねw

265:デフォルトの名無しさん
13/03/01 23:31:11.08
くそ、DPない情報に踊らされて買いそびれたじゃねぇか・・・
まだ店で売ってるのかな?

266:デフォルトの名無しさん
13/03/01 23:42:31.95
どさくさに紛れてGPUDirectも対応してないかな?

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

268:デフォルトの名無しさん
13/03/02 09:43:58.97
>>267
>エラーがでたらやり直せばいい。

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

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

269:デフォルトの名無しさん
13/03/02 11:20:49.48
そもそもエラーだと気付ければいいけどな
いつの間にか数値の一部が化けてたら気付かないぞ

270:268
13/03/02 11:37:11.88
>>269
あ、そういえばw

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

271:デフォルトの名無しさん
13/03/02 11:53:59.26
シミュレーションに使っているけど、計算ミスが問題になったことはないなぁ。
あ、そもそも精度が要るところはCPUで計算してたわ。

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

273:デフォルトの名無しさん
13/03/02 17:05:17.55
そんなに心配なら、同じ計算を別のボードでもやって比較すれば良い。
同じボードでやっても、あまり意味は無いぞ?

274:デフォルトの名無しさん
13/03/03 22:00:13.88
計算分野は狭いようで果てしなく広い分野だからな
エラー検知がどの程度重要かどうかは分野ごとに異なるから
自分の身の回りだけの話をしてもしょうがないだろ

275:デフォルトの名無しさん
13/03/03 23:45:41.01
これでTeslaと同じくらいの価格だな
URLリンク(akiba-pc.watch.impress.co.jp)

276:デフォルトの名無しさん
13/03/06 17:18:26.89
CUDAとOpenCV併用するとき、.cuファイル内でOpenCV使うことってできないの?
コンパイル時にインクルードファイル周りでエラー吐いたからとりあえずOpenCV周りだけ.cppファイルに隔離したんだけど。
後々不便そうだから対策があれば教えてほしい。

277:デフォルトの名無しさん
13/03/06 23:06:53.58
>>276
CUDAのカーネル関係はlib化して、
ビルド時にリンクする方法を取ればいいんじゃね?
NVCCってそんなに賢くないから、CPU周りで遅くなるし。

278:デフォルトの名無しさん
13/03/06 23:12:33.76
ふと、CUDAで自分で書いたカーネルを実行しつつ、OpenCV の cv::gpu:: の機能も同時に使うと、どうなっちゃうんだろう、と思った。

279:デフォルトの名無しさん
13/03/07 16:52:32.98
>>278
どちらも実行されるよ。
普通のマルチスレッドのプログラミングと
考え方は変わらん。

280:デフォルトの名無しさん
13/03/19 17:09:17.32
URLリンク(news.mynavi.jp)
【連載】
GPUの実効性能を上げるDynamic ParallelismとHyper-Q

281:デフォルトの名無しさん
13/03/19 19:27:04.30
>>280
安藤先生の記事や。
ありがたや~。

282:デフォルトの名無しさん
13/03/20 00:32:25.98
URLリンク(www.eevblog.com)

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

半田班ちょっと人柱頼む

283:デフォルトの名無しさん
13/03/20 07:16:49.30
>>282
どうやって発見したんだよ!w

284:デフォルトの名無しさん
13/03/20 08:04:17.97
英語読める割に綴り間違うんだな

285:デフォルトの名無しさん
13/03/20 09:36:22.71
>>282
レジスタってGPUのレジスタじゃないぞ。
抵抗のレジスタだ。

しかし、興味深いなあ。

286:デフォルトの名無しさん
13/03/20 09:59:06.90
>>285
半田班ってんだからわかってくれよ
>>284
携帯だったんだよ興奮してたんだよ

287:デフォルトの名無しさん
13/03/20 10:07:36.15
チップ自体は同じだから理論上可能というのは分かるんだけど、
よもやこんな方法でイネーブルにできるとはw

288:デフォルトの名無しさん
13/03/20 10:07:47.70
>>283
俺が知りたいよ

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

289:デフォルトの名無しさん
13/03/20 10:08:50.58
>>287
しかも690のがクロック高いから性能いいらしいぞ(;´Д`)
買ってみてやってみて

290:デフォルトの名無しさん
13/03/20 10:13:47.22
>>289
マジかよw

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

291:デフォルトの名無しさん
13/03/20 10:35:15.50
マジレスすると、quadroやteslaでの構成と同等なものがないから、
ドライバ側が柔軟な作りに立っていないと不具合が出るだろうな。
例えばハードウェアのモニタリング用の石が載ってなかったりするだろうから。

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

292:デフォルトの名無しさん
13/03/20 12:36:37.99
>>291
ただそれはドライバを更新しなければいいだけの話だよね。

293:デフォルトの名無しさん
13/03/20 13:02:33.89
>>292
しばらくは放置なんじゃなかな。
さすがにドライバだけでは対応しづらい気がする。
基板のリビジョンが変更になった時から使えなくなると思う。

294:デフォルトの名無しさん
13/03/20 15:20:27.90
K10化より、TITANをK20Xにしたい。

295:デフォルトの名無しさん
13/03/20 16:16:11.67
おっ
また発見されたのかよ
むかしクワドロ化の方法が発見されたのに、まだまともな対策してなかったのか

296:デフォルトの名無しさん
13/03/20 18:48:30.94
これ、カーネルドライバがボードに直接プロダクトIDを聞きにいくからカーネルフックじゃ駄目って書いてあるけど
ドライバ改造できればワンチャン有るってことじゃね。

297:デフォルトの名無しさん
13/03/20 20:54:10.54
GeForceをQuadro/Tesla化出来るからって、Quadro/Teslaの売り上げが
減るとは思えない。

298:デフォルトの名無しさん
13/03/20 22:40:19.86
tesla化するって言っても、gk110じゃないしな。gpgpuやってる人からすれば、
そんなにありがたみがないんじゃね

299:デフォルトの名無しさん
13/03/24 16:59:47.30
NVIDIA CUDA Gets Python Support | techPowerUp
URLリンク(www.techpowerup.com)

300:デフォルトの名無しさん
13/03/24 19:08:27.70
これネイティブ対応ってことなのかね。
パイソンってそんなに人気あったんだ・・・
ruby人気は国内限定?

301:デフォルトの名無しさん
13/03/24 19:23:07.90
うん

302:デフォルトの名無しさん
13/03/24 19:23:52.69
Visual StudioもPython使えるようになるパッチが出てるし
OpenCVもPython対応(一応)してるしLibreOfficeもマクロでPython使えるし
そしてCUDAまでも…ということで俺の身の周りではPython株が右肩上がり。

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

303:デフォルトの名無しさん
13/03/27 13:42:12.44
SL#(えすえるしゃーぷ)とは、GPUで実行されるプログラマブルシェーダーを、超高級言語である
C#で書けてしまうという夢のようなオープンソースのフレームワークである。
URLリンク(monobook.org)

304:デフォルトの名無しさん
13/03/27 19:32:08.19
C#イラネ

305:デフォルトの名無しさん
13/03/30 17:24:17.17
DXエラーだったのでdirectXかと思いここで聞きます
色々ググったのですが解決に至らず

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

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

306:デフォルトの名無しさん
13/03/30 23:26:35.66
スレチだし
ゲームの事は開発元に聞け
あとDX8のレンダリングにすら対応してないクソPCを捨てろ

307:デフォルトの名無しさん
13/04/02 08:57:50.31
なんでDirectXでこのスレなんだ?

308:デフォルトの名無しさん
13/04/05 14:24:31.80
本出るの久しぶりじゃない?
URLリンク(www.amazon.co.jp)

309:デフォルトの名無しさん
13/04/05 20:17:57.08
>>308
おお~、CUDA5本か!
注目だな~。

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

URLリンク(next.rikunabi.com)
つか、この人、すごい人なんだな!
原作やってたときは東大生のときだったのか・・・。

310:デフォルトの名無しさん
13/04/05 22:20:39.03
宣伝乙

311:デフォルトの名無しさん
13/04/06 08:59:19.78
もしかして、東大の初代grapeに関わった人?

312:デフォルトの名無しさん
13/04/06 09:40:17.13
そだよ!

313:デフォルトの名無しさん
13/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:デフォルトの名無しさん
13/04/12 17:39:21.97
>>157
書籍「CUDA BY EXAMPLE」の「5.3.2 ドット積の(誤った)最適化」の項には、
真逆な事が書いてあって悩んでいる。

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

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

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

315:デフォルトの名無しさん
13/04/14 00:10:01.63
>>314
多分「CUDA BY EXAMPLE」が間違っているか、CUDAのバージョンが新しくなったことによって変わったのかもしれないが、基本的に__syncthreads()はスレッド間のフェンスみたいなものだから。SIMDの概念で考えればわかると思う。

316:デフォルトの名無しさん
13/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:デフォルトの名無しさん
13/04/14 02:44:35.83
要は、_syncthreadsはwarp内の一つでも実行されたら
そのwarpの全てのthreadが実行したのと同じと看做される。
だから同一block内の全てのwarpそれぞれで、
いずれかのthreadがアクティブな状態で_syncthreadを実行すれば辻褄はあう。

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

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

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

320:デフォルトの名無しさん
13/04/24 11:40:18.49
マルチgpu環境での使い方が分からず困っています。
何か良い資料は有りませんでしょうか。

321:デフォルトの名無しさん
13/04/25 22:41:10.57
ハードウェアとCUDAのバージョン依存があるんでね?

322:デフォルトの名無しさん
13/04/26 06:03:02.19
>>320
URLリンク(www.meriken2ch.com)
マルチGPUをサポートしてて、ソースも付いてるよ。

323:デフォルトの名無しさん
13/04/26 08:14:11.43
compute capability 3.0 cuda 5.0 です。

324:デフォルトの名無しさん
13/04/26 09:28:57.17
>>320
1GPU毎に1CPUスレッドを割り当てて、それぞれのスレッドで
cudaSetDevice()を実行すればいいって話じゃないの?

325:デフォルトの名無しさん
13/04/26 14:32:31.09
>>320
デバイスドライバとの関係でメインスレッドでcudaSetDevice()すると巧くないので、
>324のようにcreate_pthread()した先でcudaSetDevice()するよろし。

326:デフォルトの名無しさん
13/04/26 16:55:32.79
cudaSetDevice()を呼び出した後はシングルgpuの場合と同様なんですよね。

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

327:デフォルトの名無しさん
13/04/26 18:34:49.73
二度目以降は無視されるよ。

328:デフォルトの名無しさん
13/04/26 21:44:04.64
各スレッドでcudaSetDevice()を呼び出して各スレッド毎に各gpuを割り当てますが、
複数のスレッドでcudaSetDevice()を呼び出さずにカーネルを実行したらどうなるのでしょうか。
また、複数のスレッドでcudaSetDevice()使って同一のデバイスを指定はできるのでしょうか。

329:デフォルトの名無しさん
13/04/26 22:55:17.00
>>328
何でそんなことしたいのかさっぱりわからん。

330:デフォルトの名無しさん
13/04/26 22:59:08.02
自分で試して確認しようとしていないようだけど、そういう環境がないのか?
何の為に学ぼうとしてるのかしらんが、そうやって無限に問い続けるつもりか?

331:デフォルトの名無しさん
13/04/27 03:09:44.84
環境がなくて試せないんです。すみません。

332:デフォルトの名無しさん
13/04/27 04:12:10.93
数千円でボードかえるのに環境がないとか甘えだろ。

333:デフォルトの名無しさん
13/04/27 07:55:25.89
>>326
CUDA C Programming Guide 5.0
3.2.6.2 Device Selection

334:デフォルトの名無しさん
13/04/28 17:13:33.38
環境がなくて試せないのに、なんでそんなに重箱の隅みたいなこと聞くのか

335:デフォルトの名無しさん
13/04/28 17:49:38.30
環境の構築(購入)に踏み切ろうか悩んでいるとか

336:デフォルトの名無しさん
13/04/29 18:00:10.69
マルチGPU環境とかこだわらなければ1万円以下でできるのにな。

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

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

339:デフォルトの名無しさん
13/05/05 13:49:16.71
可能ならディスプレイ表示用(ショボいので十分)とコンピューティング用GPUと2つ接続したいところだね。

340:デフォルトの名無しさん
13/05/05 14:00:03.24
最近はCPUオンチップGPU+CUDA用GPUという構成を組みやすい。

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

342:デフォルトの名無しさん
13/05/05 16:15:31.63
タイムアウトが最初の手段だと思うけど。。
それで不便だったらお金のかかる対策に移ればいいわけだし。

343:デフォルトの名無しさん
13/05/05 16:26:07.35
GT610なんて使っている時点でCUDAの意味がないがな。
スレッドブロック数のチューニングにすらならん。

344:デフォルトの名無しさん
13/05/05 23:29:26.34
>>340
今はiGPU搭載CPUが普通で、それがCUDAお手軽構成って感じだよね

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

345:デフォルトの名無しさん
13/05/06 20:16:35.75
linuxでやれよハゲ

346:デフォルトの名無しさん
13/05/07 07:37:55.74
>>344
iGPUでCUDAができるかよ。やるならOpenCLだ。
AMDのCPUがi7並であるわけ無いだろ。
スレ違い。

347:デフォルトの名無しさん
13/05/07 20:27:42.82
GPGPUはHSAでAMDが主流になるって感じになってきたよな
一部のスパーコンピューティングでCUDAがつかわれ、それ以外はHSAのAMDになるのかな

348:デフォルトの名無しさん
13/05/08 22:17:31.73
>>347
性能が桁違いだから仕方がない。自作のアプリではTITANは
7970の6割ぐらいの性能だし。CUDAとOpenCLって
やってることはほとんど同じだから、移行するのもそんなに
難しくなかった。

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

350:デフォルトの名無しさん
13/05/09 07:45:11.56
()

351:デフォルトの名無しさん
13/05/09 21:33:31.18
>>349
お、GCNはだいぶいいんだね。

352:デフォルトの名無しさん
13/05/16 00:41:04.04
teslaでも使える温度計測ソフトはないものか

353:デフォルトの名無しさん
13/05/19 14:04:36.33
>>352
何の温度計測なんだ?
ボードの温度なのか、何かの測定器に連動するものなのかで答えが違うと思うが。

354:デフォルトの名無しさん
13/05/26 13:54:06.19
チクショウ、GTX780は倍精度切られているらしい。
Titanに逝ってくる。

355:デフォルトの名無しさん
13/05/30 23:33:46.73
またどっかの抵抗取り替えたりするとTITAN化するような仕掛けが潜んでたりして

356:デフォルトの名無しさん
13/05/31 00:42:40.83
TITANをK20xにしたいんだけど

357:デフォルトの名無しさん
13/06/01 17:17:27.06
TITANにECCつけたもの(SMXを一つくらい品質保証のため削るかもしれないが)をK20の後継として出るんじゃないかな?
今年のSCあたりで発表されそうだ

358:デフォルトの名無しさん
13/06/01 18:12:25.50
ゲフォがFermiが倍精度1/4でKepler1/8で結局性能悪化って酷くない?
ラデへの乗り換えを検討したほうがいいかな?

359:デフォルトの名無しさん
13/06/01 18:38:27.12
CUDA使ったフリーウェアを作ってるのに、
げふぉに将来性がなくなったら詰む。

360:デフォルトの名無しさん
13/06/01 19:05:28.96
>>359
CUDA で計算している部分だけ OpenCL に変えれば済む話では?
どうせモジュール化してるんでしょ

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

スレ違いスマソ。

362:デフォルトの名無しさん
13/06/01 21:06:59.29
>>361
その辺の話は全て入門書、あるいは入門ページに載っている

363:359
13/06/01 22:07:59.04
>>362
さんクス

364:デフォルトの名無しさん
13/06/02 23:02:35.27
>>360
カーネルだけなら簡単だけど、メモリ周りとか考えると簡単にはいかないよ。

365:デフォルトの名無しさん
13/06/03 01:40:28.77
>>364
CUDAで書く時はグローバルメモリへのコアレスアクセスとか
シェアードメモリ使ってチューニングしてるけど、
OpenCLはそのへん配慮できるのかな。

366:デフォルトの名無しさん
13/06/03 04:03:55.61
>>365
普通にできるよ。

367:デフォルトの名無しさん
13/06/03 07:03:45.92
>>364
それは「げふぉに将来性がなくなったら詰む」ほどの事か?

368:デフォルトの名無しさん
13/06/03 20:38:07.29
結局、アーキテクチャ毎に最適なコードを書かなければならない。
そのように書かず、可搬性を意識しすぎるとGPU(などの並列プロセッサ)を使う旨みがない。
なんというジレンマ。

369:デフォルトの名無しさん
13/06/03 21:37:12.36
つまり、コンパイラがまだまだ馬鹿ということか

370:デフォルトの名無しさん
13/06/05 06:41:20.54
teslak20って一番安い所ではいくらで買えるんだろ?
ヤフオクで三年間分有効の保証が残ってる奴が30万なんだが安いのかな?

371:デフォルトの名無しさん
13/06/05 08:49:33.35
URLリンク(news.nvidia.com)
nvからkayla(armようcuda開発機)のおすすめメールが来た

tegra3とGPU別w

372:デフォルトの名無しさん
13/06/07 08:47:44.96
>>370
titanでいいやん。
ヤフオクで買う程度なら個人利用なんだろ?
会えてteslaを買う理由がないわ。

373:デフォルトの名無しさん
13/06/07 09:27:17.14
マルチGPUするなら、GPUDirectの有無は大きいよ>TeslaとTitan

374:デフォルトの名無しさん
13/06/08 09:03:33.73
titanでもGPUDirectはついてるよ。ないのはRDMAの機能の方だ。

375:デフォルトの名無しさん
13/06/10 13:22:19.65
CUDA5.5はどうですか

376:デフォルトの名無しさん
13/06/10 14:29:57.78
いいえ、どうではありません。

377:デフォルトの名無しさん
13/06/11 09:31:22.59
MPIプログラムのプロファイルを取れるようになったのはでかいな。
URLリンク(developer.nvidia.com)

378:デフォルトの名無しさん
13/06/15 23:05:03.06
cuda-gdbのステップ実行で、
コンソールが返ってこなくなる事がありますが、
原因を確認する方法はありますか?
(普通のgdbも使った経験が無いのですが。。。)

379:デフォルトの名無しさん
13/06/17 15:41:11.18
海外amazonでGK208なGT640かったら
やっぱSM35だった

380:デフォルトの名無しさん
13/06/18 05:37:29.42
TOP500でxeon-phiがトップなたぞ

381:デフォルトの名無しさん
13/06/18 12:06:25.31
正直微妙だ

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

382:デフォルトの名無しさん
13/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:デフォルトの名無しさん
13/06/19 09:00:43.77
>Detected Compute SM 3.5 hardware with 2 multi-processors

SMX2器ってこと?

384:デフォルトの名無しさん
13/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:デフォルトの名無しさん
13/06/19 22:01:10.70
Dynamic Parallelismも使えるの?

386:デフォルトの名無しさん
13/06/19 22:38:31.81
バス幅とL2キャッシュ典

387:デフォルトの名無しさん
13/06/20 10:08:10.18
>385
hyperQとDynamic Parallelismは使えるようだね>382

388:デフォルトの名無しさん
13/06/20 11:17:52.11
GK110でL2が1536KBだから512KBって結構でかいな

389:デフォルトの名無しさん
13/06/22 09:56:35.61
GK208きた

390:デフォルトの名無しさん
13/06/24 14:06:27.76
cudaGetDevicePropertiesで
maxThreadsPerBlockが1024とでたので、
kernel<<<32, 1024>>>()
とやったらKeplerはおKでFermiではだめだった。
基本的にはKeplerでも1ブロックあたり512スレッドが上限だったっけ?

391:デフォルトの名無しさん
13/06/30 20:50:09.84
使えるようにするまで2時間かかった

392:デフォルトの名無しさん
13/06/30 20:58:07.76
CUDAは導入がメンドイね。
HSAはよ。

393:デフォルトの名無しさん
13/07/01 NY:AN:NY.AN
>>392
Linuxだとホントしんどいよね。
何が悲しゅうてドライバインストールのために
カーネルを再コンパイルせにゃならんのか。

394:デフォルトの名無しさん
13/07/01 NY:AN:NY.AN
Linuxでの導入はなおさら OpenCL の方が楽だ

プログラムはめんどいが

395:やんやん ◆yanyan72E.
13/07/01 NY:AN:NY.AN
Linuxのカーネルのコンパイルといってもモジュールを
コンパイルするだけなのだけれど、それってそんなに面倒?
リブートも必要ないし。

396:デフォルトの名無しさん
13/07/02 NY:AN:NY.AN
GUIを使わないモードでリブートして、
カーネルコンパイルして、
もう一回リブートしないとならないと思ったけど。

397:やんやん ◆yanyan5.Xudd
13/07/02 NY:AN:NY.AN
GUIのサービス止めてモジュールをrmmodして
nVidiaが提供するスクリプトを
使ってモジュールだけをコンパイルして
insmodしてからGUIのサービス再開するだけだよ。

398:デフォルトの名無しさん
13/07/02 NY:AN:NY.AN
Bumblebee はどうなの?

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

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

399:デフォルトの名無しさん
13/07/03 NY:AN:NY.AN
>>395
適切なOSとGCCのバージョンをそろえるのが面倒。
ちょっとでもNVIDIA推奨環境と違うと
Getting Startedだけ読んででインストールするのは絶対無理
だと思う。

400:デフォルトの名無しさん
13/07/03 NY:AN:NY.AN
>>399
別に面倒でも何でもない。
CUDAならメジャーなディストリならほぼ動く。
面倒なのは単にLinuxの知識がないだけ。
例外は組み込み系で使う場合。
そもそもCUDAの開発やるのに、

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

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

401:デフォルトの名無しさん
13/07/04 NY:AN:NY.AN
>>400
その後 nouveau が邪魔だって言われて
「cuda nouveau install」でググるんですよね

402:デフォルトの名無しさん
13/07/04 NY:AN:NY.AN
多倍長整数の計算におすすめのライブラリとかある?

403:デフォルトの名無しさん
13/07/04 NY:AN:NY.AN
>>400
X79の最新のチップセット使うとUbuntu 12以上じゃないと
動かなくてですね、そいつのデフォのバージョンのgccだと
CUDAが対応しないんですわ。

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

405:デフォルトの名無しさん
13/07/06 NY:AN:NY.AN
>>404
はいはい情弱ですみませんねえ。
みんながみんなGetting Startedだけ読んでインストールできたら
「Ubuntu 12.**でCUDA 5.0入れてみた」系のブログを書く人も読む人いないですよ。
ふーんだ。

406:デフォルトの名無しさん
13/07/06 NY:AN:NY.AN
CUDA 5.0がだめなら、5.5 RCを試せばいいじゃない。

407:デフォルトの名無しさん
13/07/07 NY:AN:NY.AN
cuda5.0のgccは4.6だろ?
それよりnVidiaはFedora16のサポートが切れてることについてどう思ってるんだろう。

408:デフォルトの名無しさん
13/07/07 NY:AN:NY.AN
Ubuntuにせよ、FedoraにせよNVIDIAは最近Linuxに対してあんまりやる気ないな。
リーナスに中指立てられて批判されたからかな?

409:デフォルトの名無しさん
13/07/08 NY:AN:NY.AN
windowsが最高の開発環境だし

410:デフォルトの名無しさん
13/07/13 NY:AN:NY.AN
GPGPUはAMDになってしまったから、Nvはやる気でないだろ

411:デフォルトの名無しさん
13/07/14 NY:AN:NY.AN
はい?

412:デフォルトの名無しさん
13/07/15 NY:AN:NY.AN
適切なスレが分からなかったので、ここで質問します。
今のCUDAはCUDA CとOpenCLでバックエンドが共通になっていると聞きましたが、
今もしくは将来のCUDAで、HSA(Heterogeneous System Architecture)を
共通のバックエンドで動かすことは技術的に可能ですか?

413:デフォルトの名無しさん
13/07/17 NY:AN:NY.AN
nvidiaに聞け
公開資料にない事の予定問われても
スレの住人はnvidia関係者な訳じゃないし
関係者が居たとしても、2chで非公開の予定情報の可否なんか答える訳ないだろ

414:デフォルトの名無しさん
13/07/17 NY:AN:NY.AN
メジャーになればそれなりの対応もあるだろうが
影も形もないものを・・

415:デフォルトの名無しさん
13/07/17 NY:AN:NY.AN
技術的に可能かどうかと言われれば可能でしょ。
メモリ空間が共通化されれば、GPUの演算器がCPUのSIMD演算器のように扱えるわけだし。
ただCUDAである必要があるかどうかはNVIDIAが判断するんじゃないか?

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


次ページ
最新レス表示
レスジャンプ
類似スレ一覧
スレッドの検索
話題のニュース
おまかせリスト
オプション
しおりを挟む
スレッドに書込
スレッドの一覧
暇つぶし2ch