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

1 名前：デフォルトの名無しさん mailto:sage [2014/11/20(木) 23:14:46.66 ID:jr3oZn27.net] このスレッドは、他のスレッドでは書き込めない超低レベル、 もしくは質問者自身何が何だが分からない質問を勇気を持って書き込むスレッドです。 CUDA使いが優しくコメントを返しますが、 お礼はCUDAの布教と初心者の救済をお願いします。 CUDA・HomePage ttp://developer.nvidia.com/category/zone/cuda-zone 関連スレ GPGPU#5 ttp://peace.2ch.net/test/read.cgi/tech/1281876470/l50 前スレ 【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/ 【GPGPU】くだすれCUDAスレ part6【NVIDIA】 ttp://peace.2ch.net/test/read.cgi/tech/1348409867/ 2 名前：デフォルトの名無しさん mailto:sage [2014/11/20(木) 23:15:41.57 ID:jr3oZn27.net] 関連サイト CUDA www.nvidia.co.jp/object/cuda_home_new_jp.html CUDAに触れてみる chihara.naist.jp/people/STAFF/imura/computer/OpenGL/cuda1/disp_content CUDA のインストール blog.goo.ne.jp/sdpaninf/e/9533f75438b670a174af345f4a33bd51 NVIDIAの「GeForce 8800 GT(G92)」と次に控える64-bit GPUアーキテクチャ pc.watch.impress.co.jp/docs/2007/1031/kaigai398.htm CUDAを使う tech.ckme.co.jp/cuda.shtml NVIDIA CUDAを弄ってみた その2 dvd-r.sblo.jp/article/10422960.html CUDAベンチ wataco.air-nifty.com/syacho/2008/0 3 名前：2/cuda_2044.html KNOPPIX for CUDA http://www.yasuoka.mech.keio.ac.jp/cuda/ [] [ここ壊れてます] 4 名前： 【東電 69.8 %】 mailto:糞スレを発見したB29が焼夷弾を投下しに飛来しました [2014/11/21(金) 00:13:37.04 ID:kFWiXf0I.net] >>1 ああ、キミ！また会えたね。久しぶりだ。どうだいあの件は？どうなったか説明したまえな。 神戸市の東、芦屋西宮の知的障害者施設で未成年利用者に性的な行為をして淫行条例で逮捕された三田谷学園元職員の堂垣直人(西宮市老松町)は、結局どういう罪になったの？ 被害者家族のケアを芦屋市役所と兵庫県警はちゃんとやったのか？ 差別や虐待は環境を選べない子供には関係ない。 www.youtube.com/watch?v=JxMzW3ZlV4g&sns=em まあ、こっちに座れよ。ゆっくり話そうじゃないか。 5 名前：デフォルトの名無しさん mailto:sage [2014/11/21(金) 07:36:45.42 ID:JsFj8Vej.net] i.imgur.com/SduT7qR.jpg i.imgur.com/2f8fP3F.jpg i.imgur.com/u0ymg8c.jpg i.imgur.com/LruEssT.png i.imgur.com/AeSQqRT.jpg cisburger.com/up/bnf/6016.jpg up.pangya.tv/src/www_pangya_tv18594.jpg 6 名前：デフォルトの名無しさん mailto:sage [2014/11/21(金) 18:05:42.00 ID:qiUQrZk/.net] syncthreadsとthreadfence_blockの違いが分かりません syncthreadsだけで十分な気がしますが、どういう時に使い分けるのでしょうか？ 7 名前：デフォルトの名無しさん mailto:sage [2014/11/22(土) 13:34:27.11 ID:Ke1g3qvZ.net] >>5 shobomaru.wordpress.com/2013/09/13/synchronizing-instruction-in-direct-compute-and-cuda/ メモリの書き込み競合防止をするかどうかの違いみたいです。 これってatomic命令より軽いのか知らん? 8 名前：デフォルトの名無しさん mailto:sage [2014/11/22(土) 13:48:05.92 ID:S8C7U0PL.net] >>6 競合防止なんて書いてる？ 複数のスレッドが同じメモリに書き込む時はatomicをsyncthreadsやthreadfenceと同時に使わないといけない気がする(誰か教えてください) 9 名前：デフォルトの名無しさん mailto:sage [2014/11/22(土) 22:18:58.72 ID:Ke1g3qvZ.net] すみません、誤読しました。 「同じブロック内の全スレッドがこの命令にたどり着く」 まで待つかどうかの違いですね。 10 名前：デフォルトの名無しさん mailto:sage [2014/11/24(月) 04:18:18.83 ID:qPQKDlD6.net] 今一番コスパ高いカードってなに？ 11 名前：デフォルトの名無しさん mailto:sage [2014/11/24(月) 04:21:35.17 ID:qPQKDlD6.net] 謝罪文みても思い上がりが激しい 12 名前：デフォルトの名無しさん [2014/11/25(火) 19:26:42.72 ID:kFuypilU.net] プログラムのカーネル部分がどうしても実行されません(サンプルプログラムでは実行されていました) どなたか原因に心当たりはありませんか？(私はありません) ブレークポイントで確認したところカーネルの上下にあるクロックは実行されており、 カーネルだけが実行されていませんでした(カーネルの中へ入って行かないという意味です) 以下が呼び出しで、dim3はグローバルで定義してあります dim3 blocks((num + max - 1) / max, (num + max - 1) / max); dim3 threads(max, (1024 + max - 1) / max, 1); void calculation(void) { clock_t start, end; start = clock(); cal<<<blocks, threads>>>(con, num, points, data); end = clock(); cout << double(end - start) / CLOCKS_PER_SEC << "\n"; } 13 名前：デフォルトの名無しさん mailto:sage [2014/11/25(火) 21:02:59.83 ID:N/U8okyJ.net] >>11 dim3構造体のメンバ変数をプリントしたらどうなりますか? 14 名前：デフォルトの名無しさん mailto:sage [2014/11/25(火) 21:44:27.45 ID:kFuypilU.net] >>12 >11のcalculation()の最後の行にプリントの一文を入れてみましたが ブロックが(128,128)、スレッドが(128,8)と想定通りでした (スレッドは1ブロック当たり1024個まで配置可能なのでギリギリセーフなはずです) 忘れていましたが、呼び出し先です この中にブレークポイントを配置してもプログラムが止まらないという魔の領域となっています (もちろんnsightのcuda debuggingでデバッグしています) __global__ void cal(double con, int num, a_data *points, b_data *data) {・・・} 15 名前：デフォルトの名無しさん mailto:sage [2014/11/26(水) 18:34:20.92 ID:cpKKMAIz.net] >>13 カーネルが実行されているかどうかはどうやって確認していますか? 16 名前：デフォルトの名無しさん mailto:sage [2014/11/26(水) 19:30:35.81 ID:mOjmGjn5.net] >>14 >>13に書いた通り、ブレークポイントをカーネルの中に入れて実行されているかどうかを確認しています 描画をするプログラムなので図形が動くか動かないかでも判断できます 17 名前：デフォルトの名無しさん mailto:sage [2014/11/26(水) 20:04:14.26 ID:qey6HT7s.net] おれもカーネルに入らないケース出たわ 原因調査中・・・ 18 名前：デフォルトの名無しさん mailto:sage [2014/11/27(木) 08:04:41.88 ID:7alpN+o4.net] >>15 CUDAのデバッガ使った事がない(ひたすらprintf)ので 一般論的な事しか言えませんが、 怪しそうな処理をコメントアウトしていったらどうでしょう? 変なメモリアクセスで落ちるとか割とありがちな気が。 19 名前：デフォルトの名無しさん mailto:sage [2014/11/28(金) 22:55:16.84 ID:JeOcX4pA.net] おれもカーネルに入らない ただなぜかcygwinでコンパイルするとカーネルが起動する なんでじゃ・・・ 例のvisual studio2013もインストールしてみたいな〜 20 名前：デフォルトの名無しさん mailto:sage [2014/11/28(金) 23:06:00.72 ID:JeOcX4pA.net] >>18だけど osはwindows8.1 コンパイラはcuda6.5+vs2013 express カーネルは担当する要素を＋1するだけのもの これをコマンドプロンプトでコンパイルしてもカーネルは動かなかった(?) 動かないと判断したのは結果をmemcpyしてホスト側に返しても＋１されてなかったから もしかしたら正常にmemcpyされてないだけかもしれない カーネルに入らない人はcygwin使ってみるといいかもね 本質的な解決にはならないけど・・・ 21 名前：デフォルトの名無しさん mailto:sage [2014/11/30(日) 20:47:05.05 ID:NdicNENH.net] ビジュアルプロファイラー使ってみたら。 22 名前：デフォルトの名無しさん mailto:sage [2014/12/03(水) 14:38:19.35 ID:GXBajCbw.net] kernel実行後にcudaGetLastError()でRCを取得 そうするとkernel実行結果が分かる RCの数値は自分で調べてね 23 名前：16,21 mailto:sage [2014/12/04(木) 22:23:47.78 ID:EZ4odEf+.net] 21の情報は役だったかな？ 自分の場� 24 名前：�（LINUX）、RC=7（too many resources requested for launch）だったので コンパイルオプションに　-Xptxas -vを追加して使用レジスタ数を確認。 結果ハードウェアのレジスタ数を超過したためにカーネルの処理が行われなかったことが判明。 スレッドサイズを小さくして問題解決。 因みに使用レジスタはハードウェアによって変わる。 [] [ここ壊れてます] 25 名前：デフォルトの名無しさん mailto:sage [2014/12/05(金) 14:06:24.56 ID:+nFWXccn.net] >>22 ちょっと興味があるんですけれど、 カーネルのサイズは動的に決めてるんですか? 前にソースコードにブロック数とスレッド数をべた書きで 大きいサイズを指定したらコンパイルの段階ではじかれた事があったんで。 26 名前：>>15 mailto:sage [2014/12/05(金) 15:37:53.47 ID:fpNGtjbn.net] >>21 ありがとうすごく役に立ったよ ここ最近忙しくてｐｃ触れなかったんだ 自分も同じく「error: too many resources requested for launch」だった 原因を調べてみるよ あと、自動でエラー内容もだせるみたいだね homepage2.nifty.com/takaaki024/tips/programs/gpgpu/cuda.html >>20 ビジュアルプロファイラーも便利そうだから調べてみるよ 27 名前：デフォルトの名無しさん mailto:sage [2014/12/07(日) 18:17:39.03 ID:g9DGYGEw.net] x,y,zの3つの変数から成る構造体配列A,B(同じサイズ)があったとして BからAへそれぞれ対応するデータを転送する場合 Ax,Ay,Az,Bx,By,Bzという同じサイズの構造体でない配列が6つあったとして BからAへそれぞれ対応するデータを転送する場合(BxからAxなど) 前者と後者では後者の方が転送速度は上がりますか？ 28 名前：デフォルトの名無しさん mailto:sage [2014/12/08(月) 12:53:10.47 ID:JbpvX5Qi.net] >>25 一般論としてデータ量が同じなら一回にまとめてを転送した方が効率はよくなりますね。 29 名前：デフォルトの名無しさん mailto:sage [2014/12/11(木) 11:46:54.75 ID:kV0/O7vj.net] memcpyって同期とるもんね 30 名前：デフォルトの名無しさん mailto:sage [2014/12/12(金) 21:54:58.15 ID:PdQu+k/h.net] kernelの中で使えるタイマー関数はありますか？ 31 名前：デフォルトの名無しさん mailto:sage [2014/12/12(金) 23:43:18.27 ID:q1FKM2bt.net] clock()関数が使えるよ。 32 名前：名無しさん＠そうだ選挙に行こう mailto:sage [2014/12/13(土) 22:02:15.64 ID:B9P4oQcX.net] >>29 サンキュー 試してみるよ 33 名前：デフォルトの名無しさん mailto:sage [2014/12/16(火) 21:07:44.74 ID:6hyQD5WD.net] 自分の持っていないGPUの共有メモリの量などの詳細を知ることはできますか？ 34 名前：デフォルトの名無しさん mailto:sage [2014/12/17(水) 08:36:19.04 ID:0flByQKi.net] en.wikipedia.org/wiki/CUDA#Version_features_and_specifications の Maximum amount of shared memory per multiprocessor あたりかな。 35 名前：デフォルトの名無しさん mailto:sage [2014/12/17(水) 20:16:42.59 ID:D/43rANg.net] >>32 ありがとうございます 嬉しいことに5.0以降から容量が増えてるみたいですね 36 名前：デフォルトの名無しさん mailto:sage [2014/12/17(水) 20:55:04.06 ID:vfaS5qRM.net] なんでCCのバージョン3.5から5.0に飛んでるの？ SDKのバージョンと合わせたのか？ 37 名前：デフォルトの名無しさん mailto:sage [2014/12/19(金) 18:38:15.45 ID:dxBCSCiu.net] 二つ以上のGPUでVBOを使用する場合、データの流れはどうなっているのでしょう？ やはり一旦ディスプレイに接続側のGPU出力データが集められ出力されるのですか？ 38 名前：デフォルトの名無しさん mailto:sage [2014/12/20(土) 13:19:39.58 ID:hlsDA/2 ] [ここ壊れてます] 39 名前：G.net mailto: プロファイラ使えばわかるんじゃない？ [] [ここ壊れてます] 40 名前：デフォルトの名無しさん mailto:sage [2014/12/20(土) 13:33:50.56 ID:ARYnLzi0.net] >>34 ゲフォの800番台がスルーされたからかも? 41 名前：デフォルトの名無しさん mailto:sage [2014/12/20(土) 14:31:13.88 ID:NgIUM6cpA] メイン関数内でcudaMallocManagedを用いてユニファイドメモリを確保して その後メイン関数内でデータを代入し、GPU関数へポインタを渡したのですが、 値を利用できません。 だれか詳しく解決方法を教えてください 42 名前：デフォルトの名無しさん mailto:sage [2014/12/20(土) 14:34:48.04 ID:NgIUM6cpA] 38ですがこんな感じに参照してます __global__ void check_mem(int *file_top,int *d_test){   int x = blockIdx.x * blockDim.x + threadIdx.x;   d_test[x] = *(file_top + x); } file_topに値が入っていることはメインで確認が取れてます 43 名前：35 mailto:sage [2014/12/20(土) 22:32:39.87 ID:ovXiOWlC.net] >>36 少し気になって質問してみただけで今自分は一つしか持っていません！ もう一つGPUを購入した時の為にvisual profilerを使えるようになっておこうと思い ビルドした実行ファイルでプロファイラを使用してみたところ「Warning: No CUDA application was profiled, exiting」とエラーが出てしまいました 6.5のツールキットを使用して新しいセッションを作成→ビルドした実行ファイルを選択→設定はデフォルト、としたのですが何がダメだったのでしょう？ nvidiaの説明書を見てもさっぱりです。ヒントだけでもいいので教えてください 44 名前：デフォルトの名無しさん mailto:sage [2014/12/21(日) 10:11:52.91 ID:aUL9MF/2V] cudaDeviceReset() が必要のはず。 Visual Studioで「CUDA X.X Runtime」のプロジェクトを作ったら、 kernel.cu の return 0; の直前のコードに、以下のように書いてある。 // cudaDeviceReset must be called before exiting in order for profiling and // tracing tools such as Nsight and Visual Profiler to show complete traces. cudaStatus = cudaDeviceReset(); 45 名前：デフォルトの名無しさん mailto:sage [2014/12/21(日) 12:30:22.37 ID:C04pqXsd.net] >>40 cudaDeviceReset() が必要のはず。 Visual Studioで「CUDA X.X Runtime」のプロジェクトを作ったら、 kernel.cu の return 0; の直前のコードに、以下のように書いてある。 // cudaDeviceReset must be called before exiting in order for profiling and // tracing tools such as Nsight and Visual Profiler to show complete traces. cudaStatus = cudaDeviceReset(); 46 名前：35 mailto:sage [2014/12/21(日) 22:57:23.15 ID:McLr4XTH.net] >>42 ループしているプログラムなのでエスケープキーを押すと後処理関数をatexit関数で呼び出して終了するようになっています その後処理関数の中にcudaDeviceReset();を入れているのですがこれではダメなようです それともcudaError_t cudaStatus = cudaDeviceReset();としてcudaStatusをどこかへ渡すのでしょうか？ 47 名前：デフォルトの名無しさん mailto:sage [2014/12/22(月) 00:14:29.09 ID:6pNe5aqW.net] >>43 とりあえず、>>42を新規プロジェクト作ってプロファイラの 動作を確かめてから、あらためて自分のソースコードに反映すれば？ 48 名前：デフォルトの名無しさん [2014/12/23(火) 10:16:58.70 ID:rFJPpcq3i] 玄人志向のPCI-Express x16→PCI-Express x1変換ケーブルキット 「PCIEX16-X1/KIT」を使ってcudaやってみた人いますか？ 49 名前：デフォルトの名無しさん mailto:sage [2014/12/23(火) 15:40:43.82 ID:bsnZ8h6l.net] >>44 新規プロジェクトでサンプルプログラムが生成されるのを忘れていました おそらくサンプルのプロファイルに成功したので自分のプログラムに反映させようと思います ありがとうございました 50 名前：42 mailto:sage [2014/12/24(水) 18:51:30.92 ID:/5m6EieY.net] 一応書いておきます 調べてみた結果、必要なものはcudaDeviceReset();を呼び出すことのみでした 自分のプログラムがプロファイル出来なかった原因は.dllが.exeと同じ場所に無かったからでした 51 名前：デフォルトの名無しさん mailto:sage [2014/12/24(水) 20:32:20.58 ID:6fZpwBGv.net] 並列化についての質問です スレッドやブロックを増やしてもあまり計算速度に差が出ないのですがどのような理由が挙げられますか 52 名前：デフォルトの名無しさん mailto:sage [2014/12/24(水) 20:59:33.63 ID:SBHK+d/x.net] どう変わると思った？計算量自体は変わらんのだぞ。 53 名前：デフォルトの名無しさん mailto:sage [2014/12/24(水) 21:07:44.58 ID:6fZpwBGv.net] 最初は一つのスレッドにつき4回ほどループさせ計算をしていました その後、スレッド数を2倍にしてループ数を半分の2回しました 計算速度は2倍になるだろうと予想していましたが、あまり変わりませんでした 54 名前：デフォルトの名無しさん mailto:sage [2014/12/24(水) 21:23:52.28 ID:SBHK+d/x.net] ハード的に同時に実行できるスレッドは有限なんだから、それ以上スレッドを増やしても 物理的に速くなりようがない。 55 名前：デフォルトの名無しさん mailto:sage [2014/12/24(水) 22:24:43.11 ID:pbZqH+Xm.net] プログラム上のスレッド数とハード上のスレッド数は違うと言うことですか 56 名前：デフォルトの名無しさん mailto:sage [2014/12/24(水) 22:50:48.00 ID:3hqu78L7.net] スレッドが多ければ、メモリアクセスでスレッドが止まっている間 cudaコアは別のスレッドを実行できる 57 名前：デフォルトの名無しさん mailto:sage [2014/12/27(土) 01:57:38.03 ID:u9BI3CqV.net] 基本的にはcudaコアの数だけしか並列計算出来ないのですか だけしかと言ってもコアは何百もありますが 58 名前：デフォルトの名無しさん mailto:sage [2014/12/27(土) 06:52:06.91 ID:fxVjSbuk.net] ある瞬間、実際に並列に処理されているということと、理論上並列に扱われるということは別の話ですよ 上のレスにもありますが、計算速度的には実際に処理を行うヤツが足りていなければそこで頭打ちになるのは当然かと 59 名前：デフォルトの名無しさん mailto:sage [2014/12/27(土) 10:09:35.39 ID:W6Y2DM4+.net] cudaコアの数以上にスレッドを生成する利点は メモリアクセスの遅延の隠蔽にある 60 名前：デフォルトの名無しさん mailto:sage [2014/12/28(日) 20:39:13.97 ID:52BL0aAq.net] 550TIで使っていたプログラムを750TIで走らせて見たところ1.5倍ほど遅くなってしまいました 何故でしょうか？ 61 名前：デフォルトの名無しさん mailto:sage [2014/12/29(月) 02:19:31.23 ID:YUQudPNs.net] 腐ってやがる。早すぎたんだ 62 名前：デフォルトの名無しさん mailto:sage [2014/12/29(月) 09:47:06.28 ID:Sx0YYE+e.net] >>57 一度のカーネル実行で処理するデータ量を増やしたら改善しませんか? 63 名前：デフォルトの名無しさん mailto:sage [2014/12/29(月) 09:51:11.07 ID:Sx0YYE+e.net] >>57 maxwellは倍精度がそーとーしょぼいので、 cuda-zかなんかで性能をチェックした方がよいかもしれません。 sourceforge.jp/projects/sfnet_cuda-z/ 64 名前：デフォルトの名無しさん mailto:sage [2014/12/29(月) 12:03:33.25 ID:oV4aoJAy.net] >>60 本当にしょぼかった。陽子ビームぶち込みたい。 どうやらマクスウェルさんは演算用には向いていないようですね 65 名前：デフォルトの名無しさん mailto:sage [2014/12/29(月) 12:51:57.86 ID:oV4aoJAy.net] コア数が3倍になっていることを考慮すればそれでも遅い気がしますね 66 名前：デフォルトの名無しさん mailto:sage [2014/12/31(水) 21:08:03.80 ID:3b0Wn462.net] CUDA初学者です cudaBindTexture2D()のpitchとoffsetは何を表しているのですか？ 手元の書籍のサンプルから推測するに pitchは一次元の配列を二次元のテクスチャに入れる場合の折り返し地点のようなもの でしょうか？それならwidthとhighだけでも十分ではないかと言う疑問も出てきます。 そして一番の疑問がテクスチャメモリの存在です。 いくら二次元、三次元配列が使えるとは言え512バイトしか容量のないテクスチャメモリは64キロバイトもあるコンスタントメモリに劣るのではないでしょうか？ 長々と失礼いたしましたm(_ _)m 67 名前：デフォルトの名無しさん mailto:sage [2014/12/31(水) 23:07:39.94 ID:WrP28EMy.net] >> 68 名前：56 それじゃあストリームは何のためにあるのさ？ [] [ここ壊れてます] 69 名前：デフォルトの名無しさん mailto:sage [2015/01/01(木) 12:48:51.81 ID:82JnHkZd.net] >>64 どっちも使えるなら、実験して早い方を採用 同時に並んでいるスレッド数を増やしてcudaコアが遊ばない状況 を作り出すことが重要 ストリームの使い道は異なるカーネルの並列実行だと思ってる 70 名前：デフォルトの名無しさん mailto:sage [2015/01/02(金) 01:42:39.20 ID:aooXGYY5.net] >>65 なるほどね 71 名前：デフォルトの名無しさん [2015/01/02(金) 10:40:26.85 ID:2fT8SJ1Ez] >>65 「ストリームの使い道は異なるカーネルの並列実行」というのは 何が言いたいのか良くわからないが、捕捉すると、 ストリームはCPU-GPU間のデータ転送中（非同期のcudaMemcpy）に、GPUで演算（カーネル実行）させる為の仕組みだ。 「あるストリームがデータ転送待ち中に、別のストリームのカーネルを実行する。」 といった事を、賢くやってくれる。 （別のストリームが「異なるカーネル」でなく、「同じ関数のカーネル」でも良い。） 結果、「データ転送待ち」の時間を節約できる。 ストリームの詳細は、書籍に書いてあるが、 データ転送とカーネル実行をストリームに入れる「推奨の順番」が cc3.5からは違うので、古い本の場合は読み替える必要あり。 72 名前：67 [2015/01/02(金) 11:01:01.35 ID:2fT8SJ1Ez] まとめると、こんな感じ。 （※同じプログラムで、以下の両方を使う事も可能。） ・cudaコアの数以上にスレッドを生成する：「GPU」と「GPU側のDRAM」間のメモリアクセスの遅延の隠蔽が期待できる。 ・ストリームを使う：「CPU側のDRAM」と「GPU側のDRAM」間のメモリアクセスの時間の節約が期待できる。 73 名前：デフォルトの名無しさん mailto:sage [2015/01/03(土) 13:07:05.02 ID:yWVdPt25.net] 970/980は確かにゲームのパフォーマンスは上がってるが帯域減ってるから GPGPU用途では微妙になってしまったな 74 名前：デフォルトの名無しさん mailto:sage [2015/01/09(金) 00:33:57.02 ID:Iq4Pw+IC.net] Toolkit 6.0 + VS 2008から Toolkit 6.5 + VS 2013に移行したら 遅くなっちゃったんだけど、そういう人ほかにいる? 75 名前：デフォルトの名無しさん mailto:sage [2015/01/12(月) 00:15:12.48 ID:crrCnhEj.net] CUDAの日本語ページって4.0とかの古い情報ばっかりじゃね 6.0/6.5では全然仕様が違ってて全然使えない 76 名前：デフォルトの名無しさん [2015/01/16(金) 07:53:04.52 ID:IHSf0jGJ.net] CUDA7.0 RC 77 名前：デフォルトの名無しさん mailto:sage [2015/01/16(金) 09:02:56.43 ID:VQ2eHsT0.net] もうCUDAも成熟してしまった感があるなあ。 78 名前：デフォルトの名無しさん mailto:sage [2015/01/25(日) 10:58:33.86 ID:m2kue9j8.net] 970の影響でGPUメモリテストが流行っているね。 79 名前：デフォルトの名無しさん mailto:sage [2015/02/06(金) 21:35:53.90 ID:72/Q/UeS.net] ここ何週間かデバッグを続けているのですが原因を突き止めることが出来ません 初学者がはまりやすいミスやデバッグのこつなんかを教えてもらえませんか？ 明らかなバグなら原因を突き止めやすいのですが、かなり微妙なバグなのでなかなか見つけられず困っています 80 名前：デフォルトの名無しさん mailto:sage [2015/02/07(土) 00:48:46.19 ID:OS4q1AxS.net] printfとかで要素を表示してデバックしてみれば？ 81 名前：デフォルトの名無しさん mailto:sage [2015/02/07(土) 13:10:14.61 ID:4cvxubK6.net] syncthreadとか？ if文の中に書いてたりすると同期ずれが起こったりするなー 他には確保してないメモリへのアクセスとか？ >>76の通り、printfとかで、配列の添字とか値を表示するしかないのかな？ 82 名前：デフォルトの名無しさん mailto:sage [2015/02/08(日) 15:23:01.85 ID:E04CIgi2.net] >>76 >>77 ありがとうございます 1セットの計算量があまりにも多いのでprintfの方法は難しいです シンクロや範囲外アクセスもありませんでした 原因が分かっちゃったかも知れないので質問です 中間計算結果→atomicAd 83 名前：d 中間計算結果→配列→atomicAdd こんな風に同じ数値を使った計算でも一度配列を通してしまうとatomicAddによって追加された計算結果に差が出たりしますか？ [] [ここ壊れてます] 84 名前：デフォルトの名無しさん mailto:sage [2015/02/08(日) 15:48:00.96 ID:E04CIgi2.net] 変数に入れると精度は落ちますね お騒がせしました 85 名前：デフォルトの名無しさん mailto:sage [2015/02/08(日) 21:32:20.44 ID:BpjOkBmf.net] >>78 ちょっと面倒だけど、要素が多い場合は減らしてやってみるとか、どうだろうか？ 何はともあれ、原因判明したみたいで、おめでとう 86 名前：デフォルトの名無しさん mailto:sage [2015/02/08(日) 22:56:08.44 ID:KLuvC02r.net] >>80 それは意外な盲点でした 数を減らせば良かったのですね 87 名前：デフォルトの名無しさん mailto:sage [2015/02/09(月) 10:20:08.08 ID:pN+UjOmC.net] >>78 fpが(a+b)+c != a+(b+c)を知らないとかではないよね？ 88 名前：デフォルトの名無しさん mailto:sage [2015/02/09(月) 23:18:25.64 ID:QR2S1do8.net] volatile使うとか？ 変数の宣言とか関数の引数の型の前にvolatileをいれると・・・ 89 名前：デフォルトの名無しさん mailto:sage [2015/02/17(火) 21:40:32.15 ID:K8c74Rhe.net] >>57 750TIでGPGPUって考えていたけど、750TIって2世代前の同ランクぐらいの550TIより性能悪いのか。 一般ゲーム用VGAではGPGPU能力ってたいして要らないから落としたのかな いろいろなゲーム用VGAの単精度、倍精度の能力が載ったホームページ教えてください 90 名前：undefined mailto:undefined [2015/02/19(木) 11:23:20.97 ID:aqLRWkl1.net] 質問☆ cudaってドライバインストして、画像表示をcuda設定にするだけでは 効果ない？ 91 名前：デフォルトの名無しさん mailto:sage [2015/02/19(木) 14:30:54.02 ID:iKdaAUCi.net] >>84 FP32とFP64の一覧表ならこれとか。 ttp://www.geeks3d.com/20140305/amd-radeon-and-nvidia-geforce-fp32-fp64-gflops-table-computing/ ボトルネックになりうる点は他にもあるから、Compute Capability毎の仕様の違いも結構重要だと思う。 92 名前：デフォルトの名無しさん mailto:sage [2015/02/19(木) 23:45:36.57 ID:ngPIgbTR.net] maxwellさん自体にに倍精度が無いようだから 一世代前のkeplerさんか次世代のpascalさんを選べば良いんじゃないかな 93 名前：デフォルトの名無しさん mailto:sage [2015/02/19(木) 23:54:19.28 ID:Lt8lBsrZ.net] >>86 有難う。750TiのFP64悪すぎだな。 なんか大衆向け用でGPGPUするならFP32よ、FP64は使わないでだな。 94 名前：デフォルトの名無しさん mailto:sage [2015/02/20(金) 04:05:21.12 ID:fPdGyDpl.net] >>87 KeplerはMaxwell以上にピーキーだったような。 自分の用途がはっきりしていて、それがKeplerやMaxwellに向いているならありだろうけど。 64bit変数をほとんど使わなくても、不向きな処理ではGTX 680が570に惨敗したりする。 CUDAの開発環境とか情報量に魅力を感じて、あえて今から始めてみるという人に勧めるとすれば、個人的には 投げ売り続行中のGTX 570や580で、余裕があるなら型落ちCPU・マザボ・メモリのセット等と 組み合わせてCUDA専用マシンを用意かな。 >>88 それは差別化とか、グラフィック用途でのワットパフォーマンスとかで仕方がないかと。 95 名前：デフォルトの名無しさん [2015/02/20(金) 08:52:57.87 ID:xG3c1huj.net] keplerはinteger bit shiftが弱いGK110(tesla)以外は maxwellはkepler比で2倍のスループットになってる 96 名前：デフォルトの名無しさん mailto:sage [2015/03/04(水) 22:02:58.94 ID:krHDLIbc3] コアレッシングて、なんでハーフワープじゃないとダメなのでありますか？ フルじゃダメなのでしょうか。 #deviceQuery: # CUDA Capability Major/Minor version number: 3.2 # Warp size: 32 97 名前：デフォルトの名無しさん mailto:sage [2015/02/22(日) 18:45:23.46 ID:JhGx5uct.net] 適当なプログラム作ってみても maxwellの方がはやいね shared memoryが倍になったのも大きいなぁ 98 名前：名無し [2015/03/07(土) 10:52:38.85 ID:UBzBpgz5.net] スレチなら申し訳ない 当方、モバイルでCUDAを使用したいけど 安い方法はどれが良いと思います？ （速度はそこそこで良く、外でテストして 　パワーがいる場合はデスクトップを使用するつもり） モバイル用は安く上げたいので Chromebookかタブレットで探した方が良いですかね？ 奇をてらってJetsonのtk1にACアダプタ用のバッテリーを積むとか （可能かどうかわからないですが） ちなみにゲームはやるつもりありません。 99 名前：デフォルトの名無しさん mailto:sage [2015/03/07(土) 20:42:36.10 ID:CzdLWIdo.net] thinkpad w550sのquadro K620mはダメなの？ ノートパソコンだよ 100 名前：デフォルトの名無しさん [2015/03/08(日) 10:25:35.13 ID:TYY6zzsE.net] レスありがとうございます 安くあげたいので予算的に厳しいかと 最初だけ計算量は多いですけど 要所ごとに定数化すれば、その後は計算量がへるかと思っているので外での使用は少ないデータ量でプログラムチェックができればいいかなと考えています。 K1がのったタブレットも安いのでそこから考えてみようと思います ありがとうございました 101 名前：デフォルトの名無しさん mailto:sage [2015/03/18(水) 02:36:15.45 ID:jKTvW/7W.net] うわ、titan xの倍精度、しょぼ過ぎ・・・。 https://twitter.com/search?q=titan%20x&src=typd

---

---

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

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