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

1 ：デフォルトの名無しさん：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】
http://toro.2ch.net/test/read.cgi/tech/1314104886/
【GPGPU】くだすれCUDAスレ part6【NVIDIA】
ttp://peace.2ch.net/test/read.cgi/tech/1348409867/

2 ：デフォルトの名無しさん：2014/11/20(木) 23:15:41.57 ID:jr3oZn27.net

関連サイト
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/

3 ： 【東電 69.8 %】 ：2014/11/21(金) 00:13:37.04 ID:kFWiXf0I.net

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

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

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


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

4 ：デフォルトの名無しさん：2014/11/21(金) 07:36:45.42 ID:JsFj8Vej.net

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

5 ：デフォルトの名無しさん：2014/11/21(金) 18:05:42.00 ID:qiUQrZk/.net

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

6 ：デフォルトの名無しさん：2014/11/22(土) 13:34:27.11 ID:Ke1g3qvZ.net

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

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

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

7 ：デフォルトの名無しさん：2014/11/22(土) 13:48:05.92 ID:S8C7U0PL.net

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

8 ：デフォルトの名無しさん：2014/11/22(土) 22:18:58.72 ID:Ke1g3qvZ.net

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

9 ：デフォルトの名無しさん：2014/11/24(月) 04:18:18.83 ID:qPQKDlD6.net

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

10 ：デフォルトの名無しさん：2014/11/24(月) 04:21:35.17 ID:qPQKDlD6.net

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

11 ：デフォルトの名無しさん：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";
}

12 ：デフォルトの名無しさん：2014/11/25(火) 21:02:59.83 ID:N/U8okyJ.net

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

13 ：デフォルトの名無しさん：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)
{・・・}

14 ：デフォルトの名無しさん：2014/11/26(水) 18:34:20.92 ID:cpKKMAIz.net

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

15 ：デフォルトの名無しさん：2014/11/26(水) 19:30:35.81 ID:mOjmGjn5.net

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

16 ：デフォルトの名無しさん：2014/11/26(水) 20:04:14.26 ID:qey6HT7s.net

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

17 ：デフォルトの名無しさん：2014/11/27(木) 08:04:41.88 ID:7alpN+o4.net

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

18 ：デフォルトの名無しさん：2014/11/28(金) 22:55:16.84 ID:JeOcX4pA.net

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

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

19 ：デフォルトの名無しさん：2014/11/28(金) 23:06:00.72 ID:JeOcX4pA.net

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

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

20 ：デフォルトの名無しさん：2014/11/30(日) 20:47:05.05 ID:NdicNENH.net

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

21 ：デフォルトの名無しさん：2014/12/03(水) 14:38:19.35 ID:GXBajCbw.net

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

22 ：16,21：2014/12/04(木) 22:23:47.78 ID:EZ4odEf+.net

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

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

23 ：デフォルトの名無しさん：2014/12/05(金) 14:06:24.56 ID:+nFWXccn.net

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

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

24 ：>>15：2014/12/05(金) 15:37:53.47 ID:fpNGtjbn.net

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

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

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

25 ：デフォルトの名無しさん：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など)

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

26 ：デフォルトの名無しさん：2014/12/08(月) 12:53:10.47 ID:JbpvX5Qi.net

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

27 ：デフォルトの名無しさん：2014/12/11(木) 11:46:54.75 ID:kV0/O7vj.net

memcpyって同期とるもんね

28 ：デフォルトの名無しさん：2014/12/12(金) 21:54:58.15 ID:PdQu+k/h.net

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

29 ：デフォルトの名無しさん：2014/12/12(金) 23:43:18.27 ID:q1FKM2bt.net

clock()関数が使えるよ。

30 ：名無しさん＠そうだ選挙に行こう：2014/12/13(土) 22:02:15.64 ID:B9P4oQcX.net

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

31 ：デフォルトの名無しさん：2014/12/16(火) 21:07:44.74 ID:6hyQD5WD.net

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

32 ：デフォルトの名無しさん：2014/12/17(水) 08:36:19.04 ID:0flByQKi.net

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

33 ：デフォルトの名無しさん：2014/12/17(水) 20:16:42.59 ID:D/43rANg.net

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

34 ：デフォルトの名無しさん：2014/12/17(水) 20:55:04.06 ID:vfaS5qRM.net

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

35 ：デフォルトの名無しさん：2014/12/19(金) 18:38:15.45 ID:dxBCSCiu.net

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

36 ：デフォルトの名無しさん：2014/12/20(土) 13:19:39.58 ID:hlsDA/2G.net

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

37 ：デフォルトの名無しさん：2014/12/20(土) 13:33:50.56 ID:ARYnLzi0.net

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

38 ：デフォルトの名無しさん：2014/12/20(土) 14:31:13.88 ID:NgIUM6cpA

メイン関数内でcudaMallocManagedを用いてユニファイドメモリを確保して
その後メイン関数内でデータを代入し、GPU関数へポインタを渡したのですが、
値を利用できません。

だれか詳しく解決方法を教えてください

39 ：デフォルトの名無しさん：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に値が入っていることはメインで確認が取れてます

40 ：35：2014/12/20(土) 22:32:39.87 ID:ovXiOWlC.net

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

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

41 ：デフォルトの名無しさん：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();

42 ：デフォルトの名無しさん：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();

43 ：35：2014/12/21(日) 22:57:23.15 ID:McLr4XTH.net

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

44 ：デフォルトの名無しさん：2014/12/22(月) 00:14:29.09 ID:6pNe5aqW.net

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

45 ：デフォルトの名無しさん：2014/12/23(火) 10:16:58.70 ID:rFJPpcq3i

玄人志向のPCI-Express x16→PCI-Express x1変換ケーブルキット
「PCIEX16-X1/KIT」を使ってcudaやってみた人いますか？

46 ：デフォルトの名無しさん：2014/12/23(火) 15:40:43.82 ID:bsnZ8h6l.net

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

47 ：42：2014/12/24(水) 18:51:30.92 ID:/5m6EieY.net

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

48 ：デフォルトの名無しさん：2014/12/24(水) 20:32:20.58 ID:6fZpwBGv.net

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

49 ：デフォルトの名無しさん：2014/12/24(水) 20:59:33.63 ID:SBHK+d/x.net

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

50 ：デフォルトの名無しさん：2014/12/24(水) 21:07:44.58 ID:6fZpwBGv.net

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

51 ：デフォルトの名無しさん：2014/12/24(水) 21:23:52.28 ID:SBHK+d/x.net

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

52 ：デフォルトの名無しさん：2014/12/24(水) 22:24:43.11 ID:pbZqH+Xm.net

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

53 ：デフォルトの名無しさん：2014/12/24(水) 22:50:48.00 ID:3hqu78L7.net

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

54 ：デフォルトの名無しさん：2014/12/27(土) 01:57:38.03 ID:u9BI3CqV.net

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

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

55 ：デフォルトの名無しさん：2014/12/27(土) 06:52:06.91 ID:fxVjSbuk.net

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

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

56 ：デフォルトの名無しさん：2014/12/27(土) 10:09:35.39 ID:W6Y2DM4+.net

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

57 ：デフォルトの名無しさん：2014/12/28(日) 20:39:13.97 ID:52BL0aAq.net

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

58 ：デフォルトの名無しさん：2014/12/29(月) 02:19:31.23 ID:YUQudPNs.net

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

59 ：デフォルトの名無しさん：2014/12/29(月) 09:47:06.28 ID:Sx0YYE+e.net

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

60 ：デフォルトの名無しさん：2014/12/29(月) 09:51:11.07 ID:Sx0YYE+e.net

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

61 ：デフォルトの名無しさん：2014/12/29(月) 12:03:33.25 ID:oV4aoJAy.net

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

62 ：デフォルトの名無しさん：2014/12/29(月) 12:51:57.86 ID:oV4aoJAy.net

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

63 ：デフォルトの名無しさん：2014/12/31(水) 21:08:03.80 ID:3b0Wn462.net

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

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

64 ：デフォルトの名無しさん：2014/12/31(水) 23:07:39.94 ID:WrP28EMy.net

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

65 ：デフォルトの名無しさん：2015/01/01(木) 12:48:51.81 ID:82JnHkZd.net

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

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

66 ：デフォルトの名無しさん：2015/01/02(金) 01:42:39.20 ID:aooXGYY5.net

>>65
なるほどね

67 ：デフォルトの名無しさん：2015/01/02(金) 10:40:26.85 ID:2fT8SJ1Ez

>>65
「ストリームの使い道は異なるカーネルの並列実行」というのは
何が言いたいのか良くわからないが、捕捉すると、

ストリームはCPU-GPU間のデータ転送中（非同期のcudaMemcpy）に、GPUで演算（カーネル実行）させる為の仕組みだ。

「あるストリームがデータ転送待ち中に、別のストリームのカーネルを実行する。」
といった事を、賢くやってくれる。
（別のストリームが「異なるカーネル」でなく、「同じ関数のカーネル」でも良い。）
結果、「データ転送待ち」の時間を節約できる。

ストリームの詳細は、書籍に書いてあるが、
データ転送とカーネル実行をストリームに入れる「推奨の順番」が
cc3.5からは違うので、古い本の場合は読み替える必要あり。

68 ：67：2015/01/02(金) 11:01:01.35 ID:2fT8SJ1Ez

まとめると、こんな感じ。
（※同じプログラムで、以下の両方を使う事も可能。）

・cudaコアの数以上にスレッドを生成する：「GPU」と「GPU側のDRAM」間のメモリアクセスの遅延の隠蔽が期待できる。

・ストリームを使う：「CPU側のDRAM」と「GPU側のDRAM」間のメモリアクセスの時間の節約が期待できる。

69 ：デフォルトの名無しさん：2015/01/03(土) 13:07:05.02 ID:yWVdPt25.net

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

70 ：デフォルトの名無しさん：2015/01/09(金) 00:33:57.02 ID:Iq4Pw+IC.net

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

71 ：デフォルトの名無しさん：2015/01/12(月) 00:15:12.48 ID:crrCnhEj.net

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

72 ：デフォルトの名無しさん：2015/01/16(金) 07:53:04.52 ID:IHSf0jGJ.net

CUDA7.0 RC

73 ：デフォルトの名無しさん：2015/01/16(金) 09:02:56.43 ID:VQ2eHsT0.net

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

74 ：デフォルトの名無しさん：2015/01/25(日) 10:58:33.86 ID:m2kue9j8.net

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

75 ：デフォルトの名無しさん：2015/02/06(金) 21:35:53.90 ID:72/Q/UeS.net

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

76 ：デフォルトの名無しさん：2015/02/07(土) 00:48:46.19 ID:OS4q1AxS.net

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

77 ：デフォルトの名無しさん：2015/02/07(土) 13:10:14.61 ID:4cvxubK6.net

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

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

78 ：デフォルトの名無しさん：2015/02/08(日) 15:23:01.85 ID:E04CIgi2.net

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

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

79 ：デフォルトの名無しさん：2015/02/08(日) 15:48:00.96 ID:E04CIgi2.net

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

80 ：デフォルトの名無しさん：2015/02/08(日) 21:32:20.44 ID:BpjOkBmf.net

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

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

81 ：デフォルトの名無しさん：2015/02/08(日) 22:56:08.44 ID:KLuvC02r.net

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

82 ：デフォルトの名無しさん：2015/02/09(月) 10:20:08.08 ID:pN+UjOmC.net

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

83 ：デフォルトの名無しさん：2015/02/09(月) 23:18:25.64 ID:QR2S1do8.net

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

84 ：デフォルトの名無しさん：2015/02/17(火) 21:40:32.15 ID:K8c74Rhe.net

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

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

85 ：undefined：2015/02/19(木) 11:23:20.97 ID:aqLRWkl1.net

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

86 ：デフォルトの名無しさん：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毎の仕様の違いも結構重要だと思う。

87 ：デフォルトの名無しさん：2015/02/19(木) 23:45:36.57 ID:ngPIgbTR.net

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

88 ：デフォルトの名無しさん：2015/02/19(木) 23:54:19.28 ID:Lt8lBsrZ.net

>>86
有難う。750TiのFP64悪すぎだな。
なんか大衆向け用でGPGPUするならFP32よ、FP64は使わないでだな。

89 ：デフォルトの名無しさん：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
それは差別化とか、グラフィック用途でのワットパフォーマンスとかで仕方がないかと。

90 ：デフォルトの名無しさん：2015/02/20(金) 08:52:57.87 ID:xG3c1huj.net

keplerはinteger bit shiftが弱いGK110(tesla)以外は
maxwellはkepler比で2倍のスループットになってる

91 ：デフォルトの名無しさん：2015/03/04(水) 22:02:58.94 ID:krHDLIbc3

コアレッシングて、なんでハーフワープじゃないとダメなのでありますか？
フルじゃダメなのでしょうか。
#deviceQuery:
# CUDA Capability Major/Minor version number: 3.2
# Warp size: 32

92 ：デフォルトの名無しさん：2015/02/22(日) 18:45:23.46 ID:JhGx5uct.net

適当なプログラム作ってみても
maxwellの方がはやいね

shared memoryが倍になったのも大きいなぁ

93 ：名無し：2015/03/07(土) 10:52:38.85 ID:UBzBpgz5.net

スレチなら申し訳ない
当方、モバイルでCUDAを使用したいけど
安い方法はどれが良いと思います？
（速度はそこそこで良く、外でテストして
　パワーがいる場合はデスクトップを使用するつもり）

モバイル用は安く上げたいので
Chromebookかタブレットで探した方が良いですかね？
奇をてらってJetsonのtk1にACアダプタ用のバッテリーを積むとか
（可能かどうかわからないですが）
ちなみにゲームはやるつもりありません。

94 ：デフォルトの名無しさん：2015/03/07(土) 20:42:36.10 ID:CzdLWIdo.net

thinkpad w550sのquadro K620mはダメなの？
ノートパソコンだよ

95 ：デフォルトの名無しさん：2015/03/08(日) 10:25:35.13 ID:TYY6zzsE.net

レスありがとうございます
安くあげたいので予算的に厳しいかと
最初だけ計算量は多いですけど
要所ごとに定数化すれば、その後は計算量がへるかと思っているので外での使用は少ないデータ量でプログラムチェックができればいいかなと考えています。
K1がのったタブレットも安いのでそこから考えてみようと思います
ありがとうございました

96 ：デフォルトの名無しさん：2015/03/18(水) 02:36:15.45 ID:jKTvW/7W.net

うわ、titan xの倍精度、しょぼ過ぎ・・・。
https://twitter.com/search?q=titan%20x&src=typd

97 ：デフォルトの名無しさん：2015/03/18(水) 09:22:22.19 ID:33RtPIwm.net

単精度が7TFLOPS、倍精度が0.2TFLOPSで良いんだよな？？

詳しい事は良く分からないんだけど倍精度ってそんなに使わないものなの？
PhysXっていう物理エンジン使ったりするのに

98 ：デフォルトの名無しさん：2015/03/18(水) 10:23:22.30 ID:C11qPS4w.net

>>97
ゲームだと単精度で十分だったりするからじゃない？
もともとゲームのために作られたような物だから、倍精度の性能なんかあまりこだわってないと思う

99 ：デフォルトの名無しさん：2015/03/19(木) 07:37:41.52 ID:M0RYJxHF.net

https://developer.nvidia.com/cuda-downloads
CUDA 7 Downloads

100 ：デフォルトの名無しさん：2015/03/19(木) 07:39:40.96 ID:M0RYJxHF.net

Dear Developer,

The CUDA? 7.0 Production Release is now available to the public. Run your application faster with this latest version of the CUDA Toolkit. It features 64-bit ARM support and the simplified programming model enabled by Unified Memory. Highlights include:

New cuSOLVER library
? Accelerates key LAPACK routines, 12x faster direct sparse solvers

New C++11 language features
? Increases productivity with lambdas, auto, and more

Runtime Compilation
? Enables highly optimized kernels to be generated at runtime

Download the CUDA 7 Production Release at www.nvidia.com/getcuda

Learn more about CUDA 7 by attending these webinars:

CUDA 7 Feature Review
Date/Time: Friday, April 10th at 10:00 AM PDT
Register: https://cc.readytalk.com/r/4b0lwgeqgzrk&eom

CUDA 7 Performance Overview

Date/Time: Wednesday, April 15th at 11:30 AM PDT
Register: https://cc.readytalk.com/r/empyu1qc65l6&eom

Best regards,

Nadeem Mohammad
NVIDIA Developer Relations Team