5ちゃんねる ★スマホ版★ ■掲示板に戻る■ 全部 1- 最新50  

■ このスレッドは過去ログ倉庫に格納されています

【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/

2 :デフォルトの名無しさん:2011/08/23(火) 22:11:49.24
関連サイト
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://gpgpu.jp/article/61432191.html

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 :デフォルトの名無しさん:2011/08/23(火) 22:12:38.55
宣伝になっちゃうと難なので、一応分けて。いろいろ載っているサイト。

ttp://www.softek.co.jp/SPG/Pgi/TIPS/para_guide.html

4 :デフォルトの名無しさん:2011/08/23(火) 22:26:00.72
おい、pertだろ
まちがえんな

5 :デフォルトの名無しさん:2011/08/23(火) 22:48:26.89
>>2
個人blogのは情報古すぎなような…
「cudaさわってみた」はリンク切れだし。

6 :デフォルトの名無しさん:2011/08/25(木) 19:36:34.86
>>4
やっぱりCUDAやってる人は荒んでるんだな

7 :デフォルトの名無しさん:2011/08/26(金) 07:04:46.69
ほんとCUDAらないことで騒ぐよな

8 : 忍法帖【Lv=40,xxxPT】 :2011/08/26(金) 14:04:14.48
ケンカするのはやめてCUDAさい><

9 :デフォルトの名無しさん:2011/08/26(金) 20:59:54.29
そうだよ。喧嘩はやめ…GeForんゲフォン

10 :デフォルトの名無しさん:2011/08/26(金) 21:57:35.41
まだスレが始まっTeslaいないのに

11 :デフォルトの名無しさん:2011/08/27(土) 00:03:43.41
ときに口論になるのもみんな自分の技術にそれだKeplerイドを持ってるってことだ。

12 :デフォルトの名無しさん:2011/08/27(土) 00:09:15.96
amazon.comで検索して出てくるCUDA本ってことごとく出版がずいぶん先か延期になったりするな
なんでだろ

13 :デフォルトの名無しさん:2011/08/27(土) 21:26:38.44
無効とは情報の登録ポリシーが違う…ような気がする。
それがAmazonのポリシーなのか出版社のポリシーなのかは知らないけれど。

多分出版社のサイト直接見た方が正確。


14 :デフォルトの名無しさん:2011/08/27(土) 23:28:49.51
cudaの本が必要、って訳でも無いが出版されるなら買ってしまうと思う

15 :デフォルトの名無しさん:2011/08/31(水) 20:00:08.40
4.0になってdeviceemuが使えんのだが、CUDA載せてないPCで動かす方法ない?

16 :デフォルトの名無しさん:2011/09/01(木) 07:46:58.96
3.2を使う。

17 :デフォルトの名無しさん:2011/09/01(木) 09:05:53.29
>>16 アホしかおらんのか

18 :デフォルトの名無しさん:2011/09/01(木) 09:16:16.30
>>17
現実的じゃないか。どうせemuで正確なエミュレーションができるわけじゃないし。

19 :デフォルトの名無しさん:2011/09/01(木) 09:25:26.25
>>17
4.0でdeviceemu使う方法なんてねーよ、ってことだよ。
言わせんな恥ずかしい

20 :デフォルトの名無しさん:2011/09/01(木) 10:02:13.87
>>15
Cuda x86

21 :デフォルトの名無しさん:2011/09/01(木) 15:33:21.69
理屈の上ではKnights Ferryでも動かせるのかな?

22 :デフォルトの名無しさん:2011/09/01(木) 15:39:30.39
ubuntu 11.04 tesla c2050 cuda sdk 4.0の環境でdoubleのアトミック演算が正しく動作しません
オプションは -arch=compute_20を指定してます。

本家のプログラミングガイドに書いてある、以下のコードを実行してます。
__device__ double atomicAdd(double* address, double val)
{
unsigned long long int* address_as_ull
                  = (unsigned long long int*)address;
unsigned long long int old = *address_as_ull, assumed;
do {
assumed = old;
old = atomicCAS(address_as_ull, assumed,
             __double_as_longlong(val + __longlong_as_double(assumed)));
} while (assumed != old);
return __longlong_as_double(old);
}

引数のadrees*には共有メモリのアドレスを渡してます。
一つのスレッドのみが共有メモリにアクセスする場合は正しく加算は行われますが、二つ以上のスレッドが衝突する場合は
正しく動作しません。printfで見てみたら、二つ目のスレッドがループを抜けれずに何度も共有メモリの値に
加算しているようでした。上の文をdouble→float,long long → int に書き換えた場合は正しく動作しました。
また、atomicCASのところを二つのprintf、
printf("%d %f %f\n",threadIdx.x,__longlong_as_double(assumed) ,__longlong_as_double(old));
ではさむとなぜか正しくループを抜け、正しくsharedメモリに値を書き込みます。

原因がわかる方、教えていただけないでしょうか。

23 :デフォルトの名無しさん:2011/09/01(木) 17:40:47.67
>>22の続きです
assumed oldの値をcpu側に持ってきて表示してみたら、
oldの値(atomicCASの返す値)とassumedの値が末尾の数ビットが
異なってました。50回ほどループしてやっとビットレベルで値が一致して
ループを抜けているようでした。その間に過剰に加算が行われていたようです。

24 :デフォルトの名無しさん:2011/09/01(木) 18:47:40.63
カードの計算が正しいかチェックするベンチマークってありませんか?

四則演算と√と三角関数とかを全ビットを演算器全部チェックしろと云う無茶ぶりがきた。

計算機が計算を間違うなんてありえないのに…


25 :デフォルトの名無しさん:2011/09/01(木) 19:11:55.88
Teslaの存在意義全否定デスカ

26 :デフォルトの名無しさん:2011/09/01(木) 19:34:24.88
>>24
過去スレを読むとわかってもらえると思うが、みんなそれを求めているんだ。
そのチェックにおける最先端はおそらく長崎大学。

27 :デフォルトの名無しさん:2011/09/01(木) 19:41:57.01
1スレッドの実行がどのコアかがわかる仕組みがどのような方法か分かればいいわけだよな
誰か調べてみろよ

28 :デフォルトの名無しさん:2011/09/01(木) 23:41:13.82
>>22,23
もしかしてブロック辺りのスレッド数が32以下なら動いたりする?


29 :デフォルトの名無しさん:2011/09/01(木) 23:44:25.02
>>27
何がしたいのか分からんがそれは簡単だ。

30 :デフォルトの名無しさん:2011/09/01(木) 23:56:07.27
>>29
コアと計算結果の表が出来上がればコア間の相違が分かるじゃん
やり方知ってるなら教えてください

31 :デフォルトの名無しさん:2011/09/01(木) 23:58:10.21
煽っても教えてやらん
29じゃないけどw

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 :デフォルトの名無しさん:2011/09/02(金) 02:29:09.05
ここはおまえの日記帳だから好きに使いなさい

34 :デフォルトの名無しさん:2011/09/02(金) 11:08:38.52
>>28
ブロックあたりのスレッド数を32以下にして試してみましたが、
結果は変わらなかったです。

35 :デフォルトの名無しさん:2011/09/02(金) 18:20:08.65
>>23
とりあえず、プログラミングガイドのコードそのままだし、
実際アルゴリズム自体には問題ないから、
意図どおり動かないとしたら、ドライバかSDKのバグか
ハードの故障じゃね。
まあ、ここで詳細なバージョンとか晒してみたり、
PTXやアセンブリを貼ってみるのもいいけど、
本家のフォーラムで聞いてみたほうが早いんじゃね。

36 :デフォルトの名無しさん:2011/09/05(月) 23:02:56.52
長崎大の言う演算間違えって、具体的にどういう演算を間違えたのかの情報ってないの?
間違えますって情報だけ垂れ流して具体例出さんって、どういうことなんだべ。

37 :デフォルトの名無しさん:2011/09/05(月) 23:24:45.84
×演算間違え
○演算間違い

38 :デフォルトの名無しさん:2011/09/07(水) 18:20:00.85
CUDAを使用するプログラムで巨大なデータを扱いたいのですが、nvcc でのコンパイル時に
relocation truncated to fit: R_X86_64_32S against `.bss'
とエラーが出て完了しません。gcc ではデータ量制限を解除するのに-mcmodel=medium というオプションがあるようですが、nvcc でこれにあたるものはないでしょうか。

39 :デフォルトの名無しさん:2011/09/07(水) 19:43:00.72
巨大データってどれくらい?
CUDAは基本グローバルメモリ使用の制約があるから”巨大”は決して得意ではないのだけど

40 :38:2011/09/07(水) 22:05:04.34
>>39
今確認したところ、要素が数十万個の配列を使用していて、それらの合計は200MBほどです。
そのうち、実際にGPUに計算させるためGPUに渡すデータは50MB程度です。
上記のエラーメッセージでググったところ、データ量制限が問題だろうということだったのですが…それほど大きなデータでもないですよね。
なんなんでしょうこのエラー。

プログラミングに関してはまだ知識が浅いためチンプンカンプンなこと言ってたらすみません。

41 :デフォルトの名無しさん:2011/09/07(水) 22:36:52.16
どうせデータをソース中に書き込んでるんだろw

42 :デフォルトの名無しさん:2011/09/07(水) 22:42:26.34
>37
日本語の間違えより、長崎大の言ってる間違いの詳しい情報くれよぅ…。

43 :デフォルトの名無しさん:2011/09/07(水) 22:50:33.71
要素数が大きすぎるなら、mallocを使ってみたら?

44 :デフォルトの名無しさん:2011/09/07(水) 23:04:45.67
>40
配列は何次元?
CやC++だと配列は連続したメモリ領域と規定されてるので、
確保可能なメモリ量は物理メモリの絶対量ではなく、
アドレスを連続で確保できるメモリ量に規定されるよ。

例えば大きめの2次元配列を確保する場合は
double* pBuf = new double[row * column];

とするより
double** ppBuf = new (double*)[row];
for(int i = 0; i< row; ++i){
   ppBuf[i] = new double[column];
}

とかにしたほうがいいよ。

45 :デフォルトの名無しさん:2011/09/07(水) 23:11:07.95
ファイル読み込みで、cudaMallocHost(だったかな??)でメモリ確保すればいいじゃん

46 :38:2011/09/07(水) 23:11:33.99
>>43>>44
gccと同様にオプションでエラー回避できないかな…と思ったのですが、配列の取り方を変えるほうが良さそうですね。
特に>>44は知りませんでした。
ありがとうございました!

47 :デフォルトの名無しさん:2011/09/07(水) 23:23:20.08
良く分からんが、-Xcompiler -mcmodel=mediumじゃ駄目なん?
-Xlinkerも要るかもしれないが。

48 :デフォルトの名無しさん:2011/09/08(木) 08:01:27.80
>>44
普通はありえない。そんなことしたら、ポインタアクセスが一回余計に要るから遅くなるし、
キャッシュアウトしやすくなるから更に遅くなる。
まぁ、GPUに転送するのは1スライス分だけだろうからGPUに任せてCPUは暇こいていられるならいいかも知れんが。

49 :デフォルトの名無しさん:2011/09/08(木) 17:24:44.67
>>44
それは無い。
・メモリも無駄に消費
・アクセス性能がた落ち
・バグの温床になる

つーかC++ならvectorで初期化時に
サイズ指定したほうがいろいろ便利



50 :デフォルトの名無しさん:2011/09/08(木) 17:40:17.98
配列でやる場合、
配列の添え字がINT_MAXを越える大きさならそうするしかないと思うが・・


51 :デフォルトの名無しさん:2011/09/08(木) 20:47:04.41
キャッシャがアクセスがとドヤ顔されても・・・。
要素のデカイ配列が確保できんって話なわけだし。

52 :デフォルトの名無しさん:2011/09/08(木) 22:19:09.58
コンパイル時の話なのに
ヒープのフラグメンテーションの訳無いだろう。
nvccからgccに自由にオプション渡すやり方は
知らんが、最悪c(++)部分のコードを分割して
gccでコンパイルすりゃ良い。

仮にヒープのフラグメンテーションが酷い場合でも
>>44みたいなやり方は下策だよ。
素人におかしなやり方を覚えさせるだけ。

大きな配列確保する前に小さな物についてShrink-to-fitやればデフラグされるから
その後巨大配列を確保すれば良い。

特にGPGPUで巨大な配列扱うような用途なら
配列の配列なんて転送も困るだろうに。


53 :デフォルトの名無しさん:2011/09/08(木) 22:30:02.29
>大きな配列確保する前に小さな物についてShrink-to-fitやればデフラグされるから
>その後巨大配列を確保すれば良い。

コレ具体的にどんなコード書けばいいのん?

54 :デフォルトの名無しさん:2011/09/09(金) 00:35:55.99
メインの計算に必要な必要な小容量配列
保持用のポインタなりvectorの変数だけ始めに
ヒープに確保した後、メインの計算前の前処理をする。
この処理で、一時的なデータと最終的に使うデータの
ヒープ確保が入り乱れる所為でフラグメンテーションが発生する。

最終的に使うデータについて、大体処理した順番を守って
ひたすら同容量確保してはコピーを繰り返せば、
こういうデータは全て前方に圧縮されて配置されるようになる。


55 :デフォルトの名無しさん:2011/09/09(金) 21:48:41.98

まあスレ違いなんで他でやってよ。

56 :デフォルトの名無しさん:2011/09/09(金) 22:24:20.88
kernelへの引数に構造体は使えんの?
今やってみたらコンパイルでエラーになった

57 :デフォルトの名無しさん:2011/09/09(金) 22:30:19.36
あ、間違えてただ
構造体引数使えるべな

58 :デフォルトの名無しさん:2011/09/14(水) 11:12:56.96
>>15
http://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 :デフォルトの名無しさん:2011/09/15(木) 17:52:31.49
cuda4.0、gtx580で
Maximum dimension of a grid( 65535 , 65535 , 65535 )
ってなってるんですがz軸のブロックの取れる数は1じゃなくなったんですか?

60 :デフォルトの名無しさん:2011/09/19(月) 22:05:37.25
デバイス上の処理が長時間になる場合、途中経過を取得する方法ってないかな?
例えば100万回のモンテカルロを実行して、途中の10万回、20万回が済んだらその情報を取りたいんだけど…

61 :デフォルトの名無しさん:2011/09/19(月) 22:31:31.07
>>60
streamingでできないかな?

62 :デフォルトの名無しさん:2011/09/19(月) 22:52:18.06
streamingってなに?
SMのSじゃないよね?

63 :デフォルトの名無しさん:2011/09/20(火) 05:00:39.03
カーネル呼び出し時の<<<>>>の第3引数がstreamingの何かじゃなかったっけ?
俺は>>61を見てそういうのがあったなと思い出したレベルなので使い方は知らないけど。

64 :デフォルトの名無しさん:2011/09/20(火) 08:29:52.67
第3パラメータはsharedMemoryのサイズだよ。そこで指定しなくても使えるけど。

65 :デフォルトの名無しさん:2011/09/21(水) 22:11:25.09
ごめん第4引数だったみたい。それとAsync系関数がうんたらかんたら。

66 :デフォルトの名無しさん:2011/09/22(木) 00:50:42.07
あー、その辺だね。

67 :デフォルトの名無しさん:2011/09/22(木) 18:05:24.11
【トリップ検索】CUDA SHA-1 Tripper【GeForce】
http://hibari.2ch.net/test/read.cgi/software/1311428038/

すげえな

68 :デフォルトの名無しさん:2011/09/23(金) 11:16:26.70
個人でCUDAプログラミングやってうれしいのはその程度なのか?
と逆に不安になるんだが…

69 :やんやん ◆yanyan72E. :2011/09/23(金) 11:35:33.50
世の中なんでも膨大なデータをいかに早く処理するかが勝負の世界に
なってきているから、CUDAプログラミングの旨味はいくらでもあるだろ。

70 :デフォルトの名無しさん:2011/09/23(金) 12:04:06.11
いつかやってみようと思って半年くらいこのスレROMってるが
いまだに開発環境のセッティングの仕方すら調べていない・・・

71 :デフォルトの名無しさん:2011/09/24(土) 00:02:19.96
>>68
>>67くらいのコード書ける?

72 :デフォルトの名無しさん: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. :2011/10/04(火) 21:13:10.02
-archコマンドってなんだそら?
compute capabilityはnvccなりcl.exeのオプションに渡すものであって、
-archコマンドなんてものはないよ。

77 :デフォルトの名無しさん: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 :デフォルトの名無しさん:2011/10/04(火) 23:39:11.06
CUDA以前の話だな。コンパイラオプションの意味がわかってない。

というか、VCなら、カスタムビルド規則でCUDA Runtime API Build Ruleを追加したほうが早い

81 : 忍法帖【Lv=40,xxxPT】 :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 :デフォルトの名無しさん:2011/10/05(水) 12:14:48.16
-arch=compute_20
にすればいいんじゃねえの?

85 :デフォルトの名無しさん:2011/10/05(水) 12:17:32.55
>>84
ありがとうございます。
おっしゃるように>>75のように追加するとエラーを吐きます。
追加の仕方が違うのでしょうか。

86 :80:2011/10/05(水) 13:36:15.15
VS2010は色々変わってるんだな。
知らんかったすまん。
とりあえず「CUDA VS2010」でググって出てくる一番上のサイト参考にして、CUDA4.0のビルドカスタマイズファイルを導入した後、プロパティからCode Generationをcompute_20,sm_20にしたらいけた

87 :デフォルトの名無しさん:2011/10/05(水) 15:09:51.57
GPUのリソースを全部CUDAに割り当てるため,グラボを2枚さしてディスプレイ出力用とCUDA用に分けるというのはありますが
同じようなことをsandy bridgeとグラボ1枚でも出来るのでしょうか?

88 :デフォルトの名無しさん:2011/10/05(水) 15:36:59.85
>>87
Z68マザーなら、マザー側DVIにディスプレイつないでおけばいい話じゃ?
CUDAで計算したデータをintelで表示したいなら、
1.DirectXかOpenGLのリソースをCUDAにバインドしてデータ書き込み
2.メインメモリに転送
3.intel側に転送
って形になるかなぁ。多分。

89 :81: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: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 :デフォルトの名無しさん: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 :デフォルトの名無しさん:2011/10/06(木) 18:43:04.20
そこだけ書かれても分からないと思う。
nvccのエラーメッセージ載せないと。

94 :デフォルトの名無しさん:2011/10/06(木) 19:11:36.83
>>92
敵性語なんでよくわかんなかったけど、
VisualStudio を管理者として実行したら動く?

95 :やんやん ◆yanyan72E. :2011/10/06(木) 21:18:36.98
Visual C studioは下手にGUIに頼るよりも、コマンドラインでやった方がはやいこともある。コマンドラインでnvccやってみよう。

96 :デフォルトの名無しさん:2011/10/07(金) 16:37:49.21
単精度のmadってmulより遅いやないかー
プログラミングガイド嘘つきやないかー
ちくしょー

97 :デフォルトの名無しさん:2011/10/09(日) 01:36:18.51
>>96
そりゃ命令実行時間は同じだけど、
必要な帯域は違う。

98 :デフォルトの名無しさん:2011/10/09(日) 02:40:58.84
>>92
コマンドラインのオプションが何か1つ間違ってた気がする。
コマンドプロンプトで同じコマンド実行するとエラーが出たので、
そのエラーのひとつ前の引数が問題だった、はず。

99 :デフォルトの名無しさん:2011/10/10(月) 02:14:56.87
CUDAとRTミドルウェアを組み合わせてみた。

行った方法は、CUDAのサンプルをRTコンポーネントに
パッケージングする方法。

動作周期にあわせて、行列計算が実行されたのを確認。



100 :デフォルトの名無しさん:2011/10/10(月) 20:49:05.67
カーネル側のソース分割ってできないんだっけ?
環境はLinux cuda4.0

101 :デフォルトの名無しさん:2011/10/10(月) 22:00:17.43
できるけどコンスタントメモリが変な事になったような気がする

102 :デフォルトの名無しさん:2011/10/10(月) 22:04:05.94
このままだと5万行くらいになっちゃいそう・・・

103 :デフォルトの名無しさん:2011/10/10(月) 22:39:08.05
>>102
正直処理対象がCUDAに向いてないような。。。

104 :デフォルトの名無しさん:2011/10/14(金) 02:20:30.66
とある文献でレイテンシを表す際, 種類によってclock, cycleが使い分けられていました。

clock : texture cache, constant cache
cycle : global memory, shared memory

なんとなくメモリ, キャッシュで使い分けられているのはわかりますが, このままでは速度の比較ができません。
二つの違いは何なのか, またcycle-clock間で単位の変換が可能なのか, もしご存知の方がいたらご教授ください。

105 :デフォルトの名無しさん:2011/10/17(月) 13:50:45.89
カーネル関数内で立てたフラグをホスト側で判断する事って出来る?

106 :デフォルトの名無しさん:2011/10/17(月) 14:13:41.38
グローバルメモリ介してできるだろ

107 :デフォルトの名無しさん:2011/10/18(火) 02:35:47.28
はじめてのCUDAプログラミングの9章に載ってるcublas使うプログラム実行したんだけど、
cublasInit()あたりで「〜0x7c812afb で初回の例外が発生しました: Microsoft C++ の
例外: cudaError_enum〜」って出てうまく動かないんだけどなして?

108 :デフォルトの名無しさん: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 :デフォルトの名無しさん:2011/10/19(水) 11:56:08.38
適当なポインタで確保した後に,
そのアドレス転送してやればいいんじゃない?

110 :デフォルトの名無しさん: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: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 :デフォルトの名無しさん:2011/10/19(水) 21:10:14.17
概ねいいんでない?
前者なら、私はhoge_t hoge;で定義しちゃうけど。

113 :109:2011/10/19(水) 22:04:08.49
俺が考えてたのは後者だな。
合ってると思うよ。
でもmemcpyのsizeof(int)はsizeof(int *)にした方が良いと思う。


114 :108:2011/10/19(水) 23:57:45.47
未だにポインタが使いこなせてないと痛感しました。
このたびはアドバイス有難うございました。

115 :デフォルトの名無しさん:2011/10/20(木) 00:12:23.07
ポインタ使える男の人ってモテそうだよね

116 :デフォルトの名無しさん:2011/10/20(木) 17:52:20.01
ポインタでいろんな所を指して、
覗いたりいじくり回したりするんですね。

117 :やんやん ◆yanyan72E. :2011/10/21(金) 01:37:52.34
トラックポイントではない

118 :デフォルトの名無しさん:2011/11/03(木) 20:05:27.91
デバイスから出る「the launch timed out and was terminated.」とかのメッセージを
CPU側で受信する方法ってどうやるのかな?コンソール出力されるんだからできると思うけど

119 :デフォルトの名無しさん: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 :デフォルトの名無しさん:2011/11/07(月) 09:05:53.46
CG法と書いてあるので問題行列は対称疎行列でCUSPARSEからはcsrmv関数だけ、CUBLASからは
Level-1 BLASルーチンしか使っていないのだと思いますが、

実際にどの程度遅いのでしょうか?

ランダムアクセスを含むcsrmvはLevel-1 BLASルーチンとくらべてメモリアクセス効率が極端に下がるのは
仕様というか当然の結果です。

また古いバージョンのライブラリを使っていたりしませんか?

handle云々は基本的に関係ありません。関係するとすればよほど小さな行列を扱っている場合に限られます。



122 :デフォルトの名無しさん: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 :デフォルトの名無しさん: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 :デフォルトの名無しさん:2011/11/07(月) 18:52:46.57
構造体にポインタがあるのはダメ
理屈で考えればわかるよな

125 :デフォルトの名無しさん:2011/11/07(月) 19:13:06.14
>>124
そうでしたか
わかりました


126 :デフォルトの名無しさん:2011/11/07(月) 19:34:23.39
馬鹿でプログラミング初心者には難しいれす(^ρ^)

127 :やんやん ◆yanyan72E. :2011/11/07(月) 19:39:20.38
無駄にでっかくメモリ確保された
ただのポインタがコピーされてるだけに見えるのは気のせい?
実体をコピーしようとしている意図は読み取れるけれど、
コードはその意図を反映してないような。。。

128 :やんやん ◆yanyan72E. :2011/11/07(月) 19:50:21.46
あ、考え違いか、すまん

129 :デフォルトの名無しさん:2011/11/07(月) 19:52:32.74
>>124
わからないのでその理屈を教えていただけないでしょうか

130 :デフォルトの名無しさん:2011/11/07(月) 20:15:14.16
cudaMallocHostはあるけどcudaMallocはしてないの?書き洩れ?

131 :デフォルトの名無しさん:2011/11/07(月) 20:28:54.09
デバイス側のcudamallocが無いよね?

cudaMallocHost((void**)&vvv_d->vari,sizeof(int)*(N*N));
ってvvv_d->variをホスト側にメモリ確保しちゃってんるんだけど
必要なことは例えば
cudaMallo(reinterpret_cast<void**>(&vvv_d->vari,sizeof(int)*(N*N));
じゃないの?

132 :デフォルトの名無しさん:2011/11/08(火) 09:53:28.63
取り敢えず、ホスト側とデバイス側の二種類のメモリ空間を扱っていることを肝に銘じよう。
「構造体にポインタがある」ある場合、そのポインタがホスト側のメモリを指しているのか
デバイス側のメモリを指しているのか常に意識しないといけない。

それがいやだったら、ポインタではなく配列にしてしまえ。
cf.
struct var_str {
int var[N * N];
float var_g;
struct var_box vb[N];
};

133 :デフォルトの名無しさん:2011/11/08(火) 09:54:43.06
それ以前に、そんなデータ構造を渡しても素直に並列化できなくて速度が出ない悪寒。

134 :デフォルトの名無しさん:2011/11/08(火) 10:50:48.01
>>129
>>108-114の流れを読め.

135 :デフォルトの名無しさん:2011/11/08(火) 13:04:40.64
C# (.NET Framework 4.0)で動作するCUDAのライブラリってありますか?
CUDA.NETがそうだとは思うのですが、.NET Framework 2.0 or newerとしか書いていないので、
結局のところ.NET 4.0で動作するのかどうかわかりません。

136 :デフォルトの名無しさん:2011/11/08(火) 15:31:03.58
おまえ、そのレベルでよくこのスレに辿り着けたな。。

137 :デフォルトの名無しさん:2011/11/08(火) 18:02:17.87
とりあえず動かしてみろw

138 :デフォルトの名無しさん:2011/11/09(水) 04:38:11.56
4.1RC出てるんだな
デベロッパプログラムに登録したいけど審査あるから、
他人のソースコンパイルして遊ぶだけの俺には無理

139 :デフォルトの名無しさん:2011/11/09(水) 14:48:59.26
試しにデタラメ並べ立てて申請したら通った
んでSDKのサンプルを上から順番に走らせて遊んでたらブルスク食らった
危ねえ・・・

140 :デフォルトの名無しさん:2011/11/09(水) 22:34:20.87
>>138
審査あんの?俺普通に登録してOKでたけど

141 :デフォルトの名無しさん:2011/11/09(水) 22:45:40.52
審査ってなにを調べるの?

142 :デフォルトの名無しさん:2011/11/09(水) 23:51:20.06
ウンコの重さとか

143 :デフォルトの名無しさん:2011/11/10(木) 00:31:38.44
Parallel Nsightのアカウントとは別なのか。
面倒くさすぎだろ。

144 :デフォルトの名無しさん:2011/11/10(木) 06:20:27.98
CUDA.NETってCUDA 3.0までで止まってるよね?

145 :デフォルトの名無しさん:2011/11/10(木) 12:49:08.28
>>144
くだんねーと(CUDANET)思ったんだろ作ってる方も

146 :デフォルトの名無しさん:2011/11/11(金) 00:24:14.22
【審議中】
    ∧,,∧  ∧,,∧
 ∧ (´・ω・) (・ω・`) ∧∧
( ´・ω) U) ( つと ノ(ω・` )
| U (  ´・) (・`  ) と ノ
 u-u (l    ) (   ノu-u
     `u-u'. `u-u'

147 :デフォルトの名無しさん:2011/11/14(月) 02:32:16.94
kernelの中でどの部分でレジスタがMAXになっているか、どの値を保持しているかっていうのは簡単にわかるものですか?

148 :デフォルトの名無しさん:2011/11/14(月) 11:50:00.02
細かいことは兎も角、取り敢えずptx出力を見てみたら医院で内科医。

149 :デフォルトの名無しさん:2011/11/14(月) 18:06:51.94
shrQATest.hって4.0のtoolkitとSDK入れたらどこにあるはずですか?

150 :デフォルトの名無しさん:2011/11/15(火) 02:43:04.47
%ProgramData%\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.0\shared

151 :デフォルトの名無しさん:2011/11/15(火) 02:44:06.11
\inc

152 :デフォルトの名無しさん:2011/11/15(火) 16:34:29.61
見つかりました
ありがとうございます

153 :デフォルトの名無しさん:2011/11/15(火) 21:47:31.35
linux kernelからGPGPU使いたいんだけどKGPU以外に何かいい方法ある?

154 :デフォルトの名無しさん:2011/11/15(火) 22:55:00.82
CUDA developer用としてCUDAサイトで配布されてるドライバと普通のドライバってどんな違いがあるの?
CUDAに必須と思ったら、普通のドライバインストールしてもCUDA開発できると書いている所があった。

155 :デフォルトの名無しさん:2011/11/15(火) 23:18:16.01
一時期、developer用と一般用のバージョンが同じ時があったけど
md5sumは一致した。
おそらく、一般用ドライバのバージョンで対応が間に合っていないときに
beta版を出してるだけだと思う。

だから、両者でバージョンの新しい方を使えば問題ない。

156 :デフォルトの名無しさん:2011/11/16(水) 00:34:28.85
SC2011 OpenACC Joint Press Release - main
http://www.openacc-standard.org/announcements-1/nvidiacraypgicapsunveil%E2%80%98openacc%E2%80%99programmingstandardforparallelcomputing


157 :デフォルトの名無しさん:2011/11/16(水) 05:00:22.99
>>155
了解。レスありがと

158 :デフォルトの名無しさん:2011/11/22(火) 17:29:12.79

すいません、経験者の方教えてください。

ホストメモリをページロックメモリで確保して、
ストリームを300個くらい起動して
MemcpyAnsyncでh->d ,kenrel, h->dで非同期実行させてるんですけど、
どうやらkernelでセグって停止してしまいます。
しかし、cuda-gdbで見てみるとどうもセグるようなポジションじゃないんです。

そこで質問なんですが、MemcpyAsyncでコピーリクエストをだしてコピーが始まり、
制御がカーネルに入り、コピーが完了したストリームを実行してる途中に、
まだコピー中の領域でセグった場合、どこでセグったように見えるのでしょうか?

やはりその時実行しているカーネルの中でセグっているように見えるのでしょうか?

159 :デフォルトの名無しさん:2011/11/27(日) 22:05:36.51
全スレッドの計算結果を1コアに集約して1スレッドで処理したいんだけどいい方法あるかな?
リダクションのサンプルだと最後はCPUで合計だしこれでは都合悪いんだよね
全てGPUでやりたいのよね

160 :デフォルトの名無しさん:2011/11/27(日) 23:20:19.00
しかたないじゃん、CPUで合計出した方が速いんだから。

161 :デフォルトの名無しさん:2011/11/27(日) 23:28:38.68
>>159
バイナリツリーで足していくしか。
Ex. thread数が64のとき
if (threadIdx < 32) data[threadIdx] += data[threadIdx + 32];
__syncthreads();
if (threadIdx < 16) data[threadIdx] += data[threadIdx + 16];
__syncthreads();
if (threadIdx < 8) data[threadIdx] += data[threadIdx + 8];
__syncthreads();
if (threadIdx < 4) data[threadIdx] += data[threadIdx + 4];
__syncthreads();
if (threadIdx < 2) data[threadIdx] += data[threadIdx + 2];
__syncthreads();
if (threadIdx < 1) result = data[threadIdx] + data[threadIdx + 1];

>>160
データ量によってはそうでもないよ。

162 :デフォルトの名無しさん:2011/11/27(日) 23:35:14.30
>>161
>データ量によってはそうでもないよ。

>>159は「1スレッドで処理したい」と言っている。
それに、「計算結果を集約」と言っているだけで足し算とは言っていない。
この条件の下では、データ量が多くなれば多くなるほどCPUで処理した方が速い。
データ数が極端に少なければ、転送時間の関係でGPUでしたほうが速いかもしれないが。

163 :デフォルトの名無しさん:2011/11/27(日) 23:40:05.87
あーそうか、まさか1スレッドだけ動かす方法が判らんと言う話とは思わなかった。
単純に、「CPUで合計」の代わりをやりたいのかと思ったよ。

164 :159:2011/11/27(日) 23:55:11.26
>>160-163
レスサンクス
一度集約して最少二乗法を適用してから再度マルチスレッドで処理しようとしてるんだけど
一度CPUに戻すしかないかな・・・

165 :デフォルトの名無しさん:2011/11/28(月) 00:02:09.10
最小二乗法ならマルチスレッドでできるじゃん。

166 :デフォルトの名無しさん:2011/11/29(火) 03:07:20.82
リダクションしろよ
1ブロック分しかできないってこと?

167 :デフォルトの名無しさん:2011/11/29(火) 10:48:39.03
サイズが決まっている3次元配列どうしの演算で演算回数も決まっているんだけど、
ブロック数とスレッド数を調整して演算回数ぴったりにするとコアレッシングしなくて効率悪いんだ。
こーいう時ってどゆアプローチ?

168 :デフォルトの名無しさん:2011/11/29(火) 14:42:48.05
CPUで計算させてみる

169 :デフォルトの名無しさん:2011/11/29(火) 15:26:04.48
もちろんCPUでは計算させてみてるよぉぉん(><)
1つ目の演算対象配列はギリギリL2に乗らないサイズ、2つ目の配列はVRAMにも乗らないサイズなので悩んでる。

170 :159:2011/11/29(火) 16:52:31.44
もっと良いCPUかGPU買え

171 :デフォルトの名無しさん:2011/11/30(水) 02:20:28.60
>>159
同じブロック内の複数スレッドは共有メモリで集約できるけど、
ブロックをまたがっての集約ができないってこと?
手元のコードではブロック数の一時保存のための領域を確保し
そこにブロックごとの結果を放り込むカーネルを起動し、
次にブロック数が1、スレッド数が前のカーネルのブロック数の
カーネルを起動。その中でさらに集約して結果を1つにまとめている。
カーネル2回起動よりもCPUでやらせた方が速いかもしれないが。

172 :デフォルトの名無しさん:2011/11/30(水) 04:29:51.70
http://code.google.com/p/thrust/
のリダクション使ってみれば?

中身で何やってるかはよくわからんが

173 :デフォルトの名無しさん:2011/11/30(水) 08:40:15.94
カーネルの中で関数作って足せばいいんじゃない
threadが終わったら自作関数に投げて足すとか
バリアがいるけど

174 :デフォルトの名無しさん:2011/11/30(水) 08:52:34.48
微妙に話が噛み合っていない悪寒。
つーか、>159が混乱に輪を掛けているし。

175 :159:2011/11/30(水) 10:04:46.33
各スレッドでサンプリングされたXiとYiを使って
Σ(Ai*Xi-Yi)^2を計算して最小化するパラメータのAiを求めたいんだよ

求まったAiを使ってまたGPUでマルチスレッドで処理をするから
CPUに戻したくなかっただけ

戻せばできるのは分かるけど、戻さないでできるならその方がいいと思ったわけね

176 :やんやん ◆yanyan72E. :2011/11/30(水) 10:24:24.37
なんか式がおかしい、Ai=Yi/Xi なんて答えを求めたい訳じゃなかろう?

177 :159:2011/11/30(水) 11:22:19.91
>>176
うん
正確には多項式近似の最小二乗法
めんどいので略した

178 :171:2011/11/30(水) 14:43:43.94
>>177
非線形最小二乗法ってこと?非線形最適化法しか知らないけど、目的関数が
特殊な場合にそれを上手に利用したのが非線形最小二乗って理解。
手元では対数尤度(二乗和みたいなもの)を171の手順で計算する関数を
CUDAで書いて、MCMCの中で使って推定値を求めているけど、同じことは
最小二乗法でもできると思うんだが…。

179 :159:2011/11/30(水) 14:50:41.39
>>178
カーネルの二度起動だよね?
その方法でもできると思うけど、一度で済ませたいというのがニーズ

それよかMCMCをマルチスレッドでやるってのが分からんのだが…
あれはマルコフ連鎖だからほとんどマルチスレッドで効果上がらんだろw

180 :178:2011/11/30(水) 15:21:59.61
>>179
こちらも一度の起動で済ませたかったけど方法が分からなかったので、
分かったら是非報告よろしく。
尤度の計算対象とするデータ数があまりにばかでかくて…。マルコフ連鎖は
順次計算だけど、その中の並列化は相当のメリットがあってね。

181 :やんやん ◆yanyan72E. :2011/11/30(水) 16:50:24.20
>>179
二度起動がいやなら、ここの
ttp://neuralgorithm.org/documents/global-thread-synchronization-on-cuda-ja/
globalsyncみたいなのを使って計算→集計→計算って
やらせればいいんだろうけれど、
なんか怖いなw

182 :デフォルトの名無しさん:2011/11/30(水) 16:50:29.44
CUDA Visual Plofiler への入力ファイル(cpjファイルとcsvファイル)をコンソールを使って作成したいのですが、可能でしょうか。
GUI環境ではメモリが足りず実行できないプログラムのビジュアルプロファイルを取りたいのですが…

183 :デフォルトの名無しさん:2011/12/02(金) 18:08:55.43
linuxで4.0のtoolkitを導入した時に古いものをリムーブするかと出てきたのでyesを選んだのですが、もしかして同じ階層のものを全部削除とかしています?
X windowすら上手く開かなくなったのですが…

184 :デフォルトの名無しさん:2011/12/03(土) 18:12:31.51
unixは知らんけど、windowsだとインストールしたファイルしか消されないよ
新しいの入れる度に毎回cl.hppを手動で追加してるけど、
アンインストールしてもそれだけ残ってる

185 :デフォルトの名無しさん:2011/12/03(土) 18:46:52.04
Toolkitを入れたからじゃなくて、Driverを入れたからじゃね?
Driverを入れるときにroot権限がないとXのconfigを更新できないから。

186 :デフォルトの名無しさん:2011/12/06(火) 06:03:27.50
4.1 RC2
| NVIDIA Developer Zone
http://developer.nvidia.com/cuda-toolkit-41


187 :デフォルトの名無しさん:2011/12/06(火) 10:46:45.36
蔵人にメル栓抜きツイスターきたのか

188 :デフォルトの名無しさん:2011/12/06(火) 11:57:45.78
コンパイルどれくらい早くなってんのかな
VCでコンパイルしてて小さなcuファイルでも数秒待たされるのがガチでイライラするんだけど

189 :デフォルトの名無しさん:2011/12/06(火) 12:05:12.77
vcでコンパイルしているなら、変わるわけないだろ。
vsからnvccを起動しているなら兎も角も。

190 :デフォルトの名無しさん:2011/12/06(火) 12:18:58.00
cuファイルとcppファイルに分けるのはどういう意図があるんでしょうか?thrustって何に使えるんですか?

191 :デフォルトの名無しさん:2011/12/06(火) 12:26:24.32
>>190
・nvccに任せておきたくない部分は、.cppに書く。
・抽象的に取り敢えず書くのに便利。

192 :デフォルトの名無しさん:2011/12/06(火) 16:39:18.49
cuにkernel_foo()を書いて
cppに書いたfoo()からkernel_foo()を呼ぶのが定石かと思ってた
そうすればopenclに移行したりしてもocl_kernel_foo()を呼び出すようにすれば変更すくないし。

193 :デフォルトの名無しさん:2011/12/06(火) 17:11:07.38
>>192
その定石に則れば、>191の前半を満たせるから問題ないよ。
ただ、それだとCPU側とGPU側で共通のロジックを更に別のソースに書かないといけなくなるから適材適所だと思う。

194 :デフォルトの名無しさん:2011/12/07(水) 00:20:58.16
青木氏の本読んでも、ガイド読んでも、くすだれcudaを見ても全くわからなかったので質問します。
おそらくwarpの話で、cudaをまったく分かっていないゆえの質問だと思います。
cudaのデータの取扱いを調べるために、以下のような構造のデータをGPU側に送り、
typedef struct data{
 int i,j; //初期化でi=1,j=0 
}DATA;
Grid1コ,ThreadBlock1コ,総Thread数512コと指定して、以下のようなコードを実行させました。
__global__ void test(DATA *A){
 int i = blockDim.x*blockIdx.x + threadIdx.x;
 if(i%2==0){//threadIdの奇遇でiの値及び加算値を変更
  A->i=2; A->j+=1;
 }else{
  A->i=3; A->j+=2;
}}
5,6回実行して、iの値はthreadの総数を奇遇どちらにしても3で不変でした。
jの値は実行するたび値が異なり、j=3,5,7,9のいずれかになりました。

iの結果は各warpの32コのThreadが順次if文を実行してて、
idが奇数のときの場合が後に実行されるから、結果がi=3となるのか?という理解でよろしいのでしょうか。

また、jの結果は青木氏の言う「加算命令を実行した結果は有効なスレッドに対してのみ反映される」
の理解がいると思うのですが、そもそも有効なスレッドがどう判定されているのかわかりません。
また512コのthreadがあるのに、jの値の結果が10以下になるのはどうも腑に落ちません。
i,jの値を決めているものは何か、ご教示願います。

195 :デフォルトの名無しさん:2011/12/07(水) 00:45:36.69
+=で同じアドレスに同時書き込みしてるから

196 :デフォルトの名無しさん:2011/12/07(水) 01:09:59.92
>>195
+=で同じアドレスに同時書き込みすると、内部で何が起こるんですか?

197 :デフォルトの名無しさん:2011/12/07(水) 01:27:31.21
競合状態が発生してんじゃないの?
atomic演算とか同期が必要だと思うよ。
512スレッドで同一アドレスの変数の読み書きしてんだから。

まず512個の要素の配列作って、添え字にスレッド番号(上のi)を指定して確認してみたら?
書籍ならcuda exampleも買って読むといいかもね

198 :デフォルトの名無しさん:2011/12/07(水) 02:05:18.21
>>197
>競合状態が発生してんじゃないの?
>>195のコメと合わせて考えるに、なんとなく予想はしてましたけど、取り合いになってるんですね...

>atomic演算とか同期が必要だと思うよ。
まだザックリとしか勉強してないので、atomic演算は知らなかったです。あとで試してみます。
同期を行う場合だったら、どうすればいいのだろう。

>まず512個の要素の配列作って、添え字にスレッド番号(上のi)を指定して確認してみたら?
それは分岐条件がきちんと実行しているのか見るためのテストをしたとき確認しました。

199 :デフォルトの名無しさん:2011/12/07(水) 07:39:11.25
根本的にプログラミングの基礎が抜けている悪寒。

200 :デフォルトの名無しさん:2011/12/07(水) 09:07:22.99
>>198
atomicという排他的処理を知らないのなら並列化プログラムをやるのは早すぎる。
増してCUDAという特殊な並列化をやるのはもっと早すぎる。


201 :デフォルトの名無しさん:2011/12/07(水) 12:02:55.64
>>199
どこまで戻ればいいんでしょうか...

>>200
...これ使って課題提出しなきゃならんので、早すぎると言われても後に引けないです...

202 :デフォルトの名無しさん:2011/12/07(水) 13:13:17.16
>>198
>>201
ですけど、
>>200の"atomicという排他的処理を知らないのなら並列化プログラムをやるのは早すぎる。"
では並列化プログラムをやるにあたり、
どういったことを勉強して、どういった手順でやればいいんでしょうか?
そこがよくわからず、私はいきなりcudaに突っ込んだので四苦八苦してるわけですが...

203 :デフォルトの名無しさん:2011/12/07(水) 14:10:39.88
>>202
同期とロック

204 :デフォルトの名無しさん:2011/12/07(水) 15:51:27.47
青木本読んだんじゃないのん?

x+=1ってのは
1.xを読み出す
2.読み出した値に1を加算する
3.結果をxに格納する
って手順なんだけど以下略

205 :デフォルトの名無しさん:2011/12/07(水) 16:45:28.63
>>203
同期とロックですか。勉強します。

>>204
青木本は読みましたけど、ものにしたっていう状態じゃないです...
>x+=1ってのは...
普段、なんとなく使ってるので... 勉強不足ですみません。

206 :デフォルトの名無しさん:2011/12/07(水) 17:29:38.47
ここの一番下に載ってる資料の4-7ページにatomicの説明が少しある。
ttp://accc.riken.jp/HPC/training.html

207 :デフォルトの名無しさん:2011/12/07(水) 17:58:07.38
>>206
ありがとうございます。参照します。

208 :デフォルトの名無しさん:2011/12/07(水) 23:52:25.72
この本でも読んでみると良い。
日本語でわかりやすい。

http://www.amazon.co.jp/gp/product/4798014621/



209 :デフォルトの名無しさん:2011/12/08(木) 00:42:53.72
>>208
ありがとうございます。明日早速本屋行ってきます。

210 :デフォルトの名無しさん:2011/12/08(木) 17:37:10.02
久しぶりに新しいCUDAの本が出たようだ。
http://www.amazon.co.jp/dp/4906608000/ref=cm_sw_r_fa_dp_454Zob03VNZ25


211 :デフォルトの名無しさん:2011/12/08(木) 18:59:25.20
>>210
グラフィックをメインにしてCUDAを道具としてこれから使おうとする人にはいいかも?
3章第3項がCUDA入門。第4項の「応用プログラム」でおもしろい話が読めたらいいね。

こっち(2011/11/14発売)も気になったけど、内容説明読んだだけじゃどの程度の本なのかわからなかった。
値段とタイトルからはちょっと期待させられる。
Amazon.co.jp: CUDA Application Design and Development: Rob Farber: 洋書
http://www.amazon.co.jp/CUDA-Application-Design-Development-Farber/dp/0123884268/

212 :デフォルトの名無しさん:2011/12/09(金) 05:54:08.89
CUDAというかVisualStudioが分からない。鬱だ死のう。

213 :デフォルトの名無しさん:2011/12/09(金) 09:40:45.47
VSがわからないんじゃなく、C++がわからないんだろ

214 :デフォルトの名無しさん:2011/12/11(日) 23:39:15.74
fortranもここでいいかな?
pgiのコンパイラ買って、三重ループの前後に指示行追加してやったけどまったく速くならない。憂鬱。。

215 :デフォルトの名無しさん:2011/12/11(日) 23:56:05.05
fortran使う人って量子論とかやってる人?
どの世界の人が使うのかしら?

216 :デフォルトの名無しさん:2011/12/12(月) 08:31:48.58
fortranを使うのは過去の遺産のためだろう。
fortranを使えば速くなるわけでもないしなあ。

217 :デフォルトの名無しさん:2011/12/12(月) 09:39:06.64
fortran懐かし過ぎる
先日、二度と使うことは無いと思って本を捨てたばっかりだわ

218 :デフォルトの名無しさん:2011/12/12(月) 20:09:34.35
高レベル言語使うのは時代の趨勢じゃないかな。

LAPACKはfortranで書かれている。
fortranはnvccよりもアプリケーションに近い。
CUDAを使うのにnvccのレベルでなきゃダメということは無いと思う。
逆に、nvcc使わない人をバカにするヤツは、本物のバカだと思う。

219 :デフォルトの名無しさん:2011/12/12(月) 20:20:14.55
道具にこだわることが目的となって、肝心の結果を出せない人間を馬鹿というでは?

220 :デフォルトの名無しさん:2011/12/12(月) 21:18:03.01
TSUBAME2.0でnvccを使うのと京でfortran使うのと比べるとどうなんだろう。
京はSPARCだけで計算して、TSUBAME2.0の10倍くらいの速度みたいだ。

東工大は教育機関なのに対して、理化学研究所は結果を求められる独立行政法人。

221 :デフォルトの名無しさん:2011/12/12(月) 22:25:59.20
既存のコードを使う限りでは京のほうが速いだろう。
ただ、富士通のコンパイラがダメダメだと聞いたけど、
それは解消されたのかな?

222 :デフォルトの名無しさん:2011/12/12(月) 22:29:03.43
富士通のコンパイラは、8コアCPUへの割り当ては自動らしい。
これ、すごいことだと思うんだけど、
このコンパイラは完成しているのだろうか。

223 :デフォルトの名無しさん:2011/12/12(月) 22:46:19.38
>>222
openMPに対応してたら各CPUのコアへの割り当ては普通にコンパイラ側でやると思うけど?

59 KB
■ このスレッドは過去ログ倉庫に格納されています

★スマホ版★ 掲示板に戻る 全部 前100 次100 最新50

read.cgi ver 05.04.00 2017/10/04 Walang Kapalit ★
FOX ★ DSO(Dynamic Shared Object)