【GPGPU】NVIDIA CUDA質問スレッド

1 名前：デフォルトの名無しさん mailto:sage [2007/09/17(月) 14:54:28 ] 本家 developer.nvidia.com/object/cuda.html 892 名前：デフォルトの名無しさん mailto:sage [2008/12/18(木) 00:49:51 ] sourceforge.net/projects/cudavswizard どうぞ 893 名前：デフォルトの名無しさん mailto:sage [2008/12/18(木) 01:47:35 ] >>892 それ使ってるんですが・・・readmeを嫁ということですね 894 名前：デフォルトの名無しさん mailto:sage [2008/12/18(木) 03:39:45 ] ptx出力がどうなるように最適化していけばいいんですか？ 895 名前：デフォルトの名無しさん mailto:sage [2008/12/18(木) 07:42:38 ] >>894 >889なら、registerの数を減らしたいんだろ。ptxの出力を読んで少なくなっていればOKだw 896 名前：デフォルトの名無しさん mailto:sage [2008/12/20(土) 00:21:41 ] >>862 NVIDIA、モバイルGPUドライバを配布開始 〜ノートでPhysXなどCUDAアプリケーションが利用可能に 12月18日(現地時間)配布開始 　米NVIDIAは18日(現地時間)、モバイルGPU用の汎用ドライバを同社ホームページ上で配布開始した。 対応GPUはGeForce 8M/9M、およびQuadro NVS 100M/300Mシリーズ。 　従来、NVIDIA製モバイルGPUのドライバは、ノートPC本体のメーカーが提供していたが、 今回よりNVIDIAが汎用ドライバを提供することになった。対応OSはWindows XP/Vista。 　バージョンはデスクトップ版の180代に近い179.28 BETAで、CUDAに対応。これにより、PhysXや、 CUDA対応の動画編集ソフト、Photoshop CS4などでアクセラレーションが効くようになる。 ただし、PhysXについてはビデオメモリが256MB以上必要となる。 　なお、ソニーVAIO、レノボThinkPad、デルVostro/Latitudeシリーズ、およびHybrid SLI構成の 製品については対象外となっている。 　現バージョンはベータ版で、WHQL準拠の正式版は2009年初頭の公開を予定している。 pc.watch.impress.co.jp/docs/2008/1219/nvidia2.htm 897 名前：デフォルトの名無しさん [2008/12/21(日) 00:11:21 ] なんでvaioは対象外ですか？いじめですか？？ 898 名前：デフォルトの名無しさん mailto:sage [2008/12/21(日) 07:03:38 ] vaioとか一部のノートはリコールされてるから 899 名前：デフォルトの名無しさん mailto:sage [2008/12/21(日) 19:00:56 ] それ以前にノートだとメーカーが色々やってるから 一般的にドライバ更新さえノートのメーカー対応待ちだろ。 900 名前：デフォルトの名無しさん mailto:sage [2008/12/21(日) 22:30:31 ] シェーダプロセッサってエラーがあってもゲームでは無視されて 認識出来ないレベルで画面が崩れるだけだから 一見正常に動いてると思っててもＣＵＤＡは動かない人が多いかも ドライバが対応してるとか以前にハードが用件を満たしてない 901 名前：デフォルトの名無しさん mailto:sage [2008/12/24(水) 10:44:05 ] どこでエラーになっているのかデバッグしろよ 902 名前：デフォルトの名無しさん [2008/12/25(木) 15:53:55 ] ちょいちょいCUDAをいじり始めてるんですが、OpenCLの規格が決まったことで、 CUDAの仕様も大幅に変わってしまうんですかね。 903 名前：デフォルトの名無しさん mailto:sage [2008/12/25(木) 18:15:24 ] MacBookProなんかだと自前でドライバ突っ込めばCUDA使えたんだけど， VaioやThinkPadだとどうですか？ 使っている人がいれば情報求む． 904 名前：デフォルトの名無しさん [2008/12/25(木) 18:21:26 ] 5秒制限の条件とか回避の仕方とかがよくわからないのですが 教えてもらえないでしょうか？ 905 名前：デフォルトの名無しさん mailto:sage [2008/12/26(金) 01:44:56 ] ThinkPad T61はOpenGL関係以外は使えた。 906 名前：デフォルトの名無しさん mailto:sage [2008/12/26(金) 03:48:04 ] OpenCLはAPIのみだからCUDAとは別にライブラリが用意されるだけだと思うよ こっちは専用コンパイラも必要ない 907 名前：デフォルトの名無しさん mailto:sage [2008/12/26(金) 14:22:18 ] >>904 Gridを何回も起動するしかないのでは。だから、途中結果をメモリにのこしとかないといけない。 面倒だよね。 908 名前：デフォルトの名無しさん mailto:sage [2008/12/26(金) 14:35:42 ] >>904 Xを使わない、グラフィック機能を使わなければいけるはず。 …何のためのグラフィックカードやねんって感じだが 909 名前：デフォルトの名無しさん [2008/12/26(金) 17:45:27 ] >>907>>908 やっぱりそれぐらいしかないですか。 現状でグラフィック機能止めて使うのは無理なので Gridの分割を考えてみます。 返答ありがとうございます。 910 名前：デフォルトの名無しさん mailto:sage [2008/12/26(金) 18:15:35 ] Vistaだとレジストリ弄ればもうちょっと長くできるよーみたいなこと言ってたような、言ってなかったような 911 名前：デフォルトの名無しさん mailto:sage [2008/12/26(金) 22:25:09 ] 使うグリッド数を制限すればなんとかならなくもないんだが、効率は落ちるんだよね。 つーか、効率重視ならもう一枚挿すしかない希ガス。 912 名前：,,・´∀｀・,,）っ-○◎● mailto:sage [2008/12/27(土) 00:47:07 ] こんなの教えて貰った compview.titech.ac.jp/Members/endot/adv-app-hpc/2008schedule 913 名前：デフォルトの名無しさん mailto:sage [2008/12/27(土) 05:05:39 ] どのみちCUDA実行中は画面の描画が止まるから割り込み入れる工夫は必要でしょ 914 名前：デフォルトの名無しさん mailto:sage [2008/12/27(土) 08:47:20 ] 描画要求の方がGridよりも優先されているようにしか見えないんだけど。 止まってくれたらどんだけ効率が改善することかw 915 名前：デフォルトの名無しさん mailto:sage [2008/12/27(土) 11:52:08 ] >>912 面白そうだけど、東工大は遠いなぁ。ワープアSEに新幹線は高い。 わがまま言うと、神戸でやって欲しい。 916 名前：デフォルトの名無しさん mailto:sage [2008/12/27(土) 14:13:48 ] 大岡山の東工大なら歩いて行けるぜ！ 917 名前：,,・´∀｀・,,）っ-○◎● mailto:sage [2008/12/27(土) 23:27:30 ] >>915 六甲台か？ あの急勾配の坂は登りたくないな。 918 名前：デフォルトの名無しさん mailto:sage [2008/12/29(月) 05:06:23 ] >>912 これは紹介レベルだから、たいした内容ではないよ 919 名前：デフォルトの名無しさん mailto:sage [2008/12/29(月) 11:44:34 ] CUDAをクアッドコアで動くように出来ればいいのに 920 名前：デフォルトの名無しさん mailto:sage [2008/12/29(月) 12:15:26 ] >>919 どういう意味? CUDAをQuadCoreCPUで動かすことなら何の問題もなくできるのだけど。 921 名前：デフォルトの名無しさん mailto:sage [2008/12/29(月) 12:44:32 ] っOpenCL 922 名前：デフォルトの名無しさん [2009/01/05(月) 20:09:11 ] VIPPERの諸君　立ち上がれ！ 以前韓国を打ち破ったのは誰か知らない、小学生や、他サイトの管理人がVIPを馬鹿にしている！ 全員集結せよ。そして全員で打ち倒すのだ！ もし時間がない人は、この文章をそのまま他のスレに貼ってほしい！ 対策本部 takeshima.2ch.net/test/read.cgi/news4vip/1231149782/ 今こそ　Vipperの意地を見せつけるのだ！ 923 名前：デフォルトの名無しさん [2009/01/08(木) 17:49:55 ] nc = 100*100 bs = 50 dim3 dimBlock(bs,bs) dim3 dimGrid(sqrt(nc)/dimBlock.x,sqrt(nc)/dimBlock.y) kernel<<<dimGrid, dimBlock>>>(idata, odata, sqrt(nc) __global__ void kernel(float* idata, float*, odata, int nc) { index=blockIdx.x * blockDim.x + threadIdx.x + +(blockIdx.y * blockDim.y + threadIdx.y) * nc } この時のイメージは、Gird:2x2、Block:50x50でよいのでしょうか? それとこのままグローバルメモリで計算するのはできるのですが、一旦 シェアードメモリに退避して計算してグローバルメモリに戻す方法が サンプルを見てもうまくいきません。どういう感じになるのでしょうか? 924 名前：デフォルトの名無しさん [2009/01/08(木) 18:11:37 ] いつも疑問に思いつつ使っているのですがグローバルメモリ使うときも 共有メモリ使うときも float* data; CUDA_SAFE_CALL(cudaMalloc((void**)&data, sizeof(float) * 100) CUDA_SAFE_CALL(cudaMemcpy(data1, h_data, sizeof(float) * 100, cudaMemcpyHostToDevice)) になるのでしょうか?これはグローバルメモリにとっているのですよね。 共有メモリの話は、kernelのほうで使用するかどうかですよね? コンスタントメモリ使うときだけどうも CUDA_SAFE_CALL(cudaMemcpyToSymbol(data, h_data, sizeof(float) * 100)) としてホストからコピーするんですかね? 925 名前：デフォルトの名無しさん mailto:sage [2009/01/08(木) 23:27:14 ] >>924 cudaMemcpyToSymbolは定数メモリの他、グローバルメモリでも使えるよ。 要は、device空間で定義されたメモリと言うニュアンスでしかない。 逆に言えば、cudaMemcpyToSymbolを使わないとhost側のメモリ空間に割り付けられたアドレスを利用しようとしてしまう。 # この辺が判り難いのよね。 共有メモリはグリッド実行中しか使えないからhost側からは為す術がない。 >>923 どのサンプル見てどうやったらどう巧くいかないのか具体的に。 少なくとも、>923のような処理なら共有メモリの使いようがないんだけど。 926 名前：デフォルトの名無しさん mailto:sage [2009/01/09(金) 01:50:13 ] warp単位32スレッドで同じ命令を行って、 実際に一度に動くのは8スレッド こういう考え方でいいのですか？ 927 名前：デフォルトの名無しさん [2009/01/09(金) 18:21:54 ] >>925 結局現状では、cudaMemcpyXXXの違いは理解できそうにないです。 まあそれはよいとして(全然良くないですが) >>923 のidataは、実は、一次元配列なんです。なのにカーネルには2次元で利用しよう としているんです。 float host_idata[nc] float* idata cudaMalloc((void**) &idata, sizeof(float) * nc) cudaMemcpy(idata, host_idata, sizeof(float) * nc, cudaMemcpyHostToDevice) こんな感じのグローバル変数がidata。 としたとき、ncとbsの設定値によってはうまく処理できないことがあるのです。 これが謎です。わかる方いませんか?もちろんncはsqrtで整数になる値、設定する 値は、割り切れる値です。nc=1000*1000,bs=500で処理結果がおかしくなっています。 kernelの読んだあとエラーでinvalid configuration argument.になってます。 それと、共有メモリのほうですがこれは分かりました。 kernel<<<dimGrid, dimBlock, sizeof(float) * nc>>>(idata, odata, sqrt(nc)) 第3引数の共有メモリのサイズを定義しないとだめだった。ちゃんと説明書を 読んでいなかったというレベルでした。これは解決です。 928 名前：デフォルトの名無しさん [2009/01/09(金) 19:05:25 ] >>927 ブロックは512x512x64と出ているけどいちお3次元は分からないので512x512 としてxで512使い切るとどうもyではもう1しかとれないようです。 ncは、そのままだとエラーが出ないレベルは、bs=20つまり20x20=400ブロック 使用。21,22は割り切れないので設定できない。みたいです。 とすると一次元配列を使用するのにわざわざ、x,yの次元のブロックを使いきる テクニックは必要なさそうですね。と推察できると思います。結果からですが ただ、これがバグであるなら次のCUDAのバージョンアップで使えるようになる かもですね。 929 名前：デフォルトの名無しさん mailto:sage [2009/01/09(金) 20:53:02 ] >>928 次元とは別に1ブロックあたり512スレッドが上限じゃなかったっけ？ どっかにきちんと書いてあったはずだからバグではないよ。 930 名前：デフォルトの名無しさん mailto:sage [2009/01/09(金) 22:23:41 ] 私は1次元でやってしまうことが多いなぁ。 で、スレッド数についてはdeviceQuery参照で。 931 名前：デフォルトの名無しさん mailto:sage [2009/01/10(土) 21:32:30 ] スレッドの最大数はAPIで取得できる。全部512だと思う。 ブロックの最大は65535だから65535*512以上は1次元のグリッドでは扱えない。 あと最大は512だけど、192か256くらいのほうがパフォーマンスが出る。 932 名前：デフォルトの名無しさん mailto:sage [2009/01/14(水) 01:57:36 ] CUDA 2.1 Release ttp://forums.nvidia.com/index.php?showtopic=85832 933 名前：デフォルトの名無しさん [2009/01/14(水) 21:34:37 ] 情報サンクス。早速、Linux版ゲットした。 934 名前：デフォルトの名無しさん mailto:sage [2009/01/15(木) 01:08:35 ] 9600GTでVideoDecodeのサンプル動かそうとしたら正常に動かなくて、 結局SDK2.0+178.24まで戻さないとダメだった。 2.1にもソースは付いてるけど、2.0のままみたいだね。 935 名前：,,・´∀｀・,,）っ-●◎○ mailto:!sage [2009/01/15(木) 07:55:42 ] >>931 1 warp = 32 threadで、GeForce 8/9が24warp/block、GT200で32warp/blockが最大だから 768thread/blockと1024thread/blockが最大なんだけどね。本来は。 CUDAのドライバ側が512でリミッタかけてるんだ。罠としか言いようが無い。 逆に言うとCUDAを経由しなきゃ目いっぱい使えるかもね。 ただ、スレッドインターリーブすると１スレッドあたりで使えるレジスタ本数が減っちゃうんだよね。 メモリレイテンシを隠蔽するならスレッドを目いっぱい使ったほうがいいし 逆に一時変数を何度も再利用する場合は、thread/blockを減らして１スレッドあたりの仕事量を増やしたほうがいい。 936 名前：デフォルトの名無しさん mailto:sage [2009/01/15(木) 15:42:19 ] CUDAってシーユーディーエーって読んでいいの？ それともクダとか、なんか読み方あったりする？ 937 名前：デフォルトの名無しさん mailto:sage [2009/01/15(木) 15:57:04 ] >>936 クーダ 938 名前：,,・´∀｀・,,）っ-●◎○ mailto:sage [2009/01/15(木) 15:58:34 ] ネイティブの人は「クーダ」って発音してたよ。 Cellは「くた」だから注意な。 939 名前：デフォルトの名無しさん mailto:sage [2009/01/15(木) 16:01:48 ] ありがとう！明日から会社でクゥゥゥウーダ！！って叫びまくるわ 940 名前：デフォルトの名無しさん mailto:sage [2009/01/15(木) 16:02:44 ] きゅーだ　って発音してる俺は異端ですかね？ 941 名前：デフォルトの名無しさん mailto:sage [2009/01/15(木) 16:23:46 ] >940 そんな人いたんだ。 942 名前：デフォルトの名無しさん mailto:sage [2009/01/15(木) 16:29:33 ] 英語的に　クー　よりも　キュー　だと思わね？ 943 名前：デフォルトの名無しさん mailto:sage [2009/01/15(木) 16:56:19 ] >>941 異端だな 944 名前：デフォルトの名無しさん mailto:sage [2009/01/15(木) 16:57:46 ] どっちでもいいんじゃね ちなみにnudeはニュードとヌードの両方の発音が有 945 名前：デフォルトの名無しさん mailto:sage [2009/01/15(木) 17:13:01 ] キュもクも似てるし、まあ仮に間違って言ったとしても気づかれないんじゃない？ 946 名前：デフォルトの名無しさん [2009/01/15(木) 19:38:28 ] ホムペにはクーダと嫁 と書かれてたはず 947 名前：デフォルトの名無しさん mailto:sage [2009/01/15(木) 22:24:41 ] "The cuda and my wife" 948 名前：デフォルトの名無しさん mailto:sage [2009/01/15(木) 23:40:36 ] たしか英語版のfaqには kuh-da と書いてあったはず 949 名前：デフォルトの名無しさん mailto:sage [2009/01/15(木) 23:53:57 ] 2.1はJITサポートがミソ？ 950 名前：デフォルトの名無しさん [2009/01/16(金) 14:10:31 ] 2.1をつついた人に質問 2.1ではカーネルに渡す引数でポインタ配列は使えますか? 951 名前：デフォルトの名無しさん mailto:sage [2009/01/16(金) 18:48:04 ] 共有メモリからグローバルメモリに連続領域をコピーするにはどうすればいいでしょうか カーネルから呼び出せるmemcpyみたいなのがあれば教えてください 952 名前：デフォルトの名無しさん [2009/01/16(金) 18:58:33 ] cudaMemcpyが長いようで、the launch timed out and was terminated.が でてしまいす。これなんとかするほうほうありましたっけ? 953 名前：デフォルトの名無しさん mailto:sage [2009/01/16(金) 21:53:12 ] >>952 転送量を減らす。 >>951 cudaMemcpy 954 名前：デフォルトの名無しさん mailto:sage [2009/01/17(土) 00:16:47 ] >>951 スレッド内で連続的にメモリにアクセスするようなコードは望ましくない。 coalscedにアクセスするとなると、自分でコードを書くしかない。 共有メモリから転送を始める前に、同期を取ることも忘れずに。 >>952 参考までに、どんなハードウェアの組み合わせでどれくらい転送しようとしている? 8800GT辺りだと転送できるだけ全部転送してもタイムアウトしないのだけれど。 955 名前：デフォルトの名無しさん [2009/01/17(土) 02:52:44 ] >>954 カーネル実行後の結果データコピーでエラーになっているのですが sizeof(float)*1000バイトを100回連続でコピーしています。 100個とも別のアドレスに結果をコピーしたいためそうしてます。やはり いっぺんにメインメモリにいったんコピーしてからやったほうがいいので しょうか? 956 名前：デフォルトの名無しさん [2009/01/17(土) 03:03:49 ] >>955 sizeof(float)*10000バイトを100回連続コピーでした。 957 名前：デフォルトの名無しさん [2009/01/17(土) 07:16:40 ] >>954 ＞ >>952 ＞ どんなハードウェアの組み合わせ 958 名前：デフォルトの名無しさん mailto:sage [2009/01/17(土) 07:19:52 ] >>956 同一データ部分を100箇所にコピー？ 独立の100箇所をコピー？ 959 名前：デフォルトの名無しさん mailto:sage [2009/01/17(土) 07:55:29 ] >>956 それ、本当にメモリのコピーでタイムアウトしてる？ カーネル実行後に同期取らずにすぐにメモリのコピー始めて 100回コピーのどこかでカーネル実行がタイムアウトしてるとか。 cudaMemcpyの戻り値は Note that this function may also return error codes from previous, asynchronous launches. ということだし。 960 名前：954 mailto:sage [2009/01/17(土) 08:44:25 ] >>955=956 えーと、都合数MB程度ってことか。それだったら>959の言うようにカーネル関数のタイムアウトじゃないかな。 DEBUGつきでCUT_SAFE_CALL(間違ってたら失敬)を使うか、自前でcudaThreadsSynchronize()を呼んで戻り値を見てみて。 カーネル関数起動後は、復帰を待たずにホストの処理に戻るからカーネル関数自体のエラーはカーネル関数起動では捕捉できないのよ。 で、まさかとは思うけど同一データを100回コピーならバス越しにやってはダメ。別のデータならホスト側の用件次第かな。 # でも多分、一旦ホスト側に転送してから分配した方が速い気がする。CPUとメモリ周りがよっぽど遅くなければ。 >>957 下手な突っ込みはお郷が知れるよ。この場合、CPUが遅いんじゃないかとかバスがx1なんじゃないかとかGPUがモニタ表示か何かで よっぽど負荷が掛かっているんじゃないかとかGPU側メモリが巨大なSTELAなのかとか、組み合わせ次第で色々条件が変わってくるのよ。 961 名前：デフォルトの名無しさん [2009/01/17(土) 09:52:40 ] >>960 ＞ 下手な突っ込みはお郷が知れるよ。 質問者が答えてないのを指摘しただけなんだが… 誤解させてすまんかった 962 名前：デフォルトの名無しさん mailto:sage [2009/01/17(土) 11:09:08 ] そういう意味か。あの書き方じゃぁ、誤解されるよ。 まぁ、大量転送じゃないと言う意味で、今回は関係なさそうだね。 963 名前：デフォルトの名無しさん [2009/01/17(土) 15:38:06 ] >>952の話 >>957 ATIの統合チップセットでオンボードのATIビデオカードを無効 PCI-E 1.0 x16バスにGeforce9600GT 512MB CPUはAMDデュアルプロセッサ(結構遅いやつ) 964 名前：デフォルトの名無しさん [2009/01/17(土) 15:55:29 ] >>952の話 >>959 カーネル実行 a:cudaのエラーチェック(カーネルのエラーのありなし) cudaThreadsSynchronize for(100回 結果データをデバイスからホストにコピー b:cudaのエラーチェック(直前のcudaMemcpyのエラーのありなし) ) bのとこでthe launch timed out and was terminatedが出てるんです。 先週は時間がきてあきらめて帰ったのでここまでです。 言われたようにcudaThreadsSynchronizeの戻り値見てみたほうがよさそう ですね。 965 名前：デフォルトの名無しさん mailto:sage [2009/01/17(土) 18:45:09 ] >>951 共有メモリはカーネル起動時に動的に確保されるメモリ領域だから カーネルが終了したら消えるし、１つの実行スレッドから全部コピーなんて論理的におかしいし 共有メモリはあくまで高速動作と同期が取れる作業用のメモリとして考えたほうがいい 共有メモリを使って計算した結果は１個づつ普通のグローバルメモリに入れてやるのが正しいやり方 >>952 10000を一度に転送して実行しても 1を10000回繰り返して転送しても 実行時間は大差ないんですよ CUDAで実行する部分は出来るだけコンパクトにまとめて 呼び出す本体の方で特殊な演算関数くらいの感覚でループさせて使うのが正解 966 名前：デフォルトの名無しさん mailto:sage [2009/01/18(日) 23:59:40 ] JCublasについて おしえてエロイ人 967 名前：デフォルトの名無しさん [2009/01/19(月) 11:04:50 ] >>925の話 cudaThreadsSynchronizeの戻り値をチェックしたら the launch timed out and was terminatedが出ていました。結果コピーで エラーで落ちてたのだけどエラーデータは、その前に実行していた cudaThreadsSynchronizeの問題だったようです。 (cudaThreadsSynchronizeが正常になるまで待つとしても配列を 100x1000回、100x10000回、回すと1分待っても同期とれないようです。) つまりカーネルに同期できない処理が書かれていたのが原因だと思います。 NVIDIA_CUDA_Programming_Guide_2.0.pdfのp34-35のようなコードで かつそこのカーネルのdevPtrを参照するだけでなく書き戻す処理をや っているのでそれがだめなのだと思います。恐らくこういうピッチ処理 の場合は、参照のみが許されるのでは?と思います。 問題のコードをこのpdfの変数に直して下記に書いておきます。 __global__ void myKernel(float* devPtr, int pitch) { for (int r = 0; r < height - 1; r++) { float* row1 = (float*)((char*)devPtr + r * pitch); float* row2 = (float*)((char*)devPtr + (r + 1) * pitch); for (int c = 1; c < width - 1; c++) { row2[c] = row1[c]とrow1[c + 1]とrow1[c -1]やらの計算 } } } このp34-35を見て思うのはindex使わずにこんな処理書いてGPUでやる意味 あるの?と思いますが。自分で書いてなんなのですが 968 名前：デフォルトの名無しさん mailto:sage [2009/01/19(月) 17:28:20 ] お前はあほかｗ 969 名前：デフォルトの名無しさん mailto:sage [2009/01/19(月) 17:49:31 ] タイムアウトになる原因はそのでかいループのせい せいぜいミリ秒単位でタイムアウトを判断してるから ミリ秒以下で応答しないようなカーネルは絶対に書いては駄目 cudaThreadsSynchronizeは発行したカーネルがすべて終了するのを待つだけ グローバルメモリは読み書きは出来るが前後は保障されないので １スレッドが書き込みする箇所は限定する必要がある 共有メモリを使って他のスレッドが書き込んだ内容を参照して利用する場合に カーネル内部で___syncthreadを使う これが本来の同期の意味 970 名前：デフォルトの名無しさん [2009/01/19(月) 19:44:51 ] >>952の話 >>968のように言われるのは分かって書いてみたんだけど NVIDIA_CUDA_Programming_Guide_2.0.pdfのp34-35のようなコードが マニュアルに書いてあるのがおかしいと思う。 __global__ void myKernel(float* devPtr, int pitch){} そもそもこんな書き方じたいが書けるけど間違えな使い方だと。 この書き方しているとこにやらないようにこの部分に×印つけてほしい。 あとはコンパイラがえらくなったらfor多重ループをうまく処理する アセンブラぐらい作ってほしい。(OpenMPでパラレルfor指定すると#のタグだけ でプロセッサ使ってきって高速化なるんだけどね。そこまでぐらい将来 的にはしてほしい。) 971 名前：デフォルトの名無しさん mailto:sage [2009/01/19(月) 20:53:12 ] 文句言っても何にもならない。 そう悟る時がくるまで気長に待ちましょうよ。 そうすれば成長しまっせ。 972 名前：デフォルトの名無しさん mailto:sage [2009/01/19(月) 23:38:15 ] それよりも先ず、日本語を何とかしてくれ。 ＞間違えな使い方 ＞しているとこにやらないように ＞プロセッサ使ってきって 973 名前：デフォルトの名無しさん mailto:sage [2009/01/19(月) 23:41:25 ] そのサンプルはあくまでもcudaMallocPitch()がどういう風にメモリを確保するかの説明だからなぁ。 きっとライターは、まさかそのまま動かそうとされるは思っちゃいまいよ。 974 名前：デフォルトの名無しさん mailto:sage [2009/01/20(火) 02:02:49 ] ところでドライバ180なんだが2Dモードだかなんだかで消費電力抑える為に ドライバが自動的にクロックを下げて3Dゲーム動作させると上がる仕様なんだけど CUDA実行しても上がらないぞｗ 975 名前：デフォルトの名無しさん mailto:sage [2009/01/20(火) 02:08:30 ] ドライバを替えてくださいw つーか、それって169.09のときに指摘されて直したんじゃなかったのか?＞NVIDIA 976 名前：デフォルトの名無しさん [2009/01/20(火) 03:40:32 ] d_a[0][blockIdx.x][blockIdx.y][threadIdx.x] って使おうとしたら"expression must have arithmetic or enum type" ってでたんだけど何がいけないんすか？ 教えてください。 977 名前：デフォルトの名無しさん mailto:sage [2009/01/20(火) 07:58:55 ] >>976 d_aがポインタなら、途中で解決できない状況があるのかも。 978 名前：デフォルトの名無しさん mailto:sage [2009/01/20(火) 08:41:39 ] 共有メモリで構造体を使いたいのだけど 例えば struct data { int a; char b; } func<<<dim3(), dim3(), SIZE * sizeof(data)>>>(); __global__ kernel(){ __shared__ data d[SIZE]; } こんな感じでやるとCUDAとC側のパック構造が違うせいだと思うけどおかしくなる どうやれば出来るの？ 979 名前：デフォルトの名無しさん mailto:sage [2009/01/20(火) 08:59:56 ] パックが違うのだけが理由なら、 struct data {int a; char b; char dummy[3];}とでもしておけばいいじゃん。 980 名前：デフォルトの名無しさん mailto:sage [2009/01/20(火) 13:53:14 ] ただでさえ少ない共有メモリをそんな無駄に使えない 981 名前：デフォルトの名無しさん mailto:sage [2009/01/20(火) 20:44:37 ] >>977 世の中夜勤帰りで朝から寝てる人だっているんだよ？ 引っ越しの時ちゃんと挨拶行った？ 顔合わせたら軽く会話するとかしてちゃんとコンタクト取り続けてる？ 日頃からそういうコミニュケーションが取れてればいつ洗濯機を回していいのか いつ静かにしなければならないのか 迷惑を掛けないように生活出来るはずなんだが 982 名前：デフォルトの名無しさん mailto:sage [2009/01/20(火) 23:36:58 ] >>980 マジで言っているのなら、設計が悪い。 どうしてもパディングしたくないくらい逼迫しているなら、int一個を切り分けて3バイトと1バイトで使え。 983 名前：デフォルトの名無しさん mailto:sage [2009/01/21(水) 00:22:25 ] 共有メモリが制限されてるのに無駄な領域作って ほとんどをグローバルメモリに追いやる方がよっぽど設計が悪いでしょｗ 984 名前：デフォルトの名無しさん mailto:sage [2009/01/21(水) 00:25:51 ] GPUのメモリレイテンシって１２とかの世界だぞ CPU用のDDR2で５だからな intを内部でcharとして扱うプロセスを考慮しても共有メモリ使った方が早いんだよ 985 名前：デフォルトの名無しさん mailto:sage [2009/01/21(水) 00:35:27 ] CUDAで共有メモリを使うこと自体、無駄な同期待ちが発生するから避けたいところだと思うが。 coaxschedな読み書きができるなら、共有メモリより遅くないぞw 986 名前：デフォルトの名無しさん mailto:sage [2009/01/21(水) 02:24:30 ] >>975 (エンバグの)歴史は繰り返す。 987 名前：デフォルトの名無しさん mailto:sage [2009/01/21(水) 12:18:47 ] >>CUDAで共有メモリを使うこと自体、無駄な同期待ちが発生するから避けたいところだと思うが。 え？ 共有メモリってCPUで言う所のただのレジスタみたいなもんで同期は指定しないと発生しないと思うけど ところでローカル変数をすべて共有メモリに置いた場合のが早いんだけど ローカル変数ってデバイス上のメモリに確保されるだけで共有メモリより遅いレベルのメモリを使ってるっぽいね 988 名前：デフォルトの名無しさん mailto:sage [2009/01/21(水) 18:11:15 ] あー、表現が悪かった。共有メモリを共有目的に使うと同期を取る必要が出てくるから同期待ちが発生するということ。 レジスタみたいに使うのなら確かに関係なかったね。 で、レジスタよりも速いかどうかについてはptx見てみたいところ。 989 名前：,,・´∀｀・,,）っ-○◎● mailto:sage [2009/01/21(水) 18:18:05 ] 見た感じリードレイテンシはこれくらい レジスタ＞＞Const Memory＞Shared Memory＞＞＞＞DRAM

---

---

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

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