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

1 名前：デフォルトの名無しさん [2011/08/23(火) 22:08:06.09 ] このスレッドは、他のスレッドでは書き込めない超低レベル、 もしくは質問者自身何が何だが分からない質問を勇気を持って書き込むスレッドです。 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/ 29 名前：デフォルトの名無しさん mailto:sage [2011/09/01(木) 23:44:25.02 ] >>27 何がしたいのか分からんがそれは簡単だ。 30 名前：デフォルトの名無しさん mailto:sage [2011/09/01(木) 23:56:07.27 ] >>29 コアと計算結果の表が出来上がればコア間の相違が分かるじゃん やり方知ってるなら教えてください 31 名前：デフォルトの名無しさん mailto:sage [2011/09/01(木) 23:58:10.21 ] 煽っても教えてやらん 29じゃないけどｗ 32 名前：デフォルトの名無しさん [2011/09/02(金) 02:12:08.96 ] NvidiaのCUDAをやろうと思い、Nvidiaのカードの刺さったマシンにFedoraを 入れてCUDAのためのNvidiaのドライバを入れて使ってみたが、OSのカーネル のアップデートがあるたびに、コンソールが再起動時にブラックアウトした。 そのたびごとに、別のマシンからリモートログインをしてCUDAのための Nvidiaのドライバを再インストールする必要があった。 CUDA3.xまではそれでなんとかなったが、現在のCUDA4.0になると、 幾らやっても駄目になってしまった。どうやらFedora15ではまともに サポートされていないようなのだ。Nvidiaの提供しているCUDAのための ドライバがインストール時におかしなメッセージが出てインストールが 出来ない。 33 名前：デフォルトの名無しさん mailto:sage [2011/09/02(金) 02:29:09.05 ] ここはおまえの日記帳だから好きに使いなさい 34 名前：デフォルトの名無しさん mailto:sage [2011/09/02(金) 11:08:38.52 ] >>28 ブロックあたりのスレッド数を32以下にして試してみましたが、 結果は変わらなかったです。 35 名前：デフォルトの名無しさん mailto:sage [2011/09/02(金) 18:20:08.65 ] >>23 とりあえず、プログラミングガイドのコードそのままだし、 実際アルゴリズム自体には問題ないから、 意図どおり動かないとしたら、ドライバかSDKのバグか ハードの故障じゃね。 まあ、ここで詳細なバージョンとか晒してみたり、 PTXやアセンブリを貼ってみるのもいいけど、 本家のフォーラムで聞いてみたほうが早いんじゃね。 36 名前：デフォルトの名無しさん mailto:sage [2011/09/05(月) 23:02:56.52 ] 長崎大の言う演算間違えって、具体的にどういう演算を間違えたのかの情報ってないの？ 間違えますって情報だけ垂れ流して具体例出さんって、どういうことなんだべ。 37 名前：デフォルトの名無しさん mailto:sage [2011/09/05(月) 23:24:45.84 ] ×演算間違え ○演算間違い 38 名前：デフォルトの名無しさん mailto:sage [2011/09/07(水) 18:20:00.85 ] CUDAを使用するプログラムで巨大なデータを扱いたいのですが、nvcc でのコンパイル時に relocation truncated to fit: R_X86_64_32S against `.bss' とエラーが出て完了しません。gcc ではデータ量制限を解除するのに-mcmodel=medium というオプションがあるようですが、nvcc でこれにあたるものはないでしょうか。 39 名前：デフォルトの名無しさん mailto:sage [2011/09/07(水) 19:43:00.72 ] 巨大データってどれくらい？ CUDAは基本グローバルメモリ使用の制約があるから”巨大”は決して得意ではないのだけど 40 名前：38 mailto:sage [2011/09/07(水) 22:05:04.34 ] >>39 今確認したところ、要素が数十万個の配列を使用していて、それらの合計は200MBほどです。 そのうち、実際にGPUに計算させるためGPUに渡すデータは50MB程度です。 上記のエラーメッセージでググったところ、データ量制限が問題だろうということだったのですが…それほど大きなデータでもないですよね。 なんなんでしょうこのエラー。 プログラミングに関してはまだ知識が浅いためチンプンカンプンなこと言ってたらすみません。 41 名前：デフォルトの名無しさん mailto:sage [2011/09/07(水) 22:36:52.16 ] どうせデータをソース中に書き込んでるんだろw 42 名前：デフォルトの名無しさん mailto:sage [2011/09/07(水) 22:42:26.34 ] >37 日本語の間違えより、長崎大の言ってる間違いの詳しい情報くれよぅ…。 43 名前：デフォルトの名無しさん mailto:sage [2011/09/07(水) 22:50:33.71 ] 要素数が大きすぎるなら、mallocを使ってみたら？ 44 名前：デフォルトの名無しさん mailto:sage [2011/09/07(水) 23:04:45.67 ] >40 配列は何次元？ CやC++だと配列は連続したメモリ領域と規定されてるので、 確保可能なメモリ量は物理メモリの絶対量ではなく、 アドレスを連続で確保できるメモリ量に規定されるよ。 例えば大きめの２次元配列を確保する場合は double* pBuf = new double[row * column]; とするより double** ppBuf = new (double*)[row]; for(int i = 0; i< row; ++i){ 　　　ppBuf[i] = new double[column]; } とかにしたほうがいいよ。 45 名前：デフォルトの名無しさん mailto:sage [2011/09/07(水) 23:11:07.95 ] ファイル読み込みで、cudaMallocHost(だったかな？？)でメモリ確保すればいいじゃん 46 名前：38 mailto:sage [2011/09/07(水) 23:11:33.99 ] >>43>>44 gccと同様にオプションでエラー回避できないかな…と思ったのですが、配列の取り方を変えるほうが良さそうですね。 特に>>44は知りませんでした。 ありがとうございました！ 47 名前：デフォルトの名無しさん mailto:sage [2011/09/07(水) 23:23:20.08 ] 良く分からんが、-Xcompiler -mcmodel=mediumじゃ駄目なん？ -Xlinkerも要るかもしれないが。 48 名前：デフォルトの名無しさん mailto:sage [2011/09/08(木) 08:01:27.80 ] >>44 普通はありえない。そんなことしたら、ポインタアクセスが一回余計に要るから遅くなるし、 キャッシュアウトしやすくなるから更に遅くなる。 まぁ、GPUに転送するのは1スライス分だけだろうからGPUに任せてCPUは暇こいていられるならいいかも知れんが。 49 名前：デフォルトの名無しさん mailto:sage [2011/09/08(木) 17:24:44.67 ] >>44 それは無い。 ・メモリも無駄に消費 ・アクセス性能がた落ち ・バグの温床になる つーかC++ならvectorで初期化時に サイズ指定したほうがいろいろ便利 50 名前：デフォルトの名無しさん mailto:sage [2011/09/08(木) 17:40:17.98 ] 配列でやる場合、 配列の添え字がINT_MAXを越える大きさならそうするしかないと思うが・・ 51 名前：デフォルトの名無しさん mailto:sage [2011/09/08(木) 20:47:04.41 ] キャッシャがアクセスがとドヤ顔されても・・・。 要素のデカイ配列が確保できんって話なわけだし。 52 名前：デフォルトの名無しさん mailto:sage [2011/09/08(木) 22:19:09.58 ] コンパイル時の話なのに ヒープのフラグメンテーションの訳無いだろう。 nvccからgccに自由にオプション渡すやり方は 知らんが、最悪c(++)部分のコードを分割して gccでコンパイルすりゃ良い。 仮にヒープのフラグメンテーションが酷い場合でも >>44みたいなやり方は下策だよ。 素人におかしなやり方を覚えさせるだけ。 大きな配列確保する前に小さな物についてShrink-to-fitやればデフラグされるから その後巨大配列を確保すれば良い。 特にGPGPUで巨大な配列扱うような用途なら 配列の配列なんて転送も困るだろうに。 53 名前：デフォルトの名無しさん mailto:sage [2011/09/08(木) 22:30:02.29 ] >大きな配列確保する前に小さな物についてShrink-to-fitやればデフラグされるから >その後巨大配列を確保すれば良い。 コレ具体的にどんなコード書けばいいのん？ 54 名前：デフォルトの名無しさん mailto:sage [2011/09/09(金) 00:35:55.99 ] メインの計算に必要な必要な小容量配列 保持用のポインタなりvectorの変数だけ始めに ヒープに確保した後、メインの計算前の前処理をする。 この処理で、一時的なデータと最終的に使うデータの ヒープ確保が入り乱れる所為でフラグメンテーションが発生する。 最終的に使うデータについて、大体処理した順番を守って ひたすら同容量確保してはコピーを繰り返せば、 こういうデータは全て前方に圧縮されて配置されるようになる。 55 名前：デフォルトの名無しさん mailto:sage [2011/09/09(金) 21:48:41.98 ] ？ まあスレ違いなんで他でやってよ。 56 名前：デフォルトの名無しさん mailto:sage [2011/09/09(金) 22:24:20.88 ] kernelへの引数に構造体は使えんの？ 今やってみたらコンパイルでエラーになった 57 名前：デフォルトの名無しさん mailto:sage [2011/09/09(金) 22:30:19.36 ] あ、間違えてただ 構造体引数使えるべな 58 名前：デフォルトの名無しさん mailto:sage [2011/09/14(水) 11:12:56.96 ] >>15 code.google.com/p/gpuocelot/ Ocelot currently allows CUDA programs to be executed on NVIDIA GPUs, AMD GPUs, and x86-CPUs at full speed without recompilation. 59 名前：デフォルトの名無しさん mailto:sage [2011/09/15(木) 17:52:31.49 ] cuda4.0、gtx580で Maximum dimension of a grid( 65535 , 65535 , 65535 ) ってなってるんですがz軸のブロックの取れる数は1じゃなくなったんですか？ 60 名前：デフォルトの名無しさん mailto:sage [2011/09/19(月) 22:05:37.25 ] デバイス上の処理が長時間になる場合、途中経過を取得する方法ってないかな？ 例えば１００万回のモンテカルロを実行して、途中の１０万回、２０万回が済んだらその情報を取りたいんだけど… 61 名前：デフォルトの名無しさん mailto:sage [2011/09/19(月) 22:31:31.07 ] >>60 streamingでできないかな? 62 名前：デフォルトの名無しさん mailto:sage [2011/09/19(月) 22:52:18.06 ] streamingってなに？ SMのSじゃないよね？ 63 名前：デフォルトの名無しさん mailto:sage [2011/09/20(火) 05:00:39.03 ] カーネル呼び出し時の<<<>>>の第3引数がstreamingの何かじゃなかったっけ？ 俺は>>61を見てそういうのがあったなと思い出したレベルなので使い方は知らないけど。 64 名前：デフォルトの名無しさん mailto:sage [2011/09/20(火) 08:29:52.67 ] 第3パラメータはsharedMemoryのサイズだよ。そこで指定しなくても使えるけど。 65 名前：デフォルトの名無しさん mailto:sage [2011/09/21(水) 22:11:25.09 ] ごめん第4引数だったみたい。それとAsync系関数がうんたらかんたら。 66 名前：デフォルトの名無しさん mailto:sage [2011/09/22(木) 00:50:42.07 ] あー、その辺だね。 67 名前：デフォルトの名無しさん mailto:sage [2011/09/22(木) 18:05:24.11 ] 【トリップ検索】CUDA SHA-1 Tripper【GeForce】 hibari.2ch.net/test/read.cgi/software/1311428038/ すげえな 68 名前：デフォルトの名無しさん mailto:sage [2011/09/23(金) 11:16:26.70 ] 個人でCUDAプログラミングやってうれしいのはその程度なのか？ と逆に不安になるんだが… 69 名前：やんやん ◆yanyan72E. mailto:sage [2011/09/23(金) 11:35:33.50 ] 世の中なんでも膨大なデータをいかに早く処理するかが勝負の世界に なってきているから、CUDAプログラミングの旨味はいくらでもあるだろ。 70 名前：デフォルトの名無しさん mailto:sage [2011/09/23(金) 12:04:06.11 ] いつかやってみようと思って半年くらいこのスレROMってるが いまだに開発環境のセッティングの仕方すら調べていない・・・ 71 名前：デフォルトの名無しさん mailto:sage [2011/09/24(土) 00:02:19.96 ] >>68 >>67くらいのコード書ける？ 72 名前：デフォルトの名無しさん mailto:sage [2011/09/24(土) 15:28:34.26 ] >>71 おれがそれくらいのコード書けるようになるころには、自ずとましな使い道が浮かんでくることだろう、 っていうこと？ 73 名前：デフォルトの名無しさん [2011/10/02(日) 08:23:54.81 ] >>60 試してないけど、長い処理の途中でグローバルメモリに書き込んで、 別streamでその値を読むカーネルを起動すれば取れると思う。 キャッシュからグローバルメモリに書き込むのは特別な関数を 使うと思ったけど名前忘れた 74 名前：デフォルトの名無しさん [2011/10/02(日) 09:01:54.77 ] >>73 思い出した。atomicを使うか、書き込む側で __threadfence() 75 名前：デフォルトの名無しさん [2011/10/04(火) 20:59:53.36 ] 失礼します、先輩方お助けを。 atomicを使ったソースを"-arch compute_20"オプションでコンパイルすると次のようなエラーが出てしまいます。 「'-arch' は、内部コマンドまたは外部コマンド、操作可能なプログラムまたはバッチ ファイルとして認識されていません。 C:\Program Files (x86)\MSBuild\Microsoft.Cpp\v4.0\Microsoft.CppCommon.targets(103,5): error MSB3073: コマンド "-arch compute_20 error MSB3073: :VCEnd" はコード 9009 で終了しました。」 -archコマンドはプロパティからビルド前イベントのコマンドラインに追加してあります。 -archコマンドがあるディレクトリへパスが通ってないのかなと思うのですが、このようなコマンドがあるディレクトリの追加の方法がよくわかりません。 環境はcuda4.0 + VC2010 + teslac2070(cc2.0) + windows7 proです。 どなたかアドバイスください。 76 名前：やんやん ◆yanyan72E. mailto:sage [2011/10/04(火) 21:13:10.02 ] -archコマンドってなんだそら？ compute capabilityはnvccなりcl.exeのオプションに渡すものであって、 -archコマンドなんてものはないよ。 77 名前：デフォルトの名無しさん mailto:sage [2011/10/04(火) 21:50:45.04 ] 自分でオプションでって書いてるけど、メッセージはそれで実行した系だよな 78 名前：デフォルトの名無しさん [2011/10/04(火) 23:16:18.47 ] >>76 >>77 ありがとうございます。「cuda by example」にもatomic関数を使ったソースには-archコマンドラインオプションを指定する必要があります。 とあるのですが、みなさんはatomicを使ったソースをコンパイルするときなにかコマンドつけませんか＞＜？ 79 名前：デフォルトの名無しさん [2011/10/04(火) 23:19:48.93 ] >>78 あ、もちろん全部64bitです！ 80 名前：デフォルトの名無しさん mailto:sage [2011/10/04(火) 23:39:11.06 ] CUDA以前の話だな。コンパイラオプションの意味がわかってない。 というか、VCなら、カスタムビルド規則でCUDA Runtime API Build Ruleを追加したほうが早い 81 名前： 忍法帖【Lv=40,xxxPT】 mailto:sage [2011/10/05(水) 02:20:11.36 ] -gencode=arch=compute_20,code=\"sm_21,compute_20\" になってるなぁ俺の環境だと。VC2010ならプロジェクトのプロパティでCUDA関係の設定するところがあるんじゃない？ 82 名前：デフォルトの名無しさん [2011/10/05(水) 11:51:04.38 ] >>80 ありがとうございます。 えーと、それはおそらくVC2008の話ではないかと思います。2010にカスタムビルド規則という名前はなくなり、その代りCUDA4.0というビルドカスタマイズに統一されています。 83 名前：デフォルトの名無しさん [2011/10/05(水) 11:58:44.51 ] >>81 ご返答ありがとうございます。 そのコードはどこに書いていますか？ はい、cuda c/c++という項目にいろいろいじるところはあるのですがコンパイラオプションを設定する項目はないように思えます。 私がオプションを追加しているのはビルドイベントという項目ですね。 以前にもatomic関連の話が出ていたようです。>>22 84 名前：デフォルトの名無しさん mailto:sage [2011/10/05(水) 12:14:48.16 ] -arch=compute_20 にすればいいんじゃねえの？ 85 名前：デフォルトの名無しさん [2011/10/05(水) 12:17:32.55 ] >>84 ありがとうございます。 おっしゃるように>>75のように追加するとエラーを吐きます。 追加の仕方が違うのでしょうか。 86 名前：80 mailto:sage [2011/10/05(水) 13:36:15.15 ] VS2010は色々変わってるんだな。 知らんかったすまん。 とりあえず「CUDA VS2010」でググって出てくる一番上のサイト参考にして、CUDA4.0のビルドカスタマイズファイルを導入した後、プロパティからCode Generationをcompute_20,sm_20にしたらいけた 87 名前：デフォルトの名無しさん mailto:sage [2011/10/05(水) 15:09:51.57 ] GPUのリソースを全部CUDAに割り当てるため，グラボを2枚さしてディスプレイ出力用とCUDA用に分けるというのはありますが 同じようなことをsandy bridgeとグラボ1枚でも出来るのでしょうか？ 88 名前：デフォルトの名無しさん mailto:sage [2011/10/05(水) 15:36:59.85 ] >>87 Z68マザーなら、マザー側DVIにディスプレイつないでおけばいい話じゃ？ CUDAで計算したデータをintelで表示したいなら、 1.DirectXかOpenGLのリソースをCUDAにバインドしてデータ書き込み 2.メインメモリに転送 3.intel側に転送 って形になるかなぁ。多分。 89 名前：81 mailto:sage [2011/10/05(水) 21:07:18.54 ] >>83 ゴメン、俺もVC2010じゃないや。VC2008。 ちなみに>>81は構成プロパティの中に「CUDA Runtime API」があってそこの「コマンドライン」にある。 CUDA Runtime APIの中にGUIでGPU Architecture（sm_21）を指定するところがあるから、 コマンドを打ってるわけじゃないんだけど、結果としてnvccには>>81のようなオプションが渡ってる。 90 名前：89 mailto:sage [2011/10/05(水) 21:17:54.24 ] そもそもなぜCUDA Runtime APIという項目があるのかというとサンプルのプロジェクトをコピったから。 91 名前：デフォルトの名無しさん [2011/10/06(木) 16:35:52.62 ] >>80 >>81 ありがとうございます！code generationを>>80のように書きなおしたらオプション指定することなくコンパイル通りました。 どうやらVS2010からはcommand line指定しなくてもcode generationでオプションを渡せるようです。 本当に助かりました。ありがとうございます！ 92 名前：デフォルトの名無しさん mailto:sahud@onc.nasi.jp [2011/10/06(木) 17:22:43.25 ] 最近巷で流行ってる error MSB3721: コマンド ""C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.0\bin\nvcc.exe" -gencode=arch=compute_20,code=\"sm_20,compute_20\" --use-. ... template.cu"" はコード 2 で終了しました。 のエラーなんですけど解決法どなたか知りませんか?? CUDA FORUMにもかなり話題が上がってるっぽいんですが...私のところにも発病してしまいました。 93 名前：デフォルトの名無しさん mailto:sage [2011/10/06(木) 18:43:04.20 ] そこだけ書かれても分からないと思う。 nvccのエラーメッセージ載せないと。 94 名前：デフォルトの名無しさん mailto:sage [2011/10/06(木) 19:11:36.83 ] >>92 敵性語なんでよくわかんなかったけど、 VisualStudio を管理者として実行したら動く？ 95 名前：やんやん ◆yanyan72E. mailto:sage [2011/10/06(木) 21:18:36.98 ] Visual C studioは下手にGUIに頼るよりも、コマンドラインでやった方がはやいこともある。コマンドラインでnvccやってみよう。 96 名前：デフォルトの名無しさん mailto:sage [2011/10/07(金) 16:37:49.21 ] 単精度のmadってmulより遅いやないかー プログラミングガイド嘘つきやないかー ちくしょー 97 名前：デフォルトの名無しさん mailto:sage [2011/10/09(日) 01:36:18.51 ] >>96 そりゃ命令実行時間は同じだけど、 必要な帯域は違う。 98 名前：デフォルトの名無しさん mailto:sage [2011/10/09(日) 02:40:58.84 ] >>92 コマンドラインのオプションが何か1つ間違ってた気がする。 コマンドプロンプトで同じコマンド実行するとエラーが出たので、 そのエラーのひとつ前の引数が問題だった、はず。 99 名前：デフォルトの名無しさん mailto:sage [2011/10/10(月) 02:14:56.87 ] CUDAとRTミドルウェアを組み合わせてみた。 行った方法は、CUDAのサンプルをRTコンポーネントに パッケージングする方法。 動作周期にあわせて、行列計算が実行されたのを確認。 100 名前：デフォルトの名無しさん mailto:sage [2011/10/10(月) 20:49:05.67 ] カーネル側のソース分割ってできないんだっけ？ 環境はLinux　cuda4.0 101 名前：デフォルトの名無しさん mailto:sage [2011/10/10(月) 22:00:17.43 ] できるけどコンスタントメモリが変な事になったような気がする 102 名前：デフォルトの名無しさん mailto:sage [2011/10/10(月) 22:04:05.94 ] このままだと5万行くらいになっちゃいそう・・・ 103 名前：デフォルトの名無しさん mailto:sage [2011/10/10(月) 22:39:08.05 ] >>102 正直処理対象がCUDAに向いてないような。。。 104 名前：デフォルトの名無しさん mailto:sage [2011/10/14(金) 02:20:30.66 ] とある文献でレイテンシを表す際, 種類によってclock, cycleが使い分けられていました。 clock : texture cache, constant cache cycle : global memory, shared memory なんとなくメモリ, キャッシュで使い分けられているのはわかりますが, このままでは速度の比較ができません。 二つの違いは何なのか, またcycle-clock間で単位の変換が可能なのか, もしご存知の方がいたらご教授ください。 105 名前：デフォルトの名無しさん mailto:sage [2011/10/17(月) 13:50:45.89 ] カーネル関数内で立てたフラグをホスト側で判断する事って出来る？ 106 名前：デフォルトの名無しさん mailto:sage [2011/10/17(月) 14:13:41.38 ] グローバルメモリ介してできるだろ 107 名前：デフォルトの名無しさん mailto:sage [2011/10/18(火) 02:35:47.28 ] はじめてのCUDAプログラミングの9章に載ってるcublas使うプログラム実行したんだけど、 cublasInit()あたりで「〜0x7c812afb で初回の例外が発生しました: Microsoft C++ の 例外: cudaError_enum〜」って出てうまく動かないんだけどなして？ 108 名前：デフォルトの名無しさん mailto:sage [2011/10/19(水) 00:53:08.52 ] かなり初歩的な質問で申し訳ないのですが 構造体hoge_tにバッファを持たせPtにバッファのアドレスが入るようにしたいのですが 以下のようにするとsegmentation faultになってしまいます。 このようなことは可能なのでしょうか。 struct hoge_t{ int* Pt; int a,b,c,d; } int main(){ hoge_t *hoge_d; cudaMalloc((void**)&hoge_d,sizeof(hoge_t)); cudaMalloc((void**)&hoge_d->Pt,sizeof(int)*100); //終了処理は省略 } 109 名前：デフォルトの名無しさん mailto:sage [2011/10/19(水) 11:56:08.38 ] 適当なポインタで確保した後に， そのアドレス転送してやればいいんじゃない？ 110 名前：デフォルトの名無しさん mailto:sage [2011/10/19(水) 13:25:43.68 ] >>108 ＞cudaMalloc((void**)&hoge_d,sizeof(hoge_t)); 意味のおさらい。 & hoge_dをcudaMalloc()に渡すことで、ローカル変数hoge_dにデバイス上のアドレスが格納される。 ＞cudaMalloc((void**)&hoge_d->Pt,sizeof(int)*100); & hoge_d->Ptは、意味としては & (* hoge_d).Pt に等しい。 つまり、hoge_dがデバイス上のアドレスを保持しているにも拘らず ホスト上で* hoge_d するからsegmentation faultを引き起こす。 データ効率を考えると固定長配列にできないか検討するなどして見直すべきだが、 どうしてもこのままの構造でやりたいなら sizeof(int) * 100 確保したデバイス上の アドレスをcudaMemcpy()でhoge_d->ptにコピーする必要がある。 って、この下りは>109が書いていることと同じだね。 111 名前：108 mailto:sage [2011/10/19(水) 17:14:09.68 ] >>109,>>110 お忙しい中、レス有難うございます。 >>110 データ構造については検討しようと思いますが、 いただいたアドバイスを参考に以下のようにしてみたところsegmentation faultがなくなりました。 109,110さんの意図していたものとなっておりますでしょうか？ int main(){ hoge_t *hoge: hoge=(hoge_t*)malloc(sizeof(hoge_t)); cudaMalloc((void**)&(hoge->Pt),sizeof(int)*100); hoge_t *hoge_d; cudaMalloc((void**)&(hoge_d),sizeof(hoge_t)); cudaMemcpy(hoge_d,hoge,sizeof(hoge_t),cudaMemcpyHostToDevice); //終了処理は省略 } あるいはこういう感じでしょうか・・・ int main(){ int *tmp_pt; cudaMalloc((void**)&(tmp_pt),sizeof(int)*100); hoge_t *hoge_d; cudaMalloc((void**)&(hoge_d),sizeof(hoge_t)); cudaMemcpy(hoge_d->Pt,tmp_pt,sizeof(int),cudaMemcpyHostToDevice); //終了処理は省略 } 112 名前：デフォルトの名無しさん mailto:sage [2011/10/19(水) 21:10:14.17 ] 概ねいいんでない? 前者なら、私はhoge_t hoge;で定義しちゃうけど。 113 名前：109 mailto:sage [2011/10/19(水) 22:04:08.49 ] 俺が考えてたのは後者だな。 合ってると思うよ。 でもmemcpyのsizeof(int)はsizeof(int *)にした方が良いと思う。 114 名前：108 mailto:sage [2011/10/19(水) 23:57:45.47 ] 未だにポインタが使いこなせてないと痛感しました。 このたびはアドバイス有難うございました。 115 名前：デフォルトの名無しさん mailto:sage [2011/10/20(木) 00:12:23.07 ] ポインタ使える男の人ってモテそうだよね 116 名前：デフォルトの名無しさん mailto:sage [2011/10/20(木) 17:52:20.01 ] ポインタでいろんな所を指して、 覗いたりいじくり回したりするんですね。 117 名前：やんやん ◆yanyan72E. mailto:sage [2011/10/21(金) 01:37:52.34 ] トラックポイントではない 118 名前：デフォルトの名無しさん mailto:sage [2011/11/03(木) 20:05:27.91 ] デバイスから出る「the launch timed out and was terminated.」とかのメッセージを CPU側で受信する方法ってどうやるのかな？コンソール出力されるんだからできると思うけど 119 名前：デフォルトの名無しさん mailto:sage [2011/11/03(木) 21:13:42.56 ] >>118 デバイス側といっても、cudart辺りがドライバからのステータスを見て出していると思われ。 従って、当然リダイレクトできる。 120 名前：デフォルトの名無しさん [2011/11/07(月) 04:01:37.49 ] CUBLASとCUSPARSEを使って倍精度で計算を行うCG法のプログラムを作成したのですが、 CUSPARSEのcsrmv関数だけが異様に遅く速度が出ません。 この原因としては何が考えられますでしょうか？ 自分としてはCUBLASとCUSPARSEでそれぞれ別のhandleを使っているせいかと思っているのですが それを確かめる方法や、CUBLASとCUSPARSEでhandleを共有する方法がわかりません... 121 名前：デフォルトの名無しさん mailto:sage [2011/11/07(月) 09:05:53.46 ] CG法と書いてあるので問題行列は対称疎行列でCUSPARSEからはcsrmv関数だけ、CUBLASからは Level-1 BLASルーチンしか使っていないのだと思いますが、 実際にどの程度遅いのでしょうか？ ランダムアクセスを含むcsrmvはLevel-1 BLASルーチンとくらべてメモリアクセス効率が極端に下がるのは 仕様というか当然の結果です。 また古いバージョンのライブラリを使っていたりしませんか？ handle云々は基本的に関係ありません。関係するとすればよほど小さな行列を扱っている場合に限られます。 122 名前：デフォルトの名無しさん mailto:sage [2011/11/07(月) 15:42:01.27 ] 仰る通り対称疎行列でCUSPARSEからはcsrmv関数のみです。 行列のサイズはn×n　n=122187で、ライブラリのバージョンは4.0です。 BLASでの総計算時間が12.66sに対し、 CUSPARSEでの総計算時間が80.36sという結果です。 このCUSPARSEでの計算時間はCPUで計算を行った場合よりも遅く、 nVidiaのCUDA toolkit 4.0 Performance Reportの結果と比較しても 非常に遅いと思います。 123 名前：デフォルトの名無しさん mailto:sage [2011/11/07(月) 18:49:04.47 ] 構造体の中にある配列はどのように確保してmemコピーすればいいんでしょうか？ #define N (32) struct var_box{ int boxi; float boxf; }; struct var_str{ int *vari; // ~N*N float var_g; struct var_box *vb; // ~N }; のような構造体がある時 struct var_str *vvv,*vvv_d; vvv = (struct var_str*)malloc(sizeof(struct var_str)); vvv->vari = (int*)malloc((N*N)*sizeof(int)); vvv->vb = (struct var_box*)malloc(sizeof(struct var_box)*(N)); 値代入 cudaMallocHost((void**)&vvv_d,sizeof(struct var_str)); cudaMallocHost((void**)&vvv_d->vari,sizeof(int)*(N*N)); cudaMallocHost((void**)&vvv_d->vb,sizeof(struct var_box)*(N)); cudaMemcpy(vvv_d, vvv, sizeof(struct var_str), cudaMemcpyHostToDevice); cudaMemcpy(vvv_d->vari, vvv->vari, sizeof(int)*(N*N), cudaMemcpyHostToDevice); cudaMemcpy(vvv_d->vb, vvv->vb, sizeof(struct var_box)*(N), cudaMemcpyHostToDevice); GPUに送ってCPUに戻しただけですが float var_g;　に関しては問題なくできていますが配列にした部分が送れていないみたいです。 　そもそも確保の部分だけで0.5sかかっているのでちゃんとできてるかどうか怪しいです。 124 名前：デフォルトの名無しさん mailto:sage [2011/11/07(月) 18:52:46.57 ] 構造体にポインタがあるのはダメ 理屈で考えればわかるよな 125 名前：デフォルトの名無しさん mailto:sage [2011/11/07(月) 19:13:06.14 ] >>124 そうでしたか わかりました 126 名前：デフォルトの名無しさん mailto:sage [2011/11/07(月) 19:34:23.39 ] 馬鹿でプログラミング初心者には難しいれす（^ρ^） 127 名前：やんやん ◆yanyan72E. mailto:sage [2011/11/07(月) 19:39:20.38 ] 無駄にでっかくメモリ確保された ただのポインタがコピーされてるだけに見えるのは気のせい？ 実体をコピーしようとしている意図は読み取れるけれど、 コードはその意図を反映してないような。。。 128 名前：やんやん ◆yanyan72E. mailto:sage [2011/11/07(月) 19:50:21.46 ] あ、考え違いか、すまん 129 名前：デフォルトの名無しさん mailto:sage [2011/11/07(月) 19:52:32.74 ] >>124 わからないのでその理屈を教えていただけないでしょうか

---

---

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

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