【GPGPU】くだすれCUDAスレ part7【NVIDIA】 [転載禁止]©2ch.net

**デフォルトの名無しさん** · 2014/11/20(木) 23:14:46.66

このスレッドは、他のスレッドでは書き込めない超低レベル、
もしくは質問者自身何が何だが分からない質問を勇気を持って書き込むスレッドです。
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】
http://toro.2ch.net/test/read.cgi/tech/1314104886/
【GPGPU】くだすれCUDAスレ part6【NVIDIA】
ttp://peace.2ch.net/test/read.cgi/tech/1348409867/

**デフォルトの名無しさん** · 2014/11/20(木) 23:15:41.57

関連サイト
CUDA
http://www.nvidia.co.jp/object/cuda_home_new_jp.html

CUDAに触れてみる
http://chihara.naist.jp/people/STAFF/imura/computer/OpenGL/cuda1/disp_content

CUDA のインストール
http://blog.goo.ne.jp/sdpaninf/e/9533f75438b670a174af345f4a33bd51

NVIDIAの「GeForce 8800 GT(G92)」と次に控える64-bit GPUアーキテクチャ
http://pc.watch.impress.co.jp/docs/2007/1031/kaigai398.htm

CUDAを使う
http://tech.ckme.co.jp/cuda.shtml

NVIDIA CUDAを弄ってみたその2
http://dvd-r.sblo.jp/article/10422960.html

CUDAベンチ
http://wataco.air-nifty.com/syacho/2008/02/cuda_2044.html

KNOPPIX for CUDA
http://www.yasuoka.mech.keio.ac.jp/cuda/

【東電 69.8 %】 · 2014/11/21(金) 00:13:37.04

>>1
ああ、キミ！また会えたね。久しぶりだ。どうだいあの件は？どうなったか説明したまえな。

神戸市の東、芦屋西宮の知的障害者施設で未成年利用者に性的な行為をして淫行条例で逮捕された三田谷学園元職員の堂垣直人(西宮市老松町)は、結局どういう罪になったの？
被害者家族のケアを芦屋市役所と兵庫県警はちゃんとやったのか？
差別や虐待は環境を選べない子供には関係ない。

http://www.youtube.com/watch?v=JxMzW3ZlV4g&;sns=em

まあ、こっちに座れよ。ゆっくり話そうじゃないか。

**デフォルトの名無しさん** · 2014/11/21(金) 07:36:45.42

http://i.imgur.com/SduT7qR.jpg
http://i.imgur.com/2f8fP3F.jpg
http://i.imgur.com/u0ymg8c.jpg
http://i.imgur.com/LruEssT.png
http://i.imgur.com/AeSQqRT.jpg
http://cisburger.com/up/bnf/6016.jpg
http://up.pangya.tv/src/www_pangya_tv18594.jpg

**デフォルトの名無しさん** · 2014/11/21(金) 18:05:42.00

syncthreadsとthreadfence_blockの違いが分かりません
syncthreadsだけで十分な気がしますが、どういう時に使い分けるのでしょうか？

**デフォルトの名無しさん** · 2014/11/22(土) 13:34:27.11

>>5
http://shobomaru.wordpress.com/2013/09/13/synchronizing-instruction-in-direct-compute-and-cuda/

メモリの書き込み競合防止をするかどうかの違いみたいです。

これってatomic命令より軽いのか知らん?

**デフォルトの名無しさん** · 2014/11/22(土) 13:48:05.92

>>6
競合防止なんて書いてる？
複数のスレッドが同じメモリに書き込む時はatomicをsyncthreadsやthreadfenceと同時に使わないといけない気がする(誰か教えてください)

**デフォルトの名無しさん** · 2014/11/22(土) 22:18:58.72

すみません、誤読しました。
「同じブロック内の全スレッドがこの命令にたどり着く」
まで待つかどうかの違いですね。

**デフォルトの名無しさん** · 2014/11/24(月) 04:18:18.83

今一番コスパ高いカードってなに？

**デフォルトの名無しさん** · 2014/11/24(月) 04:21:35.17

謝罪文みても思い上がりが激しい

**デフォルトの名無しさん** · 2014/11/25(火) 19:26:42.72

プログラムのカーネル部分がどうしても実行されません(サンプルプログラムでは実行されていました)
どなたか原因に心当たりはありませんか？(私はありません)
ブレークポイントで確認したところカーネルの上下にあるクロックは実行されており、
カーネルだけが実行されていませんでした(カーネルの中へ入って行かないという意味です)

以下が呼び出しで、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";
}

**デフォルトの名無しさん** · 2014/11/25(火) 21:02:59.83

>>11
dim3構造体のメンバ変数をプリントしたらどうなりますか?

**デフォルトの名無しさん** · 2014/11/25(火) 21:44:27.45

>>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)
{・・・}

**デフォルトの名無しさん** · 2014/11/26(水) 18:34:20.92

>>13
カーネルが実行されているかどうかはどうやって確認していますか?

**デフォルトの名無しさん** · 2014/11/26(水) 19:30:35.81

>>14
>>13に書いた通り、ブレークポイントをカーネルの中に入れて実行されているかどうかを確認しています
描画をするプログラムなので図形が動くか動かないかでも判断できます

**デフォルトの名無しさん** · 2014/11/26(水) 20:04:14.26

おれもカーネルに入らないケース出たわ
原因調査中・・・

**デフォルトの名無しさん** · 2014/11/27(木) 08:04:41.88

>>15
CUDAのデバッガ使った事がない(ひたすらprintf)ので
一般論的な事しか言えませんが、
怪しそうな処理をコメントアウトしていったらどうでしょう?
変なメモリアクセスで落ちるとか割とありがちな気が。

**デフォルトの名無しさん** · 2014/11/28(金) 22:55:16.84

おれもカーネルに入らない
ただなぜかcygwinでコンパイルするとカーネルが起動する
なんでじゃ・・・

例のvisual studio2013もインストールしてみたいな～

**デフォルトの名無しさん** · 2014/11/28(金) 23:06:00.72

>>18だけど
osはwindows8.1
コンパイラはcuda6.5+vs2013 express
カーネルは担当する要素を＋1するだけのもの
これをコマンドプロンプトでコンパイルしてもカーネルは動かなかった(?)
動かないと判断したのは結果をmemcpyしてホスト側に返しても＋１されてなかったから
もしかしたら正常にmemcpyされてないだけかもしれない

カーネルに入らない人はcygwin使ってみるといいかもね
本質的な解決にはならないけど・・・

**デフォルトの名無しさん** · 2014/11/30(日) 20:47:05.05

ビジュアルプロファイラー使ってみたら。

**デフォルトの名無しさん** · 2014/12/03(水) 14:38:19.35

kernel実行後にcudaGetLastError()でRCを取得
そうするとkernel実行結果が分かる
RCの数値は自分で調べてね

**16,21** · 2014/12/04(木) 22:23:47.78

21の情報は役だったかな？
自分の場合（LINUX）、RC=7（too many resources requested for launch）だったので
コンパイルオプションに　-Xptxas -vを追加して使用レジスタ数を確認。

結果ハードウェアのレジスタ数を超過したためにカーネルの処理が行われなかったことが判明。
スレッドサイズを小さくして問題解決。
因みに使用レジスタはハードウェアによって変わる。

**デフォルトの名無しさん** · 2014/12/05(金) 14:06:24.56

>>22
ちょっと興味があるんですけれど、
カーネルのサイズは動的に決めてるんですか?

前にソースコードにブロック数とスレッド数をべた書きで
大きいサイズを指定したらコンパイルの段階ではじかれた事があったんで。

**>>15** · 2014/12/05(金) 15:37:53.47

>>21
ありがとうすごく役に立ったよ
ここ最近忙しくてｐｃ触れなかったんだ
自分も同じく「error: too many resources requested for launch」だった
原因を調べてみるよ

あと、自動でエラー内容もだせるみたいだね
http://homepage2.nifty.com/takaaki024/tips/programs/gpgpu/cuda.html

>>20
ビジュアルプロファイラーも便利そうだから調べてみるよ

**デフォルトの名無しさん** · 2014/12/07(日) 18:17:39.03

x,y,zの3つの変数から成る構造体配列A,B(同じサイズ)があったとして
BからAへそれぞれ対応するデータを転送する場合

Ax,Ay,Az,Bx,By,Bzという同じサイズの構造体でない配列が6つあったとして
BからAへそれぞれ対応するデータを転送する場合(BxからAxなど)

前者と後者では後者の方が転送速度は上がりますか？

**デフォルトの名無しさん** · 2014/12/08(月) 12:53:10.47

>>25
一般論としてデータ量が同じなら一回にまとめてを転送した方が効率はよくなりますね。

**デフォルトの名無しさん** · 2014/12/11(木) 11:46:54.75

memcpyって同期とるもんね

**デフォルトの名無しさん** · 2014/12/12(金) 21:54:58.15

kernelの中で使えるタイマー関数はありますか？

**デフォルトの名無しさん** · 2014/12/12(金) 23:43:18.27

clock()関数が使えるよ。

**名無しさん＠そうだ選挙に行こう** · 2014/12/13(土) 22:02:15.64

>>29
サンキュー
試してみるよ

**デフォルトの名無しさん** · 2014/12/16(火) 21:07:44.74

自分の持っていないGPUの共有メモリの量などの詳細を知ることはできますか？

**デフォルトの名無しさん** · 2014/12/17(水) 08:36:19.04

http://en.wikipedia.org/wiki/CUDA#Version_features_and_specifications
の
Maximum amount of shared memory per multiprocessor
あたりかな。

**デフォルトの名無しさん** · 2014/12/17(水) 20:16:42.59

>>32
ありがとうございます
嬉しいことに5.0以降から容量が増えてるみたいですね

**デフォルトの名無しさん** · 2014/12/17(水) 20:55:04.06

なんでCCのバージョン3.5から5.0に飛んでるの？
SDKのバージョンと合わせたのか？

**デフォルトの名無しさん** · 2014/12/19(金) 18:38:15.45

二つ以上のGPUでVBOを使用する場合、データの流れはどうなっているのでしょう？
やはり一旦ディスプレイに接続側のGPU出力データが集められ出力されるのですか？

**デフォルトの名無しさん** · 2014/12/20(土) 13:19:39.58

プロファイラ使えばわかるんじゃない？

**デフォルトの名無しさん** · 2014/12/20(土) 13:33:50.56

>>34
ゲフォの800番台がスルーされたからかも?

35 · 2014/12/20(土) 22:32:39.87

>>36
少し気になって質問してみただけで今自分は一つしか持っていません！

もう一つGPUを購入した時の為にvisual profilerを使えるようになっておこうと思い
ビルドした実行ファイルでプロファイラを使用してみたところ「Warning: No CUDA application was profiled, exiting」とエラーが出てしまいました
6.5のツールキットを使用して新しいセッションを作成→ビルドした実行ファイルを選択→設定はデフォルト、としたのですが何がダメだったのでしょう？
nvidiaの説明書を見てもさっぱりです。ヒントだけでもいいので教えてください

**デフォルトの名無しさん** · 2014/12/21(日) 12:30:22.37

>>38
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();

35 · 2014/12/21(日) 22:57:23.15

>>39
ループしているプログラムなのでエスケープキーを押すと後処理関数をatexit関数で呼び出して終了するようになっています
その後処理関数の中にcudaDeviceReset();を入れているのですがこれではダメなようです
それともcudaError_t cudaStatus = cudaDeviceReset();としてcudaStatusをどこかへ渡すのでしょうか？

**デフォルトの名無しさん** · 2014/12/22(月) 00:14:29.09

>>40
とりあえず、>>39を新規プロジェクト作ってプロファイラの
動作を確かめてから、あらためて自分のソースコードに反映すれば？

**デフォルトの名無しさん** · 2014/12/23(火) 15:40:43.82

>>41
新規プロジェクトでサンプルプログラムが生成されるのを忘れていました
おそらくサンプルのプロファイルに成功したので自分のプログラムに反映させようと思います
ありがとうございました

42 · 2014/12/24(水) 18:51:30.92

一応書いておきます
調べてみた結果、必要なものはcudaDeviceReset();を呼び出すことのみでした
自分のプログラムがプロファイル出来なかった原因は.dllが.exeと同じ場所に無かったからでした

**デフォルトの名無しさん** · 2014/12/24(水) 20:32:20.58

並列化についての質問です
スレッドやブロックを増やしてもあまり計算速度に差が出ないのですがどのような理由が挙げられますか

**デフォルトの名無しさん** · 2014/12/24(水) 20:59:33.63

どう変わると思った？計算量自体は変わらんのだぞ。

**デフォルトの名無しさん** · 2014/12/24(水) 21:07:44.58

最初は一つのスレッドにつき4回ほどループさせ計算をしていました
その後、スレッド数を2倍にしてループ数を半分の2回しました
計算速度は2倍になるだろうと予想していましたが、あまり変わりませんでした

**デフォルトの名無しさん** · 2014/12/24(水) 21:23:52.28

ハード的に同時に実行できるスレッドは有限なんだから、それ以上スレッドを増やしても
物理的に速くなりようがない。

**デフォルトの名無しさん** · 2014/12/24(水) 22:24:43.11

プログラム上のスレッド数とハード上のスレッド数は違うと言うことですか

**デフォルトの名無しさん** · 2014/12/24(水) 22:50:48.00

スレッドが多ければ、メモリアクセスでスレッドが止まっている間
cudaコアは別のスレッドを実行できる

**デフォルトの名無しさん** · 2014/12/27(土) 01:57:38.03

基本的にはcudaコアの数だけしか並列計算出来ないのですか

だけしかと言ってもコアは何百もありますが

**デフォルトの名無しさん** · 2014/12/27(土) 06:52:06.91

ある瞬間、実際に並列に処理されているということと、理論上並列に扱われるということは別の話ですよ

上のレスにもありますが、計算速度的には実際に処理を行うヤツが足りていなければそこで頭打ちになるのは当然かと

**デフォルトの名無しさん** · 2014/12/27(土) 10:09:35.39

cudaコアの数以上にスレッドを生成する利点は
メモリアクセスの遅延の隠蔽にある

**デフォルトの名無しさん** · 2014/12/28(日) 20:39:13.97

550TIで使っていたプログラムを750TIで走らせて見たところ1.5倍ほど遅くなってしまいました
何故でしょうか？

**デフォルトの名無しさん** · 2014/12/29(月) 02:19:31.23

腐ってやがる。早すぎたんだ

**デフォルトの名無しさん** · 2014/12/29(月) 09:47:06.28

>>53
一度のカーネル実行で処理するデータ量を増やしたら改善しませんか?

**デフォルトの名無しさん** · 2014/12/29(月) 09:51:11.07

>>53
maxwellは倍精度がそーとーしょぼいので、
cuda-zかなんかで性能をチェックした方がよいかもしれません。
http://sourceforge.jp/projects/sfnet_cuda-z/

**デフォルトの名無しさん** · 2014/12/29(月) 12:03:33.25

>>56
本当にしょぼかった。陽子ビームぶち込みたい。
どうやらマクスウェルさんは演算用には向いていないようですね

**デフォルトの名無しさん** · 2014/12/29(月) 12:51:57.86

コア数が3倍になっていることを考慮すればそれでも遅い気がしますね

**デフォルトの名無しさん** · 2014/12/31(水) 21:08:03.80

CUDA初学者です
cudaBindTexture2D()のpitchとoffsetは何を表しているのですか？
手元の書籍のサンプルから推測するに
pitchは一次元の配列を二次元のテクスチャに入れる場合の折り返し地点のようなもの
でしょうか？それならwidthとhighだけでも十分ではないかと言う疑問も出てきます。

そして一番の疑問がテクスチャメモリの存在です。
いくら二次元、三次元配列が使えるとは言え512バイトしか容量のないテクスチャメモリは64キロバイトもあるコンスタントメモリに劣るのではないでしょうか？
長々と失礼いたしましたm(_ _)m

**デフォルトの名無しさん** · 2014/12/31(水) 23:07:39.94

>>52
それじゃあストリームは何のためにあるのさ？

**デフォルトの名無しさん** · 2015/01/01(木) 12:48:51.81

>>60
どっちも使えるなら、実験して早い方を採用
同時に並んでいるスレッド数を増やしてcudaコアが遊ばない状況
を作り出すことが重要

ストリームの使い道は異なるカーネルの並列実行だと思ってる

**デフォルトの名無しさん** · 2015/01/02(金) 01:42:39.20

>>61
なるほどね

**デフォルトの名無しさん** · 2015/01/03(土) 13:07:05.02

970/980は確かにゲームのパフォーマンスは上がってるが帯域減ってるから
GPGPU用途では微妙になってしまったな

**デフォルトの名無しさん** · 2015/01/09(金) 00:33:57.02

Toolkit 6.0 + VS 2008から
Toolkit 6.5 + VS 2013に移行したら
遅くなっちゃったんだけど、そういう人ほかにいる?

**デフォルトの名無しさん** · 2015/01/12(月) 00:15:12.48

CUDAの日本語ページって4.0とかの古い情報ばっかりじゃね
6.0/6.5では全然仕様が違ってて全然使えない

**デフォルトの名無しさん** · 2015/01/16(金) 07:53:04.52

CUDA7.0 RC

**デフォルトの名無しさん** · 2015/01/16(金) 09:02:56.43

もうCUDAも成熟してしまった感があるなあ。

**デフォルトの名無しさん** · 2015/01/25(日) 10:58:33.86

970の影響でGPUメモリテストが流行っているね。

**デフォルトの名無しさん** · 2015/02/06(金) 21:35:53.90

ここ何週間かデバッグを続けているのですが原因を突き止めることが出来ません
初学者がはまりやすいミスやデバッグのこつなんかを教えてもらえませんか？
明らかなバグなら原因を突き止めやすいのですが、かなり微妙なバグなのでなかなか見つけられず困っています

**デフォルトの名無しさん** · 2015/02/07(土) 00:48:46.19

printfとかで要素を表示してデバックしてみれば？

**デフォルトの名無しさん** · 2015/02/07(土) 13:10:14.61

syncthreadとか？
if文の中に書いてたりすると同期ずれが起こったりするなー
他には確保してないメモリへのアクセスとか？

>>70の通り、printfとかで、配列の添字とか値を表示するしかないのかな？

**デフォルトの名無しさん** · 2015/02/08(日) 15:23:01.85

>>70
>>71
ありがとうございます
1セットの計算量があまりにも多いのでprintfの方法は難しいです
シンクロや範囲外アクセスもありませんでした

原因が分かっちゃったかも知れないので質問です
中間計算結果→atomicAdd
中間計算結果→配列→atomicAdd
こんな風に同じ数値を使った計算でも一度配列を通してしまうとatomicAddによって追加された計算結果に差が出たりしますか？

**デフォルトの名無しさん** · 2015/02/08(日) 15:48:00.96

変数に入れると精度は落ちますね
お騒がせしました

**デフォルトの名無しさん** · 2015/02/08(日) 21:32:20.44

>>72
ちょっと面倒だけど、要素が多い場合は減らしてやってみるとか、どうだろうか？

何はともあれ、原因判明したみたいで、おめでとう

**デフォルトの名無しさん** · 2015/02/08(日) 22:56:08.44

>>74
それは意外な盲点でした
数を減らせば良かったのですね

**デフォルトの名無しさん** · 2015/02/09(月) 10:20:08.08

>>72
fpが(a+b)+c != a+(b+c)を知らないとかではないよね？

**デフォルトの名無しさん** · 2015/02/09(月) 23:18:25.64

volatile使うとか？
変数の宣言とか関数の引数の型の前にvolatileをいれると・・・

**デフォルトの名無しさん** · 2015/02/17(火) 21:40:32.15

>>53
750TIでGPGPUって考えていたけど、750TIって2世代前の同ランクぐらいの550TIより性能悪いのか。
一般ゲーム用VGAではGPGPU能力ってたいして要らないから落としたのかな

いろいろなゲーム用VGAの単精度、倍精度の能力が載ったホームページ教えてください

**undefined** · 2015/02/19(木) 11:23:20.97

質問☆
cudaってドライバインストして、画像表示をcuda設定にするだけでは
効果ない？

**デフォルトの名無しさん** · 2015/02/19(木) 14:30:54.02

>>78
FP32とFP64の一覧表ならこれとか。
ttp://www.geeks3d.com/20140305/amd-radeon-and-nvidia-geforce-fp32-fp64-gflops-table-computing/

ボトルネックになりうる点は他にもあるから、Compute Capability毎の仕様の違いも結構重要だと思う。

**デフォルトの名無しさん** · 2015/02/19(木) 23:45:36.57

maxwellさん自体にに倍精度が無いようだから
一世代前のkeplerさんか次世代のpascalさんを選べば良いんじゃないかな