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

1 名前：デフォルトの名無しさん mailto:sage [2012/09/23(日) 23:17:47.58 ] このスレッドは、他のスレッドでは書き込めない超低レベル、 もしくは質問者自身何が何だが分からない質問を勇気を持って書き込むスレッドです。 CUDA使いが優しくコメントを返しますが、 お礼はCUDAの布教と初心者の救済をお願いします。 CUDA・HomePage ttp://developer.nvidia.com/category/zone/cuda-zone 関連スレ GPGPU#5 ttp://hibari.2ch.net/test/read.cgi/tech/1281876470/ 前スレ 【GPGPU】くだすれCUDAスレ【NVIDIA】 ttp://pc12.2ch.net/test/read.cgi/tech/1206152032/ 【GPGPU】くだすれCUDAスレ pert2【NVIDIA】 ttp://pc12.2ch.net/test/read.cgi/tech/1254997777/ 【GPGPU】くだすれCUDAスレ pert3【NVIDIA】 ttp://hibari.2ch.net/test/read.cgi/tech/1271587710/ 【GPGPU】くだすれCUDAスレ pert4【NVIDIA】 ttp://hibari.2ch.net/test/read.cgi/tech/1291467433/ 【GPGPU】くだすれCUDAスレ part5【NVIDIA】 toro.2ch.net/test/read.cgi/tech/1314104886/ 413 名前：デフォルトの名無しさん mailto:sage [2013/07/17(水) NY:AN:NY.AN ] nvidiaに聞け 公開資料にない事の予定問われても スレの住人はnvidia関係者な訳じゃないし 関係者が居たとしても、2chで非公開の予定情報の可否なんか答える訳ないだろ 414 名前：デフォルトの名無しさん mailto:sage [2013/07/17(水) NY:AN:NY.AN ] メジャーになればそれなりの対応もあるだろうが 影も形もないものを・・ 415 名前：デフォルトの名無しさん mailto:sage [2013/07/17(水) NY:AN:NY.AN ] 技術的に可能かどうかと言われれば可能でしょ。 メモリ空間が共通化されれば、GPUの演算器がCPUのSIMD演算器のように扱えるわけだし。 ただCUDAである必要があるかどうかはNVIDIAが判断するんじゃないか？ 416 名前：デフォルトの名無しさん mailto:sage [2013/07/29(月) NY:AN:NY.AN ] N×1行列とM×N行列を計算して結果をテクスチャに書き込むという単純な処理で これを合計512スレッド(Mに関して並列化)で実行しているんだけど(N=3000 M=512) 各ブロックを16×16スレッドの2ブロックよりも 各ブロックは16×1スレッドの32ブロックのほうが2〜3%速度が速いという不可解な結果が出てしまっている 何でこんなことが起こるんだろう 417 名前：デフォルトの名無しさん mailto:sage [2013/07/30(火) NY:AN:NY.AN ] 適度に粒度下がってスケジューラの効率上がったとか？ 418 名前：デフォルトの名無しさん mailto:sage [2013/07/30(火) NY:AN:NY.AN ] local memoryにレジスタが溢れているとか 419 名前：デフォルトの名無しさん mailto:sage [2013/07/30(火) NY:AN:NY.AN ] こういう予測しにくい挙動こそがGPGPUのクソなところ。 420 名前：デフォルトの名無しさん mailto:sage [2013/07/30(火) NY:AN:NY.AN ] そういえばNVIDIAはPGIを買収したらしいね ここに書くことでもないかも知れないが 421 名前：デフォルトの名無しさん mailto:sage [2013/07/31(水) NY:AN:NY.AN ] >>416 Nを並列化せずに本当に512スレッドしか使ってないんだったら、16warp*32threadより 32warp*16threadの方がわずかに効率が良いというだけの話じゃないのか？分岐やなんかで。 422 名前：デフォルトの名無しさん [2013/07/31(水) NY:AN:NY.AN ] >>416 メモリの配置次第で不可解でも何でもないと思うが。 423 名前：デフォルトの名無しさん [2013/08/07(水) NY:AN:NY.AN ] 【AMD涙目】デファクトスタンダードへの道を突き進むCUDA・・IBMやGoogleも陣営に加わる engawa.2ch.net/test/read.cgi/poverty/1375883942/ 424 名前：デフォルトの名無しさん mailto:sage [2013/08/08(木) NY:AN:NY.AN ] デファクトスタンダードへの道を突き進むCUDA　IBMやGoogleも陣営に加わる kohada.2ch.net/test/read.cgi/pcnews/1375893221/ 425 名前：デフォルトの名無しさん mailto:sage [2013/08/08(木) NY:AN:NY.AN ] 性能の低さを政治力でカバーするのか？ 426 名前：デフォルトの名無しさん mailto:sage [2013/08/08(木) NY:AN:NY.AN ] まぁニュースサイトへの宣伝広告費でカバーしてるところより良いと思うよ 427 名前：デフォルトの名無しさん mailto:sage [2013/08/08(木) NY:AN:NY.AN ] カーネル関数内で，乱数が生成されたD_c配列を用いて計算したくて， D_c配列のポインタを渡していくのが面倒なので以下としたけど，上手くいかない． __device__ double D_c[is_110][is_110][is_110]; curandGenerateUniformDouble(generator, (double*)&D_c, count); cutilSafeCall(cudaMemcpyFromSymbol(H_a, D_c, size_a));　　//ここで30：Unknown error cudaMallocの場合には動くから，乱数生成場所がおかしいんだと思うけど． でも，curandGenerateの第2引数はdouble*型だけど，&D_cで配列の先頭を示してるから， 実質同じ事だと考えてたんだけど，違うの？ 428 名前：デフォルトの名無しさん mailto:sage [2013/08/08(木) NY:AN:NY.AN ] >>427 今手元にないから記憶で書くけど、D_cの型がdouble []じゃないからだと思う。 D_cの定義をdouble D_c[is_110 * is_110 * is_110]にしてみたらどうなる? 429 名前：デフォルトの名無しさん mailto:sage [2013/08/08(木) NY:AN:NY.AN ] >>428 __device__ double D_c[is_110* is_110* is_110]; curandGenerateUniformDouble(generator, D_c, count); にしても，同じエラーだった． (double*)&D_cでdouble*型に変換してるからイケると思ってたんだけど… __device__ double *D_c;でcudamallocしたら乱数生成はできたけど， カーネル関数内でD_c[id_array]がアクセスエラーっぽい． __device__でやるのは無理があるのかなぁ？ 430 名前：デフォルトの名無しさん mailto:sage [2013/08/08(木) NY:AN:NY.AN ] >>429 うーん、よく判らんな。 少なくとも、cudaMemcpyFromSymbol()はCのコードとしては特殊(トリッキー)な仕様だから 使い難いんだよね。まるでプリプロセッサマクロのように。 だから私のところではリファクタリングの結果、使うのをやめてしまっている。 変に悩むくらいなら、代替手段を考えたほうがいいかもよ。 431 名前：デフォルトの名無しさん mailto:sage [2013/08/08(木) NY:AN:NY.AN ] >>430 特殊なのかー． アドバイス通り，普通にホストでMallocして，カーネル関数にポインタ渡す方式にします． ありがとう！ 432 名前：デフォルトの名無しさん mailto:sage [2013/08/22(木) NY:AN:NY.AN ] ドライバをアップデートすると演算性能上がりますか?落ちますか? 433 名前：デフォルトの名無しさん mailto:sage [2013/08/23(金) NY:AN:NY.AN ] __global__の関数内でグローバルメモリの内容を一気にコピーしたいんだけどそういう方法ってある？ 一要素ずつやった方が無難？ 434 名前：デフォルトの名無しさん mailto:sage [2013/08/30(金) NY:AN:NY.AN ] >>433 コピーの手前でカーネル切って、新たに __global__ copy(double *a, double *b) { int tid = blockIdx * blockDim.x + threadIdx.x; b[tid] = a[tid] } を実行するのじゃだめ? あるいはダイナミックパラレリズムで何とかなるのかな。 435 名前：デフォルトの名無しさん mailto:sage [2013/08/30(金) NY:AN:NY.AN ] カーネル分けんでもindexの生成を工夫すれば同じことでしょ。 それがDtoDなら。DtoHやHtoDなら設計を見直すべき。 436 名前：デフォルトの名無しさん mailto:sage [2013/08/31(土) NY:AN:NY.AN ] ツールキット5.5プロダクションリリースってどういう意味ですか? ベータ版とかじゃなくて正式に使えるってこと? 437 名前：デフォルトの名無しさん mailto:sage [2013/08/31(土) NY:AN:NY.AN ] curandGenerateUniformって1.0を含むけど1.0を含まない乱数生成ってないの？ 438 名前：デフォルトの名無しさん mailto:sage [2013/08/31(土) NY:AN:NY.AN ] それ使う所の式変形で対処できないの？ 439 名前：デフォルトの名無しさん mailto:sage [2013/09/03(火) 08:02:37.87 ] >>435 ブロックまたいで同期する必要がある場合は? 440 名前：デフォルトの名無しさん mailto:sage [2013/09/03(火) 13:46:17.67 ] .cubin ファイルをエディターで・・・と 公式スライドにあるのですが、本当ですか？ 441 名前：デフォルトの名無しさん mailto:sage [2013/09/04(水) 00:17:50.20 ] cudaErrorInvalidValueってカーネルがエラーはいたんだけどどういう状況でなる？ 442 名前：デフォルトの名無しさん mailto:sage [2013/09/04(水) 08:47:55.31 ] >>441 詳細はAPIマニュアル見てねだけど、 ホントにカーネルで出てる? APIの引数を間違ってるんじゃない? 443 名前：デフォルトの名無しさん mailto:sage [2013/09/05(木) 06:10:12.96 ] 実行はできるし結果もそれっぽいけど画面が不安定になるなあ 444 名前：デフォルトの名無しさん mailto:sage [2013/09/08(日) 14:31:18.78 ] >>443 カーネル実行時間が長すぎると画面が真っ暗になる場合がある。 制限時間はレジストリをいじって変えられたはず。 445 名前：デフォルトの名無しさん mailto:sage [2013/09/09(月) 21:57:16.85 ] >>444 レジストリはいじったあとなんだけどそうなる 446 名前：デフォルトの名無しさん mailto:sage [2013/09/10(火) 12:17:17.37 ] >>445 メモリアクセスが間違ってても落ちることあるですよ。 447 名前：デフォルトの名無しさん mailto:sage [2013/09/10(火) 23:29:52.04 ] Geforce GT 530 で CUDAインストーラー(5.5)が「対応デバイスが無い」とかで失敗するので 古いドライバに変えてみたり再起動繰り返したりしたのですが、 developer.nvidia.com/cuda-gpus に、GT530が載ってませんでした・・・。 マシンに「nVidia GEFORCE with CUDA」のシールあるし、GPU対応のソフトも動いているのですが 開発用としては使えないってことでしょうか？ 448 名前：デフォルトの名無しさん mailto:sage [2013/09/10(火) 23:41:30.09 ] セットアップがミスってんでしょ。 449 名前：447 mailto:sage [2013/09/11(水) 00:36:33.98 ] セットアップのミスの原因ってなにかありますか？ 展開後は高速インストールかカスタムか選ぶだけだし、 どちらを選んでも失敗します・・・。 プログレスバー見てる感じ、Toolkitのインストール中１０％くらいでエラーが出ます。 c:\NVIDIA\CUDAへの展開でかなり時間かかりますが、 インストール失敗すると、この下のインストーラー各種全部削除されて 最初からやり直す仕様なんですね・・・。（成功しても消えるのかもしれませんが） 450 名前：デフォルトの名無しさん mailto:sage [2013/09/11(水) 00:52:29.92 ] ちゃんとリリースノートやインストールノート読んでやってないところ あとはほんとに５３０のせいかどうかを他のグラボ（nvidiaね）に変えて切り分けして原因しぼっていくしかないだろ。 可能性だけなら グラボ、PCパーツの不良 DLしたソフトウエアの不良 システム不良 などたくさんあんだからさ 451 名前：デフォルトの名無しさん mailto:sage [2013/09/11(水) 02:12:53.82 ] 俺も今5.5のToolkit のインストールでずっこけてる。 Windows XP service pack 3 Quadro FX 580 Toolkit と sample以外はカスタムでインストールできたんだけど、 Toolkitが8割ぐらい行ったところで失敗しました、てなる。 インスコディレクトリにいくらかコピーできてるみたいなんだけど、 環境変数なんかは設定されてない。 リリースノート見てもGUI使うかmsiをシェルで実行しろとしか書いてなくて 困ったぽよ 452 名前：447 mailto:sage [2013/09/11(水) 04:56:20.56 ] もう朝だお。。。 Toolkitのインストーラーが失敗するから ログ取ろうとしたらなぜか成功したっぽい？？ C:\NVIDIA\CUDA\CUDAToolkit> msiexec /i "NVIDIA （略）.msi" /L*v install.log でも、サンプルのtemplateとか開いてビルドしようとしても 「error : The CUDA Toolkit v5.5 directory '' does not exist.〜」てなる。 環境変数（CudaToolkitDir？）が設定されてないのか、VisualStudioよく分かってないのか・・・。 スレチなら他行きますので・・・。 453 名前：447 mailto:sage [2013/09/11(水) 05:42:32.92 ] ここは俺の日記かお・・・。 VS2012でプロジェクト→プロパティ→構成プロパティ→デバッグ→環境 を選んで編集モードへ。 ここで「マクロ>>」をクリックすると設定されたCuda用環境変数もちらほらありますが $(CudaToolkitDir)の値がからっぽでした。 普通に動かせてる方、この辺の弄り方教えてください。 設定するパスはこれですかね？→C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.5 454 名前：デフォルトの名無しさん mailto:sage [2013/09/11(水) 10:14:42.53 ] >>453 うちでは、NVIDIAのGPUが無くてもインストールとビルドできてるよ。(当然、このPCでは実行はできないけど) コンパイラはVS2008 Standard SP1, VS2010 Professional, VS2012 Professional UP3の3つ。 マクロの値は>>453 でok。 OSの環境変数は: CUDA_PATH=C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.5 CUDA_PATH_V5_5=C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.5 NVCUDASAMPLES5_5_ROOT=C:\ProgramData\NVIDIA Corporation\CUDA Samples\v5.5 NVCUDASAMPLES_ROOT=C:\ProgramData\NVIDIA Corporation\CUDA Samples\v5.5 NVTOOLSEXT_PATH=C:\Program Files\NVIDIA Corporation\NvToolsExt\ 455 名前：447 mailto:sage [2013/09/11(水) 16:32:58.66 ] >>454 ありがとうございます、template動きました。 他のサンプルも手作業コピーで動きました。 # C:\NVIDIA\CUDA\CUDASamples\の中身を # C:\ProgramData\NVIDIA Corporation\CUDA Samples\v5.5 # にコピーして、\v5.5\Samples_vs2012.sln開いて全部ビルド成功。 VolumeRenderのfpsが2.1〜2.2の貧弱環境ですが ひと通り習得できてきちんと開発できるようになったら 新しいグラフィクスカード買わないとな・・・。 456 名前：447 mailto:sage [2013/09/11(水) 16:47:38.48 ] 追加質問です。 ビルド中に警告が大量に出ますが、手動インストールによるものなのか判断できません。 > \include\math_functions.h : > warning C4819: ファイルは、現在のコード ページ (932) で表示できない文字を > 含んでいます。データの損失を防ぐために、ファイルを Unicode 形式で保存してください。」 > というのが、たくさん出ます。 普通にインストーラーが成功した方も同じようになりますか？ エディタで開いたところ、 math_functions.hはUTF-8N、 cuda.hはSJISと表示されました。 457 名前：デフォルトの名無しさん mailto:sage [2013/09/11(水) 17:16:10.44 ] >>456 > > warning C4819: ファイルは、現在のコード ページ (932) で表示できない文字を > > 含んでいます。データの損失を防ぐために、ファイルを Unicode 形式で保存してください。」 > > というのが、たくさん出ます。 > > 普通にインストーラーが成功した方も同じようになりますか？ CUDAにかぎらず、海外のコードは、よくその警告がでる。 自分は、プロパティマネージャで、C4819を無効にしちゃった。 あと、自分が書いたコードはUTF-8(BOM無し)にしてる。 458 名前：デフォルトの名無しさん [2013/09/12(木) 01:13:44.56 ] すみません。 まだまだ初めたばかりの初心者なのですが、 質問させてください。 Tesla C2075を使っていて、DeviceQuery.exeでスペックを見たところ Total amount of global memoryが1.2GB 程度しかありませんでした。 仕様では5GB程度あるはずなのですが・・・。 OSが32bit(win7)であることが関係してたりしますか？ 459 名前：デフォルトの名無しさん mailto:sage [2013/09/12(木) 02:48:14.69 ] >>458 そうかも。 GPUのメモリはPCのメモリマップ上に割り付けられるけど、32bitでは4GBしかないので。 メモリマップは、Windowsのデバイスマネージャで「表示」→「リソース (種類別)」の「メモリ」で判る。 460 名前：デフォルトの名無しさん mailto:sage [2013/09/12(木) 10:32:10.68 ] 64bitでもVRAM全部を割り付けたりしないよ 多分、互換性の為だろうけど 461 名前：デフォルトの名無しさん mailto:sage [2013/09/12(木) 23:42:16.47 ] Windows 7 service pack 1　64bit Geforce TITAN の構成でも試してみたが5.5インストールできなかった。 Windowsでインストーラ使って5.5入れられた人いますか? 462 名前：デフォルトの名無しさん mailto:sage [2013/09/12(木) 23:54:10.36 ] >>454さんはするっとインストールできちゃった感じなんでしょうか。 463 名前：デフォルトの名無しさん mailto:sage [2013/09/13(金) 00:17:54.99 ] 連投すまん。 Visual Studioが2008 Express しか入っていないのが原因かな? 464 名前：やんやん ◆yanyan/Pails mailto:sage [2013/09/13(金) 05:10:33.37 ] >>461-463 何の問題もなくインストールできたよ 465 名前：デフォルトの名無しさん mailto:sage [2013/09/13(金) 09:34:15.52 ] >>462 >>463 toolkit 5.0, 5.5 はVS StdかProが必要だよ。 4.0, 3.0 は忘れた。 466 名前：デフォルトの名無しさん mailto:sage [2013/09/13(金) 23:20:29.29 ] おかげさまでどうにか自作ソースを5.5でコンパイルするところまでこぎつけました。 インストーラでToolkitとSample 以外を先にインストールして、 Toolkitを単独でインストール、ファイルのコピーが終わったところでこけるので、 >>454の手順で環境変数を設定して、 Sampleはインストーラを立ち上げて展開されたものを所定の位置にこぴぺしました。 467 名前：デフォルトの名無しさん mailto:sage [2013/09/14(土) 03:24:37.09 ] CUDAでホストのmain関数（グローバル関数の呼び出し込）で時間計測にclock()ってつかって問題あったっけ？ 468 名前：デフォルトの名無しさん mailto:sage [2013/09/14(土) 18:37:57.64 ] >>467 OpenMPと一緒に使うとCPUタイムがでちゃって正確な経過時間を計れない場合があります。 精度が1msecというのも場合によっては足りないかもしれません。 469 名前：デフォルトの名無しさん mailto:sage [2013/09/14(土) 22:44:07.90 ] カーネル内でmallocとかnewとかってできる？ できるとしたらどのメモリが確保されるん？ グローバル？ローカル？ 470 名前：デフォルトの名無しさん [2013/09/14(土) 22:52:42.35 ] グローカル 471 名前：デフォルトの名無しさん mailto:sage [2013/09/15(日) 01:00:15.49 ] ﾜﾛﾀｗ 472 名前：デフォルトの名無しさん mailto:sage [2013/09/15(日) 05:52:52.68 ] CUDA C Programming Guide の B.18. Dynamic Global Memory Allocation and Operations にcompute capability 2.0以上でカーネル内のmallocが使える みたいな事が書いてあるね。GPU側のグローバルメモリじゃないかな。 473 名前：デフォルトの名無しさん mailto:sage [2013/09/15(日) 10:04:46.08 ] >>472 サンクス！ GPUのグローバルメモリのヒープ領域(デフォルト6MB)に確保されるのね 474 名前：デフォルトの名無しさん mailto:sage [2013/09/15(日) 13:48:19.08 ] >>468 CPUの方でも処理（マルチスレッドではなく）してるからそういう方式では問題ないってことでいいのかな？ 475 名前：デフォルトの名無しさん mailto:sage [2013/09/15(日) 14:55:48.31 ] >>474 だったらたぶん。大丈夫 476 名前：デフォルトの名無しさん mailto:sage [2013/09/16(月) 19:37:26.22 ] カーネル関数でグローバル変数を使う方法ありますか？ 渡したい値がたくさんあるのでカーネル関数の引数がかなり多くなってしまい面倒なのですが 477 名前：デフォルトの名無しさん mailto:sage [2013/09/16(月) 21:00:31.24 ] __device__付けたらいいんですね自己解決しました 478 名前：デフォルトの名無しさん mailto:sage [2013/09/20(金) 03:32:59.30 ] グローバル化するより構造体作って引数減らしたら駄目なん？ 479 名前：デフォルトの名無しさん mailto:sage [2013/09/20(金) 22:20:48.09 ] >>476 構造体作ってポインタで渡す。 480 名前：デフォルトの名無しさん mailto:sage [2013/09/21(土) 08:08:30.99 ] 構造体渡しはつまり全メンバーコピーだからね 481 名前：デフォルトの名無しさん [2013/10/03(木) 18:57:43.41 ] cudaするのに、１１インチくらいのディスプレイのノートPCってあるの？ kakaku.comで調べても、14インチくらいのしかないけど、だれか教えてください。 482 名前：デフォルトの名無しさん mailto:sage [2013/10/03(木) 22:45:25.45 ] なぜそこでノート限定なのか 483 名前：デフォルトの名無しさん mailto:sage [2013/10/04(金) 00:18:17.93 ] ちっこいノートはインテルHDグラフィックスなんでねーの? と適当に書いてみる。 484 名前：デフォルトの名無しさん mailto:sage [2013/10/04(金) 07:25:58.57 ] いい加減、グラボを換装できるノートが欲しいよな 485 名前：!ninja mailto:sage [2013/10/04(金) 07:49:10.28 ] 普通にあるけど割高だよ。 486 名前：デフォルトの名無しさん [2013/10/04(金) 08:42:04.48 ] 以前は、IONの載ったちいさいノートPCでCUDAできたのに・・・ 487 名前：デフォルトの名無しさん mailto:sage [2013/10/04(金) 21:53:58.77 ] どこかが展示してた外付けグラボが欲しいんだが、発売する気配なしｗ 488 名前：デフォルトの名無しさん mailto:sage [2013/10/04(金) 23:18:05.53 ] 外付けグラボは過去に何製品か出たけど あまりにもニッチすぎる、発熱や性能の問題でどれもヒットしなかった 489 名前：487 mailto:sage [2013/10/04(金) 23:20:13.46 ] >>488 そうなのかぁ･･･。 残念。 490 名前：デフォルトの名無しさん mailto:sage [2013/10/05(土) 00:05:23.57 ] くだらないしつもんですが、 1デバイス内のグローバルメモリ内でデータをコピーするのに、 cudaMemcpyDeviceToDeviceを使うのと、 カーネルを書くのではどちらが速いですか? 491 名前：デフォルトの名無しさん mailto:sage [2013/10/07(月) 18:54:19.99 ] 簡単なんだから、実測してみたら? コピー専用のカーネル作ってたらロードで不利だけど、メインの処理に抱き合わせるならきっと速いよ。 492 名前：デフォルトの名無しさん mailto:sage [2013/10/08(火) 10:22:41.50 ] >>481 一台GPU積んだlinux機作ってsshできるようにすると快適だよ 493 名前：デフォルトの名無しさん mailto:sage [2013/10/08(火) 17:50:51.74 ] 最近cudaを触りはじめた者なのですが質問があります 多次元配列（2とか3次元）の中で指定した値に一番近い値を抽出する方法をcudaで処理するにはなにが一番適しているのですか？ 色々ggってみて、とりあえずソートして絶対値をとって比較すれば良いのではと思い バイトニックソートなどで試そうとしましたがソースコード付の物はどれも一次元配列の物ばかりでどうすればいいのかわかりません それにその方法をとっても総当たりより速度が余計遅くなるのではないかとも考えましたがどうなのでしょうか？ 抽象的な質問になってすみませんが誰か回答をお願いします 494 名前：デフォルトの名無しさん mailto:sage [2013/10/08(火) 18:32:54.64 ] >>493 kd木構築してNNでいいんでないの どう並列化するかは細かい状況による 495 名前：デフォルトの名無しさん mailto:sage [2013/10/08(火) 18:39:54.30 ] 総当りを（並列で）やるしかないと思うよ。あれをreductionって言うのかわからないけど。 たとえばMaxを求めるのにソートは必要なくてO(n)でできるのと同様に、 一番近い値を求める（だけ）ならばソートする必要はないはず。 496 名前：495 mailto:sage [2013/10/08(火) 18:43:15.91 ] ああ、多次元配列は変化せずに指定する値をコロコロ変えるなら最初にソートしておけば その後はO(logn)でいけるね。条件次第。 497 名前：デフォルトの名無しさん mailto:sage [2013/10/08(火) 19:37:01.90 ] >>494-496 返信ありがとうございます とりあえず自分の知識が全く足りてない事が分かったのでもう少し色々調べてみます あと、できれば496あたりの詳しい解説を知りたいのですがO(long)とはなんでしょうか？ 498 名前：デフォルトの名無しさん mailto:sage [2013/10/08(火) 19:46:50.79 ] >>497 O(long)じゃなく、O(logn)な ロガリズムnのオーダー 499 名前：495 mailto:sage [2013/10/08(火) 19:56:25.36 ] >>497 O(n)とかについては「ランダウの記号」でググるかウィキペディると説明がみつかるはずです。 （ウィキペディアの記事は「一般的なオーダー」よりも前のところは正直俺はよくわからないけど…） 500 名前：デフォルトの名無しさん mailto:sage [2013/10/09(水) 08:46:32.40 ] >>499 了解です 回答ありがとうございました また分からなくなったら来ます 501 名前：デフォルトの名無しさん mailto:sage [2013/10/09(水) 10:13:27.99 ] その多次元配列は、一次元配列に投影できないのだろうか。 502 名前：デフォルトの名無しさん mailto:sage [2013/10/09(水) 23:03:59.00 ] >>501 (i, j, k)の三次元配列だったら index = k * N^2 + j * N + i で一次元になるね。 503 名前：デフォルトの名無しさん mailto:sage [2013/10/12(土) 23:33:36.04 ] Visual Studioでcompute_20,sm_20，-maxrregcount=128でコンパイルしたら， 16 bytes stack frame, 24 bytes spill stores, 12 bytes spill loads ptxas : info : Used 63 registers, 32 bytes cmem[0], 16 bytes cmem[16] ってスピルしちゃうんですが． どこを確認すればいいんでしょうか 504 名前：デフォルトの名無しさん mailto:sage [2013/10/13(日) 04:07:02.74 ] >>503 2.xと3.0は最大レジスタ63個までだから 処理少なくするかコード見直して不要冗長な部分を削るとか あるいはレジスタ制限がゆるい1.xか3.5にする 505 名前：デフォルトの名無しさん mailto:sage [2013/10/13(日) 18:48:01.71 ] >>504 63個までだったなんて知らなかった…． 18K Shared使ってるのと，三次元グリッドなんでCP3.5にしたら，ちゃんとコンパイル通って動作しました． ありがとうございます！ 506 名前：デフォルトの名無しさん [2013/10/14(月) 04:34:33.91 ] カーネル呼び出しの行で実行がストップしてしまうのですが（次の行に進まない）、 原因がわかりません。 しかし、nvccの「-G」オプションを付けると、この事象が発生せず、 普通に動作します。 原因のわかる方、あるいは原因の特定方法がわかる方、 教えて下さい。 507 名前：デフォルトの名無しさん mailto:sage [2013/10/14(月) 10:57:27.12 ] >>506 カーネル内で不正なメモリアドレスにアクセスすると 問答無用で落ちる場合はありますね。 -Gをつけた上で、カーネル実行の後ろに printf("%s\n", cudaGetErrorString(cudaGetLastError()); と入れてみてはどうでしょうか。 508 名前：デフォルトの名無しさん mailto:sage [2013/10/14(月) 16:11:07.38 ] >>506 >>507さんが言ってるように，まずはエラー原因を拾ってみると良いのでは？ 30： Unknownが出そうな気はするけど． デバッグ情報付与するとメモリ確保位置が異なるので， デバッグ版では偶然に計算に影響を与えない領域にアクセスしていたという可能性も． まず思いつくのは，ある配列の要素Nに対してthreadを生成しているならば， そのthread数がNを超えて生成されるので（BlockDims*GridDims>N）， N以上のIdxのthreadが配列外アクセスしてるとか． あとは，カーネル関数内でPrintf()して確認していくとか． 509 名前：デフォルトの名無しさん mailto:sage [2013/10/14(月) 16:15:22.22 ] >>506 追記 CudaGetLastError()は，カーネル関数がエラー起こしてても拾わない事があるので， そのまえのカーネルが原因という可能性もありますよ． 前に起動したカーネルがあるなら，その後に主要変数をmemcpyして調べるといいのでは？ 510 名前：506 [2013/10/14(月) 16:45:39.95 ] 回答ありがとうございます。 >>507 なるほど、メモリアクセスが怪しいのですね。 確認してみましたが、エラーとななっていない事がわかりました。 「cudaGetLastError()」の代わりに、「cudaPeekAtLastError()」でも 試したのですが、エラーとなりませんでした。 >>508 メモリ確保位置が異なるということで、「-G」を付けずに試してみましたが、 やはりカーネル呼び出しの行で実行がストップしてしまい、 エラー取得の行まで進みません。 （ログを入れて確認しました。） 配列のインデックスのバグの可能性があるのですね。 細かく確認してみます。 >>509 問題の起きているカーネルが、最初に実行されるカーネルになります。 カーネル実行の前では、「cudaMemsetAsync」やら「cudaMemcpyAsync」等を 実行していたりしますが、戻り値を確認しましたが、正常値でした。 （ストリームを使っています。） 非同期関数なので、すぐにはエラーがでないそうですが。 この辺も怪しいですね。確認します。 511 名前：デフォルトの名無しさん mailto:sage [2013/10/14(月) 17:57:20.43 ] >>510 ＞配列のインデックスのバグの可能性 それはもちろんあるんだけど，ThreadIdxが計算したいNを超えていないか？という事を言いたかったわけで． 例えば8,000個の計算をしたかったときに，1,024Threads/block，8Blockでカーネル起動したとき， 192個余分にThreadを起動しているので，その処理は大丈夫？って意味だったんだけど ＞カーネル呼び出しの行で実行がストップしてしまい、エラー取得の行まで進みません ホストからカーネル起動したらすぐに次の行にステップして， カーネルの動作は投げっぱなしだったと思うんだけど…違ったっけ？ なんか，メモリアクセスエラーよりもカーネル起動すら出来ていない様な気がするので， 簡単なテストカーネル作って起動させて確認するといいのでは？ 512 名前：デフォルトの名無しさん mailto:sage [2013/10/14(月) 18:14:52.48 ] 意外と凡ミスでＨＷ制限越えた値で起動させようとしてるとか パラメータの指定のような気がする 513 名前：デフォルトの名無しさん mailto:sage [2013/10/14(月) 18:15:28.26 ] c/パラメータの指定/パラメータの指定ミス

---

---

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

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