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

1 名前：デフォルトの名無しさん mailto:sage [2010/12/04(土) 21:57:13 ] このスレッドは、他のスレッドでは書き込めない超低レベル、 もしくは質問者自身何が何だが分からない質問を勇気を持って書き込むスレッドです。 CUDA使いが優しくコメントを返しますが、 お礼はCUDAの布教と初心者の救済をお願いします。 CUDA・HomePage www.nvidia.com/cuda 関連スレ GPUで汎用コンピューティングを行うスレ hibari.2ch.net/test/read.cgi/tech/1167989627/ GPGPU#5 hibari.2ch.net/test/read.cgi/tech/1281876470/ 前スレ 【GPGPU】くだすれCUDAスレ【NVIDIA】 pc12.2ch.net/test/read.cgi/tech/1206152032/ 【GPGPU】くだすれCUDAスレ pert2【NVIDIA】 pc12.2ch.net/test/read.cgi/tech/1254997777/ 【GPGPU】くだすれCUDAスレ pert3【NVIDIA】 hibari.2ch.net/test/read.cgi/tech/1271587710/ 110 名前：デフォルトの名無しさん mailto:sage [2010/12/12(日) 16:01:53 ] 最適化しなければ、物凄い速くなるけど 最適化してしまえばそんなに性能変わらないんじゃね。 111 名前：デフォルトの名無しさん mailto:sage [2010/12/12(日) 16:45:01 ] >>110 C1060に最適化したものをそのままFermiで動かしても差はそんなにないだろうって事だと思うけど それでもコアレスの制限が緩和された影響と自動キャッシュのおかげで結構速くなると思う。 さらにFermiのストリームとかを使用してデータ転送の最適化を行えば段違いになると思う。 特に4枚挿してるということからx8*4の転送速度になって転送がボトルネックになりそうだから そこの最適化は効いてくるんじゃないかと思ってる ところで103のMBはガルフ対応だからPCI-e2.0対応だよね？ 112 名前：デフォルトの名無しさん mailto:sage [2010/12/12(日) 16:55:36 ] ごめんストリームはFermi以前から対応してたみたいね。 113 名前：デフォルトの名無しさん mailto:sage [2010/12/12(日) 17:16:59 ] ストリームは1.3だっけ? C1060って2.0だっけ? 114 名前：デフォルトの名無しさん mailto:sage [2010/12/12(日) 18:35:23 ] >>113 ストリームはCC依存でtoolkit依存のようだが？ 115 名前：デフォルトの名無しさん mailto:sage [2010/12/12(日) 18:36:43 ] >>113 ストリームはCC依存ではなくtoolkit依存のようだが？ 116 名前：デフォルトの名無しさん mailto:sage [2010/12/12(日) 20:48:29 ] メモリバウンドならその話もわかるが 多体問題って演算バウンドだったんじゃないの確か GTX580に交換するうまみはそれほど無い気がする 117 名前：デフォルトの名無しさん mailto:sage [2010/12/12(日) 21:13:38 ] ということはtesla->fermiの性能向上ってメモリ関係（キャッシュ含む）とdoubleだけってこと？ floatなら計算パワーだけで考えると大して性能向上してなかったの？ floatだけならメモリ周りのボトルネックを解消して全体的にFlopsが上がりやすくなりましたって だけだったのか・・・・ 118 名前：デフォルトの名無しさん [2010/12/12(日) 21:34:38 ] 考えるより動けということでGTX580x4枚の見積もり取ってます。 20万＋なら自分の財布でやってしまえ。 ちなみにMathematica8もCUDA対応で年始位にデリバリーとか何とか。 >>104 >MATLAB 2010bのGPUArray R+CUDAの方に感けて，まだやってませんでした。すみません。 119 名前：デフォルトの名無しさん mailto:sage [2010/12/12(日) 21:43:37 ] C1060と580のnbodyのベンチをさがしてみた C1060　443.4GFLOPS 580 447.154 single-precision GFLOP/s at 20 flops per interaction 580は過去スレからc1060はググった結果 なにこれ！？ 120 名前：デフォルトの名無しさん mailto:sage [2010/12/12(日) 21:46:09 ] ごめん↑のC2050だった 121 名前：デフォルトの名無しさん mailto:sage [2010/12/12(日) 21:46:51 ] > Compute 2.0 CUDA device: [GeForce GTX 580] 16384 bodies, total time for 10 iterations: 69.775 ms = 38.472 billion interactions per second = 769.433 single-precision GFLOP/s at 20 flops per interaction 122 名前：デフォルトの名無しさん mailto:sage [2010/12/12(日) 21:58:16 ] >>121 GTX580結構がんばるんだな… 年末のセールでELSA一枚買ってみようかな ↓480での結果（過去スレから引用） [単精度] 15360 bodies, total time for 10 iterations: 73.826 ms = 31.957 billion interactions per second = 639.149 single-precision GFLOP/s at 20 flops per interaction 123 名前：デフォルトの名無しさん [2010/12/12(日) 22:49:40 ] >>118 自分でコード書いたほうが速いし確実だと思うけどね 124 名前：デフォルトの名無しさん mailto:sage [2010/12/12(日) 22:52:47 ] 昔、GeForceをごにょごにょしたらQuadroになったように、 GeForceをごにょごにょしたらdoubleが4倍速になったりしないかな。 125 名前：デフォルトの名無しさん mailto:sage [2010/12/12(日) 22:54:13 ] 最近のはならないようになっている あきらめろん 126 名前：デフォルトの名無しさん mailto:sage [2010/12/12(日) 22:55:44 ] >>123 行列演算で、CUBLASより速いコードかける？ 127 名前：デフォルトの名無しさん mailto:sage [2010/12/12(日) 23:09:54 ] 何でそんな不毛な煽りするかね 128 名前：デフォルトの名無しさん mailto:sage [2010/12/13(月) 22:52:09 ] コンスタントメモリについて教えてください __constant__ int i; int j=1;とした場合(コンスタントメモリに配列以外を置きたい場合)になぜ cudaMemcpyToSymbol("i",&j,sizeof(int),0,cudaMemcpyHostToDevice); と文字列のようにして渡すのでしょうか また、外部ファイルにextern __const__ int i; と書くと定義が二重ですと言われます 複数ファイルにかかれたカーネルからコンスタントメモリを読み込みたいのですがどのようにすればよいのでしょうか お願いいたします 129 名前：デフォルトの名無しさん mailto:sage [2010/12/13(月) 23:24:44 ] >>128 うーん、簡単に言えば、cudaMalloc()でVRAM上のアドレスを得ることができればcudaMemcpy()できる。 一方__constant__の場合はそういうわけにはいかないので、 nvcc独自のちょっとえげつないcudaMemcopyToSymbol()を使わないといけない。 コンパイル単位の分かれたカーネル同士で共有するのは難しいかもねぇ。 130 名前：デフォルトの名無しさん mailto:sage [2010/12/13(月) 23:42:11 ] >>118 > R+CUDAの方に感けて，まだやってませんでした。すみません。 R+CUDAってどういう構成なのですが？ やはり、 R+gputools+CUDA+CULA なのですか？ それとも、Rのビルド時にCUBLASをリンクするのですか？ 当方の環境は、 Q9450 + Memory 6GB + GeForce9600GT Ubuntsu10.04(64bit) です。gotoBLASで行列演算時にコア4つ使用するところまではやったのですが、 CUDA使うともっと凄くなるのですか？ 131 名前：デフォルトの名無しさん mailto:sage [2010/12/13(月) 23:45:44 ] 速くなるかどうかはケースバイケースでアプリの内容によるだろ 132 名前：130 mailto:sage [2010/12/14(火) 23:05:18 ] >>131 >>118さんがどういう構成でやっているのか知りたいです。 未だ、コンパイルも通らない実行も出来ていないレベルの者です。 RにgotoBLAS組み込んでも、乱数の生成速度に差が出ませんでした。 BLASは乱数を生成するライブラリではないから当然ですが。 CUDAのページのサンプルに ・MersenneTwister ・Monte Carlo Option Pricing with multi-GPU support などとあるのをみると期待に胸がふくらみます。 133 名前：デフォルトの名無しさん [2010/12/15(水) 21:41:19 ] 胸は膨らみブラ破け 134 名前：デフォルトの名無しさん mailto:sage [2010/12/15(水) 21:44:27 ] だいたい乱数で時間かかるって相関乱数の相関部分じゃね〜の？ BLASがいくら早くても関係ないじゃん、相関には 135 名前：デフォルトの名無しさん [2010/12/15(水) 21:54:23 ] 乱数の作り方にもよるけど、誤り訂正符号を使った方式では符号化に O(n^2)かかるところを32*8=256ビットの乱数を生成させるのに32個の コアで16回の計算で得られるので効率的になる。 136 名前：デフォルトの名無しさん mailto:sage [2010/12/16(木) 00:26:21 ] ３流プログラマがgotoさんの名前を出して汚すなよ。プライドが少しでもあるならな。 137 名前：デフォルトの名無しさん mailto:sage [2010/12/16(木) 00:29:55 ] ・・・ 138 名前：デフォルトの名無しさん mailto:sage [2010/12/16(木) 07:15:11 ] 後藤blasだろ 釣られた？ 139 名前：デフォルトの名無しさん mailto:sage [2010/12/16(木) 11:51:04 ] 次買うのＡＴＩになりそ・・・ Compute Shaderスレないですか？ 140 名前：デフォルトの名無しさん mailto:sage [2010/12/16(木) 15:56:27 ] ATIやる人はOpenCLかCAL+ILじゃないの？ 141 名前：デフォルトの名無しさん mailto:sage [2010/12/16(木) 17:50:04 ] >>86 ものすごい遅いレスですみませんが、driverAPIなら可能です。 142 名前：デフォルトの名無しさん mailto:sage [2010/12/16(木) 19:24:21 ] >>136 BLASが乱数を生成するライブラリでないと言ったら、 後藤さんの名前を汚していることになるのですか？ 143 名前：デフォルトの名無しさん mailto:sage [2010/12/16(木) 23:05:19 ] Textureメモリってふつうのより早いの？ みんななにに使ってるの？ 144 名前：デフォルトの名無しさん mailto:sage [2010/12/17(金) 09:23:30 ] ググレ 145 名前：デフォルトの名無しさん [2010/12/17(金) 22:12:47 ] GTX580，4枚まとめ買い目指して奔走してますが，ない。 2月くらいにはなんとかならないだろか。 146 名前：デフォルトの名無しさん mailto:sage [2010/12/17(金) 22:51:22 ] CUDA Visual Profiler と Compute Visual Profilerの違いは何でしょうか 前者はバージョン2.3 後者はバージョン3.1 が入っているのですが、toolkitが3.xになってからプロファイラの名称が変わったというだけでしょうか？ 147 名前：デフォルトの名無しさん mailto:sage [2010/12/17(金) 23:04:45 ] でしょ。 148 名前：デフォルトの名無しさん mailto:sage [2010/12/17(金) 23:18:01 ] 質問です ウッカリ長時間のCUDAジョブを走らせてしまったとき 中断はどすればよいですか？ 149 名前：146 mailto:sage [2010/12/17(金) 23:34:33 ] >>147 ありがとうございます ものの本やサイトには/usr/local/cuda/cudaprof/のcudaprofiler使えって書いてあったのですが 変なWarning出て困っていたところでした 150 名前：146 mailto:sage [2010/12/17(金) 23:36:32 ] >>147 ありがとうございます ものの本やサイトには/usr/local/cuda/cudaprof/のcudaprofiler使えって書いてあったのですが 変なWarning出て困っていたところでした 151 名前：デフォルトの名無しさん mailto:sage [2010/12/18(土) 00:03:29 ] >>148 GPUはずしちゃえよ 152 名前：デフォルトの名無しさん mailto:sage [2010/12/18(土) 00:34:36 ] kill すればいいんじゃね？ 153 名前：デフォルトの名無しさん mailto:sage [2010/12/18(土) 00:44:26 ] >>145 九十九ネットストアにあったけど。あと価格コムとかで在庫もってる店あたってみたらいいんじゃない？ 154 名前：デフォルトの名無しさん mailto:sage [2010/12/18(土) 00:47:36 ] 補足、九十九にあったのはエルザのね。 155 名前：デフォルトの名無しさん mailto:sage [2010/12/18(土) 15:04:40 ] Linuxドライバにバグないか？ 最初の実行に時間がかかるのと 同様に最初の実行時にグリッド数、スレッド数が大きめだたとmemcpyでエラーが出る 小さめのグリッド数、スレッド数で最初に実行した後だと 大きめのグリッド数、スレッド数でも正常するんだが… 156 名前：デフォルトの名無しさん mailto:sage [2010/12/18(土) 15:16:14 ] >>155 ドライバのバージョンとSDK、TOOLキット、OSの情報も載せないと駄目じゃね？ あとgccか。 157 名前：デフォルトの名無しさん mailto:sage [2010/12/18(土) 15:28:07 ] OS：CentOS5.5 ドライバ、SDK、ToolKitは最新 ｇｃｃは4.1.2 今のところ原因は全く不明 ドライバのバージョンを一個前にすると初回実行時の遅延がちょっと早くなる（最新が約2.5s、一個前が約1.8ｓ） gtx470と465で試してどちらも同じ症状 158 名前：130 mailto:sage [2010/12/18(土) 16:44:30 ] ようやくgputoolsが動いた。 速いのは分かったけど、 行列の規模が、本体のメモリ容量と関係なく、ビデオカードのメモリ容量に制約される。 キャッシュイーター向けじゃなかったですね。 159 名前：デフォルトの名無しさん mailto:sage [2010/12/18(土) 17:43:36 ] >>158 行列の計算ってのはね、分割できるんですよ。 160 名前：デフォルトの名無しさん mailto:sage [2010/12/18(土) 18:13:31 ] >>159 確かに。 分割できなければコアが複数あっても意味がない。 161 名前：デフォルトの名無しさん mailto:sage [2010/12/18(土) 18:55:39 ] でもね、いったんメインメモリに移して計算するのはもったいない。 結論C2070を買え。 162 名前：130 mailto:sage [2010/12/19(日) 18:30:10 ] 結局、部分行列作成と計算結果の結合で、処理にかかる時間が逆転してしまいました。 一般論として言えるのかどうか知らないけど、 行列が（計算結果を含めて）ビデオカードのメモリにスッポリ入るなら劇的な高速化が見込めるけど、 そうでなければ、マルチコアCPU＋本体メモリのほうが速度でもメモリ消費量でも有利なようです。 163 名前：128 mailto:sage [2010/12/19(日) 20:32:38 ] コンスタントメモリはいつデバイスに確保されるのでしょうか __const__宣言時とは考えにくいので最初にcudaMemcpyToSymbolしたときでしょうか デバイスにあるということはコンスタントメモリも終了時に解放したほうがいいと思うのですが、 cudaFree()では解放できないようです コンスタントメモリはどのように解放すればいいのでしょうか。 164 名前：デフォルトの名無しさん mailto:sage [2010/12/19(日) 21:03:47 ] 行列計算なら転送を隠せるはず リンパックベンチでGPU搭載マシーンが上位に来てるんだし 165 名前：デフォルトの名無しさん mailto:sage [2010/12/19(日) 23:49:30 ] >>163 SDKのサンプルみればいいんじゃないか？ 166 名前：デフォルトの名無しさん mailto:sage [2010/12/20(月) 00:43:41 ] __constant__を気にするなら kernelコードだって似たようなものなんじゃ。 167 名前：デフォルトの名無しさん mailto:sage [2010/12/20(月) 11:19:24 ] >>163 定数メモリはcudaMalloc()のように「確保する」ということなしに使うようになっています。 その為、cudaMemcpyToSymbol()のような特殊な関数でしか転送できません。 また、「確保する」ことなしに使える以上、「解放する」手続きは用意されていません。 168 名前：デフォルトの名無しさん mailto:sage [2010/12/21(火) 15:33:13 ] 安定ソートが欲しいのですが CUDA向きの安定ソートってありますか？ 16000要素くらいです 169 名前：デフォルトの名無しさん mailto:sage [2010/12/21(火) 15:52:44 ] NVIDIAにあるバイトニック使ってみ 170 名前：デフォルトの名無しさん mailto:sage [2010/12/21(火) 15:54:21 ] SDKにmergesortってあるがそれが同じ奴だ 171 名前：168 mailto:sage [2010/12/21(火) 16:18:22 ] バイトニックは安定じゃないようなのですが、 マージソートのサンプルがSDKにあったのですね。 チェックしてみます 172 名前：デフォルトの名無しさん [2010/12/21(火) 20:58:05 ] gtx580を4基積んだ．電源が音を上げた． やれやれ． 173 名前：デフォルトの名無しさん mailto:sage [2010/12/21(火) 21:12:11 ] >>172 何Wで駄目だった？ 174 名前：デフォルトの名無しさん mailto:sage [2010/12/21(火) 21:46:21 ] >>173 1400W。ディスク系を外においてもなぁ。 175 名前：デフォルトの名無しさん mailto:sage [2010/12/21(火) 21:49:20 ] SC748あたりなら・・・ 176 名前：デフォルトの名無しさん mailto:sage [2010/12/21(火) 21:59:47 ] 7046GT-TRFね。 177 名前：デフォルトの名無しさん mailto:sage [2010/12/21(火) 22:17:39 ] 1400Wで駄目だとなると電源2つ積めるケースにして連動させるケーブルで繋いでと いったことをしないと無理そうだな。 178 名前：デフォルトの名無しさん mailto:sage [2010/12/21(火) 23:19:45 ] 無茶しやがって… 179 名前：デフォルトの名無しさん mailto:sage [2010/12/22(水) 01:04:00 ] 1400Wとか言って売っている時点で粗悪電源確定。 >>172が日本在住でないなら話は別だが。 180 名前：デフォルトの名無しさん mailto:sage [2010/12/22(水) 07:14:57 ] >>179 あれは110V/1400W。80GOLDのリタンダント。 日本じゃ実効1100W。 181 名前：デフォルトの名無しさん mailto:sage [2010/12/22(水) 22:34:50 ] 大容量PC電源はエアコン引っこ抜いて200Vから取ろう。 182 名前：デフォルトの名無しさん mailto:sage [2010/12/22(水) 23:25:26 ] >>180 なら、「1100W電源」として売るべきだろ。 排気量2000ccの車を3000ccだと言って売る自動車屋などありえない。 そのあり得ないことを長年やってきたのが、PC電源代理業者だ。 183 名前：デフォルトの名無しさん mailto:sage [2010/12/23(木) 00:59:59 ] >>182 supermicroは日本向けのサイトを持っていない。仕方あるまい。 ちなみに法人で入れたときはこの説明はあった。 www.supermicro.com/products/system/4U/7046/SYS-7046GT-TRF.cfm 184 名前：デフォルトの名無しさん mailto:sage [2010/12/24(金) 20:02:02 ] GPUの分類について教えてください 例えばGTX480は Fermi-ComputeCapability2.0-GF100-GTX480 でよいのでしょうか それともアーキテクチャとCCは1対1対応しているのでしょうか また、Fermiの分類位置はここでよいのでしょうか 185 名前：デフォルトの名無しさん mailto:sage [2010/12/24(金) 22:08:13 ] ComputeCapability2.0-GF100-GTX480 は正しいけど、FermiはC.C.2.X全部を含む GFの"F"はFermi、GTXの"T"はなぜかTeslaだったはず 186 名前：デフォルトの名無しさん mailto:sage [2010/12/24(金) 22:20:21 ] てめぇGeForce 7800GTXさんに謝れよ 187 名前：デフォルトの名無しさん mailto:sage [2010/12/24(金) 23:41:29 ] >>184 CUDA目的ならComputeCapabilityだけ意識すればよいかと。 あとはパフォーマンスが違うだけだから。 188 名前：デフォルトの名無しさん mailto:sage [2010/12/24(金) 23:43:40 ] >>186 GeForce 7800GTXさんはCUDA対象外じゃ・・・・・・ 189 名前：デフォルトの名無しさん mailto:sage [2010/12/25(土) 00:54:10 ] コンパイルされたCUDAプログラムを実行するだけの場合、ドライバのみでいけますか？ toolktのインストールは必須ですか？ LD_LIBRARY_PATH /usr/local/cuda/lib64 を通さなくてはならないということはtoolkit必須なのかな 190 名前：デフォルトの名無しさん mailto:sage [2010/12/25(土) 01:00:31 ] ライブラリ使っててもstaticならいらないはずだし、dllつかってるなら そのdllをexeと同じフォルダにおけばいいんじゃないか？ きっとlib64内のなんかのライブラリつかってんでしょ。cutilとかかな 191 名前：デフォルトの名無しさん mailto:sage [2010/12/25(土) 02:39:52 ] DLL再配布して良かったっけ 192 名前：デフォルトの名無しさん mailto:sage [2010/12/25(土) 09:47:13 ] Linuxならドライバ（カーネルモジュールを含む）も必要なので再配布だけでは解決しない。 もちろんツールキットのライブラリも必要。 193 名前：デフォルトの名無しさん [2010/12/27(月) 19:25:09 ] すみません、教えてください。 CentOS 5.5(64bit)でCUDAを導入したくて ・Developer Drivers for Linux (260.19.26) ・CUDA Toolkit for RedHat Enterprise Linux 5.5 (64bit版) をインストールし、とりあえずSDK code samplesをコンパイルして動くか確かめたのですが コンパイル後のファイルを実行しようとすると "cannot execute binary file"のメッセージが出てきてプログラムが実行できません。 この問題を解決したいのですがどこを確認してどう改善したらよいのでしょうか？ ちなみにカーネルのバージョンは2.6.9-55.0.2.ELsmpです。 194 名前：デフォルトの名無しさん mailto:sage [2010/12/27(月) 19:40:29 ] 環境設定は？ 195 名前：デフォルトの名無しさん mailto:sage [2010/12/27(月) 20:29:33 ] それはCentOS 4.5のようだが。 196 名前：デフォルトの名無しさん mailto:sage [2010/12/27(月) 22:11:09 ] >>193 例えば、ディストリビューションのnvidiaドライバなんかが残っていない？ その上に、最新ドライバをインストールしていて、 コンパイル時と実行時に別の共有ライブラリを参照していたりとか。 ていうか、自分も、ubuntuでx-updatesのドライバと最新ドライバが両方入っていて挙動不審。 ubuntuでもOS再インストールかｗ 197 名前：193 [2010/12/28(火) 00:14:59 ] >>196 返信ありがとうございます。 ちょっとその辺を調べてみます。 後日、問題が解決したらフィードバックします。 198 名前：デフォルトの名無しさん mailto:sage [2010/12/28(火) 07:48:23 ] スルーされた。勝手にしてくれ。 199 名前：デフォルトの名無しさん mailto:sage [2010/12/28(火) 16:15:18 ] いじけるなよｗ 200 名前：デフォルトの名無しさん mailto:sage [2010/12/28(火) 22:47:32 ] とある処理のスレッド数と時間を関係を調べているのですが、 スレッド数が192を超えたあたりで処理が行われていません。 スレッド数が512より小さくても処理が行われないことがあるのでしょうか？ グラボは9800GTで11レジスタ、共有メモリを36byte使用してます。 201 名前：デフォルトの名無しさん mailto:sage [2010/12/28(火) 22:48:39 ] >>200 「処理が行なわれない」ってどういう意味さ。 202 名前：デフォルトの名無しさん mailto:sage [2010/12/28(火) 23:01:14 ] >>201 説明不足でした。 処理を1000回行い、平均処理時間を計測しているのですが、 スレッド数が192を超えると0(ms)になってしまいます。 203 名前：デフォルトの名無しさん mailto:sage [2010/12/28(火) 23:19:14 ] 何かが間違っているんだろうね。 204 名前：デフォルトの名無しさん mailto:sage [2010/12/29(水) 00:34:35 ] まずはエラー確認してからだろ。 205 名前：デフォルトの名無しさん mailto:sage [2010/12/29(水) 00:38:37 ] スレッド数を増やしたことによって確保されていない領域にアクセスするような エラーが発生するような状況になっていないかどうかを確認・・・・・かな。 これが起きているとそれ以降の（タイマーを含めた）全ての処理がエラーになる。 206 名前：デフォルトの名無しさん mailto:sage [2010/12/29(水) 01:31:39 ] まずは cudaThreadSynchronize(); cudaGetLastError(); 207 名前：デフォルトの名無しさん mailto:sage [2010/12/29(水) 14:09:27 ] >>203-206 助言ありがとうございます。 まずはエラーを確認してみたいと思います。 208 名前：デフォルトの名無しさん mailto:sage [2010/12/29(水) 19:33:12 ] たぶんリソース不足でしょ。自分もよくそうなる。 カーネルの引数を構造体のポインタにまとめると解決する。 209 名前：デフォルトの名無しさん [2010/12/30(木) 20:27:57 ] GTX460ってCUDAは使えるんですかね？ lspci | grep -i nVidia でPCIバスデバイス情報確認したら VGA compatible controller: nVidia Corporation: Unknown device … ってコンピューターさんがおっしゃられたんだが… 210 名前：デフォルトの名無しさん mailto:sage [2010/12/31(金) 00:21:26 ] >>209 /sbin/update-pciids してからコンピュータさんにもう一度お伺いたててご覧なさい 211 名前：デフォルトの名無しさん mailto:sage [2010/12/31(金) 05:27:05 ] ドライバをインストールしてないんじゃないの？ 212 名前：デフォルトの名無しさん mailto:sage [2011/01/01(土) 11:59:19 ] 正月セールで狙う1万5000円くらいのグラボでオススメのものをおせーて 213 名前：デフォルトの名無しさん mailto:sage [2011/01/04(火) 09:25:33 ] GeforceとRADEONを2枚挿して GeforceをCUDA専用、RADEONをグラフィック専用とすることは出来ますか？ 214 名前：デフォルトの名無しさん [2011/01/04(火) 09:34:04 ] 日本語的に複数の意味に取れるので訂正 GeforceとRADEONの2枚を挿して GeforceをCUDA専用、RADEONをグラフィック専用とすることは出来ますか？ 215 名前：デフォルトの名無しさん mailto:sage [2011/01/04(火) 09:41:44 ] no problem 216 名前：デフォルトの名無しさん mailto:sage [2011/01/04(火) 10:36:53 ] ありがとうございます 217 名前：デフォルトの名無しさん mailto:sage [2011/01/04(火) 10:52:31 ] cudaSetDeviceFlags(cudaDeviceBlockingSync); 218 名前：デフォルトの名無しさん mailto:sage [2011/01/04(火) 10:53:37 ] ほとんど同じソースで cudaSetDeviceFlags(cudaDeviceBlockingSync); が有効になるプログラムとならないプログラムがあるのですがどこを 調べれば良いでしょうか。 219 名前：デフォルトの名無しさん mailto:sage [2011/01/04(火) 12:31:38 ] ほとんど同じソースならどこが違うのかを書いてください。 そして「有効」かどうかをCPU負荷で判定しているのかどうかも。 cudaSetDevice()を除くすべてのCUDA関数より先にcudaSetDeviceFlags()を 呼び出しているかどうか以外には特に調べるべきところはありません。 220 名前：デフォルトの名無しさん [2011/01/05(水) 18:57:39 ] マニュアル的なpdfが沢山あってどれを読めばいいのかわからないのですがどれから読むのがいいのでしょうか？ CUDA_C_Best_Practices_Guide.pdf CUDA_C_Programming_Guide.pdf Introductory CUDA Technical Training Courses VolumeI.pdf 221 名前：デフォルトの名無しさん mailto:sage [2011/01/05(水) 20:31:24 ] Introductory CUDA Technical Training Courses VolumeI.pdfで軽く概要をつかむ。 ただこれ内容が古かったと思うからどんな仕組みで動いてるかを把握するくらいでいいと思う 必要になったらまた読み返せばいい。 次にCUDA_C_Programming_Guide.pdfで基本を学ぶ。 これは読み返すことが多くなると思う。 その次にCUDA_C_Best_Practices_Guide.pdfで良い実装の仕方を学ぶ。 これは実際にプログラム組みながらでいいと思う。 あと余談だがfermi　tuning guideも読むといい。 英語が苦手なら「はじめてのCUDAプログラミング」や「CUDAプログラミング実践講座 - 超並列プロセッサにおけるプログラミング手法」 を読んでから必要なところをPDFで追加で読むと理解が早いと思う。 222 名前：デフォルトの名無しさん mailto:sage [2011/01/05(水) 21:02:09 ] フランス語のはないですかね？ 223 名前：デフォルトの名無しさん mailto:sage [2011/01/05(水) 21:15:20 ] ワロタ 224 名前：デフォルトの名無しさん mailto:sage [2011/01/05(水) 21:16:37 ] >>222 残念ながらありません 225 名前：デフォルトの名無しさん mailto:sage [2011/01/06(木) 00:39:25 ] cc2.0でコンパイルすると、cc1.3でコンパイルした時よりレジスタ使用量が 10くらい増えるのですが、なぜですか？ 226 名前：デフォルトの名無しさん mailto:sage [2011/01/06(木) 22:50:40 ] なんでだろうねー 227 名前：デフォルトの名無しさん mailto:sage [2011/01/08(土) 18:48:55 ] Fermiで各スレッドで数百バイトのワーキングエリアが必要な場合、 ローカルメモリを使うのが最良なのでしょうか？ シェアードメモリを使うと、同時に立ち上げられるスレッド数が減って 逆に遅くなりました。 228 名前：デフォルトの名無しさん mailto:sage [2011/01/09(日) 21:38:49 ] >>225 そういうのはプログラムによって変わるだろ とりあえずptxファイル嫁 >>227 >各スレッドで数百バイトのワーキングエリアが必要な場合 これが今以上減らせないならそもそもCUDAに向いてない 229 名前：デフォルトの名無しさん mailto:sage [2011/01/10(月) 17:57:08 ] グローバルメモリのアクセスは400~500clock、レジスタやシェアードメモリ、L1キャッシュはほぼ0らしいですが L2キャッシュは大体何clockかかりますか？ 230 名前：デフォルトの名無しさん mailto:sage [2011/01/10(月) 19:32:12 ] L1は大体15 L2は大体150 231 名前：デフォルトの名無しさん mailto:sage [2011/01/11(火) 09:01:34 ] 4枚発注ー。 232 名前：デフォルトの名無しさん [2011/01/13(木) 11:27:07 ] CPU向けコードとGPU向けコードを同一ファイルで作成したいと思っています。 nvccはcuでないと、cudaコードとしてコンパイルしてくれないのですが、 拡張子がcppのファイルをcudaソースファイルとしてコンパイルするオプションはありませんか? 233 名前：デフォルトの名無しさん mailto:sage [2011/01/13(木) 19:39:51 ] obj同士はnvccでリンクできるから、それぞれコンパイルすればＯＫだよ 234 名前：デフォルトの名無しさん mailto:sage [2011/01/13(木) 20:42:06 ] >>141 亀レスですが，Driver APIを用いて出来ました． ありがとうございました． 235 名前：デフォルトの名無しさん mailto:sage [2011/01/14(金) 20:47:45 ] >>234 概略をのせてくれまいか？ cudaAAA()で○○を設定 cudaBBB()でXを取得して pthread_AAA()を使ってXXXする くらいの事でもいいから 236 名前：234 mailto:sage [2011/01/15(土) 13:34:59 ] >>235 簡単にで良いなら． // 前処理 #pragma omp parallel { 　　cuInit(); // 初期化 　　cuDeviceGet(); // GPUへのハンドルを取得 　　cuCtxCreate(); // コンテキストを作成し，ホストスレッドにバインド 　　mallocとか色々 　　cuCtxPopCurrent(); // コンテキストをアンバインド } glutMainLoop(); // カーネル実行 #pragma omp parallel { 　　cuCtxPushCurrent(); // バインド 　　カーネル実行 　　cuCtxPopCurrent(); // アンバインド } 関数の引数は省略しました． 237 名前：デフォルトの名無しさん mailto:sage [2011/01/15(土) 17:57:20 ] >>236 ありがとう。参考にさせてもらいます 238 名前：デフォルトの名無しさん mailto:sage [2011/01/15(土) 21:30:19 ] Fermiもグローバルメモリアクセスはコアレッシングにしたほうがよくって キャッシュはおまけ程度って考えでおk？ ていうか各スレッドが違うグローバルメモリにアクセスするなら同じカーネル内ではキャッシュ効かないよね 239 名前：デフォルトの名無しさん mailto:sage [2011/01/15(土) 22:14:42 ] CPUのマルチコア（巨大キャッシュ、低並列）のマルチスレッドですら キャッシュ汚染はひどいからな マルチスレッドとキャッシュは相性が非常に悪い 240 名前：デフォルトの名無しさん mailto:sage [2011/01/15(土) 23:23:54 ] スレッド番号とグロバールメモリーのアドレスが並んでいないと使い物にならないな。 速度が１０倍は違う。本当に使いにくいよ。 メモリアクセスがネックの（単純な）アプリで、shared memoryの使いまわしのできないアプリでは、 マルチスレッドを上回るのは難しい。 241 名前：デフォルトの名無しさん mailto:sage [2011/01/15(土) 23:49:58 ] しかし、いままでGPUと相性の悪いといわれるプログラムでも FermiならCPUの2倍くらい出るってことは結構ある。 非Fermiならローカルメモリ使ったら即使い物にならない速度だったけど、 Fermiならキャッシュがなんとかしてくれたり。 242 名前：デフォルトの名無しさん mailto:sage [2011/01/16(日) 18:58:37 ] cublasDgemmにおいて、係数のalphaを(double)1に、 betaを(double)0に指定しても、 内部では真面目に係数をかけているのでしょうか？ 参考(PDF注意) developer.download.nvidia.com/compute/cuda/3_2_prod/toolkit/docs/CUBLAS_Library.pdf 243 名前：デフォルトの名無しさん mailto:sage [2011/01/16(日) 19:01:42 ] FORTRANのソースを自分で追っかけて確認しろ 244 名前：デフォルトの名無しさん mailto:sage [2011/01/16(日) 19:13:17 ] 速度を測ればすぐに分かりそう 245 名前：デフォルトの名無しさん mailto:sage [2011/01/16(日) 21:42:18 ] >>243 ありがとうございます。 FORTRANのコードを見る限り、ちゃんと係数の値によって 処理変えているようでした。 >>244 ありがとうございます。 実際に測定したのですが、係数の値を変えても 計算時間に明確な差は現れませんでした。 GTX480 1枚, CUDA 3.2, Windows 7 64bit A=7500x8000の行列,B=8000x7500,C=7500x7500の行列で C=alpha * A*B + beta * Cを計算 alpha = 1, beta = 0を15回で85.079秒 alpha = 10, beta = 10を15回で85.072秒 FORTRANのコードのような実装になっていない気がするのですが。 246 名前：デフォルトの名無しさん mailto:sage [2011/01/16(日) 21:45:25 ] 「はじめてのCUDAプログラミング」に出てるコード（もしくは同じ出版社からインターネット上に出てる奴） をひたすらマイナーチェンジしていくのが一番早いと思う。 スクラッチはマイナーチェンジの経験を積んでからで十分。 247 名前：デフォルトの名無しさん mailto:sage [2011/01/17(月) 05:28:45 ] PTXって機械語なの？ GPUがPTX読むみたいな事が書いてあるんだけど 248 名前：デフォルトの名無しさん mailto:sage [2011/01/17(月) 06:02:22 ] ptxはgpuのアセンブラ言語じゃないの？？ 249 名前：デフォルトの名無しさん mailto:sage [2011/01/17(月) 07:27:52 ] PTXはあくまで中間言語。 250 名前：デフォルトの名無しさん mailto:sage [2011/01/17(月) 10:47:31 ] TMPGEncで使おうと思ったらG80非対応なんだな。。 251 名前：デフォルトの名無しさん mailto:sage [2011/01/17(月) 20:15:09 ] GTX580を4枚積んでサンプルコード走らせたら電力不足来たー。 252 名前：デフォルトの名無しさん [2011/01/19(水) 09:33:50 ] CUDAツールキットを導入する際にまずやることって何？ どう頑張ってもnvccでコンパイルしたサンプルファイルが実行できませーん って開発以前の問題で精神折れそうだわ… 253 名前：デフォルトの名無しさん [2011/01/19(水) 13:17:05 ] list.cは以下の通りです． #include <stdio.h> #include <stdlib.h> #include <string.h> struct record { /* １人分のデータ */ char name[32]; /* 氏名 */ int score; /* 得点 */ struct record *next; /* 次のデータへのポインタ */ }; void printlist( struct record *list ) { if ( list != NULL ) { /* リストが空でなければ */ printf( "%16s:%4d\n", list->name, list->score ); printlist( list->next ); /* 先頭の次からリストの表示 */ } } 254 名前：デフォルトの名無しさん mailto:sage [2011/01/19(水) 13:59:28 ] >>252 実行できないならできないなりにエラーメッセージが帰ってくるのではないでしょうか CUDAコアとはスカラプロセッサのことでいいのでしょうか FermiのみCUDAコアという呼び方をするのですか？ 255 名前：デフォルトの名無しさん mailto:sage [2011/01/19(水) 14:09:09 ] >>252 実行できないならできないなりにエラーメッセージが帰ってくるのではないでしょうか CUDAコアとはスカラプロセッサのことでいいのでしょうか FermiのみCUDAコアという呼び方をするのですか？ 256 名前：デフォルトの名無しさん mailto:sage [2011/01/19(水) 21:03:35 ] Fermiで倍精度浮動小数点の三角関数計算に SFUを使えないのでしょうか 倍精度浮動小数点の計算は単精度浮動小数点演算機が2クロックで計算するのですよね 同じようにSFUが単精度計算分の2週でやってくれるということはないのでしょうか 257 名前：デフォルトの名無しさん mailto:sage [2011/01/19(水) 21:10:31 ] geforceは駄目だろ 258 名前：デフォルトの名無しさん mailto:sage [2011/01/19(水) 21:57:03 ] ドライバに　266.58 WHQL　が来たね 259 名前：デフォルトの名無しさん mailto:sage [2011/01/19(水) 22:20:50 ] >>256 fast mathを使えば？ 260 名前：デフォルトの名無しさん mailto:sage [2011/01/20(木) 00:39:31 ] fast mathは単精度専用。 そもそも倍精度使ってるくせに精度低いSFU使いたいって時点で 本当に倍精度が必要か考えるべき。 261 名前：256 mailto:sage [2011/01/20(木) 02:09:04 ] SFU遊ばしておくのはもったいないかな、と思いまして そんなに精度が低いのかー 262 名前：デフォルトの名無しさん mailto:sage [2011/01/20(木) 02:55:13 ] 応用次第だけど、SFU使うと駄目で使わなければfloatでもOKという例があった。 263 名前：デフォルトの名無しさん [2011/01/21(金) 02:52:36 ] 初歩的なこと聞いてすみません。 元々あるソースにCUDAを取り入れてるんですが構造体を丸々GPU側にコピーはできないんでしょうか？ できるのであれば誰かご教授お願いします。 264 名前：デフォルトの名無しさん mailto:sage [2011/01/21(金) 03:12:27 ] C言語の勉強をしたらできると思います。 265 名前：デフォルトの名無しさん [2011/01/21(金) 03:30:01 ] >>264 一応いろいろ試したんですが Warning : Cannot tell what pointer points to, assuming global memory space ていう警告が出ていておそらくデータの転送ができていないんです・・・。 普通のC言語を勉強すればできるんですかね・・・？ 266 名前：デフォルトの名無しさん mailto:sage [2011/01/21(金) 03:40:27 ] 構造体の中にポインタが入ってて、そのポインタがメインメモリを指してるんじゃないの？ 267 名前：デフォルトの名無しさん [2011/01/21(金) 04:09:34 ] >>266 ポインタのほかにもポインタのポインタも入っているのでそれもGPU側に送らないといけないいうことであってますか？ そうすると結構量があるので必要な部分のみを送るほうが効率がいいんですかね・・・？ 268 名前：デフォルトの名無しさん mailto:sage [2011/01/21(金) 05:42:58 ] もちろんポインタが指してるデータは送らないといけないし， 構造体を送る前に，構造体の中で指してるアドレスをGPU側のアドレスに 書き変えないといけない． 269 名前：デフォルトの名無しさん [2011/01/21(金) 06:51:09 ] 解答ありがとうございます。 なんとなくですが皆さんのおかげで理解できました。 がんばってやってみます。 270 名前：デフォルトの名無しさん mailto:sage [2011/01/21(金) 07:06:46 ] なんとなくじゃ無理だろうな 271 名前：デフォルトの名無しさん [2011/01/24(月) 01:53:29 ] 少し聞きたいんですが、 1>ptxas C:/Temp/tmpxft_00000fac_00000000-4_cuda.ptx, line 146; error : State space incorrect for instruction 'st' 1>ptxas C:/Temp/tmpxft_00000fac_00000000-4_cuda.ptx, line 147; error : State space incorrect for instruction 'st' 1>ptxas €=8・ Ptx assembly aborted due to errors いきなりこんなエラーが出たんですがどういう意味ですか？ 解決法を知っている方教えてください！！ 272 名前：デフォルトの名無しさん [2011/01/24(月) 08:07:36 ] en.wikipedia.org/wiki/Comparison_of_Nvidia_graphics_processing_units 誰かこのページでGPUの浮動小数点計算速度の比較に使われてるプログラム知ってる人いませんか？ 273 名前：デフォルトの名無しさん mailto:sage [2011/01/24(月) 10:18:04 ] >>272 ざっと見ただけではこのページのどこに浮動小数点計算速度の比較が書いてあるのか分からなかったのだが。 いずれにしてもそこに載っているスペックは全て理論ピーク性能かと。 274 名前：デフォルトの名無しさん mailto:sage [2011/01/24(月) 10:27:17 ] >>271 いきなりすぎて難しいのでとりあえず中間ファイルの.ptxファイルを保存して146行目を調べて。 あとはSDKに含まれるプログラムとかは正常にコンパイルできるのかどうかも。 275 名前：デフォルトの名無しさん mailto:sage [2011/01/24(月) 13:59:18 ] お聞きしたいのですが、プログラムを実行したときに発生するエラーで invalide device function とは何が原因で起こるエラーなんでしょうか？ 初歩的な質問ですがどなたか教えてください。 276 名前：デフォルトの名無しさん mailto:sage [2011/01/24(月) 14:38:37 ] invalide device function device で実行しているfunctionがinvalideなんだろ でもinvalideじゃなくてinvalidじゃね？ 277 名前：275 mailto:sage [2011/01/24(月) 15:13:26 ] スペルを間違えていました。invalideではなくinvalidでした。 invalidは「無効な」という意味で捉えています。 デバイス(__global__)の関数名とホストで実行する関数名は同じにしているのですが、なぜ無効と言われるのかがわからないのです。当然ですが、何かの関数名と被らないようには名前をつけています。 どなたかお願いします。 278 名前：デフォルトの名無しさん mailto:sage [2011/01/24(月) 15:38:25 ] もう少し具体的に書け 279 名前：デフォルトの名無しさん mailto:sage [2011/01/24(月) 15:56:17 ] _device_な関数を_global_と_device_な関数以外から呼び出してる予感 280 名前：275 mailto:sage [2011/01/24(月) 16:12:51 ] 詳細を書きました。長文です。すみません。 画像処理をCUDAでやっており、今は、メディアンフィルタを作っています。 OpenCVを使いたかったので、ここの方法（ttp://blogs.yahoo.co.jp/simasaki2001/18576477.html）を参考に簡単なサンプルを先生が作成し、それを基に学生の自分がプログラムを作っています。 現在、__global__ void cuda_median() というデバイスの関数をcuda_median<<<ndrange_global, ndrange_local>>>(入力画像, 入力画像のポインタ, 出力画像, 出力画像のポインタ)と呼び出しています。 ndrange_globalとndrange_localはDim3型で ndrange_local.x= atoi( argv[3] ); ndrange_local.y= 1; ndrange_local.z= 1; ndrange_global.x= 1; ndrange_global.y= (int)dev.multiProcessorCount; <-これは14となっています ndrange_global.z= 1; となっています。この部分についてよく理解できなかったので先生に質問しましたが、「人の作ったのだからよく知らない。」と言われました。現在もしっかり理解しているとは言えません。 __device__関数も一つ作成して使っていますが、__global__からしか使用していません。 281 名前：デフォルトの名無しさん mailto:sage [2011/01/24(月) 16:21:56 ] いつ、どこでのエラーなんだ？ device内のfunctionでのエラーではないのか？ 282 名前：275 mailto:sage [2011/01/24(月) 16:37:21 ] ホストでdevice関数を呼んだ後にcudaGetLastErrorを実行しており、そこでエラーが出ます。 device関数を呼ぶ前の処理におけるエラーチェックではエラーは発生していないので、device関数でエラーが起こっているのではないかと考えています。 283 名前：デフォルトの名無しさん mailto:sage [2011/01/24(月) 16:57:03 ] ttp://lzhan.wikispaces.com/file/view/gflops.png/30874460/gflops.png ttp://4.bp.blogspot.com/_nq8CjirJSCA/TDWRC5XR8vI/AAAAAAAACZY/pyHDbs2AA0E/s1600/gpgpu.png 上にあるようなCPUとGPUの浮動小数点演算性能の比較図を見るとGPUの方が圧倒的に高いように見えるのですが 性能はGPUの方が上という事でしょうか それともただ単に1コアのクロック周波数から理論的な演算性能を算出してコア数分かけた値で、 GPUはコアが多いからそう見えているのでしょうか さらにメモリの転送速度もGPUの方が速いのですよね。 他の点で総合的に見たらCPUのほうがいいということなのでしょうか 284 名前：デフォルトの名無しさん mailto:sage [2011/01/24(月) 17:01:06 ] ＞それともただ単に1コアのクロック周波数から理論的な演算性能を算出してコア数分かけた値で、 ＞GPUはコアが多いからそう見えているのでしょうか ＞さらにメモリの転送速度もGPUの方が速いのですよね。 ＞他の点で総合的に見たらCPUのほうがいいということなのでしょうか それで合ってる 285 名前：デフォルトの名無しさん mailto:sage [2011/01/24(月) 17:21:26 ] >>283 まあ、簡単に言うと計算の仕方とメモリアクセスの仕方について、GPUの方が制約がきついってこと もっとも、CPUでも高い性能を出すためには厳しい制約が付くけど 286 名前：283 mailto:sage [2011/01/24(月) 18:18:42 ] >>284-285 ありがとうございます。 つまり全てのコアを常に使い切ればこれだけの性能が出るので、 メモリアクセスやのレイテンシや転送時の待ち時間に 準備の出来た他の演算(warp)を次々にやらせていって いかに常にコアをフル回転させるか、がCUDAプログラミング。 warp内では実時間でも完全に同期して実行されるため warp内分岐は、分岐を通らないwarpは完全に遊んでしまう(コアは空いているのに他のwarpのグループを実行できない) ので遅くなってしまう(というかピーク性能が出せない)。 メモリアクセス時には他のwarpを計算すればいいが、次のwarpもメモリアクスになり、最初のwarpもまだアクセス中だったときには 3つめのwarpを計算する。このようにどんどん他のwarpを処理していくが、一度に起動できるwarpは上限がある(カーネル関数で使用するレジスタや共有メモリで決まる) ため演算に対してメモリアクセス時間が長いと、起動上限数使い切ってしまい、 それでも最初のwarpがメモリアクセス中だとコアが遊んでしまうので、 メモリアクセス時間を減らしたり(メモリアクセス回数を減らす、共有メモリを有効利用する、 バンクコンフリクトを回避して共有メモリに早くアクセスできるようにする、コアレスアクセスでグローバルメモリに早くアクセスする)、 同時起動可能なwarp数を増やす(Occupancyうんぬん) ことで見た目全コアが降る回転するようにする。 ということでしょうか 287 名前：デフォルトの名無しさん mailto:sage [2011/01/24(月) 19:19:41 ] >>280 ついこの前出たOpenCV 2.2でNPP経由でCUDA対応したから、それ使ってみたら？ exeではなくCMakeから自前ビルドしないといけないけど 288 名前：デフォルトの名無しさん [2011/01/25(火) 10:09:24 ] >>274 SDK内のプログラムはきちんと動作しています。 .ptxファイルは保存方法がぐぐっても良くわからなくて・・・。 どうやったら保存できますか？ 289 名前：アク禁長いから久々の書き込み mailto:sage [2011/01/25(火) 12:05:07 ] >>282 状況がよく判らんが、__device__な関数は__device__か__global__な関数からしか呼んじゃダメよ。 それと、<<<>>>で呼んだ関数で起きたエラーは、次にAPIを呼んだときに戻されるのでご用心。 つまり、以下の場合cudaThreadSynchronize()の戻り値がエラーになる。 -- func<<<grid, block>>>(); cudaThreadSynchronize(); >>271 取り敢えず、-ptxつけてコンパイルして味噌。 290 名前：デフォルトの名無しさん [2011/01/25(火) 12:51:53 ] >>289 コマンドプロンプトでnvcc -ptx cuda.cuであってますか？ 普段VSでコンパイルしてるんでやり方がいまいちなんですが・・・。 291 名前：283 mailto:sage [2011/01/25(火) 13:02:29 ] Fermiでコンスタントキャッシュはハードウェア上ではどこにあるのでしょうか ttp://www.4gamer.net/games/099/G009929/20100309044/images/pic2.gif のUniform Cacheでしょうか 容量と読み込み速度からしてSM内かと思ったのですが 全てのスレッドから読めるという点ではSM外のL2キャッシュなのでしょうか SM内だとしたら全SMがそれぞれ同じデータをコピーして保持するというのは少々効率が悪い気もします そもそも大容量なはずのVRAMに転送するコンスタントメモリが 64kBと小容量に決められているという事はキャッシュも同サイズ用意されていて ヒット率100％が保証されているということなのでしょうか 292 名前：デフォルトの名無しさん mailto:sage [2011/01/28(金) 08:08:16 ] ちょっと質問なんですけど、GPU上ではstaticは使えないみたいですけど、 staticの代わりになるようなものはCUDAに存在するんでしょうか？ やっぱりグローバルメモリに領域を確保して使うしかないんでしょうか？ 293 名前：デフォルトの名無しさん [2011/01/29(土) 02:47:23 ] >>292 すまん質問の意味がわからん 294 名前：デフォルトの名無しさん mailto:sage [2011/01/29(土) 21:17:57 ] >>293 すいません、ちょっと日本語がおかしいですね・・・ 要するに、ＧＰＵ上のdevice関数内で、 関数が終わっても変数の値が残るようにしたいんですけど、 それをするにはcudaMemcpyでグローバルメモリに領域確保して、 そこに値を保存するようにするしかないのでしょうか？ 295 名前：デフォルトの名無しさん mailto:sage [2011/01/29(土) 21:38:16 ] それどういうときに使うの？カウンタ？ そもそもインライン展開されるんじゃなかったっけ 296 名前：デフォルトの名無しさん mailto:sage [2011/01/29(土) 23:19:04 ] 根本的に取り組む姿勢が間違っている気がする。 297 名前：デフォルトの名無しさん mailto:sage [2011/02/01(火) 19:48:05 ] 構造体の配列って a,b,c|a,b,c|a,b,c こう置かれるからハーフワープの各スレッドがaにアクセスすると コアレスアクセスできないということですか？ 配列の構造体.a[n],.b[n],.c[n]を作った方が良いという事でしょうか 298 名前：デフォルトの名無しさん mailto:sage [2011/02/01(火) 19:56:41 ] 誰だかまるわかりw 299 名前：デフォルトの名無しさん mailto:sage [2011/02/02(水) 19:08:00 ] CUDA GPU Occupancy Calculator を使って占有率等をみるときにActive Thread Block per Multi Processor が高いほうがいい と聞くのですが、以下の二つの状況だとどちらがいいでしょうか？GTX480を用いています。 1. Enter your resource usage: Threads Per Block 256 Registers Per Thread 32 Shared Memory Per Block (bytes) 3072 GPU Occupancy Data is displayed here and in the graphs: Active Threads per Multiprocessor 1024 Active Warps per Multiprocessor 32 Active Thread Blocks per Multiprocessor 4 Occupancy of each Multiprocessor 67% 2. Enter your resource usage: Threads Per Block 1024 Registers Per Thread 32 Shared Memory Per Block (bytes) 12288 GPU Occupancy Data is displayed here and in the graphs: Active Threads per Multiprocessor 1024 Active Warps per Multiprocessor 32 Active Thread Blocks per Multiprocessor 1 Occupancy of each Multiprocessor 67% 1と2の違いはThread Per Blockを四倍にしたかどうかです。 つまりSM内のThread Block数が多いほうが単純にいいのか、BLOCK内のThread数が多いほうがいいのか、 あるいはどちらも同じなのか教えてください。 300 名前：デフォルトの名無しさん mailto:sage [2011/02/02(水) 19:38:03 ] >>299 計測してパフォーマンスのいい方にすればいいとおもうが、 １の方があるスレッドブロックが同期待ちになってるあいだに他のスレッドブロックが 実行可能な余地があるぶん性能低下はしにくいと思う 301 名前：デフォルトの名無しさん mailto:sage [2011/02/02(水) 20:23:06 ] >>300 回答ありです。 言葉足らずでした。。。 計測して実験したのですが、1のほうがパフォーマンスが上でした。 いい言葉が見つからないのですが、 1も2もブロック数は違うけど同じスレッド数ですので メモリ隠ぺいのためにスレッド数が多ければそれでいいと思っていたので 性能は同じになると思っていました。 スレッドブロックが同期待ちというのは、syncした場合におこる現象でしょうか？ 調べてみましたがやはりSMあたりのスレッド数が同じなら SMあたりのスレッドブロック数が多いほうがいいというのが いまいちわかりません。。。 よろしくお願いします 302 名前：デフォルトの名無しさん mailto:sage [2011/02/02(水) 20:35:42 ] >>301 __syncthreads(); はブロック内のスレッドを同期する。 全部のスレッドが1024として 1ブロック1024スレっどの場合、__syncthreads();すると全部のスレッドが同期される。 その間どのスレッドも実行できない。 4ブロックで1ブロック当たり256スレッドの場合、__syncthreads();しても256スレッドだけ同期されて 他の３ブロック分のスレッドは実行可能。 ＞SMあたりのスレッドブロック数が多いほうがいい スレッドブロックはSM単位に割り当てられるから「SMあたりのスレッドブロック数が多いほうがいい」 というのはActive Thread Blocks per Multiprocessorが高いほうがいいというのと同義 303 名前：デフォルトの名無しさん mailto:sage [2011/02/02(水) 20:53:25 ] Visual Profilerでlocal load localとstoreが数万回起こっているのですが何が原因でしょうか --ptxas-options=-vでlmemが出ていなかったので ローカルメモリは使っていないと思っていたのですが local〜が0でないカーネル(3種類)はすべてレジスタ/スレッドが63でした。 偶然同じ数になっていると思ってましたが レジスタは63が上限でそれ以降はローカルメモリが使われるってこと？ 304 名前：デフォルトの名無しさん mailto:sage [2011/02/02(水) 23:02:21 ] >>303 レジスタの容量を調べてみろ 例えばレジスタの本数が32768で1ブロックあたり512スレッド使うなら 1スレッドが使用できるレジスタは64以下になる、256スレッドなら128以下 ローカルメモリが使われるなんて状況は最悪だからそのプログラムは糞だ シェアードメモリ余ってんならそっちも有効に使え 305 名前：デフォルトの名無しさん mailto:sage [2011/02/02(水) 23:30:24 ] だがローカルメモリ使ったカーネルとCPU実行では， CPU実行の方が遅いこともあるから必ずしも悪ではない． 306 名前：デフォルトの名無しさん mailto:sage [2011/02/03(木) 17:37:00 ] 総和計算を何段階かにわけたときって sharedメモリ使う都合上、最後の段階は1ブロックでやるじゃん でも1ブロックのみってことは他のsm使わないから無駄してるよね 最後の段階はcpuに転送したほうがはやいってこと？ 307 名前：デフォルトの名無しさん mailto:sage [2011/02/03(木) 18:32:56 ] fermi系ならコンカレントカーネルで問題無し 308 名前：デフォルトの名無しさん mailto:sage [2011/02/04(金) 01:25:47 ] どうせ最後にcpuに持ってくるんだから、充分データ量が減ったところで転送しちゃうのも手ではあるね。 309 名前：303 mailto:sage [2011/02/04(金) 03:03:46 ] >>304 レジスタ32768ですがスレッドは192/ブロックです。 そもそもActive Thread Blocks per Multiprocessorを2以上にするように 1スレッドで使用しているレジスタ数とシェアードメモリを見て スレッド数を決めているので足りなくなるなんて事はないはず --ptxas-options=-vでlmemの項が出ないんだから使ってないと思っているのですが レジスタ/スレッドに上限なんてないですよね？ 310 名前：デフォルトの名無しさん mailto:sage [2011/02/04(金) 03:34:56 ] そのレジスタ数から考えてfermi系ではないとおもじゅんだが いっそのことスレッド数を８とか３２とかにして実行してどうなるか見てみればいいんじゃないか なにか他の原因かもよ？ 311 名前：デフォルトの名無しさん mailto:sage [2011/02/04(金) 03:35:39 ] c/おもじゅんだが/思うんだが 312 名前：303 mailto:sage [2011/02/04(金) 04:33:53 ] >>310 GF100,Fermiです 8スレッド/ブロックにして(ブロック数を増やしたので総スレッド数は変わらず) みましたがlocal load/storeは増えました。 カーネル内の変数を(なにかに使用するように)無駄に増やしてみましたが レジスタ数が63以上にならなかったので上限があるのかも、と思った次第です。 皆さんはレジスタ数が64を超える事はありますでしょうか 313 名前：デフォルトの名無しさん [2011/02/04(金) 05:01:16 ] 今環境ないからできないがnbodyのサンプルとかはプロファイルとったらどうなんだろうね。 あとフェルミならローカルメモリはキャッシュが効くんではなかったかと思ったがどうでしたっけ？ あとコンパイルオプションでレジスタの使用個数指定してたりしない？ 314 名前：303 mailto:sage [2011/02/04(金) 05:18:06 ] >>313 もちろん--maxrregcountは指定していませんでした ところで --maxrregcount 64にすると ptxas warning : Too big maxrregcount value specified 64, will be ignored とでました。 --maxrregcount 63にすると何も出ずコンパイルされます。 315 名前：デフォルトの名無しさん mailto:sage [2011/02/04(金) 05:23:42 ] ということは６３に意味はありそうだね。うーんなんだろ。即答できない 316 名前：デフォルトの名無しさん mailto:sage [2011/02/04(金) 05:42:17 ] Also, the compiler reports total local memory usage per kernel (lmem) when compiling with the --ptxas-options=-v option. Note that some mathematical functions have implementation paths that might access local memory. 317 名前：デフォルトの名無しさん mailto:sage [2011/02/04(金) 12:19:26 ] Fermiは63までしかないよ。GT200とかは127まで行けたけど。 318 名前：デフォルトの名無しさん mailto:sage [2011/02/04(金) 18:43:37 ] CC2.0でchar型の配列a[N]に int idx = blockIdx.x * blockDim.x + threadIdx.x; a[idx]; ってアクセスしたらきちんと32bitずつglobalから読みだしますか？ それとも128bitずつですか？ 319 名前：303 mailto:sage [2011/02/05(土) 08:23:45 ] >>317 そうだったんですか ちょっとずつ減らしてったら現在90くらい使ってることが判明 かきなおしかー 320 名前：デフォルトの名無しさん mailto:sage [2011/02/06(日) 13:23:49 ] 64しか使われないのに今90使ってるっておかしくないか 経験値からじゃなくて何か根拠のある数字はないのか？ 321 名前：303 mailto:sage [2011/02/06(日) 13:51:37 ] --ptxas-options=-v→63 同時に使ってるであろう変数を1個減らす→63 もういっこへらす→63 繰り返す へらす→62 やっと62になったということは減らした変数の数*2本分(doubleなので)=26本のレジスタが 63から溢れ出てローカルに保存されているんじゃないか →本来なら90くらい使われてるんだろう まあ一概に変数1個=2本とは言えないんでしょうけど 他に表示されないレジスタ数を数える方法思いつきませんでした 322 名前：デフォルトの名無しさん mailto:sage [2011/02/06(日) 15:52:57 ] 言ってることがようやくわかりました。 192スレッド/blockで動作してるなら 32768/192=170.66 で1スレッド170は理論的につかえるはずなのに他になにか制限みたいのがあるのかね もしかして待機しているactive threadの分が消費されるのか？つまり 32768/192/ActiveThreadPerBlockが1スレッドの使えるレジスタ数になるのだろうか 323 名前：303 mailto:sage [2011/02/06(日) 16:46:28 ] >>322 逆にactive threadの数がレジスタ数で決まるんじゃないでしょうか 192tread/block(= 6warp/block) → 48maxWarp/(6w/b) =8block 0shared/block → 49152maxShared/(0s/b) = ∞block 63register/tread(=12096register/block) → 32768maxRegister/(12096r/b) =2block Min(8,∞,2,8) = 2blockActive = 12warpActive → 12warp/48maxWarpActive =25% occupancy それとは関係なくFermiではregister/threadが上限63で それ以上はローカルにまわされる上に--ptxas-options=-vでlmem表示されない のだと解釈しました。 324 名前：322 mailto:sage [2011/02/06(日) 16:57:10 ] C programming guide 3.2 p83に GT is the thread allocation granularity, equal to 256 for devices of compute capability 1.0 and 1.1, and 512 for devices of compute capability 1.2 and 1.3, and 64 for devices of compute capability 2.x. ってあるがこれか？ 325 名前：303 mailto:sage [2011/02/06(日) 17:40:25 ] >>324 1ブロックで使用されるレジスタRbは2.xでは以下のように計算する Rb=ceil((Register/thread)*(thread/warp),GT)*(warp/block) すなわち Rb=ceil(r*32,64)*(warp/block) ceil(x,y)はxをyの整数倍に切り上げたもの r=18,w/b=6のとき Rb=(18*32→9*64)*6=3456 r=17,w/b=6のとき Rb=(17*32→9*64)*6=3456 つまり1warpが使用するレジスタは量子的で、その単位GT=64 →rは2の倍数で使わないと損 ということかと。r/スレッドの上限とは関係なさそうですかね。 326 名前：デフォルトの名無しさん mailto:sage [2011/02/07(月) 02:36:58 ] 上の方で出てたが、やはりCUBLASは省略すべき係数計算を 省略せずに計算しているようだ。 www.netlib.org/blas/dgemm.fに反する実装って許されるの？ 327 名前：デフォルトの名無しさん mailto:sage [2011/02/07(月) 08:00:09 ] 許せるに決まっているだろう。 328 名前：デフォルトの名無しさん mailto:sage [2011/02/07(月) 15:59:16 ] >>326 公式のフォーラムで本当にそうなのか、 もしそうならどのような理由からか聞いたほうがいいんじゃないか？ もしかしたら自分が勘違いしている部分があるかもしれないし Nvidiaの間違いなら修正があるだろうし。 329 名前：303 mailto:sage [2011/02/07(月) 16:11:41 ] ttps://people.maths.ox.ac.uk/gilesm/cuda/lec4-2x2.pdf#search='Fermi%20register%2063' 　up to 63 registers ttp://www.beyond3d.com/content/reviews/55/7 　Maximum theoretical register allocation per thread is 63 32-bit registers. このへんに63までと書いてありますが公式にはないのかな... 330 名前：デフォルトの名無しさん mailto:sage [2011/02/07(月) 23:23:08 ] CUDA VS Wizard + Fermiな環境の人いますか？ デバッグのためのprintfが使えなくて困っています。 プロジェクト＞プロパティ＞Cuda Build Rule v3.0.0＞General＞GPU Architecture の値をsm_20としたいのですが、無効な値となり設定できません。 （CUDA VS Wizard最新版がFermiより前のリリースのため） GPU Architectureの箇所を<親またはプロジェクトの規定値を継承>とし Extra Optionsに-arch sm_20と指定しても-arch sm_20 -arch sm_10となり ビルドできません。 この状況を経験＆打破された方がいらっしゃいましたらアドバイスをお願いします。。。 【環境】 Windows7 x64 Visual Studio 2008 Express CUDA Toolkit&SDK 3.1(32bit) CUDA VS Wizard 2.2(32bit) Geforce GTX 580 331 名前：デフォルトの名無しさん mailto:sage [2011/02/07(月) 23:48:13 ] >>330 あくまでCUDA VS Wizardが不具合ならという前提だけど sdkのtemplateプロジェクトを元に必要な箇所を修正してプロジェクトを 作りなおしてみるという手もある。 332 名前：330 mailto:sage [2011/02/08(火) 00:19:27 ] >>331さん ありがとうございます。頂いたアドバイスを頼りに解決しました。 templateプロジェクトを開く＞プロパティ＞あれ？CUDA Bulid Rule v3.0.14＆sm_20できる＞google先生 ＞先生「Cuda.rulesコピれ」＞！SDK導入後コピーするの忘れてた！＞�@から�AにCuda.rulesをコピー＞ヽ(´ー｀)ノ �@C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK\C\common �AC:\Program Files (x86)\Microsoft Visual Studio 9.0\VC\VCProjectDefaults ちゃんとCUDA VS Wizardのプロジェクトでもsm_20を指定できました。 CUDA VS Wizardの問題ではなく、Cuda.rulesの格納し忘れでした。失礼いたしました。 333 名前：デフォルトの名無しさん mailto:sage [2011/02/10(木) 23:09:14 ] CUBLASのcublasSetmatrix()にCUBLAS_STATUS_MAPPING_ERRORを 吐かれてしまうのですが、どういった理由が考えられますか？ 最初に最大サイズの行列サイズでallocした後、 そのdeviceメモリ領域に行列データを転送して演算、 結果をhostに書き戻して、さきほどの領域にまた別のサイズ（最大よりは小さい）の行列を hostから持ってきて演算・・・のような処理を繰り返しています。 334 名前：デフォルトの名無しさん mailto:sage [2011/02/15(火) 02:14:46 ] 卒論書き終わりました このスレのみなさんのおかげです 質問にいっぱい答えてくれてありがとう 335 名前：デフォルトの名無しさん mailto:sage [2011/02/16(水) 01:01:24 ] ちゃんと謝辞に書いたか？ 336 名前：デフォルトの名無しさん mailto:sage [2011/02/16(水) 05:21:54 ] 他のスレッドでは書き込めない超低レベルなチラ裏で恐縮ですが、 CUDA toolkitのサンプルのVecAddで、floatを全てdoubleに置き換えた上で １．ベクトル要素数Nをどこまで大きくできるか ２．CPUによるVecAddとの速度差はどの程度か を調べました。 ※GPU=GTS450 1GB、CPU=AthlonII-X4 640＠定格。CPUのVecAddは単純な3行ループで非マルチスレッド。 １．16*1000*1000までセーフ、17*1000*1000は実行時エラー(debug)。 つまりcudaMalloc()量にして384MBはセーフ、408MBはアウトでした（GTS450 1GB)。 VRAMいっぱいまで使えるわけではないのですね。 ２．clock()による計測ですが、 GPU→トータル 437 clock（内訳：cudaMalloc(ABC)=250, cudaMemcpy(AとBをhost2dev)=125, VecAdd=0, Memcpy(Cをdev2host)=62） CPU→172 clock でした(debug。releaseでも比はだいたい同じ)。 最初は内訳を出さず437と172をみてガッカリしたのですが、内訳を見てマジスンマセンwでした。 メモリ確保・転送のオーバーヘッドがそれなりにあるということは、 「host関数から高速な関数としてdevice関数を（頻繁に）呼び出す」のではなく、 極力device関数内だけで処理をさせるようにしないとうまみが少ないのですね。 337 名前：デフォルトの名無しさん mailto:sage [2011/02/17(木) 02:47:04 ] 気づいたらなんか本出てたんだけどどうあれ？ 入門書？応用書？ 338 名前：デフォルトの名無しさん mailto:sage [2011/02/17(木) 03:13:25 ] 最近？どの本のことですか？ 339 名前：デフォルトの名無しさん mailto:sage [2011/02/17(木) 04:43:56 ] CUDA by Example? 340 名前：デフォルトの名無しさん mailto:sage [2011/02/17(木) 14:56:50 ] >>339 そうそれ。2/14って書いてあったけど翻訳かな 341 名前：デフォルトの名無しさん mailto:sage [2011/02/17(木) 17:23:30 ] >>340 翻訳だよ。 内容は入門書 英文マニュアル読めるならいらない。 ストリーム、アトミック演算、OpenGLやD3Dとの連携のヒント 複数GPUを使う方法、ゼロコピーメモリ、ページロックメモリ テクスチャメモリとコンスタントメモリの使用法などが書いてある。 最適化については記述されてない。つまり ブランチダイバージェンス、GRAMのコアレスアクセス、シェアードメモリのバンクコンフリクト の説明はない 342 名前：デフォルトの名無しさん mailto:sage [2011/02/17(木) 22:51:23 ] CUDA by Example 汎用GPUプログラミング入門 ttp://www.impressjapan.jp/books/2978 NVIDIA社シニアエンジニアが執筆！　CUDAでハードウェア性能を引き出し、計算時間を劇的に減らそう 本書では、汎用GPU（GPGPU）上でのCUDA Cプログラミングの基礎からはじめ、スレッド間の連携、 コンスタントメモリ/テクスチャメモリの使用、アトミック/ストリームの活用、マルチGPUの利用、 DirectX/OpenGLとの相互運用などを解説。CUDA Cプログラミングのさまざまな重要ポイントや 注意点を取り上げています。サンプルを交えて解説しており、 CUDA Cの手法が具体的かつ明確にわかります。初級から中級への橋渡しともなる CUDA Cプログラマ必携の翻訳書です。 \2,940 (本体 \2,800＋税) 発売日：2011/02/14発売 343 名前：336 mailto:sage [2011/02/19(土) 01:05:39 ] 超低レベルな>>336の訂正で申し訳ないですが、 clock()によるGPUの実行時間は取得方法が誤っておりました。 スーパーコンピューティングニュース[東京大学情報基盤センタースーパーコンピューティング部門] www.cc.u-tokyo.ac.jp/publication/news/ にリンクが張られてある大島聡史先生のPDF： 「これからの並列計算のためのGPGPU連載講座 (�V) GPGPUプログラミング環境CUDA最適化編」 （201005_gpgpu2.pdf）に ＞funcname<<<xyz, xyz>>>(args1, args2); の形の呼び出しは非同期であるためこの前後で時刻を取得しても GPU上での計算時間は得られないと書かれてありました。 だから0clockになっていたんですね…。 （少なくとも処理後のcudaMalloc()の完了時点では処理は終わっているから 62clock未満ではあるのでしょうけど） 同pdfに、GPUの実行時間を調べられるCUDAプロファイラなるものが紹介されて いるのでこの使い方も覚えていこうと思います。 このpdfのシリーズはすごくわかりやすいと思いました。 こういうものが公開されているのはすごくありがたいです。 344 名前：デフォルトの名無しさん mailto:sage [2011/02/19(土) 08:58:53 ] このスレ読めばちゃんと書いてあるのに…… 345 名前：デフォルトの名無しさん [2011/02/19(土) 11:05:26 ] 公開されている資料といえばこんなのを見つけたが。 ttp://accc.riken.jp/HPC/training.html の下の方のCUDA。 346 名前：デフォルトの名無しさん mailto:sage [2011/02/19(土) 14:08:01 ] >>343 そういうあなたにCuda by exampleをすすめる 347 名前：デフォルトの名無しさん mailto:sage [2011/02/20(日) 01:43:34.23 ] 著者が必死に宣伝しているように見える 348 名前：デフォルトの名無しさん mailto:sage [2011/02/20(日) 01:54:08.17 ] まぁ青木本が1番 349 名前：デフォルトの名無しさん mailto:sage [2011/02/20(日) 02:46:20.30 ] SDKのドキュメントが一番だろう。 アレに必要なことは大体書いてある。 あとは実行してプロファイルとったり nvccでcubin吐いてenvytoolsで確認したり。 350 名前：デフォルトの名無しさん mailto:sage [2011/02/20(日) 05:37:37.97 ] 日本語でおｋ 351 名前：デフォルトの名無しさん mailto:sage [2011/02/20(日) 08:59:35.86 ] nVidia日本法人にまともなエンジニアがいないんだ。日本語だけじゃなんともならん 352 名前：デフォルトの名無しさん mailto:sage [2011/02/20(日) 18:02:58.21 ] SLIしてる場合、CUDAからはどういう扱いになるんでしょうか １つの大きなGPUとして見え１つのホストスレッドで扱えるのでしょうか それともSLIしていない場合と同じように複数のGPUとしてみえ、枚数分のホストスレッドで管理してやる 必要があるんでしょうか。 353 名前：デフォルトの名無しさん mailto:sage [2011/02/20(日) 18:24:47.06 ] 前に聞いたときはSLIはなんらCUDAに関与しないとのことだったけど。 354 名前：デフォルトの名無しさん mailto:sage [2011/02/20(日) 19:00:08.93 ] ということは複数GPU挿しているときと同様の扱いになるということですか。 ゲーム等ではSLIできてCUDAではできないんですね。 手軽さより自分で制御できる事を優先してるのかな 355 名前：デフォルトの名無しさん mailto:sage [2011/02/20(日) 19:07:14.32 ] もともとコアとしては遠くにあるわけだしVRAMも遠いんだから同一扱いは無理でしょ むしろSLIがうまく分担してるんじゃない 356 名前：デフォルトの名無しさん mailto:sage [2011/02/20(日) 19:39:29.66 ] SLIで充分なんだから本来グラフィックスにGTX480みたいなおばけチップはいらないんだよね 357 名前：デフォルトの名無しさん mailto:sage [2011/02/20(日) 19:44:13.09 ] でもそのおばけチップでSLIにしたらもっと性能あがるじゃん とくにCUDAスレの住民にはメリット大きいんじゃない？ 358 名前：343 mailto:sage [2011/02/20(日) 22:40:30.94 ] >>344 すみません、最初から読んだはずだったんですが、読み取れなかったか見逃していたみたいです。 >>345 情報ありがとうございます、これちょくちょく更新されているみたいですね。チェックしておきます。 >>346 まだ「はじめてのCUDAプログラミング」を買おうかどうか迷っているレベルなんですが、 行き詰まったらそれの購入も含めて考えます！ >>349 確かにProgramming Guideは英語の割にはわかりやすいと思います。 cuda cubinで検索したら「CUDA テクニカル トレーニング」という 日本語のpdf（内容はプレゼン用スライド・2008）がみつかりました。 www.nvidia.co.jp/docs/IO/59373/VolumeI.pdf nVidia公式にも日本語のドキュメントが増えてくれるとありがたいです。 一方この状況の中で翻訳したものをwikiで公開してくださっている方もいて ただただ感謝です(Fermi Compatibility GuideとFermi Tuning Guide)。 359 名前：デフォルトの名無しさん mailto:sage [2011/02/20(日) 23:40:26.16 ] >>358 まだなら「はじめてのCUDAプログラミング」を先に勧める。 英文が苦手ならさっさと本買ったほうが時間の節約になるし 実際のプログラムに時間を使った方が有益だと俺は思うよ。 360 名前：358 mailto:sage [2011/02/21(月) 00:02:16.28 ] >>359 そうですね…約2500円で省略できる時間のほうが有意義のように思います。 よし…買おう…bk1の300円ギフト券が手に入ったら…。 361 名前：デフォルトの名無しさん mailto:sage [2011/02/21(月) 10:12:48.81 ] >353 >前に聞いたときはSLIはなんらCUDAに関与しないとのことだったけど。 使えるっぽいぞ jp.slizone.com/object/sli_cuda_learn_jp.html 362 名前：デフォルトの名無しさん mailto:sage [2011/02/22(火) 00:44:44.56 ] SDｋにSLID3D10Textureてサンプルがあってソースみたんだが、 最初にDirectXとGPUのチェックがあるだけで 特にSLI独自のコーディングは見当たらないな かといって複数のGPUを制御するようなスレッド関係のコードもない それらからSLI組んでれば1つの大きなGPUとしてプログラムからは扱えるんじゃないかと 思ったんだがどうなんだろうか。 363 名前：デフォルトの名無しさん mailto:sage [2011/02/22(火) 23:08:18.71 ] >>361 なんだか滅茶苦茶な翻訳だな。 で、そこを読む限りPhysxとの連係には触れられているけど それ以上の情報はないし、英語サイトにはその記述さえ見当たらないね。 うーん、どうなんだろう 364 名前：デフォルトの名無しさん mailto:sage [2011/02/22(火) 23:39:43.06 ] 俺の友達も教務情報改造して人数出るようにしたりしてたけど日本人だったぞ 365 名前：364 mailto:sage [2011/02/22(火) 23:40:49.41 ] 間違えました ごめんなさい 366 名前：デフォルトの名無しさん mailto:sage [2011/02/24(木) 22:29:20.97 ] 最近出たインプレスの本にSLIのことが載ってた 367 名前：デフォルトの名無しさん mailto:sage [2011/02/24(木) 22:52:09.13 ] >>366 たしかに載ってるが概要の紹介程度で、SLIを使う方法や SLIからどうGPUが見え、どのように制御するか（orできるか）は書かれてない。 索引から逆引きしても１０行くらいしか書かれてないじゃん。 368 名前：デフォルトの名無しさん [2011/02/25(金) 05:40:01.78 ] 新型マックじゃ使えないの? 369 名前：デフォルトの名無しさん mailto:sage [2011/02/25(金) 06:17:44.62 ] Mac用のドライバもtoolkitもあるだろ あとはGPUが対応してるか調べろよ 370 名前：デフォルトの名無しさん [2011/02/25(金) 14:27:17.76 ] 使えない 371 名前：デフォルトの名無しさん [2011/02/28(月) 19:10:18.87 ] CUDAのアプリを作り始めたのですが、CUDAのSDKがインストールされていない環境で アプリを動かすと”cudart32_32_16.dll”が無いと怒られます。 cudart32_32_16.dll で使われている関数等々をスタティックリンクする方法があれば教えていただけないでしょうか？ CUDA 3.2を使っています。 372 名前：デフォルトの名無しさん mailto:sage [2011/02/28(月) 19:49:36.32 ] よくわからんけどコンパイル時に -L /usr/local/cuda/lib64/ -lcudart じゃだめ？ パスは適当 373 名前：デフォルトの名無しさん mailto:sage [2011/03/01(火) 00:11:54.00 ] >>372 リンカのオプションの追加の「依存ファイル」で「cudart.lib」は指定してあります。 ツールキットのcudart.lib のファイルサイズは約41KBのものが１つのみ。 cudart32_32_16.dll のファイルサイズは約375KBです。 スタティックリンク用のcudart.libが無いということでしょうか？？？ 開発環境を詳しく書きます [開発環境] OS: Win7 64bit IDE: Visual Studio C++ 2008 Express CUDA Toolkit: ver 3.2 32bitバージョン 374 名前：デフォルトの名無しさん mailto:sage [2011/03/01(火) 00:47:03.69 ] >>371 CUDA RuntimeのDLLは再配布が許可されているみたいだし 単に他の環境で動かしたいだけなら、DLL同梱でいい気がする。 ttp://developer.download.nvidia.com/compute/cuda/3_2_prod/toolkit/docs/EULA.txt 375 名前：デフォルトの名無しさん mailto:sage [2011/03/01(火) 02:29:40.71 ] >>374 ほむ！ ありがとうございます。 インストールしたときについてくるライセンスではそこら辺がよくわからなかったので助かります。 とりあえず、DLLも同梱しようと思います。（スタティックリンクしたいけど・・・） 376 名前：デフォルトの名無しさん mailto:sage [2011/03/01(火) 03:33:24.52 ] 4Gamer.net ― NVIDIA，「CUDA 4.0」の新機能を予告。より広い層の開発者にアピールする新バージョン（CUDA） www.4gamer.net/games/076/G007660/20110228018/ Fermi持ってたら恩恵受けられるのかな。なんかよくわからんけど楽しみだ。 ＞ちなみにGPUDirect 2.0とThrustをフルサポートするGPUは，Fermiアーキテクチャ限定となる。 377 名前：デフォルトの名無しさん mailto:sage [2011/03/01(火) 03:56:33.42 ] >>376 ・GPUDirect 2.0 …スパコンには良さそう？ ・Unified Virtual Addressing…どれが何のメモリなのかごっちゃになりそう？ ・Thrust…STL。。。どうでもよさそう？ 378 名前：デフォルトの名無しさん [2011/03/01(火) 23:32:35.21 ] こんにちは。cubinとptx関連の情報が欲しくてごにょごにょいじっていたのですが ちょっと行き詰まりましたのでできれば２点ご教授いただきたく… １． CUDA 3.2 SDKについてくるtemplateプロジェクトのひな形を使ってアプリを書いているのですが、 プロジェクトの[プロパティ]→[構成プロパティ]→[CUDA Runtime API]→[GPU] →[NVCC Compilation Type]で"Generate 32 bit .cubin file (-m32 -cubin)"を 選択してコンパイルしたところ "1>nvcc fatal : redefinition of argument 'machine'" と怒られました。おそらく"--machine 32" と、cubinオプションを選択したときに付く、"-m 32" の重複が怒られていると思うのですが "-m 32"はcubin形式をプルダウン選択で選ぶと自動的に付いて、 "--machine 32"はどこで付くのかわからない状況です(Win32アプリには自動で付く？？？)。 どうすれば競合を解消できるでしょうか？ ２． 仕方がないので、とりあえずコマンドラインから -Xptxas="-v" -cubin --ptxas-options -v を指定してコンパイルしたところ 無事にcubinファイルが出来たのですが、このcubinファイルはバイナリー形式で私の目では読むことが出来ませんでした。 このcubinをテキスト形式で読むにはどうしたら良いでしょうか？ (CUDA 3.0からcubinは一旦バイナリ形式で出力されるようになったようです) [環境] CUDA Toolkit, SDK 3.2 379 名前：デフォルトの名無しさん mailto:sage [2011/03/02(水) 00:46:11.36 ] ptx出力じゃ、用が足らんのけ? 380 名前：デフォルトの名無しさん [2011/03/02(水) 13:54:56.94 ] 初心者です。 今からくだ勉強したい場合はCUDA4.0がオススメですか？ 381 名前：デフォルトの名無しさん mailto:sage [2011/03/02(水) 15:15:15.23 ] >>380 むしろcuda2.2くらいから。 基礎ができてから発展にすすめ 382 名前：デフォルトの名無しさん [2011/03/02(水) 15:23:23.73 ] 381さん有り難うございます。cuda2.2からやってみます。 383 名前：デフォルトの名無しさん mailto:sage [2011/03/02(水) 15:39:44.51 ] 古いバージョンのCUDAを知る必要は普通はない。 CUDA 4.0でも基本的にこれまでの機能を全て備えている。 また古い世代のGPUは余裕がある場合だけ追うべき。 384 名前：デフォルトの名無しさん mailto:sage [2011/03/02(水) 15:51:19.20 ] >>383 同意 つか、 >cuda2.2くらい "くらい"なんて中途半端な答え方してるのかがワカラン 385 名前：デフォルトの名無しさん mailto:sage [2011/03/03(木) 23:20:08.40 ] >>379 readelf でcubinみましたが、いまいちよくわからなかったのでptxみます。 まあ、ptxもよくわからないのですが… 386 名前：デフォルトの名無しさん mailto:sage [2011/03/04(金) 11:17:18.22 ] PTXを読んで分からない人にはcubinは無理。 387 名前：デフォルトの名無しさん mailto:sage [2011/03/04(金) 20:05:59.06 ] CUDA 4,0 RC が来たよ 388 名前：デフォルトの名無しさん mailto:sage [2011/03/06(日) 02:29:43.51 ] Visual Studio 2010 まだ来ないよ 389 名前：デフォルトの名無しさん [2011/03/08(火) 02:57:27.52 ] すみません。 どなたかAMBER(Assisted Model Building with Energy Refinement)というソフトをCUDAで動かしてる人いませんか？ このソフトを Cent OS5.5 x86_64(カーネル 2.6.18-194.el5) インテルコンパイラーとの併用で動かそうと何度も試みているのですがなかなか上手くいきません。 不具合の原因はGCCのバージョンとインテルコンパイラーのバージョンの組み合わせにあると睨んでいるのですが… もし、動かせている方がいらしたらGCCのバージョンとインテルコンパイラーのバージョンの組み合わせををヒントとして教えてほしいのですが 390 名前：デフォルトの名無しさん mailto:sage [2011/03/08(火) 22:01:42.39 ] >>389 コンパイルでエラーなのかリンカでエラーなのかわからん。 取り合えずsdkのサンプルでiccで動かしてみたら？それで問題がわかるでしょ。 それが出来ないなら、 そもそもnvcc自体がgccのラッパだし、 素直にgcc使えば？ 391 名前：デフォルトの名無しさん [2011/03/09(水) 22:09:17.79 ] デバイス関数で、VC++でいうところの Sleep(0)、gccでいうところのsleep(0)？？？と同等の動作をする関数を探しているのですが 適切な関数があれば教えていただけないでしょうか。 392 名前：デフォルトの名無しさん mailto:sage [2011/03/09(水) 22:12:55.97 ] 眠らす必要がなさそうだから無いと思う 393 名前：デフォルトの名無しさん mailto:sage [2011/03/09(水) 22:49:28.90 ] >>392 フロントエンドというかデスクトップというか、とにかくグラフィックスの描画が遅くなる（当たり前ですが…）のと、 GPUの温度がかなり上がるので、 デバイス関数のループの中に、ちびっとだけSleep(0) 的なものを入れたいのですが、 関数的なものは用意されていないっぽいですかね。 例えば、どれか１つだけスレッドの処理ループを長引かせて、__SyncThreads() でもしたらいいのかしらん…？？？ 394 名前：デフォルトの名無しさん mailto:sage [2011/03/09(水) 22:53:11.30 ] 利用者の大半はカーネルを一気に計算したい派だとおも 熱いから休憩するって発想は俺含めてないんじゃね？ 395 名前：デフォルトの名無しさん mailto:sage [2011/03/09(水) 23:41:01.50 ] 自分はカーネルを終了させる必要があると思います 396 名前：デフォルトの名無しさん mailto:sage [2011/03/09(水) 23:58:40.99 ] >>395 例えば、 __global__ void test() { 　　... 　　for (i = 0; i < LOOP; i++) 　　{ 　　　　__syncThreads(); 　　　　... 　　　　// Sleep(0) 的なもの？？？ 　　　　if (blockIdx.x == 0 && blockIdx.y == 0) 　　　　{ 　　　　　　// 何か冗長な処理 　　　　} 　　} } って感じじゃSleep(0)の代わりにならないですかね？ 397 名前：デフォルトの名無しさん mailto:sage [2011/03/10(木) 00:06:55.92 ] それはビジーウェイトになんね？ 結局熱くなってしまうんだが 398 名前：デフォルトの名無しさん mailto:sage [2011/03/10(木) 00:16:02.21 ] >>397 そうなんですか… ビジーウェイト調べてみます。 399 名前：デフォルトの名無しさん mailto:sage [2011/03/10(木) 04:14:12.92 ] >>396 カーネルを分割すればいいんじゃね? わざわざ遅くする理由を思いつかないから試す気にもなれないけど。 400 名前：デフォルトの名無しさん mailto:sage [2011/03/10(木) 09:03:17.86 ] >>399 これが正解だろう。 でもそもそも、ちょっと休ませたくらいじゃ熱なんて変わらんだろうから、ドライバの設定でクロックを下げたらいい。 もしくは、放熱を強化するかだろう。 401 名前：デフォルトの名無しさん [2011/03/17(木) 13:38:21.43 ] CUDAでGTX460(1.5GHz)、4枚分程の計算能力を求めています。 GTX460を2枚刺ししたPC2台と、4枚刺ししたPC1台のどちらを組む方が良いでしょうか？ 4枚刺しですと、電源のとりまわしと冷却に不安を感じております。 402 名前：デフォルトの名無しさん mailto:sage [2011/03/17(木) 14:33:27.94 ] GPU間の通信はどれくらいしますか。 403 名前：402 mailto:sage [2011/03/17(木) 17:10:47.41 ] そういうのを意識しないなら、なるべくわけた方が安定しそうですね。 404 名前：デフォルトの名無しさん mailto:sage [2011/03/17(木) 17:47:16.68 ] >>401 単純に、書かれた内容だけで判断する限り前者の方が計算能力が高い。 但し、後者により強力なCPUとメモリを搭載する場合は一概に言えない。 405 名前：デフォルトの名無しさん mailto:sage [2011/03/21(月) 23:08:34.74 ] >>402 GPU間の通信は無いです。 >>403-404 GPU間の通信が無く、CPU側の処理も軽いので 4枚刺しによる恩恵が特になさそうなので 2台で組んだ方がよさそうですね。 ありがとうございました。 406 名前：デフォルトの名無しさん mailto:sage [2011/03/22(火) 00:55:46.70 ] GPU計算早すぎ、ワロタ。 長い目で見ると、DSPとかで食いつないでた信号処理屋的に失業の危機じゃん… 407 名前：デフォルトの名無しさん mailto:sage [2011/03/22(火) 05:48:58.16 ] GPUメモリのTLBミスが気になるのですが、何か情報ないでしょうか 408 名前：デフォルトの名無しさん mailto:sage [2011/03/22(火) 10:14:13.39 ] >>406 カスタムチップによる高速演算ボードを出していたところは尽く撤退しました。 409 名前： [―{}@{}@{}-] デフォルトの名無しさん mailto:sage [2011/03/22(火) 23:46:07.04 ] sharedメモリへのバンクコンフリクトは避けたほうが良いって話だけど、 共有メモリが16バンクに分割されているってのは�@�Aのどっちだろ。 �@マルチプロセッサのもつ16384Byteが16分割されてる �Aブロックに割り当てた共有メモリが16分割される ハードウェアが固定されてそうだから�@の気がするんだけど…。 410 名前：福盛俊明 [2011/03/23(水) 02:45:25.67 ] アハ〜♪” 411 名前：デフォルトの名無しさん mailto:sage [2011/03/23(水) 13:39:51.91 ] GPUDirect 2.0について質問です． GeForce GTX 580を二枚搭載したPC(CUDA 4.0 RC インストール済み)を使って， GPUDirect 2.0を用いたプログラムを作成しようと思っています． cudaDeviceCanAccessPeerという関数を用いて， GPUDirect 2.0が使用可能かどうか調べたのですが， 結果が0(=出来ない)ということでした． FermiでもGeForceではダメだということでしょうか？ それとも，GPUDirect 2.0は Peer-to-Peer Memory Access(←私が勝手にGPUDirect 2.0だと思っている) とは別のものなのでしょうか？ 試された方はここに居らっしゃいますでしょうか？ 412 名前：デフォルトの名無しさん mailto:sage [2011/03/23(水) 17:48:37.13 ] CanAccessPeerは無視してよい。 確かにTeslaシリーズでは1が、GeForceシリーズでは0が返る。 だけど実際にはGeForceでも通信はできる。 UVAだかなんだかに対応するGPUなら問題ないのかも。 413 名前：デフォルトの名無しさん mailto:sage [2011/03/23(水) 23:13:07.00 ] ワープってやっぱりSMの上限いっぱいまで詰め込んだほうがいいの？ 例えばSMの上限ワープが32のところに16個のワープにしても、 メモリアクセスのようなレイテンシがないなら問題ない気がするんだが。 414 名前：デフォルトの名無しさん [2011/03/24(木) 11:46:09.16 ] CUDA初心者です。 CUDAを使って、ゲームエンジンのUnrealとかcryを使ってる方いらっしゃいますか。 参考になるサイトを教えてほしいのですが 415 名前：デフォルトの名無しさん [2011/03/24(木) 20:35:29.92 ] 例えば、デバイス側の主要な処理が下のような論理演算であった場合、 論理演算部分のアセンブラ化による高速化は望めるでしょうか？ それとも、レジスタの使い方がコンパイラの方がうまくて、 より高速なアセンブラを書くのは難しいでしょうか？ __global__ void test(int *a, int *b, int *c) { 　　// 初期化処理 　　 　　// メインのループ 　　for (i = 0; i < LOOP; i++) 　　{ 　　　　a[i] = (b[i]<<10 | (b[i] >> 22)) ^ c[i]; 　　} } 416 名前：デフォルトの名無しさん mailto:sage [2011/03/24(木) 21:10:44.15 ] 少なくともその論理演算部分をアセンブラで書く必要はない。 そもそもメモリアクセスがボトルネックにならないほど多くの演算があるのかどうか。 417 名前：デフォルトの名無しさん mailto:sage [2011/03/24(木) 21:50:24.37 ] 分岐を避けるためにビット演算使いたくなるけど、いまいち貧弱なんだよなあ・・・ 418 名前：デフォルトの名無しさん [2011/03/25(金) 12:10:15.17 ] 英語か日本語でptxファイルの読み方を解説しているドキュメントがあれば教えていただけないでしょうか。 また、コンパイル時にptxファイルに明示的にコメントを挿入するコマンドがあれば教えていただけないでしょうか。 例えば、test.cu に、「__PTX_COMMENT("foo");」という行を記述すると、ptxファイルに // foo といった行が挿入されるという具合のものを探しています。 419 名前：デフォルトの名無しさん mailto:sage [2011/03/25(金) 13:36:10.03 ] PTXのマニュアルはToolkitに含まれているかと。 ptxにブロックごとに行番号情報は含まれているのでソースのどこに相当するかは分かるのでは。 特定の場所が知りたければ例えば__syncthreads()を１０個入れとくとか。 420 名前：デフォルトの名無しさん mailto:sage [2011/03/25(金) 15:16:09.98 ] perl とかでつくれそうだよね 421 名前：デフォルトの名無しさん mailto:sage [2011/03/25(金) 17:03:42.70 ] >>419 なるほど。例えば、ダミーでfoo0()とか、foo1()とか関数を作ってコメント代わりに呼び出せばよいと。 >>420 コマンドとしては無いということですね。 422 名前：デフォルトの名無しさん mailto:sage [2011/03/28(月) 20:51:27.79 ] CUFFTのハンドル確保するだけでエラーが出る…。↓エラーのでるコード。 CUT_DEVICE_INIT(argc, argv); cufftHandle plan; cufftPlan1d(&plan, 4096, CUFFT_C2C, 16); cufftDestroy(plan); ↓エラー First-chance exception at 0x000007fefd96aa7d in CudaTest.exe: Microsoft C++ exception: cudaError_enum at memory location 0x001ef830.. なずぇ・・・。cufftExecC2C走らせるとFFTは実行して変換結果も正しいのだが謎す。 423 名前：デフォルトの名無しさん mailto:sage [2011/03/29(火) 11:17:10.20 ] cufftは内部であれこれ怪しげなことしているからなぁ。 VRAMは未だ余っている筈なのにエラーになったり、作ったプランを使わずに破棄するとこけたり…… って、>422はそのパターンか? 424 名前：デフォルトの名無しさん mailto:sage [2011/03/30(水) 14:28:25.63 ] >423 cufftExecC2C実行してもやっぱでる。 海外フォーラムみてたらver 3.2が怪しくて ver3.0にしたら消えたって書いてあったんで ver3.0にしてみたがやっぱでる。うーん・・・。 425 名前：デフォルトの名無しさん mailto:sage [2011/03/30(水) 16:09:16.17 ] >>424 うーん、SDKとcufftの組み合わせで相性あるしなぁ。 あと、サイズが2の冪かどうかでも変わるかもしれない。 つーか、3.1よりは3.2の方がましだった気がするんだが…… 426 名前：デフォルトの名無しさん mailto:sage [2011/04/02(土) 17:59:42.31 ] サンプルはmain関数でCUDAを利用するのばっかだけど、 VC++でダイアログから.cuの関数使うようなサンプルってなんかない？ 427 名前：デフォルトの名無しさん mailto:sage [2011/04/02(土) 18:01:41.56 ] >>426 おまいさんは関数を分割することもできんのかね。 428 名前：デフォルトの名無しさん mailto:sage [2011/04/02(土) 19:08:36.61 ] windowsよく知らんけど それはライブラリ作んなくてもできるんかい？ 429 名前：427 mailto:sage [2011/04/02(土) 19:39:35.97 ] >427 extern "C" つけたらできたや。ごめんごめん。 430 名前：デフォルトの名無しさん mailto:sage [2011/04/02(土) 20:08:34.70 ] >>429 お前はだれ？ そんで誰にレス支店の？ 431 名前：425 mailto:sage [2011/04/02(土) 21:07:05.73 ] まちがえちゃった。テヘッ 432 名前：デフォルトの名無しさん mailto:sage [2011/04/04(月) 17:39:36.21 ] プログラミングガイドに記載されているように CU_SAFE_CALL(cuInit(0)); CU_SAFE_CALL(cuDeviceGet(&device, 0)); CU_SAFE_CALL(cuGLCtxCreate(&context, 0, device)); と初期化すると，cuGLCtxCreateで Cuda driver error 3e7 というエラーが発生してしまいます． 詳しい方居らっしゃいますでしょうか？ 当方の環境は Windows xp 32bit CUDA 4.0 RC ドライバは270.51を使用しています． 433 名前：デフォルトの名無しさん mailto:sage [2011/04/04(月) 22:25:10.79 ] CU_SAFE_CALL(cuInit(0)); CU_SAFE_CALL(cuGLInit(0)); CU_SAFE_CALL(cuDeviceGet(&device, 0)); CU_SAFE_CALL(cuGLCtxCreate(&context, 0, device)); 434 名前：デフォルトの名無しさん mailto:sage [2011/04/04(月) 22:26:11.47 ] ミスった。 1点だけ。 uGLInit(0)→uGLInit() 435 名前：432 mailto:sage [2011/04/05(火) 11:10:29.98 ] >>433-434 この通りにしましたところ，今度はcuGLInitで Cuda driver error c9 が発生しました... 436 名前：デフォルトの名無しさん [2011/04/05(火) 17:48:55.53 ] コマンドラインからnvccでコンパイルする際、普通のC++部分のコンパイルを Intelコンパイラ："icl.exe"に行わせたいのですが、どのようにコマンドを記述したら良いでしょうか？ 環境： Windows 7 CUDA 3.2 437 名前：デフォルトの名無しさん mailto:sage [2011/04/05(火) 17:59:21.74 ] nvccしか使えないところはライブラリ化してやってるけど もっといい方法あるんかいな？ 438 名前：デフォルトの名無しさん mailto:sage [2011/04/05(火) 20:34:55.57 ] 西日本も福島原発の放射能に曝される。 4/7 予測 up3.viploader.net/ippan/src/vlippan198234.jpg 発表はドイツ気象庁 www.dwd.de/ 439 名前：デフォルトの名無しさん mailto:sage [2011/04/05(火) 21:13:26.38 ] CUDAで予測出来ないんすか？ 440 名前：デフォルトの名無しさん mailto:sage [2011/04/05(火) 23:50:34.55 ] できますん 441 名前：デフォルトの名無しさん mailto:sage [2011/04/06(水) 02:41:41.96 ] どっちですか 442 名前：デフォルトの名無しさん mailto:sage [2011/04/06(水) 07:54:08.27 ] 日本語で予測できるか質問するくらい意味がないことに気付け 443 名前：デフォルトの名無しさん mailto:sage [2011/04/06(水) 10:06:08.15 ] >>442 さすがにそれはちょっと違うと思うけど 気象系シミュレーションでGPUつかってるところって結構あるよね？ 444 名前：デフォルトの名無しさん mailto:sage [2011/04/06(水) 17:09:45.35 ] ？ 445 名前：デフォルトの名無しさん mailto:sage [2011/04/06(水) 17:21:47.83 ] さすがにこれは？？？だな 446 名前：デフォルトの名無しさん mailto:sage [2011/04/06(水) 18:02:23.22 ] >>435 cuInit cuDeviceGet cuCtxCreate cuGLInit cuGLCtxCreate じゃね? 447 名前：432 mailto:sage [2011/04/07(木) 17:59:04.18 ] >>446 やはりcuGLInitで同様のエラーが出ます． 448 名前：デフォルトの名無しさん mailto:sage [2011/04/07(木) 18:05:26.23 ] 3.2でやってみれ 449 名前：432 mailto:sage [2011/04/07(木) 18:20:08.71 ] >>448 3.2でも同様です． 450 名前：デフォルトの名無しさん mailto:sage [2011/04/07(木) 19:40:03.31 ] OpenGLは初期化してるよね? 451 名前：432 mailto:sage [2011/04/08(金) 08:52:38.77 ] >>450 大変御迷惑おかけしました． 自分の無知さに絶望しました． OpenGLの初期化後に CU_SAFE_CALL(cuInit(0)); CU_SAFE_CALL(cuDeviceGet(&device, 0)); CU_SAFE_CALL(cuGLCtxCreate(&context, 0, device)); で動作させることが出来ました． 御手数おかけしました． 452 名前：デフォルトの名無しさん mailto:sage [2011/04/08(金) 20:02:33.08 ] シュミレーション 453 名前：デフォルトの名無しさん mailto:sage [2011/04/08(金) 21:02:09.88 ] あながち間違ってはいないな 454 名前：デフォルトの名無しさん mailto:sage [2011/04/09(土) 16:43:22.41 ] cudaMallocとかcudaFreeとかcudaMemcpyとか、どのヘッダに定義されてんの？ cutil.hにもcuda.hにも定義されてねーのだが。 455 名前：デフォルトの名無しさん mailto:sage [2011/04/09(土) 17:05:25.29 ] cuda_runtime.h 456 名前：デフォルトの名無しさん mailto:sage [2011/04/09(土) 17:10:51.07 ] 　　　　＿＿＿ 　　 ,;f　　　　　ヽ 　　　　 　　　 　　i:　　 　　 　　 i 　　 　　|　　　　　　　　|　　///;ﾄ, 　　|　　　　^　　^　)　////ﾞlﾞl; 　　 　　(.　　>ﾉ(､_, )ヽ､} l 　 .i .! | 　　 　 ,,∧ヽ　!-=ﾆ=- | │　 　| .| ／＼..＼＼｀ﾆﾆ´ !,　{　　 .ﾉ.ﾉ ／　 ＼ ＼￣￣￣../　　 / .| 457 名前：デフォルトの名無しさん mailto:sage [2011/04/10(日) 01:00:12.43 ] いまだにキューダと読んでしまうのだが、クーダと読む習慣が付いてる人はどれくらいいるんだろう 458 名前：デフォルトの名無しさん mailto:sage [2011/04/10(日) 11:48:25.16 ] あなた以外ほとんどがクーダだと思います 459 名前：デフォルトの名無しさん mailto:sage [2011/04/10(日) 12:13:51.20 ] main.cu メイン ドット クー？ 460 名前：デフォルトの名無しさん mailto:sage [2011/04/10(日) 12:41:02.43 ] OpenGLとの連携が遅いでござる。 フォーラム見ても遅い遅いと書いてあるし、どうしようもないのかこれ？ 461 名前：デフォルトの名無しさん mailto:sage [2011/04/10(日) 18:44:20.12 ] warpはなんて読む？ 462 名前：デフォルトの名無しさん mailto:sage [2011/04/10(日) 18:48:15.64 ] ﾜｰﾋﾟｰ 463 名前：デフォルトの名無しさん mailto:sage [2011/04/10(日) 19:12:02.26 ] ウォープ 464 名前：デフォルトの名無しさん mailto:sage [2011/04/10(日) 19:56:18.22 ] 飯野の巣 465 名前：デフォルトの名無しさん mailto:sage [2011/04/11(月) 20:09:52.50 ] warning walking working star wars それぞれなんて読む？ 466 名前：デフォルトの名無しさん mailto:sage [2011/04/11(月) 20:23:49.35 ] それ、この板で何度も繰り返し振られる話題だよな 467 名前：デフォルトの名無しさん mailto:sage [2011/04/11(月) 20:38:03.05 ] >>465 ワーニング ワーキング ウォーキング スターウォーズ 468 名前：デフォルトの名無しさん mailto:sage [2011/04/11(月) 20:41:31.72 ] スターワーズだりょ 469 名前：デフォルトの名無しさん mailto:sage [2011/04/11(月) 20:42:58.66 ] threadIdxがidのxだといつも勘違いしてしまう 470 名前：デフォルトの名無しさん mailto:sage [2011/04/11(月) 21:04:45.77 ] stdio をスタジオだといつも勘違してしまうのとはちょっと違うか 471 名前：デフォルトの名無しさん mailto:sage [2011/04/11(月) 22:09:14.04 ] カタカナ表記した時点で正しい読みから外れてる 472 名前：デフォルトの名無しさん mailto:sage [2011/04/11(月) 22:31:58.48 ] >>469 わかるわ 油断するとIdxとIdyがあると思っちゃう threadIdx.y xのy？？？ スレッドはフルスペルなんたからthreadIndexでいいじゃねえか 473 名前：デフォルトの名無しさん mailto:sage [2011/04/11(月) 23:56:43.33 ] 寧ろptx出力宜しくtid.x希望w 474 名前：デフォルトの名無しさん mailto:sage [2011/04/12(火) 22:35:25.63 ] 明示的にグローバルメモリから共有メモリへの読み込みを行う命令ってないのかね。 スレッドに分けて暗黙にコアレスアクセスさせると、どうもコードが汚くなる。 475 名前：デフォルトの名無しさん mailto:sage [2011/04/13(水) 01:04:12.90 ] ううーん、デバイス側にメモリ確保しっぱなしだとOpenGLの性能が著しく低下する…。 計算結果を最終的にOpenGLで表示するような場合だと、無駄なメモリコピーは仕方ないものとして 計算結果をホスト側に逐一戻したほうがよさげ？ 476 名前：デフォルトの名無しさん [2011/04/19(火) 21:32:44.77 ] 3.2から4.0 RC2に変更したんだけど，Cuda.rulesってないの？ 477 名前：デフォルトの名無しさん mailto:sage [2011/04/20(水) 01:10:42.96 ] つかなんでcommon\lib\Win32の中にcutil32.dllがあるんだっていう 478 名前：デフォルトの名無しさん [2011/04/20(水) 01:11:01.88 ] CUDAのアセンブラにMMXやSSEのような命令はあるでしょうか？ （CUDAのカーネル関数の中で、MMXのPADDDのような32bit整数の4並列計算のようなことができなかと思っています） 479 名前：デフォルトの名無しさん mailto:sage [2011/04/20(水) 01:16:07.37 ] >>478 CAL+IL 480 名前： 忍法帖【Lv=11,xxxPT】 mailto:sage [2011/04/20(水) 04:19:17.84 ] 2枚差しの場合、2つプログラムを走らせたら、それぞれ別のGPUを使ってくれるのですか？ 481 名前： 忍法帖【Lv=15,xxxPT】 mailto:sage [2011/04/20(水) 07:08:07.39 ] 現状では自分で何番目のGPUを使うか指定する必要があります。 482 名前：デフォルトの名無しさん mailto:sage [2011/04/20(水) 11:51:01.49 ] >>477 イミフ 483 名前：デフォルトの名無しさん [2011/04/20(水) 20:42:07.79 ] >>479 えと、CUDA用には無いということでしょうか？ add.u32 を2並列(64bitレジスタ)、あるいは4並列(128bitレジスタ)で実行するような命令セットはないでしょうか？ 484 名前：デフォルトの名無しさん mailto:sage [2011/04/20(水) 21:24:43.20 ] >>483 そのためのSIMT 485 名前：デフォルトの名無しさん [2011/04/20(水) 22:27:22.88 ] >>483 何かを誤解しているようだが、CUDAには必要ないから無い。 SSEを使わないと100%の性能が出せないとかいうCPUのような制約は無い。 add.u32という命令ならほっといてもWarp単位の３２スレッドが（ほぼ同時に）実行する。 全命令が３２並列のSIMD命令になっていると考えればよい。（３２ビットデータの場合） 486 名前：デフォルトの名無しさん [2011/04/20(水) 22:50:57.29 ] >>484-485 CUDAがSIMTなのか、SIMDなのか、SPMDなのかはさておいて… 昔の記述だとCUDAアーキテクチャのレジスタは32bitとかだったと思いますが、 今時のNVIDAのGPUには64bitレジスタ、あるいは128bitレジスタがあって、それを使って padd.u32 みたいな命令により、「1スレッド内の演算において」32bit整数の4並列処理が出来るとかないでしょうか？ ということなのですが、いかがでしょうか？ Streamでは128bitレジスタで32bit4並列の演算ができるみたいのですが・・・ 487 名前：デフォルトの名無しさん mailto:sage [2011/04/21(木) 04:36:54.19 ] そこまでやるならＣＰＵの方が早かったりしてｗ 488 名前：デフォルトの名無しさん [2011/04/21(木) 08:16:13.34 ] >>486 そんなものはないから誤解だと言ってるんだが。 AMDは１スレッドが4 or 5演算のVLIWで動作するから4並列のSIMD命令化の意味があるが CUDAは１スレッドが各種演算ユニットを１ずつしかもたない(or複数スレッドが共有）CUDA Coreで 実行されるので128bit型を用いたとしても４命令になるだけ。 489 名前：デフォルトの名無しさん mailto:sage [2011/04/21(木) 09:31:19.38 ] >>488 誤解というか、質問なのですが・・・ 今どきのCUDAでもレジスタレベルのSIMD命令は無く、 ハードウエア的に、レジスタも演算器も（今のところは）そのような用途には向いていないということですね。 ありがとうございましたゞ 490 名前：デフォルトの名無しさん mailto:sage [2011/04/21(木) 09:35:39.70 ] それよりSandyBridgeのベクタ演算機の性能はどうなんだろうな まだ試験実装だからコア数が少ないけどメインメモリ直結だからけっこう早そうなんだが 491 名前：デフォルトの名無しさん mailto:sage [2011/04/21(木) 09:59:44.72 ] その辺は簡単な報告を以前SSEスレにあげといたよ イミフな反応が多々あったけどね 492 名前：デフォルトの名無しさん mailto:sage [2011/04/21(木) 10:04:26.92 ] >489が根本的に解ってない悪寒 493 名前：デフォルトの名無しさん mailto:sage [2011/04/21(木) 21:59:22.93 ] >>492 論理レジスタとか、物理レジスタとかいう話しですか？？？ >>486への根本的な回答はどうなりますか？ 494 名前：デフォルトの名無しさん mailto:sage [2011/04/21(木) 22:29:01.49 ] とりあえず公式文書を読んでから質問してくれ。 495 名前：デフォルトの名無しさん mailto:sage [2011/04/21(木) 22:39:09.84 ] ようするに64bitレジスタに8x8のデータを入れて同時に処理出来るのかってことでしょ streamsdkには専用の命令があると？ cudaはそんなものは勝手にやると言ってる人が居るけど、どうやってやるんだ？ cudaにはそういう命令は無いよ 内部で処理することが出来るとは思えないけどな 少なくともユーザープログラムからやる方法はないよ 内部的にはそういう命令も持ってるかもしれないけど というかそういうアセンブラレベルの最適化をしなくていいようにしたのがＣＵＤＡ言語なわけで 複雑化する要因はパフォーマンスを犠牲にしてでも排除されるだろう どうしてもやりたかったら64bit変数を使って自前でやるしかないだろう 496 名前：デフォルトの名無しさん mailto:sage [2011/04/21(木) 23:07:04.51 ] >>495 >>486 への回答は”No”ということですね。 ありがとうございます。 将来的には"Yes"になると思っているのですが、そこら辺が「根本的に解ってない」と言われる所以ですかね。 「根本的に解ってない」らしいので、何がなんだかさっぱりですが・・・ # 今、スクラッチのCUDA用アプリで、SSEで書かれたCPU用アプリの約5倍（CPU側はSandyBridgeで4スレッドで動作）の速度が出ているのですが、 CPU側がきちんとAVXに対応してきたらこの差はかなり縮まりそうなんですよね・・・ 497 名前：デフォルトの名無しさん mailto:sage [2011/04/21(木) 23:11:08.25 ] ねえ、GPUの速さの要因ってCPUより圧倒的に軽量なスレッド数によるもんじゃないの まずそっからじゃないのかね？ 498 名前：デフォルトの名無しさん mailto:sage [2011/04/21(木) 23:14:37.36 ] >>496 さっきからすげー鼻につくんだけど 教わる態度じゃないよお前 教えてんの俺じゃないけどさ 499 名前：デフォルトの名無しさん mailto:sage [2011/04/21(木) 23:22:41.50 ] あれだ、自分が書いたcudaのコードが思ったはど速度が出ないから 腹癒せしているか言い訳探しているかってところだろ。 今日たまたま計測していたんだが、cufftでもfftwfの3倍程度だった。 汎用ルーチンじゃこんなもんだよ。 500 名前：デフォルトの名無しさん mailto:sage [2011/04/21(木) 23:51:10.29 ] >>498 CUDAの将来的な見通しについて知りたかった部分もあったので 失礼に見えた部分があるかもしれません。 ごめんなさい。 501 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 00:46:21.02 ] SIMDみたいなものつけて各コアを高機能にするよりは、 単純なコアの数を増やす方向に進むだろうなあ。 502 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 00:51:40.28 ] レジスタ/SM増やして欲しい… 503 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 01:02:26.93 ] GF100やGF104の倍精度（64bit）の演算って単精度（32bit）の1/4のスループットとかじゃなかったかな？ おまけにGF104は倍精度をサポートしているコア数は1/3だし。 更にHPC向けではなくグラフィック向けの物は制限かかっているらしいし。 倍精度が必要になるなら仕方ないだろうけど、32bit変数で足りるのに64bit変数を使うのはデメリットの方が圧倒的に多そう。 504 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 01:08:10.32 ] いろいろ試したがＣＰＵと比較して４倍程度しか上がらないし 労力掛かるわりには環境依存甚だしいわ ４倍くらいだったら別に我慢出来ないようなものなんてないし マルチコアＣＰＵで普通に組のが一番いいという結論に達した 505 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 01:11:09.84 ] >>504 計算したい内容を簡潔に書いてみ 計算内容によってＧＰＵに向き不向きがあるから 506 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 01:12:19.97 ] >>497 膨大なスレッド数と、その切り替えの軽量さを活かして メモリアクセスのレイテンシの隠蔽とかだっけ。 >>502 スレッドあたりの変数が多いと苦労するね。 そういう場合はシェアードメモリをうまく使って切り抜けるしかないのかな。 507 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 01:15:58.96 ] CUDAは4nbyteを1要素とした16要素SIMDって感覚だわ。(n = 1,2,3,4) メモリアクセス考えると結局16要素単位の処理の塊になってしまう。 508 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 01:19:14.16 ] >>506 共有使ってもたらん Sharedも増やせ！キャッシュも増やせ! 509 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 01:31:49.64 ] 向いてないよ。 510 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 07:57:21.85 ] CUDAとかGPGPUの普及の足引っ張ってるだけだよなぁ。 早くintel、amd,nvidiaで統一してくれよ。 機械語レベルでも、ライブラリレベルでもいいからさ。 511 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 10:14:49.72 ] OpenMPってよく知らないが_beginthreadしてるわけじゃなく CPUレベルでスイッチングしてるだけみたいだから ほとんどオーバーヘッドがないみたいだぞ 十分標準として使えそうなんだが 512 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 10:24:09.66 ] 残念ながら生成負荷は大きいよ だから何度もスレッド生成するような処理はかえって遅くなる 513 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 10:26:56.96 ] いやいや、ちゃんとスレッドプールしてるって 514 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 10:42:12.62 ] >>513 おまえ>>504だろ 試行レベルが同一 515 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 11:01:13.45 ] >>514 お前、いろいろと残念な人だな 516 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 11:09:33.71 ] と、せ〜いっぱいの切替し つか下らんこと書くならよそでやれヴぉけ 517 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 11:16:10.78 ] 切り返し な。 煽るなら突っ込まれる隙作っちゃダメだ 518 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 11:20:23.65 ] いいからスレ違いだ。 519 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 11:21:44.53 ] スレッドの生成負荷は軽いと言われているのを鵜呑みするとなかなか性能が出ないこともある気がする。 カーネル側でスレッドの生成回数を出来るだけ減らす工夫が必要だったり。 520 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 11:55:51.71 ] スレッドの生成負荷はプールすればいいから問題じゃない スレッドを切り替えるときにスタックポインタ変えたりするコンテクストスイッチの負荷が問題 521 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 13:09:56.12 ] cudaでスレッドのプールができるのかえ？ cpuの話をここに書いてる？ 522 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 13:20:03.48 ] そうそう、何度指摘してもスレ違いの話を続けているの。 523 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 13:23:34.12 ] その話、終わりました 524 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 13:55:34.91 ] どこまでをCPUで計算し、どこまでをGPUで計算するのかはGPGPUで重要 だからスレッドの性質については知っておく必要がある CUDAを使いこなすのにアップアップのアホどもには関係ないけど 525 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 14:04:51.88 ] 解決を依存性のない単純に並列化可能なアルゴリズムに置き換えられる問題を GPGPUでやるもんだと思ってたんだが違うのか？ 526 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 14:11:42.11 ] そりゃそうだが、データの転送速度、超越関数を使用するかどうかなど、 単純に並列度が高ければGPUでOKってもんじゃない Opteronの48コアマシンがお手頃価格で手に入る時代だぜ？ 527 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 14:13:49.17 ] あ、誤解の無いように書いとくが、CUDAやGPGPUを否定してるわけじゃないよ 528 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 14:15:39.88 ] 夏場に熱で吹っ飛んだのれす 529 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 17:16:29.77 ] >>526 > Opteronの48コアマシンがお手頃価格 今、おいくら万円くらい？ 530 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 17:24:19.24 ] 全部こみこみで100マソ切った 531 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 17:48:16.70 ] 高っ 532 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 17:49:59.02 ] そうか、ごめん…… 533 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 17:50:02.42 ] 48コアで何テラFlopsぐらい出んの？ 534 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 17:56:21.72 ] 高いけど、業務で使うならテスラを組み込んだサーバとたいして変わらないんじゃない？ 個人とか研究室とかだと絶対的金額から手が出せないかもしれないけど。 535 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 17:58:31.18 ] それは言えてる。 536 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 18:06:31.47 ] >>533 えーと、0.5テラFlopsくらい？ 537 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 18:59:41.35 ] 17万x4(CPU)+12万(ﾏｻﾞﾎﾞ)+HDDﾒﾓﾘｹｰｽか 欲しいけど家庭用のコンセントで大丈夫だろうか 似たようなスペックでGPGPU構築したらいくらぐらいかかるのかな 0.5テラFlopはしょぼくないか 538 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 20:43:24.14 ] >>537 家庭用コンセントじゃ無理 グラボ刺さるからGPGPUできるよ 539 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 21:43:26.64 ] >>538 このマザーで自作すると・・・ akiba.kakaku.com/pc/1103/31/131500.php 130000+100000*4+12800*16+100000+50000+α=884800+α くらい？ だとしたら100万円ほどで完成品が買えるなら良心的な価格な気がするね。 下のマザーボードだと気合いで8本GPU載せられるんだろか？ akiba.kakaku.com/pc/1104/20/213000.php 540 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 21:59:24.97 ] >>539 > 下のマザーボードだと気合いで8本GPU載せられるんだろか？ どこかに引っかかりそう。ラック筐体次第なのかな 541 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 22:17:54.30 ] マザー自作するツワモノいないかな まあ、高周波シュミレソフトで１０００万いくから 性能気にしないでいくなら自作できそうだけど・・・ 542 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 23:17:49.25 ] >>539 そのマザーのPCIeスロットはx8+x1+x8+x1+x8+x1+x8+x1だから ホスト・デバイス間の帯域不足に悩みそう。 演算性能さえあればいいプログラムなら何とかなるのかもしれないけど。 というかハイエンドグラボは複数スロット占有が一般的だし、水冷のグラボが必須になりそうｗ 543 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 23:27:34.87 ] コスト度外視で1台で最速を目指すなら P6T7 WS SuperComputer www.unitycorp.co.jp/asus/motherboard/intel/lga1366/p6t7_ws_super/index.html と EVGA GeForce GTX 580 FTW Hydro Copper 2 jp.evga.com/products/moreInfo.asp?pn=015-P3-1589-AR になるのかな。 電源の確保がものすごく大変そうだけどｗ 544 名前：デフォルトの名無しさん mailto:sage [2011/04/22(金) 23:28:08.92 ] CUDA Visual Profiler使ってるんですがなんどやっても計測後に "Error in profiler data file'(動かすディレクトリ)/temp_compute_profiler_0_0.csv' at line number 1. No column found" と出て計測結果がうまく出ません 何が原因なのでしょうか？ 545 名前：デフォルトの名無しさん mailto:sage [2011/04/23(土) 00:17:34.23 ] エラーコード見てもわからんが、俺はプログラム中にgetchar()を書いてある キーを押さなきゃ終わらないプログラムをプロファイラに渡してはまったことがあるZE 546 名前：デフォルトの名無しさん mailto:sage [2011/04/23(土) 00:24:57.17 ] >>544 Visual Profiler使ってないからよく判らんが、その作業ディレクトリはレガシーな名前かい? 空白やら何かの所為で巧くアクセスできないってことはないよね? 547 名前：デフォルトの名無しさん mailto:sage [2011/04/23(土) 00:39:41.60 ] >>545 外部から画像を読み込んでいますがキーは押さなくても終了するプログラムです >>546 作業ディレクトリと同じディレクトリに入っている計測したい実行ファイルは プロファイルを開始すると実行されるようなのでおそらくディレクトリにアクセスは出来ています プロファイラはプログラムを実行してから計測をしているようなのですがこの計測経過が100%になったら先ほどのエラーがでてしまいます temp_compute_profiler_0_0.csvは自動生成のファイルの用です 548 名前：デフォルトの名無しさん mailto:sage [2011/04/23(土) 00:56:29.31 ] >547 エラーがなんで出てるかはさっぱりわからんので、自分がこけたときの経験を列挙するぐらいしかできんのだけど、 Visual Profilerのデフォルト設定では30秒以内に処理が終わらない実行ファイルは実行を途中でうち切って 解析結果出力せず、えんえん解析を繰り返すってのはあったぜ。 549 名前：デフォルトの名無しさん mailto:sage [2011/04/23(土) 00:58:45.66 ] >>548 そこ毎回適当に300に変えてたわ しかしなんで一回でぜんぶできないかな 550 名前：デフォルトの名無しさん mailto:sage [2011/04/23(土) 01:02:49.32 ] グラフィックボードがGPGPUの処理以外にも画面表示の処理をしてる可能性が高いのだから、 画面表示による影響を平滑化するために複数回実行時間を計測してるんじゃないかと。 551 名前：デフォルトの名無しさん mailto:sage [2011/04/23(土) 14:34:55.79 ] コアレスアクセス用に半ワープ単位でアクセスしなきゃいかんのだが 16×16の二次元のスレッド作ったとして threadIdx.x = 0〜15 が半ワープに入るのか threadIdx.y = 0〜15 が半ワープに入るのか、どっちだべ。 552 名前：デフォルトの名無しさん mailto:sage [2011/04/23(土) 14:56:49.14 ] ふつうにかんがえてx 553 名前：デフォルトの名無しさん mailto:sage [2011/04/23(土) 15:03:38.31 ] アブノーマルに考えてx 554 名前：デフォルトの名無しさん mailto:sage [2011/04/23(土) 15:42:37.53 ] 真逆な意見がでてしまったわけだが、間をとってxと考えることにするよ。 555 名前：デフォルトの名無しさん mailto:sage [2011/04/23(土) 22:47:46.90 ] Larrabee はいつ出ますか？ 556 名前：デフォルトの名無しさん mailto:sage [2011/04/24(日) 00:29:55.82 ] Larrabeeはなかった事にされてるような・・・ 557 名前：デフォルトの名無しさん mailto:sage [2011/04/24(日) 01:23:47.09 ] LarrabeeってSandryのGPU部分のコア数を数百って規模に拡張したもののことだよ AVXがまさにLarrabeeのインターフェースらしい 558 名前：デフォルトの名無しさん mailto:sage [2011/04/24(日) 02:12:08.17 ] VisualStudio使ってEmuDebugの構成にしたら、 デバイス側関数もブレークポイント仕掛けられるって聞いたのにできねえ。 リンクするライブラリ変えるとかしなきゃいかんの？ 559 名前：デフォルトの名無しさん mailto:sage [2011/04/24(日) 02:25:06.20 ] NsightじゃなくてEmu使う理由は？ 560 名前：デフォルトの名無しさん mailto:sage [2011/04/24(日) 02:31:36.91 ] Nsightってグラフィックボード１枚指しでも使えんの？ Nsightを使わない理由は、NsightインストールしたらCUDA3.3がインストールされて、 でCUDA3.3のcuFFTの動作が怪しかったんで、NsightやめてCUDA3.0にしてる。 NsightでもCUDA3.0使えるのかもしれないけど、使い方調べるのがめんどくさかったので 環境構築に慣れたNsight使わないやり方でやってる。 561 名前：デフォルトの名無しさん mailto:sage [2011/04/24(日) 02:44:49.49 ] 3.3? 562 名前：デフォルトの名無しさん mailto:sage [2011/04/24(日) 03:35:32.04 ] CUDAとSandy BridgeのGPUは同時に使えますか？ 使えるとしたら、Sandy BridgeのGPU用のコーディングには何を使えば良いですか？ 563 名前：デフォルトの名無しさん mailto:sage [2011/04/24(日) 03:40:42.99 ] >>562 SBってGPGPUに対応してるの？ SDKとか出てるって聞いたことないんだけど。 564 名前：デフォルトの名無しさん mailto:sage [2011/04/24(日) 03:53:11.32 ] Sandy BridgeのGPUはOpenCL1.1に対応してるらしいけど CUDAに対応してるって話はみたことないから、CUDAとは同時に使えないと思われる。 OpenCLスレで聞いたほうがいいかも。 565 名前：デフォルトの名無しさん mailto:sage [2011/04/24(日) 03:57:39.89 ] >>564 >CUDAに対応してるって話はみたことないから、CUDAとは同時に使えないと思われる。 対応していない=同時に使えないという考えは改めた方位が良い。 566 名前：デフォルトの名無しさん mailto:sage [2011/04/24(日) 06:20:12.37 ] グラボを刺した状態でSandyのGPUコアが有効になるのかって話だよね グラフィックチップとしては完全に機能停止されるだろうけど 演算装置として稼動するんだろうか 567 名前：558 mailto:sage [2011/04/24(日) 11:04:44.80 ] EmuDebugはリンクするライブラリをcudart.lib→cudaartemu.libにかえたらできたや。 568 名前：デフォルトの名無しさん mailto:sage [2011/04/24(日) 12:25:30.29 ] 自身ぎゃああああああああああああああああああ 569 名前：デフォルトの名無しさん mailto:sage [2011/04/24(日) 16:01:57.36 ] VS2008でFirst chance Exceptionがでるが、カーネルコードの何が原因なのかさっぱりわからねえ… Nsight使うとこういうのの原因すぐに特定できるようになんの？ 570 名前：デフォルトの名無しさん mailto:sage [2011/04/24(日) 23:11:39.84 ] CUDA 4.0 RC2 をインストールしてみようと思っています。 64bit用のバイナリを作るには、CUDAも64-bit版をインストールする必要があるでしょうか？ それとも、CUDAの32-bit版でも nvcc のオプションの切り替えで32bit用と64bit用の両方のバイナリを生成することが出来るでしょうか？ 出来れば、WindowsXP(32bit) の環境で64bit用バイナリも作りたいと思っています。 571 名前：デフォルトの名無しさん mailto:sage [2011/04/24(日) 23:26:06.00 ] >>570 ＞64bit用のバイナリを作るには、CUDAも64-bit版をインストールする必要があるでしょうか たしかそうだったはず ＞CUDAの32-bit版でも nvcc のオプションの切り替えで32bit用と64bit用の両方のバイナリを生成することが出来るでしょうか？ 生成したい処理系のSDKをインストールしないといけなかったはず 切り替えはできないと思う 2つ（３２と６４）いれればOKなのかは不明 572 名前：デフォルトの名無しさん mailto:sage [2011/04/24(日) 23:42:26.64 ] >>571 レスありがとうございます。 ややこしそうなので、32bit 環境、64bit 環境それぞれを用意し、それぞれ用のバイナリを作ろうと思います。 573 名前：デフォルトの名無しさん [2011/04/25(月) 10:06:00.86 ] 64bitOS環境に64bit Toolkitを入れれば32bit&64bitの両方のバイナリを生成可能。 574 名前：デフォルトの名無しさん [2011/04/27(水) 08:28:47.19 ] GT430を使っていますが、CUDAでコンパイルすると次のようなエラーが出て コンパイルできません。 因みにC++は初心者ですがよろしくお願いします。juria_gpu.cu codepad.org/UtZeyAkV julia_gpu.cu(44): error: calling a host function("cuComplex::cuComplex") from a __device__/__global__ function("julia") is not allowed 575 名前：デフォルトの名無しさん mailto:sage [2011/04/27(水) 08:35:16.69 ] __device__/__global_　この辺良く見直せ！ 基本がわかってないと思うから基本からよく勉強だな 576 名前：デフォルトの名無しさん [2011/04/27(水) 08:43:59.77 ] エラーの意味がよくわかりません。__device__/__global__なんてどこにも 書いて無いし、サンプルコードなので動くはずなのですが、なぜか動きません。 codepad.org/oDh0YFM5 cuComplexをjulia関数から呼ぶのは許されていないという意味なのでしょうか？ 577 名前：デフォルトの名無しさん mailto:sage [2011/04/27(水) 08:50:02.31 ] globalが無いのかよ？ 全部晒せよ 578 名前：デフォルトの名無しさん [2011/04/27(水) 09:03:00.78 ] ideone.com/s9anZ ここです。よろしくお願いします。 579 名前：デフォルトの名無しさん mailto:sage [2011/04/27(水) 09:07:22.66 ] CUDA_by_exampleのだな まずそういう情報を先に出せよ 580 名前：デフォルトの名無しさん [2011/04/27(水) 09:11:49.07 ] すみません。書くのを忘れていました。 知っている方とお見受けしますので安心です。 581 名前：デフォルトの名無しさん mailto:sage [2011/04/27(水) 09:22:12.28 ] 特に問題なくコンパイルできたぞ nvccのオプションは？ 582 名前：デフォルトの名無しさん [2011/04/27(水) 09:27:35.63 ] nvcc -O julia_global.cu -lcutil32 -cutil32D 583 名前：デフォルトの名無しさん [2011/04/27(水) 09:31:28.05 ] インストールの仕方が悪いのでしょうか？ 584 名前：デフォルトの名無しさん mailto:sage [2011/04/27(水) 09:33:25.86 ] windowsはよく知らんがLinuxではnvcc -lglut julia_cpu.cuで通るんだが 上のソース名は間違いだよな？それでやってどのエラーメッセージが出るんだ？ julia_gpu.cu(44): error: calling a host function("cuComplex::cuComplex") from a __device__/__global__ function("julia") is not allowed がでるのか？ 585 名前：デフォルトの名無しさん [2011/04/27(水) 09:38:07.66 ] julia_gpu.cu です。間違いです。julia_cup.cuの場合はこちらでも通ります。 ただGPUを使おうとするとコンパイル時にエラーが出ます。 cuComplex c(-0.8, 0.156); cuComplex a(jx, jy); と a = a * a + c; の部分でエラーが出ます。原因がわかりません。 586 名前：デフォルトの名無しさん mailto:sage [2011/04/27(水) 09:40:32.60 ] 正確なコンパイルオプションとエラーメッセージを再掲よろ 587 名前：デフォルトの名無しさん [2011/04/27(水) 09:51:50.54 ] ideone.com/mZviD ここです。 588 名前：デフォルトの名無しさん mailto:sage [2011/04/27(水) 09:57:53.82 ] C:\Users\admin\CUDA\chapter04>nvcc -O julia_gpu.cu これじゃ -lcutil32とかないじゃん　なけりゃダメだろ 589 名前：デフォルトの名無しさん [2011/04/27(水) 10:01:49.76 ] リンクエラーが出るかもしれないですが、今はコンパイルエラーなので そちらでも同じようなエラーが出るか確かめてほしいです。 もしエラーが出なかったら私の環境がおかしいのかも知れません。 Linuxでの開発環境とWindowsでは少し違うと思うのでWindowsでも 確認してもらいたいです。 590 名前：デフォルトの名無しさん [2011/04/27(水) 10:05:27.84 ] ウインドウズアプリを作るのでリナックスでうまく言っても意味が 無いのですが・・・ 591 名前：デフォルトの名無しさん mailto:sage [2011/04/27(水) 10:13:19.95 ] >>589 あ、そういうこと じゃあソースが壊れてるんでなければそっちの環境がおかしそうだね 環境の再構築したほうがよさげ 592 名前：デフォルトの名無しさん mailto:sage [2011/04/27(水) 10:33:39.34 ] 因みにlinux環境で nvcc -lglut julia_gpu.cu　→　OK nvcc julia_gpu.cu　→　コンパイルOK、リンクNG 593 名前：デフォルトの名無しさん [2011/04/27(水) 11:22:50.87 ] C:\Users\admin\CUDA\cuda_by_example\chapter04>nvcc --version nvcc: NVIDIA (R) Cuda compiler driver Copyright (c) 2005-2011 NVIDIA Corporation Built on Thu_Mar_24_14:53:10_PDT_2011 Cuda compilation tools, release 4.0, V0.2.1221 こちらの開発環境ですが、nvccのバージョンを教えてください。 環境を再インストールしましたがやっぱりだめです。 構文エラーとして解釈されます。 594 名前：デフォルトの名無しさん mailto:sage [2011/04/27(水) 11:28:41.58 ] 3.2 4.0RC2？はまだ早いんでね？ 595 名前：デフォルトの名無しさん [2011/04/27(水) 11:40:43.34 ] そうですね。自分の持っているやつがGT430なのでバージョン1.0 （最新は1.3）に対応しているみたいなので古いほうをインストール してみます。 596 名前：デフォルトの名無しさん [2011/04/27(水) 11:45:26.50 ] 出来ました！ ありがとうございました。 597 名前：デフォルトの名無しさん mailto:sage [2011/04/27(水) 13:22:17.13 ] >>596 何をどうしたら解決したか書いてほしいな 598 名前：デフォルトの名無しさん mailto:sage [2011/04/27(水) 13:30:13.83 ] 流れで分かりそうなもんだけど 4.0RC→3.2で解決したようだ 599 名前：デフォルトの名無しさん mailto:sage [2011/04/27(水) 13:44:28.80 ] 4.0は大きく拡張されてるみたいだし、RCがとれても移行は様子見したほうがよさそうだね 2台とかあって3.2と4.0の環境があったりするなら切り分けしやすいから良いかもしれないけど。 600 名前：デフォルトの名無しさん [2011/04/27(水) 15:49:38.35 ] こんにちは。 16*32の行列に長さ16の配列をXORの積をしたいのですが可能でしょうか？ イメージ的には32個のブロックに16個のスレッドという感じで考えています。 行列計算なので実装はそんなに難しくないと思うのですが、具体例があまり 見当たらず苦心しています。はじめてからまだ三日なのですが、解る方が いらしたらソースを見せてほしいです。ヒントだけでも良いです。 よろしくお願いします。 601 名前：デフォルトの名無しさん mailto:sage [2011/04/27(水) 16:03:05.24 ] >>16*32の行列に長さ16の配列をXORの積をしたいのですが これどういう意味？もう少し計算を具体的に書いてくれ 602 名前：デフォルトの名無しさん [2011/04/27(水) 16:34:04.46 ] 4*8の例で考えます。 [abcd]*[{efgh}{ijkl}{mnop}[qrst}{uvwx}{yz12}{3456}{7890}]= a^e&b^f&c^g&d^h=e a^i&b^j&c^k&d^l=i ... このように計算したいです。 並列計算なら全体をばらして8個のブロックに4個のスレッドを当てれば 4サイクルで実行できると思うのですが何か間違えていますでしょうか？ よろしくお願いします。 603 名前：デフォルトの名無しさん [2011/04/27(水) 16:40:44.57 ] for(j=0;j<16;j++){ o=FG[a[j]]^GF[u1.m[j]]; p=FG[b[j]]^GF[u.m[j]]; for(i=0;i<16;i++){ d1[j]^=t[o][h1[p][i]]; d2[j]^=t[o][h2[p][i]]; } buf[j]^=d1[j]; buf[j+16]^=d2[j]; } この処理を並列化したいです。 604 名前：デフォルトの名無しさん mailto:sage [2011/04/27(水) 17:08:33.14 ] >>603 ideoneかcodepadにC版の関数をアップロードしてくれ。 そうしてくれたらcudaでどこをどう並列化するか考えてみるよ。 605 名前：デフォルトの名無しさん [2011/04/27(水) 17:11:57.15 ] 並列化したいのはループ使っている部分だけなのでそこだけです。 codepad.org/15hfT9HS 606 名前：デフォルトの名無しさん [2011/04/27(水) 18:53:48.98 ] 公開したのに反応が無い。 見ても解らないだとか汚いと罵るだけ。 どうせやる気無いんだろ？ 607 名前：デフォルトの名無しさん mailto:sage [2011/04/27(水) 18:56:58.77 ] >>606 うん。 忙しいからね(´・∀・｀ ) 608 名前：デフォルトの名無しさん mailto:sage [2011/04/27(水) 19:19:58.01 ] >>603 またお前か。 いい加減自分の手を動かしたらどうだ。 時間がないなら対価を払って時間を買え。 609 名前：デフォルトの名無しさん [2011/04/27(水) 19:50:55.96 ] だったらソース見せろとかいうな馬鹿 610 名前：デフォルトの名無しさん mailto:sage [2011/04/27(水) 19:53:32.57 ] >>606 何怒ってんだ？ その程度のコードなら普通にできるだろ バカなの？ 611 名前：デフォルトの名無しさん [2011/04/27(水) 19:57:29.62 ] 見る気も無いくせに見せろと言うほうが馬鹿 612 名前：デフォルトの名無しさん mailto:sage [2011/04/27(水) 20:06:58.92 ] 見せろって言ったのは俺じゃないし、ソース見る前からお前だって分かった。 自分でやる気が無いっていう態度が滲み出てるんだよ。 613 名前：デフォルトの名無しさん [2011/04/27(水) 20:14:43.33 ] やる気があって調べてるんだよ。 ここならもっと詳しい人が居ると思うから聞いてみただけ。 何で最初からやる気が無いなんて決め付けるの？ 614 名前：デフォルトの名無しさん mailto:sage [2011/04/27(水) 20:18:05.38 ] まぁ質問に具体性がなく丸投げだからな もう少し考えた過程が見えないと回答もしづらい罠 615 名前：デフォルトの名無しさん mailto:sage [2011/04/27(水) 20:19:27.54 ] 最初からじゃなくて俺が知っているのはここ8ヶ月くらいだが 分からない、分からない、一個教えてもらうとすぐ次が分からないで 自分で調べる気なんてなさそうに見えるんだが。 616 名前：デフォルトの名無しさん mailto:sage [2011/04/27(水) 21:13:16.83 ] >>613 Intel Parallel Studio 使って並列化。 ↓ 逆アセンブルして、コード解析してCUDA化。 ↓ (゜д゜）ｳﾏｧ 617 名前：デフォルトの名無しさん mailto:sage [2011/04/27(水) 21:18:58.59 ] VC++ 2010用のtemplateプロジェクト、どこかにないでしょうか？ CUDA 3.2 SDKにあるVC++ 2008用のtemplateプロジェクトを、VC++ 2010用に変換しようとしましたがうまくいきません。。。 618 名前：604 mailto:sage [2011/04/27(水) 21:32:18.42 ] >>605 ＞並列化したいのはループ使っている部分だけなのでそこだけです。 ＞codepad.org/15hfT9HS 悪い悪い、デートしてて遅くなった。 で、このコードのどこをcuda化したいんだ? 未だ帰り道だから、んな長いの見てられないんだ。 619 名前：デフォルトの名無しさん mailto:sage [2011/04/27(水) 22:18:04.42 ] 16人17脚はデートって言うの？ 620 名前：デフォルトの名無しさん mailto:sage [2011/04/28(木) 00:21:09.33 ] >>617 Windows版のnvccはVisual Studio 2005か2008のコンパイラが前提になっているらしいから VC++ 2010用のテンプレートだけではうまくいかないと思う。 それらのコンパイラを用意するとなったら、VS2008で作ったほうが楽な気もする・・・ 621 名前：デフォルトの名無しさん mailto:sage [2011/04/28(木) 01:56:26.69 ] >603 きたないコードだなあ・・・。 622 名前：デフォルトの名無しさん mailto:sage [2011/04/28(木) 02:06:27.26 ] >>603 俺ならこうだな。 for (i = 0 ; i < 16; i++) { 　　for (j = 0; j < 16; j++) { 　　} } 主なコーディング作法が3つくらいあるから、そのどれかにしろ。 623 名前：デフォルトの名無しさん mailto:sage [2011/04/28(木) 02:08:04.55 ] >>622 スペースきれい！ ふしぎ！ 624 名前：デフォルトの名無しさん mailto:sage [2011/04/28(木) 02:10:39.48 ] >603 あまりにもイミフなんだけど、とりあえず並列化したいなら 関数呼び出し側 　func<<<<16, 16>>>(d1, d2, t, h1, h2, buf); 処理内容 __global__ void func(hoge* d1, hogehoge* d2, hogehogehoge* t, hogehogehogehoge* h1, hogehogehogehogehoge* h2, hogehogehogehogehoge* buf){ 　　int o = FG[a[blockIdx.x]^GF[u1.m[blockIdx.x.]]; 　　int p = FG[b[blockIdx.x]^GR[u.m[blockIdx.x]]; 　　d1[blockIdx.x] ^= t[o][h1[p][threadIdx.x]]; 　　d2[blockIdx.x] ^= t[o][h2[p][threadIdx.x]]; 　　buf[blockIdx.x] ^= d1[blockIdx.x]; 　　buf[blockIdx.x + 16] ^= d2[blockIdx.x]; } 酔っ払いながらコード書いたから間違いある気がするけど、 CUDAの根本的な文法も理解しようとせず使おうとしてるから、バレないし問題ないよね。 625 名前：デフォルトの名無しさん mailto:sage [2011/04/28(木) 20:52:23.73 ] >>617 VC2010はCUDA 4.0 RC2から対応で、SDKにテンプレートが付属してる。 うちの環境だとMSBuild関連が正しくインストールされなかったのでひと手間必要だったけど。 626 名前：デフォルトの名無しさん mailto:sage [2011/04/29(金) 23:49:38.06 ] CUDAのコンパイラってレジスタ不足でもコンパイル通っちゃうのか？ VSで開発してるんだが、どうもレジスタが足りないとFirst chance exceptionエラーが出る気がしてきた。 627 名前：デフォルトの名無しさん mailto:sage [2011/04/29(金) 23:58:25.76 ] >48 祈るってアプローチが間違いとは言わないけど、 同じ計算を2並列で行うとか、間違いなく構成を変えられる方法があれば 演算器を入れ替えて2回行うとかって手もあるよ。 2倍時間がかかってもまだCPUには勝てるでしょ。 628 名前：デフォルトの名無しさん mailto:sage [2011/04/30(土) 00:29:11.48 ] PowerDirectorなどCUDAを使ってエンコードを高速化出来るアプリがありますは、SLI環境ではさらに高速化出来るのでしょうか？ ググると出来ないというソースが目立ちますが... X58マザー使っています。 629 名前：デフォルトの名無しさん mailto:sage [2011/04/30(土) 00:35:04.10 ] レジスタ不足ってローカルメモリ使われるんじゃないの? 630 名前：デフォルトの名無しさん mailto:sage [2011/04/30(土) 01:12:10.48 ] >>628 cuda4はどうか知らないが、それ以前のcudaではSLIだからより速くできると言うことはない。 勿論GPU2枚挿しに対応したアプリなら、2枚刺した方が速くはなるがSLIには関係ないようだ。 631 名前：デフォルトの名無しさん mailto:sage [2011/04/30(土) 03:38:04.89 ] >629 引数はsharedメモリに確保されるってどっかで見たことあるけど… sharedメモリつかってないプログラムだしなあ。 使ってるのはGeForce9800 - G92コアなんだけど、 コアに上限の24ワープ割り当てるVSがエラー吐いて動かなくなるんだけど、 控え見えに16ワープ放りこむと問題なく動く。 cuFFTとかの標準ライブラリでも動作が怪しかったりするのは レジスタ周りの取り扱いのせいなんじゃって気がしてきてるのですよ。 632 名前：デフォルトの名無しさん mailto:sage [2011/04/30(土) 08:15:07.06 ] 引数はCC2.0からはコンスタントメモリで それ以前はシェアードだったかな 関数内変数はレジスタで不足したら勝手にグローバルにいくんでしょ ワープ数が上限ぴったりならエラーは吐かないと思うんだけど 633 名前：デフォルトの名無しさん mailto:sage [2011/04/30(土) 08:23:40.36 ] オーバーしても、普通は遅くなるだけだと思うんだけどなぁ。 634 名前：デフォルトの名無しさん mailto:sage [2011/04/30(土) 15:46:08.21 ] >>631 そんな曖昧な知識で推測する前にPDFちゃんと読むんだ 635 名前：デフォルトの名無しさん mailto:sage [2011/04/30(土) 16:50:47.84 ] >634 了解したぜ 636 名前：デフォルトの名無しさん mailto:sage [2011/04/30(土) 23:22:46.13 ] >>630 ありがとう。 SLI環境はもう少し様子見します。。 637 名前：デフォルトの名無しさん mailto:sage [2011/05/02(月) 12:42:12.87 ] 最近のグラボはATX電源から電源引っ張ってくんだな。 Fermiに変えたが早すぎワロタ。 638 名前：デフォルトの名無しさん mailto:sage [2011/05/02(月) 13:14:07.57 ] それ以前からfermiって劇的な変化じゃね キャッシュはでかい 639 名前：デフォルトの名無しさん mailto:sage [2011/05/08(日) 01:40:31.32 ] FlashPlayerもCUDAつかえるの？ 640 名前：デフォルトの名無しさん mailto:sage [2011/05/08(日) 05:13:16.28 ] Tesla S2070ってなんで抹消されてしまったのでしょうか。 nVIDIAのページにもELSAのページにもなくなってしまったのですが。 641 名前：CUDA初心者 [2011/05/09(月) 11:25:03.07 ] XPでFFTをかけるソースがあったのでvisualstudio2008で実行してみたところ FFT.exe の 0x7c812afb で初回の例外が発生しました: Microsoft C++ の例外: cudaError_enum (メモリの場所 0x0012ae3c)。 というエラーが出てしまい困っています。ホストからデバイスへのメモリのコピー等 の簡単なプログラムは動きます。ちなみにGeforce210です。 642 名前：デフォルトの名無しさん mailto:sage [2011/05/09(月) 21:02:35.89 ] debugを人に頼むような人間は一つ解決しても次に同じような所で引っかかるからきりがないわな。 643 名前：デフォルトの名無しさん mailto:sage [2011/05/10(火) 09:55:41.04 ] >>641は何も頼んでない。日記を書いてるだけじゃないか？ 644 名前：デフォルトの名無しさん mailto:sage [2011/05/10(火) 22:29:43.44 ] >641 多分それ、cuFFTのバグ。そのエラーコードはメモリアクセス例外のときに出る。 俺の場合9800GTから550Tiに買い換えたらエラーが消えた。 cuFFTがデバイスメモリ使い過ぎて、デバイスが積んでるメモリ量によっては エラーを吐くとかが原因の気がしてるんだが、本当の理由はわからん。 NVidia側のバグフィックスを待つぐらいしか現状で対策はないと思う。 645 名前：デフォルトの名無しさん mailto:sage [2011/05/11(水) 00:23:29.41 ] cuFFTのバグなら、サイズが2の冪かどうかでも変わるね。 cuFFTの内部で結構デバイスメモリを使うようだから、ボードの半分以下のFFTしかかけられない。 646 名前：デフォルトの名無しさん mailto:sage [2011/05/11(水) 19:42:26.40 ] >>642 FFTなんて本来ライブラリの側で対処すべき次元のもので、 ユーザー側がデバッグに煩わされるべきものではないがな。 647 名前：デフォルトの名無しさん mailto:sage [2011/05/11(水) 19:57:25.35 ] >>646 もしバグだとおもうならフォーラムで問い合わせたり報告したりしたほうが良いんじゃないか？ バグなら修正してもらえる可能性があるし、そうじゃないとしたら自分の間違いだとわかる 648 名前：デフォルトの名無しさん mailto:sage [2011/05/11(水) 23:13:45.45 ] ubuntu 10.04、CUDA4.0でSDKをコンパイルした後に deviceQueryを実行しようとすると Error: API mismatch: the NVIDIA kernel module has version 270.29, but this NVIDIA driver component has version 270.40. Please make sure that the kernel module and all NVIDIA driver components have the same version. と出てしまい実行出来ません。どうしたら良いでしょうか？ 649 名前：デフォルトの名無しさん mailto:sage [2011/05/11(水) 23:20:38.40 ] >>648 www.alc.co.jp/ 650 名前：デフォルトの名無しさん mailto:sage [2011/05/11(水) 23:33:32.11 ] nvidia-current カーネルを削除してSDKを再コンパイルで行けました ありがとうございます 651 名前：デフォルトの名無しさん mailto:sage [2011/05/11(水) 23:35:55.99 ] nvidia-current カーネルを削除してSDKを再コンパイルで行けました ありがとうございます 652 名前：デフォルトの名無しさん mailto:sage [2011/05/12(木) 00:05:11.00 ] >>647 フォーラムには既にあったと思う。 653 名前：CUDA初心者 [2011/05/12(木) 15:13:05.97 ] >>644 ありがとうございます。 550Ti持ってるのでそっちでやってみようと思います。 654 名前：デフォルトの名無しさん [2011/05/16(月) 10:53:35.28 ] Developer Drivers for Linuxをインストールすると、 Xwindowの解像度が640*480固定になるのはどうしてなのかしら？ 655 名前：デフォルトの名無しさん mailto:sage [2011/05/16(月) 10:57:58.45 ] おまえがばかだからじゃないか？ 656 名前：デフォルトの名無しさん mailto:sage [2011/05/16(月) 11:00:18.75 ] xの設定書き換えちゃってるんじゃないの? 657 名前：デフォルトの名無しさん [2011/05/18(水) 12:36:20.91 ] Ubuntu 11.04 CUDA 4.0でドライバをインストールしようとすると、 The Nouveau kernel driver is currently in use by your system. というエラーが出て、先に進めません。 どうしたらNouveau kernel driverをとりのぞけますか? 658 名前：デフォルトの名無しさん mailto:sage [2011/05/18(水) 12:42:47.54 ] これ読んでみ https://help.ubuntu.com/community/BinaryDriverHowto/Nvidia 659 名前：デフォルトの名無しさん [2011/05/18(水) 13:18:46.93 ] >>658 早速の返信ありがとうございます。 sudo apt-get --purge remove xserver-xorg-video-nouveau と打って、再起動もしてみましたが、やはり同じメッセージが帰ってきます。 ほかに、何か手立てはないものでしょうか? 660 名前：デフォルトの名無しさん mailto:sage [2011/05/18(水) 21:39:19.83 ] >>657 俺がLive CDで試したときは kernelの引数にnouveau.modeset=0として起動後に modprobe -r nouveauした。 そもそもnouveauがモジュールではなくて組み込みになってるなんてことはないよね？ 661 名前：デフォルトの名無しさん [2011/05/19(木) 04:06:30.55 ] インテルコンパイラには対応していないの? 662 名前：デフォルトの名無しさん mailto:sage [2011/05/19(木) 04:12:46.19 ] Coalesced Accessって簡単にいうとどゆうこと 663 名前：デフォルトの名無しさん mailto:sage [2011/05/19(木) 08:14:50.08 ] >>662 １６人１７脚で走ること ただし、一番端っこの人の番号に制限あり 664 名前：デフォルトの名無しさん mailto:sage [2011/05/19(木) 08:28:38.02 ] >>663 なんだかすごく遅くなりそうな例えだなｗ 665 名前：デフォルトの名無しさん mailto:sage [2011/05/19(木) 09:19:43.15 ] 論文書くときにさ、遠くの図書館に借りに行くわけよ 16冊まで借りられるから自分の欲しいタイトル1冊と隣に並んでるの15冊借りて帰って来るの そん中にフラットメイト15人が欲しい本が入ってたらラッキー 何人か足りなかったらまたまた誰か派遣して16冊借りてこなくちゃなんない 一回で16人全員満足したり16冊全部が有効に使われたらうれしいな 俺ら棚のブックエンドの右からまとめてとってく癖があるから 司書さんはいいかんじに本並べといてね 家にあればこっちのもんだから欲しい本被っててもおkおk てのはどうだろう 666 名前：デフォルトの名無しさん mailto:sage [2011/05/19(木) 09:33:17.23 ] >>663はCoalesingの説明になってないぞ。 食堂のカウンターに並んでいる先頭の１６人が皆「カレーライスください」という。 食堂のおばちゃんが「皆カレーかい？鍋ごとあげるから皆で分けてね。」 もちろん正確ではないが速くなりそうな例ではある。 667 名前：デフォルトの名無しさん [2011/05/19(木) 15:22:10.76 ] 初歩的な質問で恐縮なのですが、 皆さんCUDA版プログラムとCPU版プログラムの速度比較は どのようにされてますでしょうか？ CUDA3.0以降はエミュレーションモードがなくなったようなので、 別にCPU版プログラムも作って比較するのがよいでしょうか？ 668 名前： 忍法帖【Lv=25,xxxPT】 mailto:sage [2011/05/19(木) 16:17:52.26 ] エミュレーションと比較して何の意味があるの？ NVIDIAの回し者でない限り、iccでカリカリにチューンしたプログラムと比較するべし。 669 名前： 忍法帖【Lv=26,xxxPT】 mailto:sage [2011/05/19(木) 18:19:14.33 ] CPU版と比較しても仕方が無い。 iccでチューニングした上で、更にその先の為にcudaを使うのだから。 まして、エミュレーションだなんてなんの悪夢だか。 670 名前：デフォルトの名無しさん mailto:sage [2011/05/19(木) 21:39:40.22 ] >662 メモリ読み書きするなら32byte, 64byte, 128byte単位にしてください。 これだけ。 671 名前：デフォルトの名無しさん [2011/05/20(金) 14:19:17.96 ] SDK,toolkitを再インストールしたところ サンプルを実行すると cudart32_31_9.dllが見つからなかったため、このアプリケーションを開始できませんでした 。cudartを間違えて消してしまったと思われるのですが、どうやったら解消されるのでしょうか？ ぐぐってみたのですが、わかりませんでした。 672 名前：デフォルトの名無しさん mailto:sage [2011/05/20(金) 14:31:37.09 ] PCを再起動させてみたらどう 673 名前：デフォルトの名無しさん mailto:sage [2011/05/20(金) 15:41:59.99 ] 最悪でも再インストールで済むだろ。 674 名前：デフォルトの名無しさん [2011/05/20(金) 17:07:36.68 ] 再インストールしてみたのですが変わりませんでした泣 675 名前：デフォルトの名無しさん mailto:sage [2011/05/20(金) 17:37:59.88 ] >>674 よかったな、cudartを間違って消してしまったわけじゃないことが判ったじゃないか。 676 名前：デフォルトの名無しさん mailto:sage [2011/05/27(金) 03:06:11.39 ] CUDA 4.0 677 名前：デフォルトの名無しさん [2011/05/27(金) 11:09:49.93 ] __CUDACC__ってもう定義されていないの? 678 名前：デフォルトの名無しさん mailto:sage [2011/05/28(土) 08:21:44.26 ] みんな業務でCUDAつかってんの？ 679 名前：デフォルトの名無しさん mailto:sage [2011/05/28(土) 08:56:58.21 ] 一部使ってますが。 680 名前：デフォルトの名無しさん mailto:sage [2011/05/28(土) 21:17:04.74 ] CUDA4.0、GTX480M + Notebook Developer Drivers for WinVista and Win7 (270.61)では device query含め、CUDA Cのサンプルが全部動かない。 OpenCLとDirectComputeは動く。 なんでだろう。 681 名前： 忍法帖【Lv=1,xxxP】 mailto:sage [2011/05/29(日) 11:17:41.19 ] 坊やだからさ 682 名前：デフォルトの名無しさん mailto:sage [2011/05/29(日) 11:48:30.13 ] VS2010とCUDA4.0の環境でrulesファイルを使えた方おりませんか？ $(CUDA_PATH)\extras\visual_studio_integration\rules にあるファイルをコピーしたのですが、これを設定してプロジェクトを読み込みなおすと 「要素 <UsingTask> 内の属性 "AssemblyFile" の値 "$(CudaBuildTasksPath)" を評価した結果 "" は無効です」 といったエラーが出てプロジェクトが読み込めないようです…。 683 名前：デフォルトの名無しさん mailto:sage [2011/05/29(日) 21:12:48.05 ] 完全な逐次処理（どう頑張っても並列化できない）のを、 CUDAで高速化してちょ、といわれたんだけど殴っていい？ 684 名前：デフォルトの名無しさん mailto:sage [2011/05/29(日) 21:18:05.50 ] >>683 いや無理にでもやれよ それが仕事ってもんだ 685 名前：デフォルトの名無しさん mailto:sage [2011/05/29(日) 21:27:29.77 ] ゼロコピーってキャッシュされないでしょ。 んで、ゼロコピーを使うと遅くなるケースが出てきたんで、 仕方なく使わないことにしたんだけど、 やっぱりcudaMallocやcudaMemcpyの時間がもったいないんで どうにかしたいんだけど、なんか知恵ある？ 686 名前：デフォルトの名無しさん mailto:sage [2011/05/29(日) 21:53:26.07 ] >>683 どう頑張っても並列化できないを証明すれば良いじゃん． なかなか難しいよ． 687 名前：デフォルトの名無しさん mailto:sage [2011/05/29(日) 22:08:28.81 ] >>685 ゼロコピーって何よ。 何をしたいのか判らんけど、cudaMemset()も使えない内容なの? 688 名前：682 mailto:sage [2011/05/29(日) 22:53:30.43 ] 解決しました。ちょうど公式フォーラムで同じ問題に遭遇してた人が居たようです。 先ほどあった投稿で解決することができました forums.nvidia.com/index.php?showtopic=201433 689 名前：デフォルトの名無しさん [2011/05/29(日) 23:36:05.38 ] もふ。 マジレスするとゼロコピーとは、cudaHostGetDevicePointer()を使うやつのことだす。 cudaMallocHost()でホスト側のメモリを確保しとかないとダメだけど。 690 名前：デフォルトの名無しさん mailto:sage [2011/05/30(月) 01:01:06.08 ] ゼロコピーと言うのは知らんかった。 処で、勿体無いと言うけどcudaMalloc()って時間掛かる? 691 名前：デフォルトの名無しさん mailto:sage [2011/06/01(水) 01:01:48.66 ] 演算速度について質問です。 プログラミングガイドによるとfloatにおけるmulとmadの速度は同一だとあるのですが、 for(i=0;i<N;i++) a = b*c; と for(i=0;i<N;i++) a = d*(b+c); では明らかに前者のほうが短い時間で処理を行っています。 dをコンスタンとメモリや定数にしても同じでした。 いったいどういうことなんでしょうか？ 助けてください。 692 名前：デフォルトの名無しさん mailto:sage [2011/06/01(水) 01:07:21.88 ] >>691 どこのプログラミングガイドだよ教えろよ教えてくださいお願いします 693 名前：デフォルトの名無しさん mailto:sage [2011/06/01(水) 01:12:04.19 ] madって(A*B)+Cじゃないの？ 694 名前：デフォルトの名無しさん mailto:sage [2011/06/01(水) 01:14:43.17 ] すみませんまったく環境とか書いてませんでした。 GPU : GeForce GTX 285M CUDA Toolkit : 2.2 です。 >>692 NVIDIA CUDA Programming Guide 2.2です。 第5章Instruction Performance の 5.1.1.1 Arithmetic Instructions より Throughput of single-precision floating-point add, multiply, and multiply-add is 8 operations per clock cycle. だそうです。 695 名前：デフォルトの名無しさん mailto:sage [2011/06/01(水) 01:27:04.68 ] >>693 ごめんなさい書き間違えました。 for(i=0;i<N;i++) a = d*b+c; です。 696 名前： 忍法帖【Lv=2,xxxPT】 mailto:sage [2011/06/01(水) 02:09:02.07 ] まぁ、ptxを貼ってみろ。 697 名前：デフォルトの名無しさん mailto:sage [2011/06/02(木) 01:22:02.76 ] 積和演算なら a = b*c+d*e とかじゃないのん？畳込みとか行列積とかで頻出する計算です。 698 名前：デフォルトの名無しさん mailto:sage [2011/06/02(木) 01:30:56.84 ] DSPの命令見てみたら、積和演算は a = a + (b*c) だったわ。 a+= (b*c) でもええかもしれんが。 699 名前：デフォルトの名無しさん mailto:sage [2011/06/02(木) 03:02:35.87 ] >>680 270.81でないとCUDAランタイムAPIが変な動作をするみたい つまり現時点ではPC用のドライバしかない 700 名前：デフォルトの名無しさん mailto:sage [2011/06/02(木) 11:07:00.65 ] 270.61に 「Unified Virtual Addressing (UVA) やGPUDirect? v2.0を特徴とするCUDA 4.0を使ったアプリケーション用サポートの追加」 と書いておきながら動かないなんてふざけた話だわ 701 名前：デフォルトの名無しさん mailto:sage [2011/06/02(木) 22:17:32.06 ] CUDA勉強はじめました。 でも世の中に出てる日本語のCUDAの本って4冊ともFermi以前のアーキテクチャでの説明だからsharedメモリが16kbyteだったり、1ブロックあたりのスレッド数が512だったりするんで困ります。 なんとかして。 702 名前：691 mailto:sage [2011/06/02(木) 22:33:23.38 ] コンスタントメモリや定数の値をレジスタに読み込む時間をまったく無視していました。 たぶんそのせいで遅くなっているんだと思います。 ご迷惑かけてすみません。 703 名前：デフォルトの名無しさん mailto:sage [2011/06/02(木) 22:59:17.85 ] >>700 ちょうど新しいドライバが出たみたいです 704 名前：デフォルトの名無しさん mailto:sage [2011/06/02(木) 23:33:53.87 ] >>701 英語も勉強汁 705 名前：デフォルトの名無しさん mailto:sage [2011/06/02(木) 23:44:11.15 ] oh! i can't read english. :-P 706 名前：デフォルトの名無しさん mailto:sage [2011/06/03(金) 00:27:22.61 ] >>701 Fermi用の正しい値が分かる人なら何も困らないような気が・・・・。 707 名前：デフォルトの名無しさん mailto:sage [2011/06/03(金) 22:26:32.06 ] Compute Visual Profilerを使いこなしてる人っている？ あれ、結局、どう活用していかわからん。 708 名前：デフォルトの名無しさん mailto:sage [2011/06/04(土) 10:50:31.43 ] >>701 Fermiの方が簡単になってるからいんじゃね >>707 ローカルメモリ使われてたりバンクコンフリトあるかをチェックすんじゃね 709 名前：デフォルトの名無しさん mailto:sage [2011/06/04(土) 11:02:21.92 ] 北川景子かわゆす 710 名前：デフォルトの名無しさん mailto:sage [2011/06/04(土) 11:18:35.39 ] そうか？ 711 名前：デフォルトの名無しさん mailto:sage [2011/06/04(土) 20:06:16.65 ] あれ？4.0ってcutil64D.libなくなっちゃった？ 712 名前：デフォルトの名無しさん mailto:sage [2011/06/04(土) 23:19:09.39 ] 本当だなくなってる 713 名前：デフォルトの名無しさん mailto:sage [2011/06/05(日) 00:30:29.73 ] 　 714 名前：デフォルトの名無しさん mailto:sage [2011/06/05(日) 00:56:21.39 ] 自前でビルドしなさい 715 名前：デフォルトの名無しさん mailto:sage [2011/06/05(日) 04:04:47.93 ] 今Cを勉強中ですが、CUDAと並行して勉強できますか？ 716 名前：デフォルトの名無しさん mailto:sage [2011/06/05(日) 10:23:06.12 ] >>715 少なくとも小中で算数（数学）や社会、英語などを並行して勉強してきたはず 717 名前：デフォルトの名無しさん mailto:sage [2011/06/05(日) 10:42:05.41 ] >>715 CUDAのプログラム記法以外に知るべき事が多い。と言うかむしろそちらが大半じゃないか。並行してできる、がんばれ 718 名前：デフォルトの名無しさん mailto:sage [2011/06/05(日) 18:19:38.33 ] どっちもものにならないと思う。どっちかだけやればものになるものでもないけど。 719 名前：デフォルトの名無しさん mailto:sage [2011/06/05(日) 18:43:41.04 ] 4.0でコンパイルが通らないの結構多いね CUDA BY EXAMPLEのサンプルをいくつか試したがエラーをはく。 まだ3.2使うわ 720 名前：デフォルトの名無しさん mailto:sage [2011/06/05(日) 19:14:22.50 ] 4.0 から cutil32.lib とかって自分でビルドするようになったのね。 インストール後に検索しても見つからないから、変だな〜と思った。 Direct Compute の Example が動かないのもよくわからん。 721 名前：デフォルトの名無しさん mailto:sage [2011/06/06(月) 09:41:55.58 ] ある程度C(というかプログラミング)しってなきゃ百パー無理だろ 722 名前：デフォルトの名無しさん mailto:sage [2011/06/06(月) 13:37:59.32 ] %を略すな、池沼 723 名前：デフォルトの名無しさん mailto:sage [2011/06/06(月) 13:54:50.09 ] 。 724 名前：デフォルトの名無しさん mailto:sage [2011/06/06(月) 21:16:41.99 ] >>721 それを知らなかったので、PATHの設定やmake実行さえ一苦労でした。 人によっては、ディレクトリの概念が難しい場合もあるかもしれません。 725 名前：デフォルトの名無しさん mailto:sage [2011/06/06(月) 21:26:48.74 ] cudaはただでさえ難しいぞ 初心者に手にを得る代物じゃね〜べよ 726 名前：デフォルトの名無しさん mailto:sage [2011/06/06(月) 21:32:48.38 ] 資料も英語しかないしかなりきつかったわ 727 名前：デフォルトの名無しさん mailto:sage [2011/06/08(水) 14:58:17.39 ] OpenCLとどっちが楽ですか？ 728 名前：デフォルトの名無しさん mailto:sage [2011/06/08(水) 16:43:07.02 ] Cudaの方が楽 OpenCLはOpenGLと同様、自助努力が基本になる 729 名前：デフォルトの名無しさん mailto:sage [2011/06/09(木) 05:38:12.86 ] OpenCLとどっちが速いですか？ 730 名前：デフォルトの名無しさん mailto:sage [2011/06/09(木) 06:45:26.86 ] 自助努力って具体的に何？ 731 名前：デフォルトの名無しさん mailto:sage [2011/06/09(木) 10:03:14.93 ] >>423 亀だけど、cufftPlanは対象とするデータのサイズと同じサイズをバッファとして確保するみたい。だから、二倍以上メモリが空いてないとこける。 732 名前：デフォルトの名無しさん mailto:sage [2011/06/09(木) 17:23:54.07 ] >>730 そういうのを自分で考えたり調べたりする姿勢や行動 733 名前：デフォルトの名無しさん mailto:sage [2011/06/09(木) 18:41:25.08 ] OpenGL de プログラミングというサイトや CUDAプログラミング入門とかフリーであるもの 見たほうがそこらの書籍よりかなり役立つよ OpenCLは少し大変だけどCPUならdouble型 オプションあるからいいんだよね ＣＵＤＡにもあるのかな? 734 名前：デフォルトの名無しさん mailto:sage [2011/06/09(木) 18:58:20.29 ] あった、ごめん 735 名前：デフォルトの名無しさん mailto:sage [2011/06/11(土) 18:37:21.68 ] visual studio pro, radeon 6000 台　で ati stream ないし　open cl 使っての　並列FPU高速化ってどのくらい大変ですか？ 一般のプログラムと英語ドキュメント読解に問題はないです 736 名前：デフォルトの名無しさん mailto:sage [2011/06/11(土) 18:40:16.57 ] スレタイ読む能力はなかったようだな 737 名前：デフォルトの名無しさん mailto:sage [2011/06/11(土) 19:04:07.23 ] 間違えました。すいません。 738 名前：デフォルトの名無しさん mailto:sage [2011/06/12(日) 06:56:49.05 ] 漫才かw 739 名前：デフォルトの名無しさん mailto:sage [2011/06/12(日) 19:39:36.69 ] >>732 CUDAを勉強してみようかと少し取り組んでみましたが、 C言語の基本も知らぬプログラミング初心者なので C言語から勉強しようと思ってます。 740 名前：デフォルトの名無しさん mailto:sage [2011/06/12(日) 19:49:59.75 ] Rやoctaveに触れると、Cの世界に戻れない。 故にCUDAに触れることも出来なくなる。 741 名前：デフォルトの名無しさん mailto:sage [2011/06/12(日) 19:55:24.78 ] Rでモンテカルロ出来るんだっけ？ 742 名前：デフォルトの名無しさん mailto:sage [2011/06/12(日) 20:12:36.74 ] >>739 C言語の文法を知ることは必須としても、まずプログラムがどういうもので何が出来るものなのか、って視点からも眺めた方がいいよ。 文法を覚えることに溺れないように 743 名前：デフォルトの名無しさん mailto:sage [2011/06/12(日) 20:32:53.39 ] >>741 統計解析ソフトだから乱数の生成は得意だと思うけど、 ループが苦手だから本当にモンテカルロが必要な場面でが役立たずかも。 Rユーザから見ると、curandライブラリがRから呼び出せたら嬉しい。 744 名前：デフォルトの名無しさん [2011/06/13(月) 16:13:05.08 ] カーネルを起動するのにかかるコストってNVIDIAの資料等に載ってるんでしょうか？ 745 名前：デフォルトの名無しさん mailto:sage [2011/06/13(月) 17:09:26.42 ] 資　料　読　め　よ 746 名前：デフォルトの名無しさん mailto:sage [2011/06/13(月) 20:52:40.31 ] 　　　　　　　　　　　YES　→　【見つかった？】　─　YES → じゃあ聞くな死ね 　　　　　　　　　／　　　　　　　　　　　　　　　　　 ＼ 【資料見た？】　　　　　　　　　　　　　　　　　　　　　　　 NO → なら、ねぇよ 　　　　　　　　　＼ 　　　　　　　　　　　 NO → 死ね 747 名前：デフォルトの名無しさん mailto:sage [2011/06/13(月) 22:35:02.62 ] 日本語の試料ないですか？ 頭痛いです。 748 名前：デフォルトの名無しさん mailto:sage [2011/06/13(月) 22:46:41.93 ] 　　　　　　　　　　　YES　→　【見つかった？】　─　YES → じゃあ聞くな死ね 　　　　　　　　　／　　　　　　　　　　　　　　　　　 ＼ 【探した？】　　　　　　　　　　　　　　　　　　　　　　　 NO → それより僕と踊りませんか？ 　　　　　　　　　＼ 　　　　　　　　　　　 NO → 死ね 749 名前：デフォルトの名無しさん mailto:sage [2011/06/13(月) 22:56:54.15 ] >>746 くだすれなんだし並列に書いてみてはどうか 750 名前：デフォルトの名無しさん mailto:sage [2011/06/13(月) 23:03:14.51 ] >>749 ifが多すぎてパフォーマンスが落ちると思う 751 名前：デフォルトの名無しさん mailto:sage [2011/06/13(月) 23:07:55.03 ] >>750 嫁>>1 752 名前：デフォルトの名無しさん mailto:sage [2011/06/14(火) 09:21:34.64 ] >>1の序文が全く無視されているな 753 名前：デフォルトの名無しさん mailto:sage [2011/06/14(火) 09:45:27.25 ] >>747 日本語の書籍購入するのが早いとおもうよ。 ・ASCII.technologies (アスキードットテクノロジーズ) 2010年 08月号 ・はじめてのCUDAプログラミング ・CUDA by Example 汎用GPUプログラミング入門 この辺がお勧め。自分はこの順で読んだ。 でもプログラミングガイドとか読むのは避けられないはず。 ・CUDAプログラミング実践講座 - 超並列プロセッサにおけるプログラミング手法 これは買ってないけど良さそう。 754 名前：デフォルトの名無しさん mailto:sage [2011/06/14(火) 09:47:57.13 ] >>753 上から三つ目は一つ目と二つ目が参考資料だっからな アスキーは触りとしてはわかりやすかった どれにしろFermiについて書いてないけどね 755 名前：デフォルトの名無しさん mailto:sage [2011/06/14(火) 09:55:59.01 ] 4.0になると、もうどうしましょ 756 名前：デフォルトの名無しさん mailto:sage [2011/06/14(火) 20:59:49.06 ] ■後藤弘茂のWeekly海外ニュース■ AMDが発表したメインストリームAPU「Llano」のアーキテクチャ ttp://pc.watch.impress.co.jp/docs/column/kaigai/20110614_452948.html >もう1つのバスは「Fusion Compute Link (Onion)」で、こちらはCPUとのコヒーレントバスになっている。 >このバスを使うと、GPUコアが、これまでできなかったCPUキャッシュへのスヌープをできるようになる。 >このOnionバスを使うことで、GPUでの汎用コンピューティング時に、 >CPUとGPUの間での無駄なメモリコピーを排除するゼロコピーが実現できる。 757 名前：デフォルトの名無しさん mailto:sage [2011/06/14(火) 22:14:25.10 ] あとはメインメモリがGDDR5になるのを待つだけだな． 758 名前：デフォルトの名無しさん mailto:sage [2011/06/14(火) 22:50:31.08 ] また妙なことを始めたなAMDw 759 名前：デフォルトの名無しさん mailto:sage [2011/06/14(火) 22:59:19.26 ] メモリコピーがネックにならない計算ってむしろおかしい？ 計算の方が時間かかってるんだけど もっと最適化できそうとか 760 名前：デフォルトの名無しさん mailto:sage [2011/06/15(水) 01:56:49.38 ] 利用しているデータ量が少ないなら別におかしくはないと思うけど。 761 名前：デフォルトの名無しさん [2011/06/15(水) 14:58:40.10 ] SDK4.0(x32)のtemplateを別フォルダにコピーしてVS2008でビルドしようとしたらcutil32D.libが開けませんって出ます。 助けてください。 762 名前：デフォルトの名無しさん mailto:sage [2011/06/15(水) 16:01:33.65 ] >>761 720 763 名前：デフォルトの名無しさん mailto:soab_b_nb@yahoo.co.jp [2011/06/15(水) 17:10:58.62 ] >>762 レスありがとうございます。 cutilとshrUtilsは既にビルドはしているんですが、開けないとでるんですよ。。 764 名前：デフォルトの名無しさん mailto:sage [2011/06/15(水) 21:41:20.65 ] >759 メモリコピーの時間がかかってでも並列演算したいって問題がほとんどだから 別におかしくないんじゃない。 並列計算が早すぎてメモリコピーがネックのように見えるってことはあるけど。 765 名前：デフォルトの名無しさん mailto:sage [2011/06/15(水) 22:05:59.38 ] メモリのコピー時間は0でも 同時にはアクセスできないよね、きっと。 でも大した問題でもないか。 766 名前：デフォルトの名無しさん mailto:sage [2011/06/16(木) 03:21:57.35 ] Microsoft Going All-in on GPU Computing ≪ NVIDIA blogs.nvidia.com/2011/06/microsoft-going-all-in-on-gpu-computing/?sf1642229=1 767 名前：デフォルトの名無しさん mailto:sage [2011/06/16(木) 06:03:44.90 ] >>754 アスキーは2冊あって2010/08はFermiについても掲載されてるよ そのほかにOpenCL、Direct compute, ATI streamについても載ってる 768 名前：デフォルトの名無しさん mailto:sage [2011/06/16(木) 15:34:19.10 ] 物凄く基礎的な質問ですいません。 今大学でCUDAの勉強を始めた者です。 現在貸し出されたPCに入っていたのがGeForce7600GSで、 CUDAの公式を見る限りだとCUDAに対応していないようですが どうにかして動かすことは出来ないでしょうか？ 一応一番初期のVer1.0〜最新の4.0までひと通り試してみたのですが サンプルのdeviceQueryを動かしてもCUDA Driver Versionは0.0のままでした。 もし可能でしたら少しでも助言を・・・ 無理なら無理と言ってくだされば大学側に申請して 可能なGPUの購入を考えてます。 769 名前：デフォルトの名無しさん mailto:sage [2011/06/16(木) 15:37:06.49 ] CUDAの勉強するならCUDAに対応したグラボを用意して貰え 対応してない物が仮に動いたとしても、それが正確な動作かわからんから勉強にならんぞ 770 名前：デフォルトの名無しさん mailto:sage [2011/06/16(木) 15:54:45.30 ] cuda x86 Amazon EC2 GPU買わなくてもおｋ 771 名前：デフォルトの名無しさん mailto:sage [2011/06/16(木) 16:13:17.23 ] >>768 安いのなら1万切るから自分でグラボ買ってもいいんじゃない？ 772 名前：デフォルトの名無しさん mailto:sage [2011/06/16(木) 16:48:57.04 ] >>771 大学研究室ならキッチリ金出してくれて580くらいはポンと買ってくれるだろうから自分で買う意味ないだろう 773 名前：デフォルトの名無しさん mailto:sage [2011/06/16(木) 18:37:25.67 ] ゲーム用と思われるのもあれだからTesla 20 Cシリーズ買ってもらいなよ 774 名前：デフォルトの名無しさん mailto:sage [2011/06/16(木) 22:59:07.51 ] >>768 授業？独学？ 研究のわけないよね GPGPUの研究でDX9のGPUとか遠回しなアカハラだｗ 775 名前：デフォルトの名無しさん mailto:sage [2011/06/17(金) 04:29:44.26 ] はじめてのCUDA読んでソースみたほうが早いかもよ それかここのpdfがお勧め ttp://accc.riken.jp/HPC/training.html GPUと大量の配列使うならCUDAが楽だよねえ 自作のテキストだけど、インストール部分だけなら役に立つかもしれない ttps://docs.google.com/viewer?a=v&pid=explorer&chrome=true&srcid=0B3RsNc5-fK5OZWY4MDVjYTItOTM2MS00NDc3LWE3NjEtZTc5YmQyMjg4Y2Q4&hl=ja&authkey=CJHji7AH 776 名前：デフォルトの名無しさん mailto:sage [2011/06/17(金) 07:42:07.25 ] >>775 横からだが、これは面白そう 読んでみるわ 777 名前：デフォルトの名無しさん mailto:sage [2011/06/17(金) 10:15:21.27 ] 同じく横からだが、ブクマした 全部印刷しておいても良いかもしれん 778 名前：デフォルトの名無しさん mailto:sage [2011/06/19(日) 02:32:29.65 ] >>775 参考になりそうだ。DLできないの？ 779 名前：デフォルトの名無しさん mailto:sage [2011/06/19(日) 10:39:21.50 ] できますん 780 名前：デフォルトの名無しさん mailto:sage [2011/06/19(日) 22:33:54.54 ] 二つ質問があります �@felmiにはキャッシュが搭載されているみたいですがプログラム中で明示しなくてもキャッシュにより高速化するのですか？ �AGPUでcuda化して高速化が見込めるアプリケーションが載ってるサイトとか知りませんか？ソースコードつきで。できればC言語だとうれしいです 781 名前：デフォルトの名無しさん mailto:sage [2011/06/19(日) 22:41:52.94 ] めんどく だれかこたえてあげれや 782 名前：デフォルトの名無しさん mailto:sage [2011/06/19(日) 22:54:32.14 ] この程度の事を自分で調べられず、聞いてしまう人間には 使いこなすのは無理だよ。 783 名前：デフォルトの名無しさん mailto:sage [2011/06/19(日) 23:07:34.93 ] >>782 一応は調べました せめて�@だけでも答えてもらえると助かります �Aに関しては自分もいくつか見つけましたが他の方がどういったサイトを使ってるか気になったので聞いてみました 784 名前：デフォルトの名無しさん mailto:sage [2011/06/19(日) 23:08:08.80 ] 機種依存文字-10 綴り間違い-10 sage+3 投稿時間+10 内容に言及しなくても-7点だ もうちょいがんばれ 785 名前：デフォルトの名無しさん [2011/06/19(日) 23:31:02.35 ] 教えるきないなら書き込むなよ、暇人 786 名前：デフォルトの名無しさん mailto:sage [2011/06/19(日) 23:52:40.95 ] >>785 ﾌﾟｯ 787 名前：デフォルトの名無しさん mailto:sage [2011/06/19(日) 23:52:54.57 ] キャッシュは勝手に使われるが、 キャッシュによって高速化(何と比べて？)するかはアプリによる。 ttp://journal.mycom.co.jp/articles/2010/07/21/fermi_cache/index.html こんな話もあるから、キャッシュが如何使われるか位は 意識しないと駄目だし、究極を目指せば結局コアレスアクセスに なってキャッシュの意味は無くなるかも知れない。 で、そこまで頑張れば微妙にteslaの方が速くなる。 788 名前：デフォルトの名無しさん mailto:sage [2011/06/20(月) 02:53:19.67 ] ランダムアクセスでテストしてみたが（sm1.1でコンパイル）、コアレスと比較して、260では1/5まで速度が落ちたが、470は2/5だった びっくらこいた 高速化できる例はCUDA SDKのサンプルみればよろし 789 名前：デフォルトの名無しさん mailto:sage [2011/06/20(月) 02:55:44.71 ] 補足 G80/G92のように総レジスタ数の少ないGPUはもっと遅くなる。大体1/10 790 名前：デフォルトの名無しさん mailto:sage [2011/06/20(月) 07:37:03.23 ] キャッシャが要らなきゃ共有メモリを増やせるし、Fermiは便利だね。 791 名前：デフォルトの名無しさん mailto:sage [2011/06/22(水) 16:10:37.96 ] fermiで２ワープを２サイクルで処理していると記載を見かけるのですが、なぜ１サイクル１ワープでなく2サイクル２ワープなのでしょうか？ どういう原理なのかご存知の方がいらっしゃいましたら教えていただけませんか。 また、ご存知でしたらその記載がどこにあるか教えていただけませんか。 お願いします。 792 名前：デフォルトの名無しさん mailto:sage [2011/06/22(水) 18:40:23.89 ] 「fermi half warp」でググればいくらでも出てくる．

---

---

[ 新着レスの取得/表示 (agate) ] / [ 携帯版 ]

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