- 1 名前：デフォルトの名無しさん mailto:sage [2009/10/08(木) 19:29:37 ]
- このスレッドは、他のスレッドでは書き込めない超低レベル、
もしくは質問者自身何が何だが分からない質問を勇気を持って書き込むスレッドです。
CUDA使いが優しくコメントを返しますが、
お礼はCUDAの布教と初心者の救済をお願いします。
CUDA・HomePage
www.nvidia.com/cuda
関連スレ
GPUで汎用コンピューティングを行うスレ
pc11.2ch.net/test/read.cgi/tech/1167989627/
GPGPU#3
pc12.2ch.net/test/read.cgi/tech/1237630694/
- 150 名前：デフォルトの名無しさん mailto:sage [2009/11/18(水) 18:43:48 ]
- CUDAでCPU、GPUを並列に動作させられますか？
CPU GPU 並列 CUDA あたりで検索しても出てきません。
原理的にできそうな気もするのですが、
GPUで操作させている間はドライバを働かせているので無理なのでしょうか？
- 151 名前：デフォルトの名無しさん [2009/11/18(水) 18:52:14 ]
- スレッドを使えばできますよ。GPUのプログラムは殆んどCPUに負荷をかけません。
- 152 名前：デフォルトの名無しさん mailto:sage [2009/11/18(水) 22:50:37 ]
- >>150
GPUを起動したスレッドを待機状態のまま放置しておけば、他のスレッドでCPU資源を遣り繰りできます。
但し、SPの個数を超えるようなグリッドを食わせると途中で転送処理が入るので要注意。
- 153 名前：デフォルトの名無しさん [2009/11/19(木) 06:44:52 ]
- CentOSで動きますか？
- 154 名前：デフォルトの名無しさん mailto:sage [2009/11/19(木) 14:51:15 ]
- FermiでC言語が使えるってどーいうこと？
- 155 名前：デフォルトの名無しさん mailto:sage [2009/11/19(木) 17:13:26 ]
- >>153
動きます。
- 156 名前：デフォルトの名無しさん mailto:sage [2009/11/19(木) 19:37:25 ]
- fermi用のCコンパイラを用意しましたってことじゃねーの？
- 157 名前：デフォルトの名無しさん mailto:sage [2009/11/20(金) 19:04:36 ]
- スレッド内部でレジスタがどのように使われているかわかりません
変数の数＝レジスタの数という認識をしていたのですが、そうではないようです
どなたかご教授ください
- 158 名前：デフォルトの名無しさん [2009/11/21(土) 10:30:39 ]
- すいません、テクスチャメモリで疑問があります。
よろしければどなたか教えて頂けませんか？
コードは
gpu.fixstars.com/index.php/%E3%83%86%E3%82%AF%E3%82%B9%E3%83%81%E3%83%A3%E3%83%A6%E3%83%8B%E3%83%83%E3%83%88%E3%82%92%E4%BD%BF%E3%81%86
のテクスチャメモリを参照するていう項目のソースをとりあえず勉強がてらコンパイルしてみました。
配列に値を3ずつ入れていって、バインドして、それにtex1Dでアクセスするというだけのものです
13行目にインデクスに3.14fを用いている事から線形補間なしの場合、小数切捨てで3になり
配列[3]の値を読んで、9が返ってくるのは理解できたんですが、線形補間を有りにすると
9と12の間を補間してるはずなのに7.921875なんていう値が帰ってきています。不思議に思ってv2.3のPrograming guideを読んだところ、P.139に記述があり、どうも与えた値が-0.5されているようです。
ので、はじめから+0.5シフトしてやれば正しい値になりそうですが、そもそも何故-0.5されているのかがわかりません。
検索してもぜんぜん出てこないので当たり前な事なのかもしれませんがドツボにハマってしまってわかりません。よろしければどなたか教えて頂けませんか？長文失礼しました
- 159 名前：デフォルトの名無しさん [2009/11/21(土) 10:32:29 ]
- 158です。
上記質問で訂正があります。programming guideのP.139ではなくP.137でした。すいません。
- 160 名前：デフォルトの名無しさん mailto:sage [2009/11/21(土) 11:49:04 ]
- >>158
テクスチャだからとしか言い様が無いような。
配列の[0]の値は0 〜 1が守備範囲で中心は0.5
...
配列の[i]の値はi 〜 i+1が守備範囲で中心はi+0.5
...
配列の[N-1]の値はN-1 〜 Nが守備範囲で中心はN-0.5
N個の値を0〜Nの範囲に均等にマッピングするためにこうなっている。
- 161 名前：デフォルトの名無しさん mailto:sage [2009/11/21(土) 12:40:17 ]
- >>160
なるほど、こんな感じか。
buffer[2.5] = 6.0
buffer[3.0] = (buffer[3.5]-buffer[2.5])*0.5 + buffer[2.5] = 7.5
buffer[3.5] = 9.0
buffer[4.0] = (buffer[4.5]-buffer[3.5])*0.5 + buffer[3.5] = 10.5
buffer[4.5] = 12.0
buffer[3.141592653589] = (buffer[3.5]-buffer[2.5])*0.641592653589 + buffer[2.5] = 7.924777960767
あれ？ちょっと違う値だね
- 162 名前：158 [2009/11/22(日) 15:50:57 ]
- >>159,160
ありがとうございます。
なんとなく理由が理解できた気がします。
が、こちらでも計算しましたが
buffer[3.141592653589]=7.924777960767
となって件のサイトの値7.921875とは合わないですね・・・
buffer[9.80665]=27.91995となって27.925781ではないですね・・・
低精度の線形補間というのは有効数字3桁って事なんですかね？
- 163 名前：デフォルトの名無しさん mailto:sage [2009/11/23(月) 05:07:25 ]
- リニアフィルタつきのテクスチャでは、
座標(0,0)てのは左上端のドットが持つ四角い領域の左上端を意味するからね。
ドットの真ん中は座標(0.5, 0.5)になる。
真ん中にそのドットの本来の色（値）がくる。
精度はシラネ。リファレンスマニュアルに何か書いてあると思うけど。
- 164 名前：デフォルトの名無しさん mailto:sage [2009/11/23(月) 14:39:37 ]
- CUDA3.2上で動くTesla1070Sを使っているんだが
doubleでは、sqrtとか動作しない？
//#define DOUBLE double
#define DOUBLE float
__global__ void sqrtTest(DOUBLE *A)
{
int x=threadIdx.x;
(*(A+x))=(DOUBLE)sqrt(*(A+x));
}
みたいにして、実験したんだが、doubleにすると
値がそのまんま帰ってくるんだが orz
- 165 名前：デフォルトの名無しさん mailto:sage [2009/11/23(月) 18:16:06 ]
- >164
コンパイル時に -arch=sm_13 はつけてる？
- 166 名前：デフォルトの名無しさん mailto:sage [2009/11/23(月) 19:09:46 ]
- >>165
それだっ！動作しました。
cublasは倍精度で動いたんで、いろいろ調べたんですが。
ありがとうございます。
- 167 名前：デフォルトの名無しさん mailto:sage [2009/11/24(火) 05:01:59 ]
- >>164
CUDAって3．2までいっているの？
つい最近2．3が出たような気がする。
ひょっとして2．3の間違い？
- 168 名前：デフォルトの名無しさん mailto:sage [2009/11/24(火) 12:40:41 ]
- おそらく。
今は3.0のβだね。
- 169 名前：デフォルトの名無しさん [2009/11/25(水) 03:38:40 ]
- newでメモリを確保するのは反則ですか？
- 170 名前：デフォルトの名無しさん mailto:sage [2009/11/25(水) 06:36:25 ]
- >>169
どうやってnewで確保された領域をGPUに転送するの？
- 171 名前：デフォルトの名無しさん mailto:sage [2009/11/25(水) 10:28:10 ]
- >>169
CPU側のメモリでしたら反則ではありません。但し、VCで使う場合は*.cuでnewしてもmsvcrtにリンクできません。
- 172 名前：デフォルトの名無しさん mailto:sage [2009/11/25(水) 19:49:02 ]
- >>169
パフォーマンス求めるなら、論外です。
- 173 名前：デフォルトの名無しさん mailto:sage [2009/11/25(水) 22:00:21 ]
- newでメモリ確保するのがだめならどうやってCPU側のメモリ確保するの?
全部静的確保?
それともcudaHostAllocを使えって話?
- 174 名前：デフォルトの名無しさん mailto:sage [2009/11/26(木) 06:26:18 ]
- メモリ転送ってＳＳＥで高速化されたりするんかな？
だとしたらnewより専用で用意されたものを使ったほうがいいかもね
- 175 名前：デフォルトの名無しさん mailto:sage [2009/11/26(木) 06:58:15 ]
- >>174
メモリ転送ってCPUメモリ間？それともHOSTｰGPU間？
前者だったら高速化はされるけど、後者はDMAでPCIEにダイレクトに転送されるだろうから、
SSEは関係ないんじゃない？
- 176 名前：デフォルトの名無しさん mailto:sage [2009/11/26(木) 17:12:50 ]
- これは・・・
108 名前：デフォルトの名無しさん[sage] 投稿日：2008/07/14(月) 19:24:11
>>106
GPUはWarp単位で同じインストラクションが走るから、要は16人17脚みたいに考えればいい。
メモリアクセスを16人17脚によるパン喰い競争みたいに考えると、自分のパンが目の前にある状態がcoalesced。
そのとき、2,3人パンを喰う必要がなくてもスルーするだけだから問題ない。
処が、二人のパンが入れ違っていたらそこで入れ替える間、みんなが待たされることになるって感じ。
# 判っている人には判るけど、判っていない人には判らない説明だなw
>>107
共有メモリを使うかどうか違うだけだと思うけど。ptx出力させて較べてみたら?
109 名前：デフォルトの名無しさん[sage] 投稿日：2008/07/15(火) 00:26:16
>>108
その説明、いただいてもいいですか？
110 名前：デフォルトの名無しさん[sage] 投稿日：2008/07/15(火) 01:56:27
>>109
本にするならもっと書かせてくれw
Vipのwikiに載せるなら是非やってくれ
金取って講習するのに使うのなら分け前よこせw
- 177 名前：デフォルトの名無しさん mailto:sage [2009/11/26(木) 19:45:48 ]
- 対称行列になるものをGPUに送信したいのだが、うまい方法はない？
一般の場合は、GPU上では対称ではないとみたいのだが
送る際対称になる場合が結構あって、転送時間無駄だなあと。
- 178 名前：,,・´∀｀・,,）っ-○○○ mailto:sage [2009/11/27(金) 00:42:44 ]
- >>175
これじゃないの？
gpu.fixstars.com/index.php/WriteCombine%E3%83%A1%E3%83%A2%E3%83%AA%E3%82%A2%E3%82%AF%E3%82%BB%E3%82%B9%E3%82%92%E9%AB%98%E9%80%9F%E5%8C%96%E3%81%99%E3%82%8B
- 179 名前：,,・´∀｀・,,）っ-○○○ mailto:sage [2009/11/27(金) 00:47:50 ]
- このコードはいろいろ酷いからそのまま使えると思っちゃ駄目よ。
srcのアドレスはこの場合64バイトでアラインメントされてないといけない。
destのほうも最低限16バイトアライン
- 180 名前：デフォルトの名無しさん mailto:sage [2009/11/27(金) 01:21:10 ]
- >>176
なかなか良い説明だなｗ
- 181 名前：デフォルトの名無しさん mailto:sage [2009/11/27(金) 01:24:20 ]
- >>176
あと、競争するコースに分岐があったら、
その分岐に用のある人が２，３人だったとしても全員一応付き合わされた後
本流に戻るというか、そんな感じだな。
- 182 名前：デフォルトの名無しさん mailto:sage [2009/11/27(金) 14:02:37 ]
- >>180,181
あれ、みんな読んでないのかな。
最近ようやくCUDA本がでたわけだけど、
まんまこの文章書いてあるんだよね。
青木先生乙
- 183 名前：デフォルトの名無しさん mailto:sage [2009/11/27(金) 14:41:15 ]
- >>182
絵もついてるしなｗ
- 184 名前：デフォルトの名無しさん mailto:sage [2009/11/27(金) 19:44:26 ]
- 地球シミュレータを蹴散らし一躍脚光を浴びたというのに
中身スカスカなスレだなｗ
- 185 名前：デフォルトの名無しさん mailto:sage [2009/11/28(土) 02:15:17 ]
- CUDAスパコンってそのときだけのものなんじゃない？
研究機関のスパコンは定期的に更新されるもので、更新したら前の代の
ソフトウェア資産はさっぱり使えなくなりましたじゃとても困るだろう。
- 186 名前：,,・´∀｀・,,）っ-○○○ mailto:sage [2009/11/28(土) 03:50:50 ]
- まあG80〜GT200世代のコードも一応次のFermiでは動くし
NVIDIAが父さんしない限りはずっと続くんじゃない？
- 187 名前：デフォルトの名無しさん mailto:sage [2009/11/28(土) 12:16:54 ]
- >>182
もうある程度把握しちゃってるからいまさら入門書買ってもとも
思ったけど、どうなんだろ。
- 188 名前：デフォルトの名無しさん mailto:sage [2009/11/28(土) 12:19:29 ]
- >>185
IBMがHPC向けにはCELLやめてOpenCLに舵を切った的なことが書いてある。
ま、TheInquirerの記事は眉唾ではあるけど。
www.theinquirer.net/inquirer/news/1563659/cell-hpc-material
あ、ここCUDAスレだけどまあ似たようなもんということで。
- 189 名前：,,・´∀｀・,,）っ-○○○ mailto:sage [2009/11/28(土) 14:03:35 ]
- ×CellやめてOpenCL
○Cellベッタリのコード書くのやめてOpenCL
段階的には切り捨てることも考えられるが
いきなりOpenCLで他のデバイスとかいっても、資産が無いじゃん
- 190 名前：デフォルトの名無しさん mailto:sage [2009/11/28(土) 20:47:07 ]
- はじめてのCUDAプログラミング
買った人いる？
どうだった？
- 191 名前：デフォルトの名無しさん mailto:sage [2009/11/28(土) 22:01:41 ]
- 図書館に頼んだら陳列は来月からと言われてしまった
- 192 名前：デフォルトの名無しさん mailto:sage [2009/11/28(土) 23:23:30 ]
- 本なんか読まなくても分かるだろ
- 193 名前：デフォルトの名無しさん mailto:sage [2009/11/29(日) 00:55:57 ]
- 日販は使えない会社だ。
- 194 名前：デフォルトの名無しさん mailto:sage [2009/11/29(日) 01:10:09 ]
- >>190
思ったより安かったのでぽちってきた。
ASCII.techの特集も買ったけどアクセスの最適化あたりで苦戦中なレベルなので、
どうだった？とか言われても答えられないかもしれない。
- 195 名前：デフォルトの名無しさん mailto:sage [2009/11/29(日) 12:11:38 ]
- >>184
これ、とんでもない誤報
浮動小数点演算を理解していないバカコミの馬鹿記事
科学技術立国 日本の恥を世界に晒したもの
「ふざけたやつがペンもつな、馬鹿野郎」だ。
- 196 名前：デフォルトの名無しさん mailto:sage [2009/11/29(日) 14:01:50 ]
- 浮動小数点数
- 197 名前：デフォルトの名無しさん mailto:sage [2009/11/29(日) 15:13:49 ]
- GTX295と電源買った〜
さて、何に使おうw
- 198 名前：デフォルトの名無しさん mailto:sage [2009/11/29(日) 15:40:35 ]
- っBOINC
- 199 名前：デフォルトの名無しさん mailto:sage [2009/11/29(日) 15:46:59 ]
- >>176
おにいさん、もっと！
- 200 名前：197 mailto:sage [2009/11/29(日) 18:00:24 ]
- smokeParticles.exe の動作が 8600GTS より遅く感じるんだがなぜだー
- 201 名前：デフォルトの名無しさん mailto:sage [2009/11/29(日) 18:14:08 ]
- >>200
２個同時使用に対応してないんだろ
- 202 名前：197 mailto:sage [2009/11/29(日) 18:19:17 ]
- >>201
それでも8600GTSよりは早くなるはずじゃない？
あまり詳しくないんで間違ってたらすまん。
- 203 名前：デフォルトの名無しさん mailto:sage [2009/11/29(日) 18:26:30 ]
- >>202
言われて見ればそうだな、約２倍に性能アップしてるはずだし
いつからか分からんけど古いＣＵＤＡと最近のＣＵＤＡのサンプルプロジェクトが入れ替わってるからな
パーティクル関係の数字が増えてるけど同じプログラムでやってみた？
- 204 名前：197 mailto:sage [2009/11/29(日) 18:37:06 ]
- >>203
同じプログラムでやってる。
ベンチマークでもやってみるかな。
- 205 名前：デフォルトの名無しさん mailto:sage [2009/11/29(日) 19:47:47 ]
- 文字列処理をさせてはみたものの
遅すぎて使い物にならねーぞこんちくしょー
- 206 名前：デフォルトの名無しさん mailto:sage [2009/11/29(日) 19:49:56 ]
- >>205
単精度な数値計算に変換すれば良いんでないかい
- 207 名前：デフォルトの名無しさん mailto:sage [2009/11/29(日) 20:06:22 ]
- >>206
UTF-8をどうやって数値計算にすればいいぉか？
- 208 名前：デフォルトの名無しさん mailto:sage [2009/11/29(日) 20:14:33 ]
- UTFの全領域を使う分けじゃなければ、必要な部分だけを数値にマップするとか。
- 209 名前：デフォルトの名無しさん mailto:sage [2009/11/29(日) 20:57:44 ]
- 文字列処理って言ってもいろいろあるだろ
- 210 名前：197 mailto:sage [2009/11/29(日) 22:02:50 ]
- ベンチマークやったら電源が落ちた・・・
780Wじゃ足りないのかな
- 211 名前：デフォルトの名無しさん mailto:sage [2009/11/29(日) 22:28:55 ]
- >>210
よっぽどの詐欺電源でも買ってない限りは、さすがに足りないってことは無いと思うけど
初期不良じゃないか？
- 212 名前：デフォルトの名無しさん mailto:sage [2009/11/29(日) 22:35:29 ]
- 熱落ちじゃないの？
- 213 名前：デフォルトの名無しさん mailto:sage [2009/11/29(日) 22:39:18 ]
- Geforceは80 PLUS シルバー以上の
電源じゃないとまともに高負荷に耐えられないぞ
- 214 名前：197 mailto:sage [2009/11/30(月) 06:48:15 ]
- 80 Plusって書いてないわ・・・
CORAZON ってやつです。
- 215 名前：デフォルトの名無しさん mailto:sage [2009/11/30(月) 08:07:19 ]
- >>214
www.keian.co.jp/products/products_info/kt_780as_sli/kt_780as_sli.html
12Vが25AだなHDDとかいっぱい着いてたら微妙だな
最小構成なら問題ない範囲だと思うけど
てかこれ安いな、これ買えばよかったｗ
- 216 名前：197 mailto:sage [2009/11/30(月) 17:40:28 ]
- 25Aが2つ書いてあるのはどういう意味？
- 217 名前：デフォルトの名無しさん mailto:sage [2009/11/30(月) 20:42:55 ]
- 定員25人のエレベーター2基と、定員50人のエレベーター1基は違うというのは分かるな？
- 218 名前：デフォルトの名無しさん mailto:sage [2009/11/30(月) 20:46:53 ]
- ＨＤＤとかPCIE用の電源が２本電源の中にあるってことだ
ちゃんと分割して接続すれば12x25=300WあるからGTX295のTDP300Wは支えられる
HDDとかつけるともう足りない
ちなみにうちの700W表記の電源は一本36Aある
- 219 名前：197 mailto:sage [2009/11/30(月) 21:20:06 ]
- IDE電源にハードディスクは接続しているけれども、
PCIEの電源は2本ともGTX295に挿してるぜ。
そういう意味じゃない？
ごめん詳しくないんだ。
大人しくこのあたり買っておくかな。
www.scythe.co.jp/images/energia/energia-label0800.jpg
- 220 名前：デフォルトの名無しさん mailto:sage [2009/11/30(月) 21:25:34 ]
- >>219
検索すりゃわかるが今の電源でGTX295を動かしてる人が居るから動くって
ＨＤＤ１個にしてみて周辺機器もはずしてＯＣしてるならデフォルトにして
配線を入れ替えたりしてみてダメなら初期不良だろう
- 221 名前：デフォルトの名無しさん mailto:sage [2009/11/30(月) 21:30:13 ]
- つまり、安いもんじゃないし保障が切れたらもともこもないから
電源が原因だとしても保障があるうちに買った店に持っていって確認してもらったほうがいい
お店の人が電源がだめだと言うなら電源を買えばいいし
お店にあるちゃんとした電源でもダメだったら初期不良で交換してもらえるからね
- 222 名前：197 mailto:sage [2009/11/30(月) 21:50:02 ]
- そうしてみる。
電源も一緒に買ったものだし、持って行ってみるわ。
めっちゃ勉強になった。ありがとう。
- 223 名前：197 mailto:sage [2009/12/01(火) 06:40:06 ]
- 追記
HDを1つにしたら、落ちるまでの時間が長くなりました
- 224 名前：デフォルトの名無しさん mailto:sage [2009/12/01(火) 07:27:11 ]
- 排熱もやばいんだろw
きちんとしたケースとFAN買えw
- 225 名前：デフォルトの名無しさん [2009/12/01(火) 17:12:58 ]
- 多次元配列の領域確保、コピーってcudaMallocとcudaMemcpyでできる？
- 226 名前：デフォルトの名無しさん mailto:sage [2009/12/01(火) 18:48:14 ]
- できる。
てか普通のCみたいにコピーできる。
ちなみにCudaだと多次元配列だと面倒だから1次元配列として扱うことがおおい。
cudaMemcpyAsyncってのもある
- 227 名前：デフォルトの名無しさん mailto:sage [2009/12/01(火) 20:22:02 ]
- 面倒っていうか１次元しか扱えないし
- 228 名前：デフォルトの名無しさん mailto:sage [2009/12/01(火) 21:14:19 ]
- 1次元だけだったか
自分の技量が足りなくてできないのかと思ってた。
- 229 名前：デフォルトの名無しさん mailto:sage [2009/12/02(水) 19:12:31 ]
- 今日からCUDA触ってみたのですが、全然速くない…
device側で1MB×2(dest, src)をアロケートして、hostからデータをコピー
for (int n = 0; n < 1024*1024; n += 512)
{
CUDA_Func<<<1, 512, 0>>>(dest, src, nPos);
}
hostからdeviceへコピー
=====
__global__ static void CUDA_Func(BYTE* pDest, BYTE* pSrc, int nPos)
{
int i = blockIdx.x * blockDim.x + threadIdx.x + nPos;
pDest[i] = ((int)pDest[i] + (int)pSrc[i]) >> 1;
}
=====
なんて事をやっているのですが、CPUの方が速いです
Visual Profilerを見ると、各CUDA_FuncのGPU Time は 8〜9us で終わってますが、CPU Timeが80〜150us になってます
こんなものでしょうか？アドバイス頂けると嬉しいです
Win7/GF8800/SDK 2.3
- 230 名前：デフォルトの名無しさん mailto:sage [2009/12/02(水) 19:56:18 ]
- >>229
CUDA_Func<<<1, 512, 0>>>(dest, src, nPos);を2ｋ回も呼んでるのがまず悪いんじゃね？
あとは詳しい人に任せた。
俺も勉強中。共有メモリのバンクコンフリクトがわけわからねえ。
- 231 名前：デフォルトの名無しさん mailto:sage [2009/12/02(水) 21:07:48 ]
- >>229
そりゃカーネルをキックするコストはかなりでかいから、ループで何回も呼んだらCPUにまけるだろ。。。。
>>230
shared memoryのbank conflictは
ものすごーーーーく大雑把にいうと、thread_id順でshared memoryのアドレスにアクセスすると各バンクのチャネルがぶつからなくて、パラレルに出来るよってお話
- 232 名前：229 mailto:sage [2009/12/02(水) 21:33:41 ]
- レスありがとうございます。
リニアに1次元配列を処理するような事は意味がないと言う事でしょうか？
例えば、ある程度の長さの、サンプリング単位のPCMの演算や、ピクセル単位の画像の演算とか…
- 233 名前：229 mailto:sage [2009/12/02(水) 21:40:38 ]
- 連投スマソ
なんかレスを書いてて、やっとピンと来たんですが、
例えば各スレッドでさらにループで回して、CUDA_Funcを減らせば良い的な話だったりします？
>>229 のコードで言うと
CUDA_Func内で1KB分ループさせて、各スレッドへは1KBのオフセットを渡す。
その分、CUDA_Funcの起動回数を減らす。
違う…？
- 234 名前：初心者 mailto:sage [2009/12/02(水) 21:55:19 ]
- >>299
通常CPUなら、forで何回もやるような処理を
CUDAのカーネルを一発たたくことによって処理させるっていうのが基本的な考え方じゃないの？
あと、メモリは一度になるべく大きくとって転送したほうが効率がいいらしいよん
- 235 名前：デフォルトの名無しさん mailto:sage [2009/12/02(水) 22:17:38 ]
- >>233
そのコードを見た感じ、1024*1024*512スレッドつかってることになってるけど、
何をしてるの？
- 236 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 06:31:18 ]
- いや>>229のやり方は正しいよ
画面描画とＣＵＤＡは同期処理だから大きい単位でやると画面がタイムアウト起こす
これ以上の最適化はＣＵＤＡでは不可能
これで遅いというならそれがそのカードの性能限界だと考えるしかないな
- 237 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 06:39:52 ]
- ちなみに>>229のCPUとカードの具体的な名前と周波数と
PCIEの速度とx16 gen2とかね
遅いって実際にどれくらい遅かったのは知りたいね
上位のクアッドCPUと8400GSなんかじゃ勝負にならないのは当たり前だから
- 238 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 06:59:13 ]
- >>229
ごめん元のソースが何したいのかよく分からないんだけど、こんなことしたいわけ？
1024ピクセル×1024ピクセルの二つの領域の明るさ平均を取るみたいな？
CPUなら、縦座標が外側ループで1024回×横座標が内側ループで1024回、回ると。
//device側で1MB×3(dest, src1, src2)をアロケートして、hostからsrc1,src2にデータをコピー
//512スレッドを起動するなら、外側ループの回る回数はCPUだと1024、GPUだと1024÷スレッド数512
for (int n = 0; n < 1024 / 512; n ++)
{
CUDA_Func<<<1, 512, 0>>>(dest, src1, src2, nPos);
}
//deviceからHostへdestをコピー
=====
__global__ static void CUDA_Func(BYTE* pDest, BYTE* pSrc1, BYTE* pSrc2, int nPos)
{
int i = nPos*512 + threadIdx.x * 1024; //各GPUスレッドが動き出す起点、縦にずれてるわけ
for (k = 0; k < 1024; k++) //各GPUスレッドは呼ばれると横向きに1024回ループする。内側ループをGPUスレッドで処理。
{
pDest[i + k] = ((int)pSrc1[i + k] + (int)pSrc2[i + k]) >> 1;
}
}
=====
ちなみにこれだといちいちグローバルメモリへのアクセスになるんで、
//Sharedにスクラッチコピー
//スクラッチコピー分だけループ処理
//SharedからGlobalに書き出し
するともっと速くなる
- 239 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 07:11:45 ]
- >>238
早けりゃいいってもんじゃないぞ
そんなもん低クラスのカードで動かしたら一発で画面真っ暗だわ
- 240 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 07:16:32 ]
- あとsharedメモリはそんな使い方するもんじゃないだろう
毎回コピーしてたらそのコストの方がデカイわな
- 241 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 07:22:16 ]
- sharedメモリってあれだ
１スレッドでやる計算が複雑な時に頻繁に変数の値を更新するだろ
そういう時にグローバルメモリよりもアクセスが早い一時領域として利用するもんだ
こういう計算自体が単純なケースでは効果はない
- 242 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 07:30:33 ]
- >>239
別人だけど8400GSくらいだとそうなの？経験上何msを超えるとハングする？
ググルとOSにより2秒や5秒でタイムアウトとあるがギリギリまでやるのはまずそうな気はする。
- 243 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 07:38:09 ]
- あ、失礼 coalescedになってなかった。こうかな??
//4回CUDA_Funcを呼び出す方向で。
for (int n = 0; n < 1024 / 256; n ++)
{
CUDA_Func<<<1, 512, 0>>>(dest, src1, src2, n);
}
//nは「縦」の分割数
//512スレッドが連続した512バイトを取り込む。二回動くと、1ピクセル×横に1024ピクセルを処理。
//上に向かって縦256回回る（k)
__global__ static void CUDA_Func(BYTE* pDest, BYTE* pSrc1, BYTE* pSrc2, int n)
{
for (k=0; k < 256; k++)
{
//動き出す起点は各スレッドで1バイトずつずれてる。
int address = n*1024*256 + k*1024 + threadIDx.x;
//1024バイトを512スレッドで処理するので、二回。
pDest[address] = ((int)pSrc1[address]+(int)pSrc2[address])>>1;
pDest[address+512] = ((int)pSrc1[address+512]+(int)pSrc2[address+512])>>1;
}
}
実際書いて動かさないと良く分からないすな。グレーのビットマップ二つ用意してやってみる形かな。
- 244 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 08:17:42 ]
- >>242
タイムアウトが何秒とか議論することですらない
マウスすら動かない状態が２、３秒も続くようなアプリはアプリ失格だろ
- 245 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 09:05:21 ]
- CPUより遅いくらいの8400GSで動かそうとしたなら2,3秒のフリーズ程度なら止む無し。
クラッシュしてデータを失わせるかもしれないリスクを犯すよりは
起動時に簡単なベンチ走らせて遅いGPUはハネちゃうのもありかな。
8400GSを考慮したせいでミドル以上のグラボの足を引っ張るとか馬鹿すぎる。
- 246 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 09:14:35 ]
- プログレスバー表示したら遅くなるから表示しないで画面を固まらせるなんて
そんなものは個人で使うだけにするんだなｗ
- 247 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 10:20:49 ]
- ちょっと待て、みんな一台のGPUであれこれやろうとしているのか?
それじゃ出るスピードも出ないぞ。
- 248 名前：229 mailto:sage [2009/12/03(木) 10:22:20 ]
- 皆様おはようございます。そして、レス感謝です。
朝イチで打ち合わせがあるので、結果だけ取り急ぎ報告します。
前のコード
0 memcpyHtoD 332.928 2155.51
1867.52 memcpyHtoD 332.512 1848.49
3403.26 CUDA_Func 10.624 1158.18
3588.86 CUDA_Func 8.864 119.289
(略)
767008 memcpyDtoH 289.504 997.333
CUDA_Funcでループ
0 memcpyHtoD 332.864 2149.65
1815.04 memcpyHtoD 332.512 1792.27
3264.26 CUDA_Func 11235.1 12351.3
28136.2 memcpyDtoH 286.368 1402.62
満足行く結果ではありませんが、速くはなりました。CPUでリニアに処理した方が速いです。AthlonX2 @1GHz〜3GHz
あと、気づいたのですが、当方の環境ではRDP経由でCUDAが動きませんでした。ちょっとヤバイかも…
詳細は追ってフォローさせて下さい。
- 249 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 10:25:19 ]
- cudaMemcpyは同期を取ってから転送するから、結果の利用のタイミングぎりぎりまで実行を遅らせられれば
見掛け上の処理時間を短縮できるよ。
- 250 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 11:27:12 ]
- >>248
RDP経由でCUDAが動かないのは仕様です。
|
 |
