【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

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