- 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/
- 89 名前：81 mailto:sage [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 mailto:sage [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 名前：デフォルトの名無しさん mailto:sahud@onc.nasi.jp [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 名前：デフォルトの名無しさん mailto:sage [2011/10/06(木) 18:43:04.20 ]
- そこだけ書かれても分からないと思う。
nvccのエラーメッセージ載せないと。
- 94 名前：デフォルトの名無しさん mailto:sage [2011/10/06(木) 19:11:36.83 ]
- >>92
敵性語なんでよくわかんなかったけど、
VisualStudio を管理者として実行したら動く？
- 95 名前：やんやん ◆yanyan72E. mailto:sage [2011/10/06(木) 21:18:36.98 ]
- Visual C studioは下手にGUIに頼るよりも、コマンドラインでやった方がはやいこともある。コマンドラインでnvccやってみよう。
- 96 名前：デフォルトの名無しさん mailto:sage [2011/10/07(金) 16:37:49.21 ]
- 単精度のmadってmulより遅いやないかー
プログラミングガイド嘘つきやないかー
ちくしょー
- 97 名前：デフォルトの名無しさん mailto:sage [2011/10/09(日) 01:36:18.51 ]
- >>96
そりゃ命令実行時間は同じだけど、
必要な帯域は違う。
- 98 名前：デフォルトの名無しさん mailto:sage [2011/10/09(日) 02:40:58.84 ]
- >>92
コマンドラインのオプションが何か1つ間違ってた気がする。
コマンドプロンプトで同じコマンド実行するとエラーが出たので、
そのエラーのひとつ前の引数が問題だった、はず。
- 99 名前：デフォルトの名無しさん mailto:sage [2011/10/10(月) 02:14:56.87 ]
- CUDAとRTミドルウェアを組み合わせてみた。
行った方法は、CUDAのサンプルをRTコンポーネントに
パッケージングする方法。
動作周期にあわせて、行列計算が実行されたのを確認。
- 100 名前：デフォルトの名無しさん mailto:sage [2011/10/10(月) 20:49:05.67 ]
- カーネル側のソース分割ってできないんだっけ？
環境はLinux cuda4.0
- 101 名前：デフォルトの名無しさん mailto:sage [2011/10/10(月) 22:00:17.43 ]
- できるけどコンスタントメモリが変な事になったような気がする
- 102 名前：デフォルトの名無しさん mailto:sage [2011/10/10(月) 22:04:05.94 ]
- このままだと5万行くらいになっちゃいそう・・・
- 103 名前：デフォルトの名無しさん mailto:sage [2011/10/10(月) 22:39:08.05 ]
- >>102
正直処理対象がCUDAに向いてないような。。。
- 104 名前：デフォルトの名無しさん mailto:sage [2011/10/14(金) 02:20:30.66 ]
- とある文献でレイテンシを表す際, 種類によってclock, cycleが使い分けられていました。
clock : texture cache, constant cache
cycle : global memory, shared memory
なんとなくメモリ, キャッシュで使い分けられているのはわかりますが, このままでは速度の比較ができません。
二つの違いは何なのか, またcycle-clock間で単位の変換が可能なのか, もしご存知の方がいたらご教授ください。
- 105 名前：デフォルトの名無しさん mailto:sage [2011/10/17(月) 13:50:45.89 ]
- カーネル関数内で立てたフラグをホスト側で判断する事って出来る？
- 106 名前：デフォルトの名無しさん mailto:sage [2011/10/17(月) 14:13:41.38 ]
- グローバルメモリ介してできるだろ
- 107 名前：デフォルトの名無しさん mailto:sage [2011/10/18(火) 02:35:47.28 ]
- はじめてのCUDAプログラミングの9章に載ってるcublas使うプログラム実行したんだけど、
cublasInit()あたりで「〜0x7c812afb で初回の例外が発生しました: Microsoft C++ の
例外: cudaError_enum〜」って出てうまく動かないんだけどなして？
- 108 名前：デフォルトの名無しさん mailto:sage [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 名前：デフォルトの名無しさん mailto:sage [2011/10/19(水) 11:56:08.38 ]
- 適当なポインタで確保した後に，
そのアドレス転送してやればいいんじゃない？
- 110 名前：デフォルトの名無しさん mailto:sage [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 mailto:sage [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 名前：デフォルトの名無しさん mailto:sage [2011/10/19(水) 21:10:14.17 ]
- 概ねいいんでない?
前者なら、私はhoge_t hoge;で定義しちゃうけど。
- 113 名前：109 mailto:sage [2011/10/19(水) 22:04:08.49 ]
- 俺が考えてたのは後者だな。
合ってると思うよ。
でもmemcpyのsizeof(int)はsizeof(int *)にした方が良いと思う。
- 114 名前：108 mailto:sage [2011/10/19(水) 23:57:45.47 ]
- 未だにポインタが使いこなせてないと痛感しました。
このたびはアドバイス有難うございました。
- 115 名前：デフォルトの名無しさん mailto:sage [2011/10/20(木) 00:12:23.07 ]
- ポインタ使える男の人ってモテそうだよね
- 116 名前：デフォルトの名無しさん mailto:sage [2011/10/20(木) 17:52:20.01 ]
- ポインタでいろんな所を指して、
覗いたりいじくり回したりするんですね。
- 117 名前：やんやん ◆yanyan72E. mailto:sage [2011/10/21(金) 01:37:52.34 ]
- トラックポイントではない
- 118 名前：デフォルトの名無しさん mailto:sage [2011/11/03(木) 20:05:27.91 ]
- デバイスから出る「the launch timed out and was terminated.」とかのメッセージを
CPU側で受信する方法ってどうやるのかな？コンソール出力されるんだからできると思うけど
- 119 名前：デフォルトの名無しさん mailto:sage [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 名前：デフォルトの名無しさん mailto:sage [2011/11/07(月) 09:05:53.46 ]
- CG法と書いてあるので問題行列は対称疎行列でCUSPARSEからはcsrmv関数だけ、CUBLASからは
Level-1 BLASルーチンしか使っていないのだと思いますが、
実際にどの程度遅いのでしょうか？
ランダムアクセスを含むcsrmvはLevel-1 BLASルーチンとくらべてメモリアクセス効率が極端に下がるのは
仕様というか当然の結果です。
また古いバージョンのライブラリを使っていたりしませんか？
handle云々は基本的に関係ありません。関係するとすればよほど小さな行列を扱っている場合に限られます。
- 122 名前：デフォルトの名無しさん mailto:sage [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 名前：デフォルトの名無しさん mailto:sage [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 名前：デフォルトの名無しさん mailto:sage [2011/11/07(月) 18:52:46.57 ]
- 構造体にポインタがあるのはダメ
理屈で考えればわかるよな
- 125 名前：デフォルトの名無しさん mailto:sage [2011/11/07(月) 19:13:06.14 ]
- >>124
そうでしたか
わかりました
- 126 名前：デフォルトの名無しさん mailto:sage [2011/11/07(月) 19:34:23.39 ]
- 馬鹿でプログラミング初心者には難しいれす（^ρ^）
- 127 名前：やんやん ◆yanyan72E. mailto:sage [2011/11/07(月) 19:39:20.38 ]
- 無駄にでっかくメモリ確保された
ただのポインタがコピーされてるだけに見えるのは気のせい？
実体をコピーしようとしている意図は読み取れるけれど、
コードはその意図を反映してないような。。。
- 128 名前：やんやん ◆yanyan72E. mailto:sage [2011/11/07(月) 19:50:21.46 ]
- あ、考え違いか、すまん
- 129 名前：デフォルトの名無しさん mailto:sage [2011/11/07(月) 19:52:32.74 ]
- >>124
わからないのでその理屈を教えていただけないでしょうか
- 130 名前：デフォルトの名無しさん mailto:sage [2011/11/07(月) 20:15:14.16 ]
- cudaMallocHostはあるけどcudaMallocはしてないの?書き洩れ？
- 131 名前：デフォルトの名無しさん mailto:sage [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 名前：デフォルトの名無しさん mailto:sage [2011/11/08(火) 09:53:28.63 ]
- 取り敢えず、ホスト側とデバイス側の二種類のメモリ空間を扱っていることを肝に銘じよう。
「構造体にポインタがある」ある場合、そのポインタがホスト側のメモリを指しているのか
デバイス側のメモリを指しているのか常に意識しないといけない。
それがいやだったら、ポインタではなく配列にしてしまえ。
cf.
struct var_str {
int var[N * N];
float var_g;
struct var_box vb[N];
};
- 133 名前：デフォルトの名無しさん mailto:sage [2011/11/08(火) 09:54:43.06 ]
- それ以前に、そんなデータ構造を渡しても素直に並列化できなくて速度が出ない悪寒。
- 134 名前：デフォルトの名無しさん mailto:sage [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 名前：デフォルトの名無しさん mailto:sage [2011/11/08(火) 15:31:03.58 ]
- おまえ、そのレベルでよくこのスレに辿り着けたな。。
- 137 名前：デフォルトの名無しさん mailto:sage [2011/11/08(火) 18:02:17.87 ]
- とりあえず動かしてみろｗ
- 138 名前：デフォルトの名無しさん mailto:sage [2011/11/09(水) 04:38:11.56 ]
- 4.1RC出てるんだな
デベロッパプログラムに登録したいけど審査あるから、
他人のソースコンパイルして遊ぶだけの俺には無理
- 139 名前：デフォルトの名無しさん mailto:sage [2011/11/09(水) 14:48:59.26 ]
- 試しにデタラメ並べ立てて申請したら通った
んでSDKのサンプルを上から順番に走らせて遊んでたらブルスク食らった
危ねえ・・・
- 140 名前：デフォルトの名無しさん mailto:sage [2011/11/09(水) 22:34:20.87 ]
- >>138
審査あんの？俺普通に登録してOKでたけど
- 141 名前：デフォルトの名無しさん mailto:sage [2011/11/09(水) 22:45:40.52 ]
- 審査ってなにを調べるの？
- 142 名前：デフォルトの名無しさん mailto:sage [2011/11/09(水) 23:51:20.06 ]
- ウンコの重さとか
- 143 名前：デフォルトの名無しさん mailto:sage [2011/11/10(木) 00:31:38.44 ]
- Parallel Nsightのアカウントとは別なのか。
面倒くさすぎだろ。
- 144 名前：デフォルトの名無しさん mailto:sage [2011/11/10(木) 06:20:27.98 ]
- CUDA.NETってCUDA 3.0までで止まってるよね？
- 145 名前：デフォルトの名無しさん mailto:sage [2011/11/10(木) 12:49:08.28 ]
- >>144
くだんねーと（CUDANET）思ったんだろ作ってる方も
- 146 名前：デフォルトの名無しさん mailto:sage [2011/11/11(金) 00:24:14.22 ]
- 【審議中】
∧,,∧ ∧,,∧
∧ (´・ω・) (・ω・｀) ∧∧
( ´・ω) Ｕ) ( つと ﾉ(ω・｀ )
｜ Ｕ ( ´・) (・｀ ) と ﾉ
ｕ-ｕ （l ) ( ﾉu-ｕ
`ｕ-ｕ'. `ｕ-u'
- 147 名前：デフォルトの名無しさん mailto:sage [2011/11/14(月) 02:32:16.94 ]
- kernelの中でどの部分でレジスタがMAXになっているか、どの値を保持しているかっていうのは簡単にわかるものですか？
- 148 名前：デフォルトの名無しさん mailto:sage [2011/11/14(月) 11:50:00.02 ]
- 細かいことは兎も角、取り敢えずptx出力を見てみたら医院で内科医。
- 149 名前：デフォルトの名無しさん mailto:sage [2011/11/14(月) 18:06:51.94 ]
- shrQATest.hって4.0のtoolkitとSDK入れたらどこにあるはずですか？
- 150 名前：デフォルトの名無しさん mailto:sage [2011/11/15(火) 02:43:04.47 ]
- %ProgramData%\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.0\shared
- 151 名前：デフォルトの名無しさん mailto:sage [2011/11/15(火) 02:44:06.11 ]
- \inc
- 152 名前：デフォルトの名無しさん mailto:sage [2011/11/15(火) 16:34:29.61 ]
- 見つかりました
ありがとうございます
- 153 名前：デフォルトの名無しさん mailto:sage [2011/11/15(火) 21:47:31.35 ]
- linux kernelからGPGPU使いたいんだけどKGPU以外に何かいい方法ある？
- 154 名前：デフォルトの名無しさん mailto:sage [2011/11/15(火) 22:55:00.82 ]
- CUDA developer用としてCUDAサイトで配布されてるドライバと普通のドライバってどんな違いがあるの？
CUDAに必須と思ったら、普通のドライバインストールしてもCUDA開発できると書いている所があった。
- 155 名前：デフォルトの名無しさん mailto:sage [2011/11/15(火) 23:18:16.01 ]
- 一時期、developer用と一般用のバージョンが同じ時があったけど
md5sumは一致した。
おそらく、一般用ドライバのバージョンで対応が間に合っていないときに
beta版を出してるだけだと思う。
だから、両者でバージョンの新しい方を使えば問題ない。
- 156 名前：デフォルトの名無しさん mailto:sage [2011/11/16(水) 00:34:28.85 ]
- SC2011 OpenACC Joint Press Release - main
www.openacc-standard.org/announcements-1/nvidiacraypgicapsunveil%E2%80%98openacc%E2%80%99programmingstandardforparallelcomputing
- 157 名前：デフォルトの名無しさん mailto:sage [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 名前：デフォルトの名無しさん mailto:sage [2011/11/27(日) 22:05:36.51 ]
- 全スレッドの計算結果を1コアに集約して1スレッドで処理したいんだけどいい方法あるかな？
リダクションのサンプルだと最後はCPUで合計だしこれでは都合悪いんだよね
全てGPUでやりたいのよね
- 160 名前：デフォルトの名無しさん mailto:sage [2011/11/27(日) 23:20:19.00 ]
- しかたないじゃん、CPUで合計出した方が速いんだから。
- 161 名前：デフォルトの名無しさん mailto:sage [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 名前：デフォルトの名無しさん mailto:sage [2011/11/27(日) 23:35:14.30 ]
- >>161
＞データ量によってはそうでもないよ。
>>159は「1スレッドで処理したい」と言っている。
それに、「計算結果を集約」と言っているだけで足し算とは言っていない。
この条件の下では、データ量が多くなれば多くなるほどCPUで処理した方が速い。
データ数が極端に少なければ、転送時間の関係でGPUでしたほうが速いかもしれないが。
- 163 名前：デフォルトの名無しさん mailto:sage [2011/11/27(日) 23:40:05.87 ]
- あーそうか、まさか1スレッドだけ動かす方法が判らんと言う話とは思わなかった。
単純に、「CPUで合計」の代わりをやりたいのかと思ったよ。
- 164 名前：159 mailto:sage [2011/11/27(日) 23:55:11.26 ]
- >>160-163
レスサンクス
一度集約して最少二乗法を適用してから再度マルチスレッドで処理しようとしてるんだけど
一度CPUに戻すしかないかな・・・
- 165 名前：デフォルトの名無しさん mailto:sage [2011/11/28(月) 00:02:09.10 ]
- 最小二乗法ならマルチスレッドでできるじゃん。
- 166 名前：デフォルトの名無しさん mailto:sage [2011/11/29(火) 03:07:20.82 ]
- リダクションしろよ
1ブロック分しかできないってこと？
- 167 名前：デフォルトの名無しさん mailto:sage [2011/11/29(火) 10:48:39.03 ]
- サイズが決まっている3次元配列どうしの演算で演算回数も決まっているんだけど、
ブロック数とスレッド数を調整して演算回数ぴったりにするとコアレッシングしなくて効率悪いんだ。
こーいう時ってどゆアプローチ？
- 168 名前：デフォルトの名無しさん mailto:sage [2011/11/29(火) 14:42:48.05 ]
- CPUで計算させてみる
- 169 名前：デフォルトの名無しさん mailto:sage [2011/11/29(火) 15:26:04.48 ]
- もちろんCPUでは計算させてみてるよぉぉん(><)
1つ目の演算対象配列はギリギリL2に乗らないサイズ、2つ目の配列はVRAMにも乗らないサイズなので悩んでる。
- 170 名前：159 mailto:sage [2011/11/29(火) 16:52:31.44 ]
- もっと良いCPUかGPU買え
- 171 名前：デフォルトの名無しさん mailto:sage [2011/11/30(水) 02:20:28.60 ]
- >>159
同じブロック内の複数スレッドは共有メモリで集約できるけど、
ブロックをまたがっての集約ができないってこと？
手元のコードではブロック数の一時保存のための領域を確保し
そこにブロックごとの結果を放り込むカーネルを起動し、
次にブロック数が1、スレッド数が前のカーネルのブロック数の
カーネルを起動。その中でさらに集約して結果を1つにまとめている。
カーネル2回起動よりもCPUでやらせた方が速いかもしれないが。
- 172 名前：デフォルトの名無しさん mailto:sage [2011/11/30(水) 04:29:51.70 ]
- code.google.com/p/thrust/
のリダクション使ってみれば？
中身で何やってるかはよくわからんが
- 173 名前：デフォルトの名無しさん mailto:sage [2011/11/30(水) 08:40:15.94 ]
- カーネルの中で関数作って足せばいいんじゃない
threadが終わったら自作関数に投げて足すとか
バリアがいるけど
- 174 名前：デフォルトの名無しさん mailto:sage [2011/11/30(水) 08:52:34.48 ]
- 微妙に話が噛み合っていない悪寒。
つーか、>159が混乱に輪を掛けているし。
- 175 名前：159 mailto:sage [2011/11/30(水) 10:04:46.33 ]
- 各スレッドでサンプリングされたXiとYiを使って
Σ（Ai*Xi-Yi）^2を計算して最小化するパラメータのAiを求めたいんだよ
求まったAiを使ってまたGPUでマルチスレッドで処理をするから
CPUに戻したくなかっただけ
戻せばできるのは分かるけど、戻さないでできるならその方がいいと思ったわけね
- 176 名前：やんやん ◆yanyan72E. mailto:sage [2011/11/30(水) 10:24:24.37 ]
- なんか式がおかしい、Ai=Yi/Xi なんて答えを求めたい訳じゃなかろう？
- 177 名前：159 mailto:sage [2011/11/30(水) 11:22:19.91 ]
- >>176
うん
正確には多項式近似の最小二乗法
めんどいので略した
- 178 名前：171 mailto:sage [2011/11/30(水) 14:43:43.94 ]
- >>177
非線形最小二乗法ってこと？非線形最適化法しか知らないけど、目的関数が
特殊な場合にそれを上手に利用したのが非線形最小二乗って理解。
手元では対数尤度（二乗和みたいなもの）を171の手順で計算する関数を
CUDAで書いて、MCMCの中で使って推定値を求めているけど、同じことは
最小二乗法でもできると思うんだが…。
- 179 名前：159 mailto:sage [2011/11/30(水) 14:50:41.39 ]
- >>178
カーネルの二度起動だよね？
その方法でもできると思うけど、一度で済ませたいというのがニーズ
それよかMCMCをマルチスレッドでやるってのが分からんのだが…
あれはマルコフ連鎖だからほとんどマルチスレッドで効果上がらんだろw
- 180 名前：178 mailto:sage [2011/11/30(水) 15:21:59.61 ]
- >>179
こちらも一度の起動で済ませたかったけど方法が分からなかったので、
分かったら是非報告よろしく。
尤度の計算対象とするデータ数があまりにばかでかくて…。マルコフ連鎖は
順次計算だけど、その中の並列化は相当のメリットがあってね。
- 181 名前：やんやん ◆yanyan72E. mailto:sage [2011/11/30(水) 16:50:24.20 ]
- >>179
二度起動がいやなら、ここの
ttp://neuralgorithm.org/documents/global-thread-synchronization-on-cuda-ja/
globalsyncみたいなのを使って計算→集計→計算って
やらせればいいんだろうけれど、
なんか怖いなw
- 182 名前：デフォルトの名無しさん mailto:sage [2011/11/30(水) 16:50:29.44 ]
- CUDA Visual Plofiler への入力ファイル(cpjファイルとcsvファイル)をコンソールを使って作成したいのですが、可能でしょうか。
GUI環境ではメモリが足りず実行できないプログラムのビジュアルプロファイルを取りたいのですが…
- 183 名前：デフォルトの名無しさん mailto:sage [2011/12/02(金) 18:08:55.43 ]
- linuxで4.0のtoolkitを導入した時に古いものをリムーブするかと出てきたのでyesを選んだのですが、もしかして同じ階層のものを全部削除とかしています？
X windowすら上手く開かなくなったのですが…
- 184 名前：デフォルトの名無しさん mailto:sage [2011/12/03(土) 18:12:31.51 ]
- unixは知らんけど、windowsだとインストールしたファイルしか消されないよ
新しいの入れる度に毎回cl.hppを手動で追加してるけど、
アンインストールしてもそれだけ残ってる
- 185 名前：デフォルトの名無しさん mailto:sage [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
developer.nvidia.com/cuda-toolkit-41
- 187 名前：デフォルトの名無しさん mailto:sage [2011/12/06(火) 10:46:45.36 ]
- 蔵人にメル栓抜きツイスターきたのか
- 188 名前：デフォルトの名無しさん mailto:sage [2011/12/06(火) 11:57:45.78 ]
- コンパイルどれくらい早くなってんのかな
VCでコンパイルしてて小さなcuファイルでも数秒待たされるのがガチでイライラするんだけど
- 189 名前：デフォルトの名無しさん mailto:sage [2011/12/06(火) 12:05:12.77 ]
- vcでコンパイルしているなら、変わるわけないだろ。
vsからnvccを起動しているなら兎も角も。
|
 |
