【GPGPU】NVIDIA CUDA質問スレッド

1 名前：デフォルトの名無しさん mailto:sage [2007/09/17(月) 14:54:28 ] 本家 developer.nvidia.com/object/cuda.html 152 名前：142 mailto:sage [2008/02/17(日) 15:08:14 ] ブロック数を1つから2つに増やして、スレッド数を100から50に減らして実行してみたのですが、 速くなるどころか逆に2倍くらい遅くなってしまいました。そういうもんなのでしょうか？ dim3  grid( 1, 1, 1); dim3  threads(100, 1, 1); ↓ dim3  grid( 2, 1, 1); dim3  threads(50, 1, 1); 153 名前：デフォルトの名無しさん mailto:sage [2008/02/17(日) 15:44:38 ] プログラムにもよると思うけどありえる話だろうね 154 名前：デフォルトの名無しさん mailto:sage [2008/02/17(日) 15:53:10 ] 何スレッドが同時に動くかはGPUによっても違うので一概には言えませんが、 少なくとも8の倍数にはなるのでスレッド数50や100は効率が落ちます。 例えばスレッド数を64にしても50のときと所要時間は変わらないかも知れません。 又、デバイス側の関数呼び出しは完了していても演算自体は終わっていないかも知れません。 cudaThreadSynchronize()してから時間を計った方がいいかも知れません。 155 名前：デフォルトの名無しさん mailto:sage [2008/02/19(火) 17:46:37 ] 研究室のPCにLinux入れてCUDA動かしてますが、 ユーザ権限実行でもOSが死ぬプログラムを余裕で書けてしまいますね これはCUDA使う以上、仕方ないことなんですよねきっと バグのあるデバイスドライバ使ってるのと同じことと、あきらめるしか ないのでしょうか 156 名前：デフォルトの名無しさん mailto:sage [2008/02/19(火) 18:03:53 ] vista対応と自動回復が働いてくれることを渇望しましょう 157 名前：デフォルトの名無しさん mailto:sage [2008/02/19(火) 19:31:36 ] そして危険なプログラムが走らないような検証系に進むのであった。 といってもDirectX動かす権限があればOS死ぬのは容易いような。 158 名前：デフォルトの名無しさん mailto:sage [2008/02/19(火) 19:36:50 ] vista対応マダー？ 159 名前：デフォルトの名無しさん mailto:sage [2008/02/20(水) 05:07:39 ] vistaは要らない子。 *BSD対応マダーー？？ 160 名前：デフォルトの名無しさん mailto:sage [2008/02/20(水) 21:32:15 ] いやむしろTRON対応 161 名前：デフォルトの名無しさん mailto:sage [2008/02/21(木) 13:54:22 ] *BSD対応は確かに欲しいな・・・。 っていうか、オプソにするくらいの気持ちじゃないと この規格は普及せんような・・・。 将来的にnVIDIAのGPUが携帯電話などに使われて行った時に なかなか面倒じゃないかなぁ。 162 名前：デフォルトの名無しさん mailto:sage [2008/02/24(日) 00:15:33 ] NVIDIAはしばらく内部仕様公開しないだろうね。 ptxは読めてもその先の動作がわからないから辛いわ。 AMDには圧倒的に差をつけていそうだし。 GeForce9600出ましたな。 8800はStreamProcessorが128個×1.35GHzだったけど，今度は64個×1.625GHzみたい。 並列度が低い場合でも早くなるっていうこと？ 163 名前：デフォルトの名無しさん mailto:sage [2008/02/24(日) 00:31:43 ] だめっぽい ttp://www.4gamer.net/games/047/G004716/20080220030/ 164 名前：デフォルトの名無しさん mailto:sage [2008/02/24(日) 00:44:42 ] >>157-158 ユーザースペースで動作させた場合の性能が気になりますね。 >>159-161 *BSDやgentoo等への対応も欲しいですね。 欲をいえばVista以外でのHybrid SLIのHybrid Power対応も w 自動切換えとかGeForce Boostとかは無理でも、コマンド一発でオンボとカードを切り替えられるようになるといいのですが。 165 名前：デフォルトの名無しさん mailto:sage [2008/02/24(日) 09:31:22 ] doubleにはいったいいつ対応するのか 166 名前：デフォルトの名無しさん mailto:sage [2008/02/24(日) 09:38:24 ] pc11.2ch.net/test/read.cgi/tech/1188374938/137 167 名前：デフォルトの名無しさん mailto:sage [2008/02/24(日) 14:10:04 ] www.nvidia.com/object/cuda_learn_products.html 9600GT，CUDAも実行できなければ遅い。doubleが使える世代への場繋ぎでしょう。 168 名前：デフォルトの名無しさん mailto:sage [2008/02/24(日) 14:17:49 ] >>167 NVIDIAは伝統的に新しいGPUでCUDAが使えるようになるまでに少なくともドライバを更新してくるから、 今のところ未だ、9600GTでCUDAが使えるようにならないとは言えない。 169 名前：デフォルトの名無しさん mailto:sage [2008/02/24(日) 14:28:20 ] 普通に使えるようになるでしょ。つかヘッダとか弄ったら使えたりして。 で、VC2008+cuda1.1の食べ合わせって駄目なのかよ。1.0はいけるらしいけど orz forums.nvidia.com/index.php?showtopic=53494 170 名前：デフォルトの名無しさん mailto:sage [2008/02/26(火) 07:10:16 ] Vistaのbeta版が３月の終わりごろに出るらしい forums.nvidia.com/index.php?s=70658fb133a8b4a52524f6ec7f2f3408&showtopic=60431 171 名前：デフォルトの名無しさん mailto:sage [2008/02/26(火) 11:05:55 ] >>170 �d 172 名前：デフォルトの名無しさん [2008/03/02(日) 23:09:30 ] ずっと上の方で誰か質問しててまだ回答無いけど、 ゲフォ8400GSでCUDA動きますか？ WinXPでは動いたんだけど、Fedora Linuxでは 認識しなくてエミュレーションに成ってしまった。 何かと金の掛かるWinではCUDAやりたくないです。 173 名前：デフォルトの名無しさん mailto:sage [2008/03/02(日) 23:27:05 ] >>172 一覧にはあるから使えるんじゃないかというレスがあった希ガス。 問題は、ドライバをインストールできるかどうかでそれは端末とBIOS次第かな。 ToolKitはFedora7用のがあるそうだから、逆にそれ以前のはダメかもしれない。 あーでも、エミュレーションでコンパイルできているなら大丈夫か。 コンパイルオプションを指定すれば、ちゃんとGPUコードのモジュールもできるんじゃないかな。 デバイス認識させるのは癖があるから要注意で。 Xが動く状況じゃないと、巧く認識できなくて苦労した記憶がある。 巧く認識できていれば/dev/nvidia0と/dev/nvidiactlができている筈なんで、 それがないならroot権限で正しくコンパイルできたGPUコードのサンプルを動かせば ドライバの.koを使ってデバイスのエントリができる筈。 Linuxのデバイス周りは詳しくないんで、頓珍漢なことを書いているようなら失礼。 174 名前：デフォルトの名無しさん mailto:sage [2008/03/03(月) 10:30:40 ] しかし、VisualStuido2005EEと組み合わせれば無料でできるのに「金の掛かる」とはまた意味不明な…… WindowsならGPUのプロファイリングツールも使えるし、それはそれで悪くない選択だと思うんだけどね。 # といいつつ、未だ自宅で環境整備してないけどw 175 名前：デフォルトの名無しさん mailto:sage [2008/03/03(月) 10:32:18 ] 8600GT　WIN2Kの俺に謝れ！！！ XPが無い… 176 名前：デフォルトの名無しさん mailto:sage [2008/03/03(月) 10:36:43 ] >>175 Win2kで動かないのであれば、WinXP買うのも馬鹿馬鹿しいね。ゴメン、配慮が足りなかった。 177 名前：デフォルトの名無しさん mailto:sage [2008/03/03(月) 18:37:47 ] XPが嫌いってだけのもあるんだけどなー ムダに重たくて 178 名前：デフォルトの名無しさん mailto:sage [2008/03/03(月) 18:50:25 ] >>177 テーマをクラシックにするか、システムプロパティでパフォーマンスを優先にすればWin2Kに較べて重くないと思いますよ。 179 名前：デフォルトの名無しさん mailto:sage [2008/03/03(月) 18:58:12 ] 構ってちゃんの相手するなよ・・・ 180 名前：デフォルトの名無しさん mailto:sage [2008/03/04(火) 02:14:10 ] そもそもWindowsが糞高いし、ソース付いて来ないし。 不自由過ぎてハックする気が萎える。 181 名前：デフォルトの名無しさん mailto:sage [2008/03/04(火) 17:32:33 ] こんちわす。 .NET環境からCUDAとデータの受け渡しをする例、なんてのは 無いもんでしょうかね。既存のVB.NET、C#.NETの統計解析を 高速化致したく。 相関係数＝Σ(xj-xk)^2÷√Σ(xj-xjの平均)^2*(xk-xkの平均) なんてのが多発し、j,kは1〜1000、それぞれ1万点ずつ、x1〜x1000の全ての 組み合わせで計算、という具合です。でも「総和を取る」とかはGPUは 苦手でしょうかね・・・。 182 名前：デフォルトの名無しさん mailto:sage [2008/03/04(火) 17:44:36 ] 総和は苦手なので、適当に分割して小計を求めてから最後に総計を求めることになると思う。 それはさておき、先ずはその相関係数の式自体を変形してみてはどうだろう。 分子も分母のそれぞれも平均値を使わない形にできるはずだが。 # ついでに言えば分母の最後は2乗が抜けている希ガス。 それをやると、平均を使う2パス処理ではなくΣxj, Σxj^2, Σxk, Σxk^2,...などを求める1パスですむ筈。 あ、肝腎のVB/C#からの呼び出し方はわかんね。CUDA側は"*.obj"になるからリンクさえできればなんとかなるのかな? まぁ、計算量が充分大きいならCUDA側を"*.exe"にして起動するようにしても大差ない気もするけど。 183 名前：デフォルトの名無しさん mailto:sage [2008/03/05(水) 12:43:43 ] どうもどすえ。自分なりにお勉強してみました。 .NET←→CUDAは、UnmanagedMemoryStream/共有メモリ/セマフォ、とかで なんとかなりそうです。典型的な処理は以下のようにすればよいのかなと。 ・データは、(128×8)センサー{測定データ64個単位×256｝のように CPU側で整形する。平均0、分散1に正規化し、欠損・異常値の前処理も しておく。これはNオーダーの処理なので十数秒ですむ。 ・GPU側には、C(j,k)＝(A(j,k)−B(j,k))^2　(j,kは0-7) つまり8×8画素の画像の、差の二乗？を繰り返し計算させる。 D(j,k)＝D(j,k)＋C(j,k)で集計する。 センサー(0〜1023番)対センサー(0〜1023番)なんで、N^2オーダー。 今はここに数時間掛かってます。 ・終わったらホスト側でD(j,k)を取り出して、画素を全部足す。 数時間が数分になってくれると嬉しいのですがｗ 「GPU Gem」第三版を注文したので、もうちょっとお勉強してみるす。 184 名前：デフォルトの名無しさん mailto:sage [2008/03/05(水) 13:27:36 ] だから、相関係数を求める式を変形しろって。 CUDA云々以前だぞ。 185 名前：デフォルトの名無しさん mailto:sage [2008/03/05(水) 14:02:20 ] え、相関係数を求める式にはもう「平均」は無いすよ。 ごめんどこ怒られてるのかわかんね・・・ 186 名前：デフォルトの名無しさん mailto:sage [2008/03/05(水) 14:29:04 ] 「相関係数」でぐぐれかす 187 名前：デフォルトの名無しさん mailto:sage [2008/03/05(水) 14:33:41 ] お前ら両方とももういい加減にしろ 188 名前：デフォルトの名無しさん mailto:sage [2008/03/05(水) 16:49:01 ] 統計の基本: 相関係数は普通ピアソンの積率相関係数を指す。 これは、r = Sxy / √(Sxx * Syy)で求められる。 このとき、Sxyはx,yの共分散、Sxxはxの分散、Syyはyの分散。 分散はSxx = Σ(xi - x￣)^2。Syyについても同様。 # ここではxの平均を便宜上x￣で現わす。 この式は、Sxx = Σxx - Σx*Σx/nに変形できる。 このとき、Σxxはxiの平方和、Σxはxiの総和。 また、共分散はSxy = Σ(xi - x￣)*(yi - y￣)。 この式は、Sxy = Σxy - Σx*Σy/nに変形できる。 このとき、Σxyはxiとyiの積の総和。 つまり、相関係数はr = (Σxy - Σx*Σy/n)/√((Σxx - Σx*Σx/n)*(Σyy - Σy*Σy/n))。 従って、線形時間で求められる。また、1回のスキャンで充分なので、各要素を保存する必要がない。 189 名前：デフォルトの名無しさん mailto:sage [2008/03/06(木) 01:21:05 ] 上のほうにもあるけど、総和の計算って実際どーすればいいんだ？ スレッドブロックごとに小計を求めて、各ブロックの結果をまとめるとかいう 方法になりそうなのは予想できるが、実際どんな風に書けばいいのかわからん。 具体的にはベクトルの長さが出したいんだが、各要素の２乗の和がうまく出せない。 ヒントでもいいので分かる方教えてください。 190 名前：デフォルトの名無しさん mailto:sage [2008/03/06(木) 02:29:35 ] 1つのブロックを総和の計算に割り当てる…とか そんなことしか思い浮かばん 191 名前：上の方の人 mailto:sage [2008/03/06(木) 07:58:00 ] 後で暇があったら実際のコードの集計部分だけ晒してみるか。 192 名前：デフォルトの名無しさん mailto:sage [2008/03/06(木) 19:32:29 ] CUDAを使ったjpegの展開ライブラリはありませんか？ 193 名前：デフォルトの名無しさん mailto:sage [2008/03/06(木) 19:55:22 ] IDCTはそこそこ効率よさそうだけど前段のビットストリーム分解が不向きっぽいぞ。 194 名前：デフォルトの名無しさん mailto:sage [2008/03/06(木) 22:20:39 ] >>191 そうしてくれると助かります。 よろしくお願いします。 195 名前：デフォルトの名無しさん mailto:sage [2008/03/06(木) 23:43:14 ] デイビッド・カーク氏来日，CUDAカンファレンス2008開催 ttp://www.4gamer.net/games/032/G003263/20080306051/ 196 名前：デフォルトの名無しさん mailto:sage [2008/03/07(金) 04:38:20 ] やっとできたらしい日本語版のCUDAZONEとプログラミングガイドの日本語版。 ttp://www.nvidia.co.jp/object/cuda_develop_jp.html しかし、機械翻訳そのままなのかな。余りに訳が酷くて原文読まないと意味が判らんとは…… 197 名前：デフォルトの名無しさん [2008/03/07(金) 15:34:38 ] CUT_EXIT(argc, argv)はなんでargc, argvを与えなくちゃいけないのでしょうか？ 198 名前：デフォルトの名無しさん mailto:sage [2008/03/07(金) 15:54:16 ] >>197 引き数にnopromptが指定されていない場合は "Press ENTER to exit..."のメッセージを出して入力待ちにするためです。 詳しくは、cutil.h参照で。 199 名前：デフォルトの名無しさん [2008/03/07(金) 18:23:26 ] floatであった変数をいくつかまとめてfloat4を使うようにしました。 そうしたところプログラムの実行速度が10%ほど落ちたのですが、float4よりfloatの方が 実行効率がよいのでしょうか？ 200 名前：デフォルトの名無しさん mailto:sage [2008/03/07(金) 23:32:28 ] よく知らないが、float4ってSIMD用の型じゃないの？ 201 名前：デフォルトの名無しさん mailto:sage [2008/03/08(土) 00:20:55 ] 単にまとめただけじゃ速くならないと思う。一番いいのは-ptxして出力眺めることなんだけどね。 202 名前：デフォルトの名無しさん [2008/03/08(土) 09:24:11 ] なんで、nvidiaのフォーラムには韓国語はあるのに日本語はないの？ 203 名前：デフォルトの名無しさん mailto:sage [2008/03/08(土) 11:42:50 ] 流石にそれはnvidiaに聞いてくれ。需要が少ないと言うか、要望が少ないのだろ。 半島人と違って日本人は奥床しいからw 204 名前：デフォルトの名無しさん mailto:sage [2008/03/08(土) 11:57:24 ] アメリカの地下鉄とかの案内板や自販機にハングルはあるけど 日本語はないとかいう話を思い出しました w 東大キャンパスであった、CUDAカンファレンス2008の動画とかは公開されないのでしょうかね？ 205 名前：デフォルトの名無しさん mailto:sage [2008/03/08(土) 21:17:16 ] ずいぶん前の話だが、20-26あたりで議論されてる問題って結局どんな風に 並列処理するのが最適なんだ？ 総和の計算の話が出てたから思い出したのだが。 前に試行錯誤したが、高速化できなかったり、値がおかしくなったりして あきらめてしまっていた。 どうぞ皆様知恵をわけてください。 この前のカンファレンスでも言われていたがチューニングにかける時間を 惜しんではいけないのだね。 206 名前：デフォルトの名無しさん mailto:sage [2008/03/08(土) 21:27:59 ] 総和みたいな計算は向いてないとも言ってたけどね 207 名前：デフォルトの名無しさん mailto:sage [2008/03/08(土) 21:59:03 ] >>205 結局>20は、>24で書いたようにインデックスで並列にした。 要はこんな感じ。 -- 大きな配列tmpがあるとき、下図の+で区切られた範囲ごとに集計したい。 tmp+-----+-----+--------+-----+---+... つまり、1番目のセクションは最初の(先頭の)+から次の+まで、2番目のセクションはその次の+まで…… ここで、こんな構造体を考える。 struct sections {short begin, end;} こいつの配列を作って次のように値をセットする。 {{0, 6}, {6, 12}, {12, 21}, {21, 27}, {27, 31}, ...} これをデバイス側に転送しておいて、一つのデータスレッドが一つのセクションを担当する形で集計した。 この方法の問題点： ・セクションのサイズが不均衡なので、ワープ内でも不均衡だと無駄なからループが発生してしまう。 # 1ワープ内では同じインストラクションが走ってしまうため。つまり、使用効率が落ちる状態。 ・セクションサイズがワープ内では極力等しくなるようにソートすると、今度はアクセス場所がランダムになる。 # 先程の配列が、例えば{{0, 6}, {6, 12}, {21, 27}, {301, 307}, ...}のようになってしまう。 いずれにしても、綺麗に並べることができない。 但し、今回は前段のtmp配列への演算結果の格納が所要時間の大半を占めたために適当にチューニングして放置。 尚、予告した集計のロジックは、総和ではなく↑とも違う小計算出だったので条件が違うと言うことでパス。 総計は興味があるので、その内テストできたら改めて晒そうと思う。 208 名前：デフォルトの名無しさん mailto:sage [2008/03/09(日) 13:44:37 ] いきなりですが質問です xp64でVisualStudio2005と8800GTS(G92)がありCUDAドライバー、 ツールキット、SDK、と入れたんですが、結局ビルド方法とか どのPDFに書いてあんだろ、というのがよくワカンネですよ。 VS2005のプロジェクト設定をx64に変えて、cutil32D.libを 拾ってきたZIPから解凍して入れたら、mandelbrotとか一応 動いたけど。他のmandelbrot描画ソフトより劇速なんで、 動いてはいるみたいです。 209 名前：デフォルトの名無しさん mailto:sage [2008/03/09(日) 16:10:18 ] うーん、IDEは使ってないから判らないなぁw pdfじゃなくて、ReleaseNoteか何かにビルド方法かいてなかったっけ? 後で見てみるわ。 210 名前：205 mailto:sage [2008/03/09(日) 23:15:54 ] >>206 確かに言っていた。まぁそれは言われなくとも自明なことなんだが。 ついでに、CUDAがあるからGPUに向いていない処理も容易に書けるようになってしまった とも言っていたな。 だが、膨大なデータをホストに書き戻して、総和を計算し、その結果をまたデバイスに 送って使うということは出来れば避けたかった。 GPUが苦手な処理でもデバイス側でやらせるというのがいいのか、転送コストを覚悟の上で ホストに戻して計算したほうがいいのか実際確かめるために試行錯誤してみている。 >>207 説明ありがとう。 おかげで24でかかれていたことが理解できたよ。 その考え方は思いつかなかったな。 俺ももうちょい考えてみていい感じにできたら晒そうかと思う。 211 名前：デフォルトの名無しさん mailto:sage [2008/03/10(月) 02:56:03 ] ベクトルの総和を出すscanlarge.pdfを見てみると、 elementsが65536ならCPU−GPUでほぼ一緒、 100万elementsでGPUがCPUの5倍速い、ようですね。 なるほどGPU向きではないな・・・ 212 名前：デフォルトの名無しさん mailto:sage [2008/03/10(月) 03:54:26 ] HostがCore2Duo辺りなら未だいいんだけど、Woodcrest辺りだと厳しいんだよねw 超越関数を大量に使う演算だと、GPUが随分有利になるんだけど。 # SSEには超越関数がない→ベクタ化のためには関数テーブルが必要→メモリ消費→キャッシュミス 213 名前：デフォルトの名無しさん mailto:sage [2008/03/10(月) 11:09:00 ] cudaでfloat4*float4ってできないの？ 各要素ごとに4行に分けて書かないとだめ？ 214 名前：デフォルトの名無しさん mailto:sage [2008/03/10(月) 18:33:07 ] >>213 float4 * float4 はないみたいだ。 ちなみに、 -- __device__ float4 a, b, c; __device__ static void func() { c = make_float4(a.x * b.x, a.y * b.y, a.z * b.z, a.w * b.w); } -- こんなコードを書いてみたけど、 -- .global .align 16 .b8 a[16]; .global .align 16 .b8 b[16]; .global .align 16 .b8 c[16]; : : ld.global.v4.f32 {$f1,$f4,$f7,$f10}, [a+0]; // ld.global.v4.f32 {$f2,$f5,$f8,$f11}, [b+0]; // .loc 4 12 0 mul.f32 $f3, $f1, $f2; // mul.f32 $f6, $f4, $f5; // mul.f32 $f9, $f7, $f8; // mul.f32 $f12, $f10, $f11; // st.global.v4.f32 [c+0], {$f3,$f6,$f9,$f12}; // -- こうなった。 215 名前：デフォルトの名無しさん mailto:sage [2008/03/11(火) 19:17:03 ] FreeBSDのLinuxエミュでCUDAを使っているんだけど、 Linuxと比較してどれくらい速度が落ちているのかな? ほぼネィティブと同程度ならこのまま使っていこうかと 思っているんだけど。 216 名前：デフォルトの名無しさん [2008/03/12(水) 05:34:03 ] Windows上でコンパイルしているのですが、nmakeは使えないのでしょうか？ cygwinのmakeコマンドであればコンパイルできるようなのですが、nmakeを使うと NMAKE : fatal error U1077: 'C:\CUDA\bin\nvcc.EXE' : リターン コード '0xc0000005' Stop. となってコンパイルできません。このエラーは何を意味しているのでしょうか？ 217 名前：デフォルトの名無しさん mailto:sage [2008/03/12(水) 05:43:10 ] float4ってsimdで計算してくれないの？ 218 名前：デフォルトの名無しさん mailto:sage [2008/03/12(水) 07:57:11 ] >>217 そもそもsimdなんてありませんが何か。つーか、>214で回答ついてるっしょ。 >>216 nvccが終了ステータスに5を積んでいるようですが、 ちゃんと動いた結果ならnmake側で無視することができると思います。 そもそもちゃんと動いてますか? >>215 4種類ほどボード別に速度を測ったデータがあるからそれでも載せてみましょうか。 219 名前：216 mailto:sage [2008/03/12(水) 08:30:07 ] >>217 早速のご回答ありがとうございます。 Makefileの内容ですが、非常に初歩的な test.exe: test.cu test_kernel.cu     nvcc test.cu です。直接、コマンドプロンプトから打ち込んだ場合、問題なくコンパイルできます。 220 名前：デフォルトの名無しさん mailto:sage [2008/03/12(水) 08:49:24 ] >>219 レス番違いますぜ。 取り敢えず、nvccがエラーステータスを返しているかどうかと、 nmakeでそれを無視できないか調べてみてはいかがでしょう。 221 名前：デフォルトの名無しさん mailto:sage [2008/03/12(水) 12:49:19 ] そういうわけで、各ボードの比較表を作ったので貼ってみる。 --前半 "name","Tesla C870","GeForce 8800 GTX","GeForce 8800 GTS","GeForce 8800 GT","GeForce 8600 GTS","GeForce 8600 GT","使用コマンド" "totalGlobalMem","0x5ffc0000","0x2ff50000","0x27f50000","0x1ffb0000","0xffb0000","0xffb0000","deviceQuery他" "[MiB]",1535.75,767.31,639.31,511.69,255.69,255.69, "sharedMemPerBlock","0x4000","0x4000","0x4000","0x4000","0x4000","0x4000", "[KiB]",16,16,16,16,16,16, "regsPerBlock",8192,8192,8192,8192,8192,8192, "warpSize",32,32,32,32,32,32, "memPitch","0x40000","0x40000","0x40000","0x40000","0x40000","0x40000", "[KiB]",256,256,256,256,256,256, "maxThreadsPerBlock",512,512,512,512,512,512, "maxThreadsDim",512,512,512,512,512,512, ,512,512,512,512,512,512, ,64,64,64,64,64,64, "maxGridSize",65535,65535,65535,65535,65535,65535, ,65535,65535,65535,65535,65535,65535, ,1,1,1,1,1,1, "totalConstMem","0x10000","0x10000","0x10000","0x10000","0x10000","0x10000", "[KiB]",64,64,64,64,64,64, "major:minor",1:00,1:00,1:00,1:01,1:01,1:01, "clockRate","1350000[kHz]","1350000[kHz]","1188000[kHz]","1512000[kHz]","1458000[kHz]","1188000[kHz]", "textureAlignment","0x100","0x100","0x100","0x100","1x100","0x100", 222 名前：デフォルトの名無しさん mailto:sage [2008/03/12(水) 12:49:57 ] --後半 "with Xeon",,,,,,,"使用コマンド" "convolutionFFt2d","145.05[MPix/s]","107.99[MPix/s]","99.57[MPix/s]","107.00[MPix/s]",,,"convolutionFFT2D" "convolutionRowGPU","1100.00[MPix/s]","1071.34[MPix/s]","717.34[MPix/s]","1457.11[MPix/s]",,,"convolutionTexture" "cudaMemcpyToArray","1209.69[MPix/s]","1141.85[MPix/s]","1049.30[MPix/s]","1483.92[MPix/s]",,, "simpleTexture","1180.83[MPix/s]","1125.08[MPix/s]","799.22[MPix/s]","910.22[MPix/s]",,,"simpleTexture" "clockTest",24766,24930,25346,22396,,,"clock" "with Core2Duo",,,,,,,"使用コマンド" "convolutionFFt2d",,,,"106.64[MPix/s]","41.21[MPix/s]","32.93[MPix/s]","convolutionFFT2D" "convolutionRowGPU",,,,"1433.39[MPix/s]","467.43[MPix/s]","375.54[MPix/s]","convolutionTexture" "cudaMemcpyToArray",,,,"1475.25[MPix/s]","893.47[MPix/s]","596.14[MPix/s]", "simpleTexture",,,,"873.01[MPix/s]","370.00[MPix/s]","285.30[MPix/s]","simpleTexture" "clockTest",,,,24066,65006,71736,"clock" 223 名前：デフォルトの名無しさん mailto:sage [2008/03/12(水) 13:58:27 ] これってプログラミングガイドの最後にあるやつ？ 224 名前：デフォルトの名無しさん mailto:sage [2008/03/12(水) 14:29:37 ] CUT_EXITを実行せずにプログラムを終了した場合、どのような悪影響がありますか？ 225 名前：デフォルトの名無しさん mailto:sage [2008/03/12(水) 15:01:43 ] >>221 SharedMemoryって16KBしか割り当てられないのですか？？ 226 名前：デフォルトの名無しさん mailto:sage [2008/03/12(水) 15:10:28 ] 幾ら何でも同じスレの中くらい検索してくれたっていいじゃないか(TT >>223 いいえ、サンプルプログラムの出力です。 >>224 なーーんにも。どうしても気になるならcutil.hを眺めてください。 >>225 この表では、1block辺りのSharedMemoryのサイズを示しています。 227 名前：デフォルトの名無しさん mailto:sage [2008/03/12(水) 18:15:17 ] Cのmallocの場合、NULLかどうかでメモリが確保できたか判別できますが、 cudaMallocでメモリが確保されたかどうかを調べる方法を教えてください 228 名前：デフォルトの名無しさん mailto:sage [2008/03/12(水) 18:31:03 ] CUT_CHECK_ERROR 229 名前：デフォルトの名無しさん mailto:sage [2008/03/12(水) 18:31:16 ] cudaMalloc() cudaError_t cudaMalloc(void** devPtr, size_t count); allocates count bytes of linear memory on the device and returns in *devPtr a pointer to the allocated memory. The allocated memory is suitably aligned for any kind of variable. The memory is not cleared. cudaMalloc() returns cudaErrorMemoryAllocation in case of failure. 230 名前：228 mailto:sage [2008/03/12(水) 18:42:05 ] ああ、あとデバッグモードじゃないとだめだよ 231 名前：デフォルトの名無しさん mailto:sage [2008/03/13(木) 03:10:17 ] warpサイズって簡単に言うと何？ 232 名前：デフォルトの名無しさん mailto:sage [2008/03/13(木) 06:48:58 ] 同一インストラクションが同時に走る、演算ユニットの集まり? 233 名前：デフォルトの名無しさん mailto:sage [2008/03/13(木) 13:50:16 ] CUDA初心者です。 CUDA難しいよ 難しいよCUDA ・<<<Dg, Db, sizeof(float)*num_threads*4>>> て書いたら動かないんすかね。 　int sharemem_size = sizeof(float)*num_threads*4; 　<<<Dg, Db, sharemem_size >>> て書いたら動いた。 ・畳み込み演算させたら、CPUでやるより1.7倍しか速くないの(´・ω・`) ・動いてる間、システム全部固まってる。最初焦った。 CUDAの道険しいﾅﾘ 234 名前：デフォルトの名無しさん mailto:sage [2008/03/13(木) 15:37:14 ] 初歩的な質問かとは思いますが、各ストリームプロセッサにそれぞれ1つずつスレッドが割り当てられるのですか？ 235 名前：デフォルトの名無しさん mailto:sage [2008/03/13(木) 18:22:14 ] __syncthreads()ではブロック内でしか同期できませんが、全ブロックを同期したいときは カーネルを抜けていったんCPU側に戻すしかないのでしょうか？ 236 名前：デフォルトの名無しさん mailto:sage [2008/03/14(金) 02:24:24 ] cudaMallocHostとCのmallocの違いを教えてください 237 名前：デフォルトの名無しさん mailto:sage [2008/03/14(金) 05:10:46 ] >>235 確かカーネルは非同期で実行されたはずだから、 カーネルを連続で呼んだら同期されない・・・と思う。 <<<カーネルの呼び出し>>> cudaThreadSynchronize(); <<<カーネルの呼び出し>>> ならいけるんじゃないの？自信ないけど。 238 名前：デフォルトの名無しさん mailto:sage [2008/03/14(金) 11:06:04 ] >>236 cudaMallocHostでメモリ確保すると、ページングなしスワッピングなし 物理メモリに連続して常駐、のメモリ領域が確保されるはずなのだ！ と思ってcudaMallocHost使ってます。mallocとmlockの合せ技？ 239 名前：デフォルトの名無しさん mailto:sage [2008/03/14(金) 16:02:07 ] Windows, VC2005なんですが、PTXを出すためのオプションはどこに 書けばいいんどすか( ・ω・`) 苦労してPhenomの6倍速いまで持ってきたんですが、さすがに この先はアセンブラ見てみないと何をどうすればいいか分からんどす。 モノは信号処理用の畳み込み(4096)どす。 240 名前：デフォルトの名無しさん mailto:sage [2008/03/14(金) 16:41:32 ] >>239 プロパティでカスタムビルドであることをチェックできない? オプションは言うまでもなく -ptx 241 名前：デフォルトの名無しさん mailto:sage [2008/03/14(金) 16:56:28 ] グラフィックカードを2枚挿した状態で、メインメモリからGPU側に転送する際に どのようにしてGPUを指定すればよいのでしょうか？ カーネルから見られるのは、1枚のGPUだけという認識でよろしいのでしょうか？ 242 名前：デフォルトの名無しさん mailto:sage [2008/03/14(金) 23:48:11 ] >>241 cudaGetDeviceCount()からcudaChooseDevice()あたりでは。プログラミング ガイドのD.1。二枚刺しﾃﾗｳﾗﾔﾏｼｽ >>240 プロジェクトのプロパティのカスタムビルドステップ、等にもコマンドライン 指定が何もなかとです。次はVSディレクトリに潜って.cuのRuleファイルを 探してみるです。 243 名前：デフォルトの名無しさん [2008/03/15(土) 03:32:18 ] RuntimeAPIとDriverAPIの違いを教えてください。 244 名前：デフォルトの名無しさん mailto:sage [2008/03/15(土) 12:17:00 ] CUDA素晴らしすぎです。ようやくグリッドとブロックとスレッドと globalとsharedとconstantを理解して、今「Xeonより100倍速いぜ！」 な速度を叩き出しています。もっとチューニングがんばってみよう。 >>243 RuntimeよりDriverの方が低レベルとして書いてありますけど、 似たようなの一杯有るし、違い良く分からないですよね。 245 名前：デフォルトの名無しさん mailto:sage [2008/03/15(土) 13:20:14 ] 恐らくは、Driverの方はC++で使うことを想定してない。 実際問題として、nvccはC++としてコンパイルするのでRuntimeAPIの方が使いやすいと思う。 246 名前：デフォルトの名無しさん [2008/03/15(土) 17:14:57 ] >>244 segdmm のC800では120GFlops程度と言われていますが 4coreのXeonでも75GFlopsです。 どのくらい違いますか教えてください。 247 名前：244 mailto:sage [2008/03/16(日) 00:22:24 ] >>246 やってみたのが4096単位データ×128個に対して4096のフィルタで コンボリューション掛ける、なんですが、Xeon2GHzの1スレッド だと250M操作/秒位しか出来ないですよ。キャッシュミス連発？ GeForce8800ですと、28G操作/秒で動きます。 248 名前：244 mailto:sage [2008/03/16(日) 04:08:51 ] あ、そんで、Xeonでやるとこうだってのはもれじゃ無くて既存の なので、弄れないし、作りはよく分からんのです。とりあえず 「このルーチン速くするライブラリ作れんか？CUDAとか言う奴で」 と言われてやってみたら100倍速いので怒りが有頂天になった所です。 でも冷静に考えると、CPU側もカリカリに書けばあと10倍位速くなる 可能性はありますわな。たぶんSSEとか使ってないし。 でもそこから10倍にするには、つまりもう何台かパソコン追加しないと ならんわけなので、やっぱりCUDA可愛いよ可愛いよCUDAｗ 249 名前：デフォルトの名無しさん mailto:sage [2008/03/16(日) 08:32:33 ] 私のところじゃ、Xeonの3-10倍だなぁ。単体テストで。 # 他の演算と組み合わせると、1.5-2倍が限度（TT 250 名前：デフォルトの名無しさん mailto:sage [2008/03/16(日) 09:07:29 ] 「Xeon」てみなさんが言ってるのは、DPマシン8コアのこと？それとも1コアのこと？ 1コアに負けるんじゃ確かにちょっと悲しいよね。 251 名前：デフォルトの名無しさん mailto:sage [2008/03/16(日) 09:14:14 ] 1スッドレ、って書いてるじゃん 252 名前：デフォルトの名無しさん mailto:sage [2008/03/16(日) 09:34:49 ] つまりこういう暑苦しい感じが最強だと。 喧嘩してる場合じゃねぇ！Quadコア×デュアルCPU＋8800GTXの2枚刺し、 の友情で最強を目指そうじゃないかお前ら！夕日に向かって海まで走れ！ 敵は巨大行列のLU分解だ！手ごわいぞO(N^3)！

---

---

[ 続きを読む ] / [ 携帯版 ]

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