- 1 名前：デフォルトの名無しさん mailto:sage [2008/03/22(土) 11:13:52 ]
- このスレッドは、他のスレッドでは書き込めない超低レベル、
もしくは質問者自身何が何だが分からない質問を勇気を持って書き込むスレッドです。
CUDA使いが優しくコメントを返しますが、
お礼はCUDAの布教と初心者の救済をお願いします。
CUDA・HomePage
www.nvidia.com/cuda
関連スレ
【GPGPU】NVIDIA CUDA質問スレッド
pc11.2ch.net/test/read.cgi/tech/1190008468/
GPUで汎用コンピューティングを行うスレ
pc11.2ch.net/test/read.cgi/tech/1167989627/
GPGPU#2
pc11.2ch.net/test/read.cgi/tech/1188374938/
- 224 名前：デフォルトの名無しさん mailto:sage [2008/08/25(月) 01:01:51 ]
- 間違った。>>219だ。
- 225 名前：デフォルトの名無しさん mailto:sage [2008/08/25(月) 01:06:43 ]
- >>219
openofficeはいったい何なんだ？マルチプラットフォームのソフトウェアじゃないのか？
- 226 名前：デフォルトの名無しさん [2008/08/25(月) 02:40:58 ]
- お前は仕事でopenofficeにかかわったのか？
違うだろ
派遣先のプラットフォームだろ
お前の会社がマルチプラットフォームの製品を開発してんのか？
違うだろ
派遣先のプラットフォームだろ
現実見ようぜ
- 227 名前：デフォルトの名無しさん [2008/08/25(月) 02:51:05 ]
- >>226
まず、仮定からおかしい。
・マルチプラットフォームのソフトウェアはopenofficeだけではない。
・>>225の仕事の内容を断定する根拠が見当たらない。
・なぜ無数にある可能性から>>226の仮定を導き出したのか不明。
説明をお願いします。
- 228 名前：デフォルトの名無しさん [2008/08/25(月) 02:56:05 ]
- >>227
>・マルチプラットフォームのソフトウェアはopenofficeだけではない。
まず仮定からおかしい。
実際に仕事でマルチプラットフォームを開発しているのかを問題にしている。
>>225がopenofficeを開発したというのなら謝る。おみそれした。
>・>>225の仕事の内容を断定する根拠が見当たらない。
まず自分で開発してそうにないopenofficeを出してくるあたり、
ただのバカ＝派遣ってこと
>・なぜ無数にある可能性から>>226の仮定を導き出したのか不明。
派遣に可能性なんてない
- 229 名前：デフォルトの名無しさん mailto:sage [2008/08/25(月) 03:04:17 ]
- 荒らしは放置の方向で。
- 230 名前：あぼーん [2008/08/25(月) 08:29:00 ]
- 219のほうがバカだと思います。
マルチぷラットフォームの仕事はあると思います。
- 231 名前：デフォルトの名無しさん mailto:sage [2008/08/25(月) 09:16:36 ]
- だからお前はマルチプラットフォームの仕事をしたことあるのかと
- 232 名前：デフォルトの名無しさん mailto:sage [2008/08/25(月) 11:23:15 ]
- サターンとプレステで同じゲームを同時に作ったぞ
同時に作っただけだけど
- 233 名前：デフォルトの名無しさん mailto:sage [2008/08/25(月) 12:41:48 ]
- >>231
Firefoxはマルチプラットフォームじゃないのかね。
- 234 名前：デフォルトの名無しさん mailto:sage [2008/08/25(月) 12:43:22 ]
- マルチプラットフォームの要件はオープンソースではかなり多い。
オープンソース系の企業ならそういう仕事していてもおかしくないね。
- 235 名前：デフォルトの名無しさん mailto:sage [2008/08/25(月) 12:48:10 ]
- 日本語読めないの？
聞かれてるのは「お前は仕事でマルチプラットフォームをやったことがあるのか？」だぞ
- 236 名前：デフォルトの名無しさん mailto:sage [2008/08/25(月) 12:51:51 ]
- >>235
話のすりかえ。
俺自身は学生時代にFirefoxの開発に携わったが、
今の話題は「マルチプラットフォームの仕事があるかないか」。
お前が間違っていることに気づけ。
- 237 名前：デフォルトの名無しさん mailto:sage [2008/08/25(月) 12:53:16 ]
- そもそも、ここの住人が開発に携わったことがあるかどうかなんて誰も興味を持っていないし、
信憑性もない。
もっと有益は議論をすべきだと思うけどね。
- 238 名前：デフォルトの名無しさん mailto:sage [2008/08/25(月) 12:54:24 ]
- スレ違い お前らネットイナゴは佃煮にされて死んでしまえ
- 239 名前：デフォルトの名無しさん [2008/08/25(月) 12:58:38 ]
- 派遣先企業で派遣先企業のプラットフォームに合わせたものを実装するのがおまえら派遣プログラマの仕事だろ
派遣先正社員からマルチプラットフォーム対応なんて指示されてないだろ
くだらないこと考えて余計な工数使うなよ
- 240 名前：デフォルトの名無しさん mailto:sage [2008/08/25(月) 13:13:32 ]
- >派遣先企業で派遣先企業のプラットフォームに合わせたものを実装するのがおまえら派遣プログラマの仕事だろ
派遣先がパッケージベンダってこともあるだろう
なにムキになってるの？
- 241 名前：デフォルトの名無しさん mailto:sage [2008/08/25(月) 13:23:54 ]
- >>239
何故派遣にこだわる？
- 242 名前：デフォルトの名無しさん mailto:sage [2008/08/25(月) 13:49:24 ]
- ID出すとしょうもない素人以下の奴がいろんなスレに粘着してた
- 243 名前：デフォルトの名無しさん mailto:sage [2008/08/25(月) 14:17:22 ]
- >>242
この板はID出ないじゃん
- 244 名前：デフォルトの名無しさん mailto:sage [2008/08/25(月) 14:18:21 ]
- なんだみんな派遣だったのか
- 245 名前：デフォルトの名無しさん mailto:sage [2008/08/25(月) 15:59:39 ]
- 派遣に異常なコンプレックスを持っている人がいるようですが、
池沼の方なので基本放置でおねがいします。
- 246 名前：デフォルトの名無しさん mailto:sage [2008/08/25(月) 16:06:10 ]
- 夏の間だけ2chもアカウント認証式にすればいいと思う。
- 247 名前：デフォルトの名無しさん mailto:sage [2008/08/25(月) 21:59:49 ]
- 正社員にコンプレックスでしょ？
派遣のどこにコンプレックスをもてと
- 248 名前：デフォルトの名無しさん mailto:sage [2008/08/25(月) 23:45:38 ]
- ひどい脱線ぶりだな。
- 249 名前：デフォルトの名無しさん mailto:sage [2008/08/25(月) 23:49:36 ]
- 派遣野郎はスレ違い
- 250 名前：デフォルトの名無しさん mailto:sage [2008/08/26(火) 00:13:24 ]
- >>232
それありそう。PS3 と Xbox360 とか。どう似せるかに力が注がれる・・・
- 251 名前：デフォルトの名無しさん mailto:sage [2008/08/26(火) 00:59:55 ]
- ここはなんのスレなんだ
- 252 名前：デフォルトの名無しさん mailto:sage [2008/08/26(火) 04:09:43 ]
- くだすれ
- 253 名前：デフォルトの名無しさん mailto:sage [2008/08/26(火) 06:23:56 ]
- CUDAはマルチプラットフォーム。
- 254 名前：デフォルトの名無しさん mailto:sage [2008/08/26(火) 07:35:34 ]
- アホかい。
- 255 名前：デフォルトの名無しさん [2008/09/25(木) 14:29:37 ]
- ━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━
■ 今週の名言
──────────────────────────────────────
「HPCは“ハイ・プロダクティビティ・コンピューティング”」
日本SGIのHPC・サービス事業本部で本部長を務める田坂隆明執行役員）
- 256 名前：デフォルトの名無しさん [2008/09/28(日) 22:50:51 ]
- CUDAってSIMD並列といいうことは、過去のスパコン用のコード
移植したら激速ってことでOK?
- 257 名前：デフォルトの名無しさん mailto:sage [2008/09/29(月) 11:43:54 ]
- >>256
SIMDじゃないよ。仮にそうだとしても、そんなに単純な話じゃない。
- 258 名前：デフォルトの名無しさん mailto:sage [2008/09/30(火) 05:16:33 ]
- SIMDだよｗ
- 259 名前：デフォルトの名無しさん mailto:sage [2008/09/30(火) 16:35:15 ]
- SIMDじゃないなら、何なのか問いたい。
- 260 名前：デフォルトの名無しさん mailto:sage [2008/09/30(火) 17:38:27 ]
- SIMDだったりMIMDだったり
- 261 名前：デフォルトの名無しさん mailto:sage [2008/09/30(火) 21:12:21 ]
- 実行はSIMDだけど
コードはスカラだから面倒臭いよ
- 262 名前：デフォルトの名無しさん [2008/10/01(水) 10:37:52 ]
- NVIDIAはSIMTと言っているね。SIMD＋αくらい。
オンボメモリはSX8i並だからやっぱ速いんじゃねーの。
- 263 名前：デフォルトの名無しさん mailto:sage [2008/10/01(水) 13:58:23 ]
- 理屈はいいからなんか作れよ
＞＞過去のスパコン用のコード移植したら激速ってことでOK?
OKの一言で終わることをなにぐだぐだやってんだか
- 264 名前：デフォルトの名無しさん mailto:sage [2008/10/01(水) 18:46:54 ]
- 要するに、「移植できるもんならしてみろ」ってことだろｗ
- 265 名前：デフォルトの名無しさん mailto:sage [2008/10/01(水) 23:50:07 ]
- だね。単純に移植できるとは思わない方がいい。
- 266 名前：デフォルトの名無しさん mailto:sage [2008/10/02(木) 00:27:50 ]
- たしかにブロック切り分けの所もう少しスマートにならないのかと思うのだが
あんなもんデバイスの種類で最適化した組み合わせを自動算出すりゃいいだろ
- 267 名前：デフォルトの名無しさん mailto:sage [2008/10/02(木) 00:29:52 ]
- >>266
CUDA2でストリームプロセッサ数なんかも取得できるようになったから、
是非とも自動算出関数を作ってみてくれ。
- 268 名前：デフォルトの名無しさん mailto:sage [2008/10/02(木) 01:08:48 ]
- いや作るのは簡単なんだけどCUDAってクラスも使えないしラップ出来ないからだめじゃん
- 269 名前：デフォルトの名無しさん mailto:sage [2008/10/02(木) 01:11:17 ]
- あいや出来るのか
使うのなら作るぞ
- 270 名前：デフォルトの名無しさん mailto:sage [2008/10/06(月) 20:37:38 ]
- CUDAに興味があっていろいろ調べてるんだけど
slashdot.jp/hardware/comments.pl?sid=421471&cid=1432609
こんなにメモリアクセスがきついのか…
へたれの俺はCPUで頑張ろうって気になってきた
このあたりの条件が和らぐのはどれくらい後だろう？
- 271 名前：デフォルトの名無しさん mailto:sage [2008/10/06(月) 21:10:43 ]
- それでもビデオカードは一世代速いメモリを積んでいるから、レイテンシを誤魔化せれば、CPUより多少はいいんじゃないかな。
CPUもGPUも、演算能力よりメモリのスピードがネックなんだよねえ。
- 272 名前：デフォルトの名無しさん mailto:sage [2008/10/06(月) 22:18:57 ]
- >>270
そうそう、グローバルメモリはコヒーレントな読み出しなら4クロックなのに
ランダムアクセスすると100倍のクロックを喰ってしまう。
スラッシュドットの記事の書き方はちょっと言葉が足りてなくて、
実際にはいかにメモリをコヒーレントな読み書きで済ませるかが鍵ということになる。
例えば、CUFFTの2Dを1000回回すのと1Dバッチと自前の転置を2回ずつ動かすのとでは、
64x64辺りだと余り変わらないのに128x128になると途端に2Dの方が速くなる。
プロファイラで調べると、CUFFTの内部で用意されている転置関数はコヒーレントな読み書きしかしていない。
一方、自前の転置はコヒーレントな読み出しとインコヒーレントな書き出しになっている。
つまり、1000回呼び出すコストをインコヒーレントな書き出しコストが上回ってしまうということだ。
- 273 名前：デフォルトの名無しさん mailto:sage [2008/10/07(火) 01:53:43 ]
- CPUより遅くなるということはない
電気代の割りに効果が少ないというのはあるｗ
- 274 名前：デフォルトの名無しさん mailto:sage [2008/10/11(土) 06:33:45 ]
- TMPGEncで軽い処理はCPUより遅い。
VRAMとのメモリ転送あるし。
- 275 名前：デフォルトの名無しさん mailto:sage [2008/10/15(水) 22:02:44 ]
- メインメモリと共有するタイプのGPUで処理した方が速い事も状況によってはあるのかな
- 276 名前：デフォルトの名無しさん mailto:sage [2008/10/15(水) 22:26:03 ]
- CUDAでは有り得ないからスレ違い。GPGPUスレにでもどうぞ。
- 277 名前：デフォルトの名無しさん mailto:sage [2008/10/26(日) 18:43:49 ]
- すいません、超低級な質問です。
dim3 threads(128, 1, 1);
dim3 grid(128, 1, 1);
hoge_kernel <<< grid, threads >>> (d_ptr, 128);
とかで関数を呼び出したんですが、ホスト側のスレッド数が128個生成
されるみたいです。これって、GPU内にスレッドが出来るんだと思って
たんですが、違うんでしょうか？
ちなみに、-deviceemu はつけてません。コンソールに以下のように
出るので、GPUにはいってると思います。
Using device 0: GeForce 8600 GT
- 278 名前：デフォルトの名無しさん mailto:sage [2008/10/26(日) 20:58:06 ]
- >>277
＞ホスト側のスレッド数が128個生成されるみたいです。
そう判断した根拠は?
それとは別になるが、128x128が妥当かどうかは検証が必要だと思う。
- 279 名前：デフォルトの名無しさん mailto:sage [2008/10/27(月) 17:11:09 ]
-
質問です。
__device__ void hoge_kernel(void)
{
}
extern "C" __global__ void hoge(void)
{
dim3 threads(16, 1, 1);
dim3 grid(16, 1, 1);
hoge_kernel <<< grid, threads >>> ();
}
上をコンパイルすると、
「error: call can not be configured」
となるんですが、これはどういうことなんでしょうか。
- 280 名前：デフォルトの名無しさん mailto:sage [2008/10/27(月) 17:47:00 ]
- __global__な関数から__device__の関数を<<<>>>で呼ぶことはできない。
- 281 名前：デフォルトの名無しさん mailto:sage [2008/11/01(土) 09:40:52 ]
- ちょっと質問させてくらさい。。
８ストリーミングプロセッサ（SE)単位で構成されている
ストリーミングマルチプロセッサ（SM)でのことなのですが、
グローバル（ローカル）メモリ上へのデータのロード、
ストア命令は100サイクル以上かかってしまいますよね。
その間、8つのストリーミングプロセッサはおお休みしている
のでしょうか？
それとも、他のWarpのインストラクションが割り当てられたり
するのでしょうか。そうであるといいのですが、
そうするとSharedメモリがあとからのWARPにのっとられてしまって
まずそうですよね???
レジスタ的には32bit×2048本あるようなので(280の場合)
OKそうですが、Sharedメモリは16KBytesしかないし。
実際LD、STに160サイクルぐらいくったら、
160命令ぐらい無駄にしちゃって効率をかなり落としてる気がするので、
なんかやです。
24WARPぐらい、バッファリングしていて、割り当てられるWARPを
out of oderで実行するという記事も、どっかで見たような
気がするのですがSharedメモリがじゃまして無理があるような・・・
なんか良い方法があるのかな???
っていうか、なにか根本的な部分で俺の勘違い？
- 282 名前：デフォルトの名無しさん mailto:sage [2008/11/01(土) 10:12:52 ]
- ・グローバルメモリアクセスは、最大400(?)クロック掛かるが、最短では4クロックで済む。
# そのためには、coalescedにアクセスできるように工夫する必要がある。
・各ストリーミングプロセッサは、独立して動作する。Sharedメモリも同様。
例えば、行列の転置のような処理の場合、普通に書くとcoalescedに読んでincoherentに書かざるを得ない。
# 或いはその逆か。
そこで、CUFFT内で行なっている転置処理では、(プロファイルで見る限り)一旦共有メモリにおいて同期を取ることで、
読み書き共にcoalescedアクセスを維持しているようだ。
- 283 名前：デフォルトの名無しさん mailto:sage [2008/11/01(土) 18:40:52 ]
- >>282
早速のレス、さんくす。
なるほど、最短４クロックですか。coalescedにしてもレイテンシー
だからだめかなって思ったけど、よくよく考えると、DRAMへの直接
アクセスに数百クロックってのはおかしいことに気づいた。
各ストリーミングプロセッサは、独立だけど１インストラクション単位で
同期してるんじゃないんかな（んなことぁない？）
・・・っておもったけどブランチがあるから、TPXレベルから見ると独立か???
最初見たときSharedメモリも他のSPのregも０レイテンシで使える
ようだったんで、演算と独立にグローバルメモリとのロード、ストアができん
１６K程度のSharedメモリ、あまり意味ないじゃんとかおもったけど、
そんなことなさそだね。
なかなか面白そうなチップだ。
分岐粒度が32って実際、汎用でどうなんだろ。
ベクトルつかったことないんでわからんが、
適当にセルオートマタでも乗っけて遊んでみまつ。
- 284 名前：デフォルトの名無しさん [2008/11/12(水) 22:42:34 ]
- CUDA-Zなんて便利なものがありました。
forum.nvidia.co.jp
- 285 名前：デフォルトの名無しさん mailto:sage [2008/11/12(水) 22:55:32 ]
- >>284
kwsk
- 286 名前：デフォルトの名無しさん mailto:sage [2008/11/12(水) 23:19:24 ]
- ttp://cuda-z.sourceforge.net/
- 287 名前：デフォルトの名無しさん [2008/11/14(金) 19:39:30 ]
- あげてすまん。
CUDAのためにGTX280つきのPCかったのだが
ほかにもう一枚グラボ買わないといけんのかな？
あとPCIバスしか残ってないのだが・・・
LINUXで使う予定。
- 288 名前：デフォルトの名無しさん mailto:sage [2008/11/14(金) 22:57:10 ]
- >>287
別に一枚でも大丈夫。二枚挿しても(開発には)却って面倒だったり。
まぁ、二枚あれば表示に足引っ張られる心配なくなるから安定はするけどね。
- 289 名前：デフォルトの名無しさん mailto:sage [2008/11/15(土) 00:41:03 ]
- >>288
そうなのか、安心した。
ありが㌧
- 290 名前：デフォルトの名無しさん mailto:sage [2008/11/15(土) 17:39:36 ]
- 質問です。
float にてゴリゴリ計算して、結果を返すプログラムを書いてみたのですが、
普通のアルゴリズム、CUDA+GPU、CUDA+CPU(deviceemu) の2つで比較して
みたところ、思ったより差が大きいのですが、こんなもんですか？
もしくは、機種依存があったりするんでしょうか？
- 291 名前：290 mailto:sage [2008/11/15(土) 17:40:23 ]
- >>290
すいません、間違えました。
「2つで比較して」->「3つで比較して」
- 292 名前：デフォルトの名無しさん mailto:sage [2008/11/15(土) 19:43:01 ]
- 何の差?
演算精度なら、そりゃぁあるさ、floatなんだもの。
- 293 名前：デフォルトの名無しさん mailto:sage [2008/11/15(土) 20:51:43 ]
- ホストは、デフォルトではfloatをfloatで計算してくれないからな。
SSEを使うなりで、ホストがちゃんとfloatで計算すれば、結果は一致するんジャマイカ？
- 294 名前：デフォルトの名無しさん mailto:sage [2008/11/15(土) 20:57:47 ]
- >>293
一致する保証はないよ。CUDAのドキュメントにもあるけれど、超越関数はGPU内部の組み込み版を使うと若干誤差が残る。
いずれにしてもfloatの想定の範囲内だから、実用上は問題にならないけどね。
- 295 名前：,,・´∀｀・,,）っ-●◎○ mailto:sage [2008/11/15(土) 22:26:10 ]
- x87 SSE CellB.E. CUDA の浮動小数サポートの対比表みたいなのがCUDAのマニュアルにあったな。
確かに完全じゃない。
糞まじめに準拠してるのはx87くらいだ
- 296 名前：290 mailto:sage [2008/11/16(日) 02:56:28 ]
- >>292-295
なるほど。出てきた結果が、CPU上の double で計算した結果と、
GPUの float で計算した結果が、最大1%程度違ったから、正直
驚きました。普段doubleしか使って無くて、誤差なんかほとんど
気にしなくてよかったので。
誤差が蓄積してくようなタイプのアルゴリズムではないと思っていた
だけに、少し驚きました。
- 297 名前：デフォルトの名無しさん [2008/11/20(木) 18:53:51 ]
- Cで使っていた自作ライブラリは、
nvccでコンパイルし直さなきゃダメなの？
- 298 名前：,,・´∀｀・,,）っ-○◎- mailto:sage [2008/11/20(木) 18:58:40 ]
- 食べてみたけどおいしかったよ。
一つどうぞ、あーんして
- 299 名前：,,・´∀｀・,,）っ-●◎○ mailto:sage [2008/11/20(木) 19:44:47 ]
- 飾りだ、食うなボケ。
- 300 名前：,,・´∀｀・,,）っ---- mailto:sage [2008/11/20(木) 21:36:00 ]
- 僕のもうないです
>>299から貰ってね
- 301 名前：デフォルトの名無しさん mailto:sage [2008/11/20(木) 23:10:32 ]
- >>297
ホスト(CPU)上で実行させるものならVC++と一緒にリンクできる。
（cppIntegrationサンプル参照）
GPU上で実行させたいならそれなりにいじくらないと無理。
- 302 名前：デフォルトの名無しさん mailto:sage [2008/12/01(月) 15:16:29 ]
- -deviceemuでは動くプログラムが、GPUを使うと誤動作します。
__device__関数に引数が一部しか渡らなくなります(float4のz成分だけしか渡せない)
ループ回数を極端に減らすと改善されるのですが、これはregisterメモリがパンクしてしまったということでしょうか？
- 303 名前：デフォルトの名無しさん mailto:sage [2008/12/01(月) 17:15:56 ]
- ソースを見ないとなんとも言えませんが、ループ回数とregisterは依存関係にはありません。
nvcc -ptxでptxファイルを出力して眺めてみては如何でしょう。
- 304 名前：デフォルトの名無しさん mailto:sage [2008/12/01(月) 23:02:30 ]
- ptxファイルですか。。。
正直どこをどう見たらいいのかわからないので敬遠していましたが、
遂に避けられないとこまで来てしまいましたかねえ。
見るべきポイントとかありましたら、よろしければ教えてください。
- 305 名前：,,・´∀｀・,,）っ-●◎○ mailto:sage [2008/12/01(月) 23:09:56 ]
- ptxもネイティブコードじゃなしに中間コードの断片だからな。
レジスタが何本割り当てられるかとかそういった情報すら持ってないから困る。
その点Larrabeeは16本+16本+αとかレジスタ本数が決まってて
コード生成時点で割り当てが決まってしまう。
結果的に静的な最適化がしやすい。
良し悪し。
- 306 名前：デフォルトの名無しさん [2008/12/02(火) 13:02:50 ]
- ﾘｮｳﾃﾆ●ヽ( ・∀・)ﾉ● CUDA！
- 307 名前：デフォルトの名無しさん mailto:sage [2008/12/03(水) 17:58:11 ]
- .oとかにしてgccでリンクすることってできますか？
- 308 名前：デフォルトの名無しさん mailto:sage [2008/12/04(木) 00:09:51 ]
- >>307
何を? どこで? なんで質問もろくにできないの?
ちなみに、エスパー募集ならお門違い。
- 309 名前：デフォルトの名無しさん mailto:sage [2008/12/04(木) 00:52:30 ]
- >>307
普通にできますよ。
-l/usr/local/cuda/lib -lcudart
等のオプションは当然自分で追加する必要があります。
- 310 名前：デフォルトの名無しさん mailto:sage [2008/12/04(木) 06:13:34 ]
- Windowsで、とかいう話だったりしてね。
- 311 名前：デフォルトの名無しさん mailto:sage [2008/12/04(木) 10:05:39 ]
- .objじゃなくて.oって書いてあるんだからそれはないんじゃない？
- 312 名前：デフォルトの名無しさん mailto:sage [2008/12/04(木) 22:31:14 ]
- それだったらわざわざ聞くかねぇ。まぁいいか。
- 313 名前：デフォルトの名無しさん mailto:sage [2008/12/12(金) 18:47:36 ]
-
【GPGPU】CUDA/ATI STREAM 速度・画質検証スレ
pc11.2ch.net/test/read.cgi/jisaku/1228764782/
- 314 名前：デフォルトの名無しさん mailto:sage [2008/12/14(日) 19:30:29 ]
- これからCUDAを始めようといろいろサイトを巡回しているのですが、
グリッドサイズとブロックサイズはどのようにして決めたらいいのでしょうか?
同期をとるかとらないかで変わると思いますが、とりあえずとらない場合でお願いします。
- 315 名前：デフォルトの名無しさん mailto:sage [2008/12/14(日) 20:41:30 ]
- >>314
BlockSize(スレッド数)は32の倍数で余り大きくない辺り。GridSize(ブロック数)は充分大きく。
- 316 名前：デフォルトの名無しさん mailto:sage [2008/12/14(日) 21:31:49 ]
- >>315
ありがとうございます。
もう一つ質問させていただきたいのですが、たとえば、512x512のようにブロックサイズ、グリッドサイズ
ともに最大値未満の場合、<16, 32, 1>のように複数の次元に分けるのと、<512, 1, 1>のように1つの次元
にまとめるのとどちらがいいのでしょうか?
多分通常はブロックサイズは各次元の最大値<512, 512, 64>をオーバーすることが多いと思うので、こちらは
分けることになると思いますが、グリッドサイズの最大値は<65536, 65536, 1>もあるので、
次元を分けなくてもいいときは分けない方がいいのでしょうか?
- 317 名前：デフォルトの名無しさん mailto:sage [2008/12/14(日) 22:01:51 ]
- threadサイズが実際の分散処理の単位になる
つまりthread数だけ並列演算が行われる
例えばthread(1)にしてblock(10)とかだと順次処理と変わらない
ただしthreadは256までしか出来ないのでblockという単位が容易されてる
blockは単にプログラムがやりやすいように用意されただけのもので
実際に分散処理には影響しない
- 318 名前：デフォルトの名無しさん mailto:sage [2008/12/14(日) 22:15:42 ]
- >>316
言いたいことがよく判らんが、演算量が512x512ならブロックを8x8にしてグリッドを64x64でいいんでない?
或いはxyの区別が重要でないならブロックを64、グリッドを4096とか。
私は後者の方針でptxファイルの出力とプロファイルの結果を睨みながら決定することが多いけど。
# 横着するときは前者だなぁw
- 319 名前：デフォルトの名無しさん mailto:sage [2008/12/14(日) 22:17:30 ]
- >>316
単に内部ループが
for(x=0;x<?;x++){}
_global_
か
for(x=0;x<?;x++){
for(y=0;y<?;y++){
_global_
}
}
になるかの違い
- 320 名前：デフォルトの名無しさん mailto:sage [2008/12/14(日) 22:24:34 ]
- 元の位置を特定する為にglobal関数でインデックスを計算することになるから
もしxやyのどちらかで収まるなら１つだけ使うのが良い
global関数内部の計算コストが一番ネックになる部分だからね
インデックスの計算部分が一番邪魔で削れるだけ削った方が良い
- 321 名前：デフォルトの名無しさん mailto:sage [2008/12/14(日) 22:35:25 ]
- >>317,319
ありがとうございます。ようやく喉のつっかえが取れました。
>>318
試しに昔書いた画像処理のコードを移植しようとしていたのですが、入力によって
いろいろサイズが変わるので、512x512で固定するのはちょっとやりづらいかなと…
ただ、ここまでの話からするとchannel(RGB)*height*widthの三次元配列に画像をつっこんで、
grid = <height, channel, 1>
block = <32, width / 32, 1>
で処理を回せそうですね。画像の幅をを32の倍数に合わせる必要はありそうですが。
>>320
ありがとうございます。実際に使うときは上のような感じにブロックとグリッドの数を決めようかなと
思いますが、どうでしょうか?
- 322 名前：デフォルトの名無しさん mailto:sage [2008/12/14(日) 22:59:47 ]
- heightは超えてもいいの？
範囲は内部でインデックスが範囲かどうか判定するよりは
端数を含めたメモリを確保して必要無い部分も計算させてしまった方が速いよ
- 323 名前：デフォルトの名無しさん mailto:sage [2008/12/14(日) 23:22:38 ]
- >>322
そこまで大きな画像を入れるかどうかは分かりませんが、確かに高さが65536を
超える場合は分割してやる必要がありますね。
>端数を含めたメモリを確保して必要無い部分も計算させてしまった方が速いよ
ということは、配列にコピーするときに幅を合わせてコピーするのではなく、後ろに空白部分を
まとめた方がいいということでしょうか?
#これって多次元配列を引数に渡すのってどうするんだろう…
- 324 名前：デフォルトの名無しさん mailto:sage [2008/12/15(月) 00:01:31 ]
- いや、最終的には全て一次元の配列でそれが適切な間隔でアクセスされるように並んでいるのが一番いいことになる。
だから、画像なんかの場合だと座標値に意味がないなら詰め込んで構わない。
実際には、コンボリューションフィルターなどで隣の画素を参照しないといけないことが多いだろうけど、
その場合はどうせ共有メモリを使ってアクセスを減らすなどの工夫が必要だからどの途32の倍数に拘る必要はない気がする。
|
 |
