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

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/ 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が動かないのは仕様です。 251 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 12:36:54 ] 俺もやってみたらこうなりましたが。 Using device 0: GeForce 9500 GT GPU threads : 512 Processing time GPU: 5.406832 (ms) Processing time CPU: 18.742046 (ms) Test PASSED GPU: 78.0000 87.0000 177.0000 1077.0000 CPU: 78.0000 87.0000 177.0000 1077.0000 Press ENTER to exit... カーネルはこう。 __global__ void testKernel( float* g_idata1, float* g_idata2, float* g_odata, int n) { 　// access thread id 　const unsigned int tid = threadIdx.x; 　// access number of threads in this block 　const unsigned int num_threads = blockDim.x; 　__syncthreads(); 　unsigned int startaddress = n * 1024 * num_threads; 　for (int j = 0; j < num_threads; j++) 　{ 　　for (int k = 0; k < 1024; k = k + num_threads) 　　{ 　　　unsigned int accessAddress = startaddress + k + tid; 　　　g_odata[accessAddress] = (g_idata1[accessAddress] + g_idata2[accessAddress]) / 2.0; 　　} 　　__syncthreads(); 　} } 252 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 12:40:39 ] //ホストでこうして、 for( unsigned int i = 0; i < 1024 * 1024; ++i) { h_idata1[i] = (float) (66.0 + i); h_idata2[i] = (float) (88.0 + i); } // こう実行。 for (int i = 0; i < (int)(1024 / num_threads) ; i++) { testKernel<<< grid, threads >>>( d_idata1, d_idata2, d_odata, i); } //CPUはこう。Athlon 2.3GHz。 computeCPU(float* idata1, float* idata2, float* reference) { for( unsigned int i = 0; i < 1024 * 1024; i++) { reference[i] = (idata1[i] + idata2[i]) / 2.0; } } これじゃサイズが小さすぎてあんまり比較にならないと思うっすよ。 景気よく4096×4096でやってみるといいかな？ 253 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 12:46:40 ] 4096x4096でやってみたらだいぶ差が出てきましたよ。 GPU threads : 2048 Processing time GPU: 33.380219 (ms) Processing time CPU: 260.214355 (ms) 254 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 13:12:34 ] accessAddressの計算おかしくありません？ 255 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 13:34:08 ] コードちゃんと見て無いけど、floatならGPUでパラで動かした方が速いに決まってる athlon系の浮動小数点演算のコストは、整数演算のざっと平均70倍 intel系は知らないにゃぁ 256 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 13:49:01 ] computeCPU(float* idata1, float* idata2, float* reference) { for( unsigned int i = 0; i < 1024 * 1024; i+=4) { static const __m128 div2 = { 0.5f, 0.5f, 0.5f, 0.5f }; __m128 tmp = _mm_load_ps(&idata1[i]); tmp = _mm_add_ps(tmp, _mm_load_ps(&idata2[i])); tmp = _mm_mul_ps(tmp, div2); _mm_store_ps(&reference[i], tmp); } } あとどっかでprefetchnta噛ませるといいかも。 CPU側は最低限SSE使おうや。 大学関係者も含めて比較用のCPU側コードが酷いのが多すぎる。 257 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 13:57:09 ] ぶたぎりですが、 geforce gtx285 と　tesla c1060 って どのくらい違います？ c2050がどうせでるから、c1060と285が 対してかわらないなら、285がいいかと思うのですけど。 ところで www.nvidia.co.jp/object/personal_supercomputing_jp.html いつのまにかnvidiaのページに 「Tesla GPU コンピューティングプロセッサは発売されています: - Tesla C2050/C2070 - Tesla C1060」 こんなことが書いてあるけど、まだ出てないですよね。 258 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 14:25:16 ] >>257 そんなこと言うと、NVIDIAの営業に「そんなアキバ的発想はダメですよ」って言われちゃいますよw 259 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 14:27:55 ] CPU側は、やっぱりCore i7で8スレッド並列とか動かしてあげないとだめじゃね？ そんでGTX285と勝負するみたいな。 260 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 14:29:52 ] 率直に言って、なんで値段が5倍のTeslaが売れるのか良く分からないのです＜アキバ的発想ですが Fermiアーキテクチャも、GeForceが先に出るんじゃないですか？　来年2月くらいでしょうかね。

---

---

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

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