11/08/23 22:11:49.24
関連サイト
CUDA
URLﾘﾝｸ(www.nvidia.co.jp)
CUDAに触れてみる
URLﾘﾝｸ(chihara.naist.jp)
cudaさわってみた
URLﾘﾝｸ(gpgpu.jp)
CUDA のインストール
URLﾘﾝｸ(blog.goo.ne.jp)
NVIDIAの「GeForce 8800 GT(G92)」と次に控える64-bit GPUアーキテクチャ
URLﾘﾝｸ(pc.watch.impress.co.jp)
CUDAを使う
URLﾘﾝｸ(tech.ckme.co.jp)
NVIDIA CUDAを弄ってみた その2
URLﾘﾝｸ(dvd-r.sblo.jp)
CUDAベンチ
URLﾘﾝｸ(wataco.air-nifty.com)
KNOPPIX for CUDA
URLﾘﾝｸ(www.yasuoka.mech.keio.ac.jp)
3:デフォルトの名無しさん
11/08/23 22:12:38.55
宣伝になっちゃうと難なので、一応分けて。いろいろ載っているサイト。
URLﾘﾝｸ(www.softek.co.jp)
4:デフォルトの名無しさん
11/08/23 22:26:00.72
おい、pertだろ
まちがえんな
5:デフォルトの名無しさん
11/08/23 22:48:26.89
>>2
個人blogのは情報古すぎなような…
「cudaさわってみた」はリンク切れだし。
6:デフォルトの名無しさん
11/08/25 19:36:34.86
>>4
やっぱりCUDAやってる人は荒んでるんだな
7:デフォルトの名無しさん
11/08/26 07:04:46.69
ほんとCUDAらないことで騒ぐよな
8: 忍法帖【Lv=40,xxxPT】
11/08/26 14:04:14.48
ケンカするのはやめてCUDAさい＞＜
9:デフォルトの名無しさん
11/08/26 20:59:54.29
そうだよ。喧嘩はやめ…GeForんｹﾞﾌｫﾝ
10:デフォルトの名無しさん
11/08/26 21:57:35.41
まだスレが始まっTeslaいないのに
11:デフォルトの名無しさん
11/08/27 00:03:43.41
ときに口論になるのもみんな自分の技術にそれだKeplerイドを持ってるってことだ。
12:デフォルトの名無しさん
11/08/27 00:09:15.96
amazon.comで検索して出てくるCUDA本ってことごとく出版がずいぶん先か延期になったりするな
なんでだろ
13:デフォルトの名無しさん
11/08/27 21:26:38.44
無効とは情報の登録ポリシーが違う…ような気がする。
それがAmazonのポリシーなのか出版社のポリシーなのかは知らないけれど。
多分出版社のサイト直接見た方が正確。
14:デフォルトの名無しさん
11/08/27 23:28:49.51
cudaの本が必要、って訳でも無いが出版されるなら買ってしまうと思う
15:デフォルトの名無しさん
11/08/31 20:00:08.40
4.0になってdeviceemuが使えんのだが、CUDA載せてないPCで動かす方法ない？
16:デフォルトの名無しさん
11/09/01 07:46:58.96
3.2を使う。
17:デフォルトの名無しさん
11/09/01 09:05:53.29
>>16 アホしかおらんのか
18:デフォルトの名無しさん
11/09/01 09:16:16.30
>>17
現実的じゃないか。どうせemuで正確なエミュレーションができるわけじゃないし。
19:デフォルトの名無しさん
11/09/01 09:25:26.25
>>17
4.0でdeviceemu使う方法なんてねーよ、ってことだよ。
言わせんな恥ずかしい
20:デフォルトの名無しさん
11/09/01 10:02:13.87
>>15
Cuda x86
21:デフォルトの名無しさん
11/09/01 15:33:21.69
理屈の上ではKnights Ferryでも動かせるのかな？
22:デフォルトの名無しさん
11/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:デフォルトの名無しさん
11/09/01 17:40:47.67
>>22の続きです
assumed oldの値をcpu側に持ってきて表示してみたら、
oldの値(atomicCASの返す値)とassumedの値が末尾の数ビットが
異なってました。50回ほどループしてやっとビットレベルで値が一致して
ループを抜けているようでした。その間に過剰に加算が行われていたようです。
24:デフォルトの名無しさん
11/09/01 18:47:40.63
カードの計算が正しいかチェックするベンチマークってありませんか？
四則演算と√と三角関数とかを全ビットを演算器全部チェックしろと云う無茶ぶりがきた。
計算機が計算を間違うなんてありえないのに…
25:デフォルトの名無しさん
11/09/01 19:11:55.88
Teslaの存在意義全否定デスカ
26:デフォルトの名無しさん
11/09/01 19:34:24.88
>>24
過去スレを読むとわかってもらえると思うが、みんなそれを求めているんだ。
そのチェックにおける最先端はおそらく長崎大学。
27:デフォルトの名無しさん
11/09/01 19:41:57.01
1スレッドの実行がどのコアかがわかる仕組みがどのような方法か分かればいいわけだよな
誰か調べてみろよ
28:デフォルトの名無しさん
11/09/01 23:41:13.82
>>22,23
もしかしてブロック辺りのスレッド数が32以下なら動いたりする？
29:デフォルトの名無しさん
11/09/01 23:44:25.02
>>27
何がしたいのか分からんがそれは簡単だ。
30:デフォルトの名無しさん
11/09/01 23:56:07.27
>>29
コアと計算結果の表が出来上がればコア間の相違が分かるじゃん
やり方知ってるなら教えてください
31:デフォルトの名無しさん
11/09/01 23:58:10.21
煽っても教えてやらん
29じゃないけどｗ
32:デフォルトの名無しさん
11/09/02 02:12:08.96
NvidiaのCUDAをやろうと思い、Nvidiaのカードの刺さったマシンにFedoraを
入れてCUDAのためのNvidiaのドライバを入れて使ってみたが、OSのカーネル
のアップデートがあるたびに、コンソールが再起動時にブラックアウトした。
そのたびごとに、別のマシンからリモートログインをしてCUDAのための
Nvidiaのドライバを再インストールする必要があった。
CUDA3.xまではそれでなんとかなったが、現在のCUDA4.0になると、
幾らやっても駄目になってしまった。どうやらFedora15ではまともに
サポートされていないようなのだ。Nvidiaの提供しているCUDAのための
ドライバがインストール時におかしなメッセージが出てインストールが
出来ない。
33:デフォルトの名無しさん
11/09/02 02:29:09.05
ここはおまえの日記帳だから好きに使いなさい
34:デフォルトの名無しさん
11/09/02 11:08:38.52
>>28
ブロックあたりのスレッド数を32以下にして試してみましたが、
結果は変わらなかったです。
35:デフォルトの名無しさん
11/09/02 18:20:08.65
>>23
とりあえず、プログラミングガイドのコードそのままだし、
実際アルゴリズム自体には問題ないから、
意図どおり動かないとしたら、ドライバかSDKのバグか
ハードの故障じゃね。
まあ、ここで詳細なバージョンとか晒してみたり、
PTXやアセンブリを貼ってみるのもいいけど、
本家のフォーラムで聞いてみたほうが早いんじゃね。
36:デフォルトの名無しさん
11/09/05 23:02:56.52
長崎大の言う演算間違えって、具体的にどういう演算を間違えたのかの情報ってないの？
間違えますって情報だけ垂れ流して具体例出さんって、どういうことなんだべ。
37:デフォルトの名無しさん
11/09/05 23:24:45.84
×演算間違え
○演算間違い
38:デフォルトの名無しさん
11/09/07 18:20:00.85
CUDAを使用するプログラムで巨大なデータを扱いたいのですが、nvcc でのコンパイル時に
relocation truncated to fit: R_X86_64_32S against `.bss'
とエラーが出て完了しません。gcc ではデータ量制限を解除するのに-mcmodel=medium というオプションがあるようですが、nvcc でこれにあたるものはないでしょうか。
39:デフォルトの名無しさん
11/09/07 19:43:00.72
巨大データってどれくらい？
CUDAは基本グローバルメモリ使用の制約があるから”巨大”は決して得意ではないのだけど
40:38
11/09/07 22:05:04.34
>>39
今確認したところ、要素が数十万個の配列を使用していて、それらの合計は200MBほどです。
そのうち、実際にGPUに計算させるためGPUに渡すデータは50MB程度です。
上記のエラーメッセージでググったところ、データ量制限が問題だろうということだったのですが…それほど大きなデータでもないですよね。
なんなんでしょうこのエラー。
プログラミングに関してはまだ知識が浅いためチンプンカンプンなこと言ってたらすみません。
41:デフォルトの名無しさん
11/09/07 22:36:52.16
どうせデータをソース中に書き込んでるんだろw
42:デフォルトの名無しさん
11/09/07 22:42:26.34
>37
日本語の間違えより、長崎大の言ってる間違いの詳しい情報くれよぅ…。
43:デフォルトの名無しさん
11/09/07 22:50:33.71
要素数が大きすぎるなら、mallocを使ってみたら？
44:デフォルトの名無しさん
11/09/07 23:04:45.67
>40
配列は何次元？
CやC++だと配列は連続したメモリ領域と規定されてるので、
確保可能なメモリ量は物理メモリの絶対量ではなく、
アドレスを連続で確保できるメモリ量に規定されるよ。
例えば大きめの２次元配列を確保する場合は
double* pBuf = new double[row * column];
とするより
double** ppBuf = new (double*)[row];
for(int i = 0; i< row; ++i){
ppBuf[i] = new double[column];
}
とかにしたほうがいいよ。
45:デフォルトの名無しさん
11/09/07 23:11:07.95
ファイル読み込みで、cudaMallocHost(だったかな？？)でメモリ確保すればいいじゃん
46:38
11/09/07 23:11:33.99
>>43>>44
gccと同様にオプションでエラー回避できないかな…と思ったのですが、配列の取り方を変えるほうが良さそうですね。
特に>>44は知りませんでした。
ありがとうございました！
47:デフォルトの名無しさん
11/09/07 23:23:20.08
良く分からんが、-Xcompiler -mcmodel=mediumじゃ駄目なん？
-Xlinkerも要るかもしれないが。
48:デフォルトの名無しさん
11/09/08 08:01:27.80
>>44
普通はありえない。そんなことしたら、ポインタアクセスが一回余計に要るから遅くなるし、
キャッシュアウトしやすくなるから更に遅くなる。
まぁ、GPUに転送するのは1スライス分だけだろうからGPUに任せてCPUは暇こいていられるならいいかも知れんが。
49:デフォルトの名無しさん
11/09/08 17:24:44.67
>>44
それは無い。
・メモリも無駄に消費
・アクセス性能がた落ち
・バグの温床になる
つーかC++ならvectorで初期化時に
サイズ指定したほうがいろいろ便利
50:デフォルトの名無しさん
11/09/08 17:40:17.98
配列でやる場合、
配列の添え字がINT_MAXを越える大きさならそうするしかないと思うが・・
51:デフォルトの名無しさん
11/09/08 20:47:04.41
キャッシャがアクセスがとドヤ顔されても・・・。
要素のデカイ配列が確保できんって話なわけだし。
52:デフォルトの名無しさん
11/09/08 22:19:09.58
コンパイル時の話なのに
ヒープのフラグメンテーションの訳無いだろう。
nvccからgccに自由にオプション渡すやり方は
知らんが、最悪c(++)部分のコードを分割して
gccでコンパイルすりゃ良い。
仮にヒープのフラグメンテーションが酷い場合でも
>>44みたいなやり方は下策だよ。
素人におかしなやり方を覚えさせるだけ。
大きな配列確保する前に小さな物についてShrink-to-fitやればデフラグされるから
その後巨大配列を確保すれば良い。
特にGPGPUで巨大な配列扱うような用途なら
配列の配列なんて転送も困るだろうに。
53:デフォルトの名無しさん
11/09/08 22:30:02.29
>大きな配列確保する前に小さな物についてShrink-to-fitやればデフラグされるから
>その後巨大配列を確保すれば良い。
コレ具体的にどんなコード書けばいいのん？
54:デフォルトの名無しさん
11/09/09 00:35:55.99
メインの計算に必要な必要な小容量配列
保持用のポインタなりvectorの変数だけ始めに
ヒープに確保した後、メインの計算前の前処理をする。
この処理で、一時的なデータと最終的に使うデータの
ヒープ確保が入り乱れる所為でフラグメンテーションが発生する。
最終的に使うデータについて、大体処理した順番を守って
ひたすら同容量確保してはコピーを繰り返せば、
こういうデータは全て前方に圧縮されて配置されるようになる。
55:デフォルトの名無しさん
11/09/09 21:48:41.98
？
まあスレ違いなんで他でやってよ。
56:デフォルトの名無しさん
11/09/09 22:24:20.88
kernelへの引数に構造体は使えんの？
今やってみたらコンパイルでエラーになった
57:デフォルトの名無しさん
11/09/09 22:30:19.36
あ、間違えてただ
構造体引数使えるべな
58:デフォルトの名無しさん
11/09/14 11:12:56.96
>>15
URLﾘﾝｸ(code.google.com)
Ocelot currently allows CUDA programs to be executed on NVIDIA GPUs, AMD GPUs, and x86-CPUs at full speed without recompilation.
59:デフォルトの名無しさん
11/09/15 17:52:31.49
cuda4.0、gtx580で
Maximum dimension of a grid( 65535 , 65535 , 65535 )
ってなってるんですがz軸のブロックの取れる数は1じゃなくなったんですか？
60:デフォルトの名無しさん
11/09/19 22:05:37.25
デバイス上の処理が長時間になる場合、途中経過を取得する方法ってないかな？
例えば１００万回のモンテカルロを実行して、途中の１０万回、２０万回が済んだらその情報を取りたいんだけど…
61:デフォルトの名無しさん
11/09/19 22:31:31.07
>>60
streamingでできないかな?
62:デフォルトの名無しさん
11/09/19 22:52:18.06
streamingってなに？
SMのSじゃないよね？
63:デフォルトの名無しさん
11/09/20 05:00:39.03
カーネル呼び出し時の<<<>>>の第3引数がstreamingの何かじゃなかったっけ？
俺は>>61を見てそういうのがあったなと思い出したレベルなので使い方は知らないけど。
64:デフォルトの名無しさん
11/09/20 08:29:52.67
第3パラメータはsharedMemoryのサイズだよ。そこで指定しなくても使えるけど。
65:デフォルトの名無しさん
11/09/21 22:11:25.09
ごめん第4引数だったみたい。それとAsync系関数がうんたらかんたら。
66:デフォルトの名無しさん
11/09/22 00:50:42.07
あー、その辺だね。
67:デフォルトの名無しさん
11/09/22 18:05:24.11
【トリップ検索】CUDA SHA-1 Tripper【GeForce】
ｽﾚﾘﾝｸ(software板)
すげえな
68:デフォルトの名無しさん
11/09/23 11:16:26.70
個人でCUDAプログラミングやってうれしいのはその程度なのか？
と逆に不安になるんだが…
69:やんやん ◆yanyan72E.
11/09/23 11:35:33.50
世の中なんでも膨大なデータをいかに早く処理するかが勝負の世界に
なってきているから、CUDAプログラミングの旨味はいくらでもあるだろ。
70:デフォルトの名無しさん
11/09/23 12:04:06.11
いつかやってみようと思って半年くらいこのスレROMってるが
いまだに開発環境のセッティングの仕方すら調べていない・・・
71:デフォルトの名無しさん
11/09/24 00:02:19.96
>>68
>>67くらいのコード書ける？
72:デフォルトの名無しさん
11/09/24 15:28:34.26
>>71
おれがそれくらいのコード書けるようになるころには、自ずとましな使い道が浮かんでくることだろう、
っていうこと？
73:デフォルトの名無しさん
11/10/02 08:23:54.81
>>60
試してないけど、長い処理の途中でグローバルメモリに書き込んで、
別streamでその値を読むカーネルを起動すれば取れると思う。
キャッシュからグローバルメモリに書き込むのは特別な関数を
使うと思ったけど名前忘れた
74:デフォルトの名無しさん
11/10/02 09:01:54.77
>>73
思い出した。atomicを使うか、書き込む側で __threadfence()
75:デフォルトの名無しさん
11/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.
11/10/04 21:13:10.02
-archコマンドってなんだそら？
compute capabilityはnvccなりcl.exeのオプションに渡すものであって、
-archコマンドなんてものはないよ。
77:デフォルトの名無しさん
11/10/04 21:50:45.04
自分でオプションでって書いてるけど、メッセージはそれで実行した系だよな
78:デフォルトの名無しさん
11/10/04 23:16:18.47
>>76
>>77
ありがとうございます。「cuda by example」にもatomic関数を使ったソースには-archコマンドラインオプションを指定する必要があります。
とあるのですが、みなさんはatomicを使ったソースをコンパイルするときなにかコマンドつけませんか＞＜？
79:デフォルトの名無しさん
11/10/04 23:19:48.93
>>78
あ、もちろん全部64bitです！
80:デフォルトの名無しさん
11/10/04 23:39:11.06
CUDA以前の話だな。コンパイラオプションの意味がわかってない。
というか、VCなら、カスタムビルド規則でCUDA Runtime API Build Ruleを追加したほうが早い
81: 忍法帖【Lv=40,xxxPT】
11/10/05 02:20:11.36
-gencode=arch=compute_20,code=\"sm_21,compute_20\"
になってるなぁ俺の環境だと。VC2010ならプロジェクトのプロパティでCUDA関係の設定するところがあるんじゃない？
82:デフォルトの名無しさん
11/10/05 11:51:04.38
>>80
ありがとうございます。
えーと、それはおそらくVC2008の話ではないかと思います。2010にカスタムビルド規則という名前はなくなり、その代りCUDA4.0というビルドカスタマイズに統一されています。
83:デフォルトの名無しさん
11/10/05 11:58:44.51
>>81
ご返答ありがとうございます。
そのコードはどこに書いていますか？
はい、cuda c/c++という項目にいろいろいじるところはあるのですがコンパイラオプションを設定する項目はないように思えます。
私がオプションを追加しているのはビルドイベントという項目ですね。
以前にもatomic関連の話が出ていたようです。>>22
84:デフォルトの名無しさん
11/10/05 12:14:48.16
-arch=compute_20
にすればいいんじゃねえの？
85:デフォルトの名無しさん
11/10/05 12:17:32.55
>>84
ありがとうございます。
おっしゃるように>>75のように追加するとエラーを吐きます。
追加の仕方が違うのでしょうか。
86:80
11/10/05 13:36:15.15
VS2010は色々変わってるんだな。
知らんかったすまん。
とりあえず「CUDA VS2010」でググって出てくる一番上のサイト参考にして、CUDA4.0のビルドカスタマイズファイルを導入した後、プロパティからCode Generationをcompute_20,sm_20にしたらいけた
87:デフォルトの名無しさん
11/10/05 15:09:51.57
GPUのリソースを全部CUDAに割り当てるため，グラボを2枚さしてディスプレイ出力用とCUDA用に分けるというのはありますが
同じようなことをsandy bridgeとグラボ1枚でも出来るのでしょうか？
88:デフォルトの名無しさん
11/10/05 15:36:59.85
>>87
Z68マザーなら、マザー側DVIにディスプレイつないでおけばいい話じゃ？
CUDAで計算したデータをintelで表示したいなら、
1.DirectXかOpenGLのリソースをCUDAにバインドしてデータ書き込み
2.メインメモリに転送
3.intel側に転送
って形になるかなぁ。多分。
89:81
11/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
11/10/05 21:17:54.24
そもそもなぜCUDA Runtime APIという項目があるのかというとサンプルのプロジェクトをコピったから。
91:デフォルトの名無しさん
11/10/06 16:35:52.62
>>80
>>81
ありがとうございます！code generationを>>80のように書きなおしたらオプション指定することなくコンパイル通りました。
どうやらVS2010からはcommand line指定しなくてもcode generationでオプションを渡せるようです。
本当に助かりました。ありがとうございます！
92:デフォルトの名無しさん
11/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:デフォルトの名無しさん
11/10/06 18:43:04.20
そこだけ書かれても分からないと思う。
nvccのエラーメッセージ載せないと。
94:デフォルトの名無しさん
11/10/06 19:11:36.83
>>92
敵性語なんでよくわかんなかったけど、
VisualStudio を管理者として実行したら動く？
95:やんやん ◆yanyan72E.
11/10/06 21:18:36.98
Visual C studioは下手にGUIに頼るよりも、コマンドラインでやった方がはやいこともある。コマンドラインでnvccやってみよう。
96:デフォルトの名無しさん
11/10/07 16:37:49.21
単精度のmadってmulより遅いやないかー
プログラミングガイド嘘つきやないかー
ちくしょー
97:デフォルトの名無しさん
11/10/09 01:36:18.51
>>96
そりゃ命令実行時間は同じだけど、
必要な帯域は違う。
98:デフォルトの名無しさん
11/10/09 02:40:58.84
>>92
コマンドラインのオプションが何か1つ間違ってた気がする。
コマンドプロンプトで同じコマンド実行するとエラーが出たので、
そのエラーのひとつ前の引数が問題だった、はず。
99:デフォルトの名無しさん
11/10/10 02:14:56.87
CUDAとRTミドルウェアを組み合わせてみた。
行った方法は、CUDAのサンプルをRTコンポーネントに
パッケージングする方法。
動作周期にあわせて、行列計算が実行されたのを確認。
100:デフォルトの名無しさん
11/10/10 20:49:05.67
カーネル側のソース分割ってできないんだっけ？
環境はLinux cuda4.0
101:デフォルトの名無しさん
11/10/10 22:00:17.43
できるけどコンスタントメモリが変な事になったような気がする
102:デフォルトの名無しさん
11/10/10 22:04:05.94
このままだと5万行くらいになっちゃいそう・・・
103:デフォルトの名無しさん
11/10/10 22:39:08.05
>>102
正直処理対象がCUDAに向いてないような。。。
104:デフォルトの名無しさん
11/10/14 02:20:30.66
とある文献でレイテンシを表す際, 種類によってclock, cycleが使い分けられていました。
clock : texture cache, constant cache
cycle : global memory, shared memory
なんとなくメモリ, キャッシュで使い分けられているのはわかりますが, このままでは速度の比較ができません。
二つの違いは何なのか, またcycle-clock間で単位の変換が可能なのか, もしご存知の方がいたらご教授ください。
105:デフォルトの名無しさん
11/10/17 13:50:45.89
カーネル関数内で立てたフラグをホスト側で判断する事って出来る？
106:デフォルトの名無しさん
11/10/17 14:13:41.38
グローバルメモリ介してできるだろ
107:デフォルトの名無しさん
11/10/18 02:35:47.28
はじめてのCUDAプログラミングの9章に載ってるcublas使うプログラム実行したんだけど、
cublasInit()あたりで「～0x7c812afb で初回の例外が発生しました: Microsoft C++ の
例外: cudaError_enum～」って出てうまく動かないんだけどなして？
108:デフォルトの名無しさん
11/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:デフォルトの名無しさん
11/10/19 11:56:08.38
適当なポインタで確保した後に，
そのアドレス転送してやればいいんじゃない？
110:デフォルトの名無しさん
11/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
11/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:デフォルトの名無しさん
11/10/19 21:10:14.17
概ねいいんでない?
前者なら、私はhoge_t hoge;で定義しちゃうけど。
113:109
11/10/19 22:04:08.49
俺が考えてたのは後者だな。
合ってると思うよ。
でもmemcpyのsizeof(int)はsizeof(int *)にした方が良いと思う。
114:108
11/10/19 23:57:45.47
未だにポインタが使いこなせてないと痛感しました。
このたびはアドバイス有難うございました。
115:デフォルトの名無しさん
11/10/20 00:12:23.07
ポインタ使える男の人ってモテそうだよね
116:デフォルトの名無しさん
11/10/20 17:52:20.01
ポインタでいろんな所を指して、
覗いたりいじくり回したりするんですね。
117:やんやん ◆yanyan72E.
11/10/21 01:37:52.34
トラックポイントではない
118:デフォルトの名無しさん
11/11/03 20:05:27.91
デバイスから出る「the launch timed out and was terminated.」とかのメッセージを
CPU側で受信する方法ってどうやるのかな？コンソール出力されるんだからできると思うけど
119:デフォルトの名無しさん
11/11/03 21:13:42.56
>>118
デバイス側といっても、cudart辺りがドライバからのステータスを見て出していると思われ。
従って、当然リダイレクトできる。
120:デフォルトの名無しさん
11/11/07 04:01:37.49
CUBLASとCUSPARSEを使って倍精度で計算を行うCG法のプログラムを作成したのですが、
CUSPARSEのcsrmv関数だけが異様に遅く速度が出ません。
この原因としては何が考えられますでしょうか？
自分としてはCUBLASとCUSPARSEでそれぞれ別のhandleを使っているせいかと思っているのですが
それを確かめる方法や、CUBLASとCUSPARSEでhandleを共有する方法がわかりません...
121:デフォルトの名無しさん
11/11/07 09:05:53.46
CG法と書いてあるので問題行列は対称疎行列でCUSPARSEからはcsrmv関数だけ、CUBLASからは
Level-1 BLASルーチンしか使っていないのだと思いますが、
実際にどの程度遅いのでしょうか？
ランダムアクセスを含むcsrmvはLevel-1 BLASルーチンとくらべてメモリアクセス効率が極端に下がるのは
仕様というか当然の結果です。
また古いバージョンのライブラリを使っていたりしませんか？
handle云々は基本的に関係ありません。関係するとすればよほど小さな行列を扱っている場合に限られます。
122:デフォルトの名無しさん
11/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:デフォルトの名無しさん
11/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:デフォルトの名無しさん
11/11/07 18:52:46.57
構造体にポインタがあるのはダメ
理屈で考えればわかるよな
125:デフォルトの名無しさん
11/11/07 19:13:06.14
>>124
そうでしたか
わかりました
126:デフォルトの名無しさん
11/11/07 19:34:23.39
馬鹿でプログラミング初心者には難しいれす（^ρ^）
127:やんやん ◆yanyan72E.
11/11/07 19:39:20.38
無駄にでっかくメモリ確保された
ただのポインタがコピーされてるだけに見えるのは気のせい？
実体をコピーしようとしている意図は読み取れるけれど、
コードはその意図を反映してないような。。。
128:やんやん ◆yanyan72E.
11/11/07 19:50:21.46
あ、考え違いか、すまん
129:デフォルトの名無しさん
11/11/07 19:52:32.74
>>124
わからないのでその理屈を教えていただけないでしょうか
130:デフォルトの名無しさん
11/11/07 20:15:14.16
cudaMallocHostはあるけどcudaMallocはしてないの?書き洩れ？
131:デフォルトの名無しさん
11/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:デフォルトの名無しさん
11/11/08 09:53:28.63
取り敢えず、ホスト側とデバイス側の二種類のメモリ空間を扱っていることを肝に銘じよう。
「構造体にポインタがある」ある場合、そのポインタがホスト側のメモリを指しているのか
デバイス側のメモリを指しているのか常に意識しないといけない。
それがいやだったら、ポインタではなく配列にしてしまえ。
cf.
struct var_str {
int var[N * N];
float var_g;
struct var_box vb[N];
};
133:デフォルトの名無しさん
11/11/08 09:54:43.06
それ以前に、そんなデータ構造を渡しても素直に並列化できなくて速度が出ない悪寒。
134:デフォルトの名無しさん
11/11/08 10:50:48.01
>>129
>>108-114の流れを読め．
135:デフォルトの名無しさん
11/11/08 13:04:40.64
C# (.NET Framework 4.0)で動作するCUDAのライブラリってありますか？
CUDA.NETがそうだとは思うのですが、.NET Framework 2.0 or newerとしか書いていないので、
結局のところ.NET 4.0で動作するのかどうかわかりません。
136:デフォルトの名無しさん
11/11/08 15:31:03.58
おまえ、そのレベルでよくこのスレに辿り着けたな。。
137:デフォルトの名無しさん
11/11/08 18:02:17.87
とりあえず動かしてみろｗ
138:デフォルトの名無しさん
11/11/09 04:38:11.56
4.1RC出てるんだな
デベロッパプログラムに登録したいけど審査あるから、
他人のソースコンパイルして遊ぶだけの俺には無理
139:デフォルトの名無しさん
11/11/09 14:48:59.26
試しにデタラメ並べ立てて申請したら通った
んでSDKのサンプルを上から順番に走らせて遊んでたらブルスク食らった
危ねえ・・・
140:デフォルトの名無しさん
11/11/09 22:34:20.87
>>138
審査あんの？俺普通に登録してOKでたけど
141:デフォルトの名無しさん
11/11/09 22:45:40.52
審査ってなにを調べるの？
142:デフォルトの名無しさん
11/11/09 23:51:20.06
ウンコの重さとか
143:デフォルトの名無しさん
11/11/10 00:31:38.44
Parallel Nsightのアカウントとは別なのか。
面倒くさすぎだろ。
144:デフォルトの名無しさん
11/11/10 06:20:27.98
CUDA.NETってCUDA 3.0までで止まってるよね？
145:デフォルトの名無しさん
11/11/10 12:49:08.28
>>144
くだんねーと（CUDANET）思ったんだろ作ってる方も
146:デフォルトの名無しさん
11/11/11 00:24:14.22
【審議中】
∧,,∧ ∧,,∧
∧ (´・ω・) (・ω・｀) ∧∧
( ´・ω) Ｕ) ( つと ﾉ(ω・｀ )
｜ Ｕ ( ´・) (・｀ ) と ﾉ
ｕ-ｕ （l ) ( ﾉu-ｕ
`ｕ-ｕ'. `ｕ-u'
147:デフォルトの名無しさん
11/11/14 02:32:16.94
kernelの中でどの部分でレジスタがMAXになっているか、どの値を保持しているかっていうのは簡単にわかるものですか？
148:デフォルトの名無しさん
11/11/14 11:50:00.02
細かいことは兎も角、取り敢えずptx出力を見てみたら医院で内科医。
149:デフォルトの名無しさん
11/11/14 18:06:51.94
shrQATest.hって4.0のtoolkitとSDK入れたらどこにあるはずですか？
150:デフォルトの名無しさん
11/11/15 02:43:04.47
%ProgramData%\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.0\shared
151:デフォルトの名無しさん
11/11/15 02:44:06.11
\inc
152:デフォルトの名無しさん
11/11/15 16:34:29.61
見つかりました
ありがとうございます
153:デフォルトの名無しさん
11/11/15 21:47:31.35
linux kernelからGPGPU使いたいんだけどKGPU以外に何かいい方法ある？
154:デフォルトの名無しさん
11/11/15 22:55:00.82
CUDA developer用としてCUDAサイトで配布されてるドライバと普通のドライバってどんな違いがあるの？
CUDAに必須と思ったら、普通のドライバインストールしてもCUDA開発できると書いている所があった。
155:デフォルトの名無しさん
11/11/15 23:18:16.01
一時期、developer用と一般用のバージョンが同じ時があったけど
md5sumは一致した。
おそらく、一般用ドライバのバージョンで対応が間に合っていないときに
beta版を出してるだけだと思う。
だから、両者でバージョンの新しい方を使えば問題ない。
156:デフォルトの名無しさん
11/11/16 00:34:28.85
SC2011 OpenACC Joint Press Release - main
URLﾘﾝｸ(www.openacc-standard.org)
157:デフォルトの名無しさん
11/11/16 05:00:22.99
>>155
了解。レスありがと
158:デフォルトの名無しさん
11/11/22 17:29:12.79
すいません、経験者の方教えてください。
ホストメモリをページロックメモリで確保して、
ストリームを300個くらい起動して
MemcpyAnsyncでh->d ,kenrel, h->dで非同期実行させてるんですけど、
どうやらkernelでセグって停止してしまいます。
しかし、cuda-gdbで見てみるとどうもセグるようなポジションじゃないんです。
そこで質問なんですが、MemcpyAsyncでコピーリクエストをだしてコピーが始まり、
制御がカーネルに入り、コピーが完了したストリームを実行してる途中に、
まだコピー中の領域でセグった場合、どこでセグったように見えるのでしょうか？
やはりその時実行しているカーネルの中でセグっているように見えるのでしょうか？
159:デフォルトの名無しさん
11/11/27 22:05:36.51
全スレッドの計算結果を1コアに集約して1スレッドで処理したいんだけどいい方法あるかな？
リダクションのサンプルだと最後はCPUで合計だしこれでは都合悪いんだよね
全てGPUでやりたいのよね
160:デフォルトの名無しさん
11/11/27 23:20:19.00
しかたないじゃん、CPUで合計出した方が速いんだから。
161:デフォルトの名無しさん
11/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:デフォルトの名無しさん
11/11/27 23:35:14.30
>>161
＞データ量によってはそうでもないよ。
>>159は「1スレッドで処理したい」と言っている。
それに、「計算結果を集約」と言っているだけで足し算とは言っていない。
この条件の下では、データ量が多くなれば多くなるほどCPUで処理した方が速い。
データ数が極端に少なければ、転送時間の関係でGPUでしたほうが速いかもしれないが。
163:デフォルトの名無しさん
11/11/27 23:40:05.87
あーそうか、まさか1スレッドだけ動かす方法が判らんと言う話とは思わなかった。
単純に、「CPUで合計」の代わりをやりたいのかと思ったよ。
164:159
11/11/27 23:55:11.26
>>160-163
レスサンクス
一度集約して最少二乗法を適用してから再度マルチスレッドで処理しようとしてるんだけど
一度CPUに戻すしかないかな・・・
165:デフォルトの名無しさん
11/11/28 00:02:09.10
最小二乗法ならマルチスレッドでできるじゃん。
166:デフォルトの名無しさん
11/11/29 03:07:20.82
リダクションしろよ
1ブロック分しかできないってこと？
167:デフォルトの名無しさん
11/11/29 10:48:39.03
サイズが決まっている3次元配列どうしの演算で演算回数も決まっているんだけど、
ブロック数とスレッド数を調整して演算回数ぴったりにするとコアレッシングしなくて効率悪いんだ。
こーいう時ってどゆアプローチ？
168:デフォルトの名無しさん
11/11/29 14:42:48.05
CPUで計算させてみる
169:デフォルトの名無しさん
11/11/29 15:26:04.48
もちろんCPUでは計算させてみてるよぉぉん(><)
1つ目の演算対象配列はギリギリL2に乗らないサイズ、2つ目の配列はVRAMにも乗らないサイズなので悩んでる。
170:159
11/11/29 16:52:31.44
もっと良いCPUかGPU買え
171:デフォルトの名無しさん
11/11/30 02:20:28.60
>>159
同じブロック内の複数スレッドは共有メモリで集約できるけど、
ブロックをまたがっての集約ができないってこと？
手元のコードではブロック数の一時保存のための領域を確保し
そこにブロックごとの結果を放り込むカーネルを起動し、
次にブロック数が1、スレッド数が前のカーネルのブロック数の
カーネルを起動。その中でさらに集約して結果を1つにまとめている。
カーネル2回起動よりもCPUでやらせた方が速いかもしれないが。
172:デフォルトの名無しさん
11/11/30 04:29:51.70
URLﾘﾝｸ(code.google.com)
のリダクション使ってみれば？
中身で何やってるかはよくわからんが
173:デフォルトの名無しさん
11/11/30 08:40:15.94
カーネルの中で関数作って足せばいいんじゃない
threadが終わったら自作関数に投げて足すとか
バリアがいるけど
174:デフォルトの名無しさん
11/11/30 08:52:34.48
微妙に話が噛み合っていない悪寒。
つーか、>159が混乱に輪を掛けているし。
175:159
11/11/30 10:04:46.33
各スレッドでサンプリングされたXiとYiを使って
Σ（Ai*Xi-Yi）^2を計算して最小化するパラメータのAiを求めたいんだよ
求まったAiを使ってまたGPUでマルチスレッドで処理をするから
CPUに戻したくなかっただけ
戻せばできるのは分かるけど、戻さないでできるならその方がいいと思ったわけね
176:やんやん ◆yanyan72E.
11/11/30 10:24:24.37
なんか式がおかしい、Ai=Yi/Xi なんて答えを求めたい訳じゃなかろう？
177:159
11/11/30 11:22:19.91
>>176
うん
正確には多項式近似の最小二乗法
めんどいので略した
178:171
11/11/30 14:43:43.94
>>177
非線形最小二乗法ってこと？非線形最適化法しか知らないけど、目的関数が
特殊な場合にそれを上手に利用したのが非線形最小二乗って理解。
手元では対数尤度（二乗和みたいなもの）を171の手順で計算する関数を
CUDAで書いて、MCMCの中で使って推定値を求めているけど、同じことは
最小二乗法でもできると思うんだが…。
179:159
11/11/30 14:50:41.39
>>178
カーネルの二度起動だよね？
その方法でもできると思うけど、一度で済ませたいというのがニーズ
それよかMCMCをマルチスレッドでやるってのが分からんのだが…
あれはマルコフ連鎖だからほとんどマルチスレッドで効果上がらんだろw
180:178
11/11/30 15:21:59.61
>>179
こちらも一度の起動で済ませたかったけど方法が分からなかったので、
分かったら是非報告よろしく。
尤度の計算対象とするデータ数があまりにばかでかくて…。マルコフ連鎖は
順次計算だけど、その中の並列化は相当のメリットがあってね。
181:やんやん ◆yanyan72E.
11/11/30 16:50:24.20
>>179
二度起動がいやなら、ここの
URLﾘﾝｸ(neuralgorithm.org)
globalsyncみたいなのを使って計算→集計→計算って
やらせればいいんだろうけれど、
なんか怖いなw
182:デフォルトの名無しさん
11/11/30 16:50:29.44
CUDA Visual Plofiler への入力ファイル(cpjファイルとcsvファイル)をコンソールを使って作成したいのですが、可能でしょうか。
GUI環境ではメモリが足りず実行できないプログラムのビジュアルプロファイルを取りたいのですが…
183:デフォルトの名無しさん
11/12/02 18:08:55.43
linuxで4.0のtoolkitを導入した時に古いものをリムーブするかと出てきたのでyesを選んだのですが、もしかして同じ階層のものを全部削除とかしています？
X windowすら上手く開かなくなったのですが…
184:デフォルトの名無しさん
11/12/03 18:12:31.51
unixは知らんけど、windowsだとインストールしたファイルしか消されないよ
新しいの入れる度に毎回cl.hppを手動で追加してるけど、
アンインストールしてもそれだけ残ってる
185:デフォルトの名無しさん
11/12/03 18:46:52.04
Toolkitを入れたからじゃなくて、Driverを入れたからじゃね?
Driverを入れるときにroot権限がないとXのconfigを更新できないから。
186:デフォルトの名無しさん
11/12/06 06:03:27.50
4.1 RC2
| NVIDIA Developer Zone
URLﾘﾝｸ(developer.nvidia.com)
187:デフォルトの名無しさん
11/12/06 10:46:45.36
蔵人にメル栓抜きツイスターきたのか
188:デフォルトの名無しさん
11/12/06 11:57:45.78
コンパイルどれくらい早くなってんのかな
VCでコンパイルしてて小さなcuファイルでも数秒待たされるのがガチでイライラするんだけど
189:デフォルトの名無しさん
11/12/06 12:05:12.77
vcでコンパイルしているなら、変わるわけないだろ。
vsからnvccを起動しているなら兎も角も。
190:デフォルトの名無しさん
11/12/06 12:18:58.00
cuファイルとcppファイルに分けるのはどういう意図があるんでしょうか？thrustって何に使えるんですか？
191:デフォルトの名無しさん
11/12/06 12:26:24.32
>>190
・nvccに任せておきたくない部分は、.cppに書く。
・抽象的に取り敢えず書くのに便利。
192:デフォルトの名無しさん
11/12/06 16:39:18.49
cuにkernel_foo()を書いて
cppに書いたfoo()からkernel_foo()を呼ぶのが定石かと思ってた
そうすればopenclに移行したりしてもocl_kernel_foo()を呼び出すようにすれば変更すくないし。
193:デフォルトの名無しさん
11/12/06 17:11:07.38
>>192
その定石に則れば、>191の前半を満たせるから問題ないよ。
ただ、それだとCPU側とGPU側で共通のロジックを更に別のソースに書かないといけなくなるから適材適所だと思う。
194:デフォルトの名無しさん
11/12/07 00:20:58.16
青木氏の本読んでも、ガイド読んでも、くすだれcudaを見ても全くわからなかったので質問します。
おそらくwarpの話で、cudaをまったく分かっていないゆえの質問だと思います。
cudaのデータの取扱いを調べるために、以下のような構造のデータをGPU側に送り、
typedef struct data{
int i,j; //初期化でi=1,j=0
}DATA;
Grid1コ,ThreadBlock１コ,総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:デフォルトの名無しさん
11/12/07 00:45:36.69
+=で同じアドレスに同時書き込みしてるから
196:デフォルトの名無しさん
11/12/07 01:09:59.92
>>195
+=で同じアドレスに同時書き込みすると、内部で何が起こるんですか？
197:デフォルトの名無しさん
11/12/07 01:27:31.21
競合状態が発生してんじゃないの？
atomic演算とか同期が必要だと思うよ。
512スレッドで同一アドレスの変数の読み書きしてんだから。
まず５１２個の要素の配列作って、添え字にスレッド番号（上のi）を指定して確認してみたら？
書籍ならcuda exampleも買って読むといいかもね
198:デフォルトの名無しさん
11/12/07 02:05:18.21
>>197
>競合状態が発生してんじゃないの？
>>１９５のコメと合わせて考えるに、なんとなく予想はしてましたけど、取り合いになってるんですね...
>atomic演算とか同期が必要だと思うよ。
まだザックリとしか勉強してないので、atomic演算は知らなかったです。あとで試してみます。
同期を行う場合だったら、どうすればいいのだろう。
>まず５１２個の要素の配列作って、添え字にスレッド番号（上のi）を指定して確認してみたら？
それは分岐条件がきちんと実行しているのか見るためのテストをしたとき確認しました。
199:デフォルトの名無しさん
11/12/07 07:39:11.25
根本的にプログラミングの基礎が抜けている悪寒。
200:デフォルトの名無しさん
11/12/07 09:07:22.99
>>198
atomicという排他的処理を知らないのなら並列化プログラムをやるのは早すぎる。
増してCUDAという特殊な並列化をやるのはもっと早すぎる。
201:デフォルトの名無しさん
11/12/07 12:02:55.64
>>199
どこまで戻ればいいんでしょうか...
>>200
...これ使って課題提出しなきゃならんので、早すぎると言われても後に引けないです...
202:デフォルトの名無しさん
11/12/07 13:13:17.16
>>198
>>201
ですけど、
>>200の"atomicという排他的処理を知らないのなら並列化プログラムをやるのは早すぎる。"
では並列化プログラムをやるにあたり、
どういったことを勉強して、どういった手順でやればいいんでしょうか？
そこがよくわからず、私はいきなりcudaに突っ込んだので四苦八苦してるわけですが...
203:デフォルトの名無しさん
11/12/07 14:10:39.88
>>202
同期とロック
204:デフォルトの名無しさん
11/12/07 15:51:27.47
青木本読んだんじゃないのん？
x+=1ってのは
1.xを読み出す
2.読み出した値に1を加算する
3.結果をxに格納する
って手順なんだけど以下略
205:デフォルトの名無しさん
11/12/07 16:45:28.63
>>203
同期とロックですか。勉強します。
>>204
青木本は読みましたけど、ものにしたっていう状態じゃないです...
>x+=1ってのは...
普段、なんとなく使ってるので... 勉強不足ですみません。
206:デフォルトの名無しさん
11/12/07 17:29:38.47
ここの一番下に載ってる資料の4-7ページにatomicの説明が少しある。
URLﾘﾝｸ(accc.riken.jp)
207:デフォルトの名無しさん
11/12/07 17:58:07.38
>>206
ありがとうございます。参照します。
208:デフォルトの名無しさん
11/12/07 23:52:25.72
この本でも読んでみると良い。
日本語でわかりやすい。
URLﾘﾝｸ(www.amazon.co.jp)
209:デフォルトの名無しさん
11/12/08 00:42:53.72
>>208
ありがとうございます。明日早速本屋行ってきます。
210:デフォルトの名無しさん
11/12/08 17:37:10.02
久しぶりに新しいCUDAの本が出たようだ。
URLﾘﾝｸ(www.amazon.co.jp)
211:デフォルトの名無しさん
11/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: 洋書
URLﾘﾝｸ(www.amazon.co.jp)
212:デフォルトの名無しさん
11/12/09 05:54:08.89
CUDAというかVisualStudioが分からない。鬱だ死のう。
213:デフォルトの名無しさん
11/12/09 09:40:45.47
VSがわからないんじゃなく、C++がわからないんだろ
214:デフォルトの名無しさん
11/12/11 23:39:15.74
fortranもここでいいかな？
pgiのコンパイラ買って、三重ループの前後に指示行追加してやったけどまったく速くならない。憂鬱。。
215:デフォルトの名無しさん
11/12/11 23:56:05.05
fortran使う人って量子論とかやってる人？
どの世界の人が使うのかしら？
216:デフォルトの名無しさん
11/12/12 08:31:48.58
fortranを使うのは過去の遺産のためだろう。
fortranを使えば速くなるわけでもないしなあ。
217:デフォルトの名無しさん
11/12/12 09:39:06.64
fortran懐かし過ぎる
先日、二度と使うことは無いと思って本を捨てたばっかりだわ
218:デフォルトの名無しさん
11/12/12 20:09:34.35
高レベル言語使うのは時代の趨勢じゃないかな。
LAPACKはfortranで書かれている。
fortranはnvccよりもアプリケーションに近い。
CUDAを使うのにnvccのレベルでなきゃダメということは無いと思う。
逆に、nvcc使わない人をバカにするヤツは、本物のバカだと思う。
219:デフォルトの名無しさん
11/12/12 20:20:14.55
道具にこだわることが目的となって、肝心の結果を出せない人間を馬鹿というでは？
220:デフォルトの名無しさん
11/12/12 21:18:03.01
TSUBAME2.0でnvccを使うのと京でfortran使うのと比べるとどうなんだろう。
京はSPARCだけで計算して、TSUBAME2.0の10倍くらいの速度みたいだ。
東工大は教育機関なのに対して、理化学研究所は結果を求められる独立行政法人。
221:デフォルトの名無しさん
11/12/12 22:25:59.20
既存のコードを使う限りでは京のほうが速いだろう。
ただ、富士通のコンパイラがダメダメだと聞いたけど、
それは解消されたのかな？
222:デフォルトの名無しさん
11/12/12 22:29:03.43
富士通のコンパイラは、8コアCPUへの割り当ては自動らしい。
これ、すごいことだと思うんだけど、
このコンパイラは完成しているのだろうか。
223:デフォルトの名無しさん
11/12/12 22:46:19.38
>>222
openMPに対応してたら各CPUのコアへの割り当ては普通にコンパイラ側でやると思うけど？
224:デフォルトの名無しさん
11/12/13 11:47:52.86
>>223
OpenMPがそんな事やるかよ。やるのはOSだぞ。
>>222が言っているのはNUMAのことだ。
現時点でも普通のLinuxならnumactlを使えばかなりのケースでノードの割り当てができるようになる。
CUDAを使う場合でも、複数GPUを使う場合に、MPIを使う場合に有効だ。
225:やんやん ◆yanyan72E.
11/12/13 14:16:05.38
なんだ？OpenMPやNUMAを何のことだと思ってるんだ？
というか釣られたか？
226:デフォルトの名無しさん
11/12/13 17:55:43.92
>>222
> これ、すごいことだと思うんだけど、
スパコンの世界では、ずうううっと前からやってるんじゃない？
227:デフォルトの名無しさん
11/12/13 20:56:29.90
じゃあ、そろそろ俺らも本気出すか
228:デフォルトの名無しさん
11/12/14 09:03:43.27
>>227
本気だしたら、負けだと思うんだ。
229:デフォルトの名無しさん
11/12/14 20:16:13.92
4Gamer.net ― NVIDIA，CUDA 4.1をリリース。CUDAコンパイラのソースコード公開も
URLﾘﾝｸ(www.4gamer.net)
中国時間2011年12月14日，NVIDIAは「CUDA 4.1」をリリースした。
最大の注目点は, LLVMベースとなるCUDAコンパイラが搭載されそのソースコードが公開された点だ。
LLVM（Low Level Virtual Machine）はAppleなどが参加する“言語非依存”のコンパイラ環境だ。
NVIDIAは同時に「CUDAプラットフォームのオープン化」を明言し，
LLVMベースとなるCUDAコンパイラのソースコードを研究者やツールベンダーに公開している。
要するに，NVIDIAのGPUでしか利用できなかったCUDA環境がほかのCPUやGPUに広がっていく可能性が出てきたのだ。
たとえばCPUや，それこそRadeonなどの他社製GPUでもCUDAを利用できる可能性が出てきたわけで，
CUDAの標準化をさらに推し進める起爆剤となり得るのである。
230:デフォルトの名無しさん
11/12/14 22:38:03.27
これって、個人でもソースもらえるのかな
231:デフォルトの名無しさん
11/12/14 23:11:41.54
へえ、ソースを公開したんだ。
でも思いっきりOpenCLとかぶるな。
OpenCLが死亡か。
232:デフォルトの名無しさん
11/12/15 08:04:39.43
なんで、OpneCLが死亡するんだ？
233:デフォルトの名無しさん
11/12/15 11:54:57.42
まだ来てなくない？
234:デフォルトの名無しさん
11/12/15 11:59:38.42
Radeon用CUDA作っても、特定の機種用になりそうだな
235:デフォルトの名無しさん
11/12/18 02:35:32.30
そもそもRADEON上でCUDAってマトモに動くのか？
一応HD3000世代のRV670からはFP64対応らしいが
236:デフォルトの名無しさん
11/12/19 02:14:51.68
動かないよ
動かそうとする人も現時点ではいないだろ
まだまだGPGPU向けには問題ありそうだし
237:デフォルトの名無しさん
11/12/20 01:48:20.74
cudaスレだが、
研究室や自分の周りでけで高速化すればいいんだったらcudaでいいんだろうけど
製品に組みこむならintelの普及度やなるべく多くの人に使ってもらえることを
考えるとOpenCLのほうがいいんじゃないかと思ってる。
photoshopのGPU対応もOpenCLでやってるみたいだし。
238:デフォルトの名無しさん
11/12/20 03:10:58.69
Photoshopは現状CUDAとATi Streamです
高速化しようとすると結局はデバイス毎に最適化したプログラムを書くことになるので、
デバイスに特化したライブラリの方が融通が利きます
Parallel NsightでのデバッグはOpenCLではできません
239:デフォルトの名無しさん
11/12/20 03:38:50.36
NVIDIA Opens Up CUDA, Compiler Source Code is Out - Bright Side Of News*
URLﾘﾝｸ(www.brightsideofnews.com)
240:デフォルトの名無しさん
11/12/20 18:56:37.44
CUDAのハードウェア種別を問わない展開が可能になったという事は、
CUDAの上に乗っかる形で実装されているPhysXやOptiX等もまた
ハードウェア種別を問わず動作可能になる可能性があるという事だよな
期待したい
241:デフォルトの名無しさん
11/12/20 22:23:45.55
別にソースコードなくたって、ある程度の知識を持った人なら
CUDAを移植できたと思うけど
242:デフォルトの名無しさん
11/12/20 22:49:46.93
RADEONでPhysXはマーケティングの理由で使えないんだろ？
243:デフォルトの名無しさん
11/12/20 23:19:03.65
AMDがCUDAを実装されるの嫌がってPhysX蹴ったんだっけ
244:デフォルトの名無しさん
11/12/21 19:57:27.52
質問：
GPU側に2次元配列を送るには、
１.2次元を1次元に直してからcudaMallocでメモリ確保→cudaMemcpyで送信
２.cudaMallocPitchでメモリ確保→cudaMemcpy2Dで送信
という2つの方法がありますけど、
構造体のメンバーに2次元配列があるものをGPU側に送る場合はどのようにしたらいいんでしょうか？
245:デフォルトの名無しさん
11/12/21 20:20:10.20
ポインタじゃなく配列ならそのままでうまくいかんか？
246:デフォルトの名無しさん
11/12/21 21:23:30.55
>>244
再度検証してたけど、サジ投げた。
最終的には以下の構造体をGPU側に送信して処理した後、CPU側に戻したい。
typedef struct test{
int a; int b[5];
}TEST;
typedef struct data{
int c; TEST d[1][10];
}DATA;
上をcudaMallocでメモリ確保＆cudaMemcpy送信しようとしてるけど、違うのかな。
教示お願いします。
247:デフォルトの名無しさん
11/12/21 21:36:40.66
そのまま転送すればいいじゃん
248:デフォルトの名無しさん
11/12/21 21:49:19.70
>>246です。
すいません。できました。お騒がせしました。
249:デフォルトの名無しさん
11/12/22 08:21:40.56
>>229
このコンパイラって、謎のCをPTXに変換する部分だけじゃないの？
今のフロントエンド(cudafe)は、他社製品のEDGベースだから、出てこないと思うが。
cudafe -> PTXへのコンパイラ(今だとnvopencc) -> ptxas
のまんなかが出来るだけなので、この部分のソースがあったところで、CUDAコンパイラを自分で作れるとか、
CUDAを他のプロセッサで動かせるようになるとかいうものではない。
250:デフォルトの名無しさん
11/12/22 13:07:41.69
.cファイルがコンパイルできません
たすけて
251:デフォルトの名無しさん
11/12/22 13:13:00.92
んｖｃｃじゃできんだよ
252:デフォルトの名無しさん
11/12/26 15:17:58.63
あの、OpenCLってオワコンなの？IntelとAMDが支持してるのに全然盛り上がってないんだが。
253:デフォルトの名無しさん
11/12/26 21:25:42.32
科学技術計算向けがtesla一択だから他を使う理由もないってっところなのかね。
わざわざcudaで組んだものを他に移植する理由もないだろうし。
254:デフォルトの名無しさん
11/12/26 21:42:01.58
OpenCLで書くと長くなってしまう。
それだけのこと
255:デフォルトの名無しさん
11/12/27 02:04:51.35
C++ Bindings使えよ。
256:デフォルトの名無しさん
11/12/27 10:05:27.99
ちょっとしたラッパーかくのも嫌なのか
257:デフォルトの名無しさん
11/12/27 11:26:29.65
超素人の質問です。
コマンドラインでnvccを使いコンパイルを行っているのですが、OpenGLのサンプルコードをコンパイルできません。
freeglutを入れてます。
gcc -lglut program.cではコンパイル可能でした。
しかしnvcc -lglut program.cuとするとコンパイル不可となります。
-Lや-lによってライブラリやヘッダファイル先を指定したところで、コンパイルできませんでした。
阿呆な質問かもしれませんが、自己解決できそうにないのでよろしくおねがいします。
258:デフォルトの名無しさん
11/12/27 11:33:58.46
凄いアフォだな
エラーメッセージすべてコピペし
259:帝徒＝繪璃奈＝啓北商業の野島えり
11/12/27 14:53:52.03
主犯 少頭劣一族＝蔗冽一族とは
中国 華喃の山の梺の村八分の家。
二間位の横長で玄関の右側がお勝手。
大正に猿のままで生まれたのが
鈴木あゆみ(网(アミ) 范襤の子)
フィリピン人の范襤と 同じくフィリピン人のモンゴルに逃げた『シバ』との間の男児。
日本名 鈴木ひろしと聞いた。
その後、親戚の鈴木大樹を殺し 戸籍を使用。
今は一文字 雉。
260:デフォルトの名無しさん
11/12/29 21:08:02.00
母体になるプログラムなしで、
CUDAのコード単体で動かせるツールって無いかねぇ。
外出先でCUDAのサンプル見せるのがメンドイ。
例えば ./cuda-run < source.cuda みたいな感じで実験できるといい感じ。
codepadみたいな感じで動かせればなおよしだけど。
261:デフォルトの名無しさん
11/12/29 22:01:29.67
そういうスクリプトかけばよい
262:デフォルトの名無しさん
12/01/05 10:51:17.15
>>260
＞母体になるプログラムなしで、
それってなんのこと? 外出先のPCで動かしたいってこと?
コンパイラのことなら、事前にコンパイルしておくかリモートでアクセスすればいい。
ドライバのことなら、外出先の端末にインストールさせて貰うしかない。
外出先のPCなんか頼らず、自前のPCを持っていけばいいじゃん。
263:デフォルトの名無しさん
12/01/05 13:36:51.62
>>257
手元に開発環境がなく、エラーメッセージをコピー忘れしてました。すみません。
# cc Drawtest.c -lglut
# ./a.out
（空のウィンドウが製作される）
（Drawtest.cをDrawtest.cuに中身は同じで拡張子だけ変えて）
# nvcc Drawtest.cu -lglut
/usr/bin/ld: skipping incompatible /usr/local/cuda/bin/../lib64/libglut.a when searching for -lglut
（-lglutを見付かられなかった？）
# nvcc -L /usr/local/cuda/lib64/ -l /usr/local/cuda/include/ Drawtest.cu -lglut
/usr/bin/ld: cannot find -l/usr/local/cuda/include/
collect2: ld はステータス 1 で終了しました
となりました。
264:デフォルトの名無しさん
12/01/05 14:25:08.28
インクルードパスの指定は -l じゃなくて -I な
265:デフォルトの名無しさん
12/01/05 15:05:37.06
/usr/bin/ld: skipping incompatible～って64/32ビット混在時のエラーメッセージじゃなかったかな
266:デフォルトの名無しさん
12/01/05 15:09:35.86
>>264
なんというミス。
目で見て違いが分からない。
L（小文字）→i（大文字）に変えて同じようにしてみました。
# nvcc -L /usr/local/cuda/lib64/ -I /usr/local/cuda/include/ Drawtest.cu -lglut
/usr/bin/ld: skipping incompatible /usr/local/cuda/lib64//libglut.a when searching for -lglut
/usr/bin/ld: skipping incompatible /usr/local/cuda/bin/../lib64/libglut.a when searching for -lglut
結局パス指定なしと同じ見たいですね。
267:デフォルトの名無しさん
12/01/05 15:37:13.27
おれも64/32の違いかと思ったが、パス指定から指摘してみた
実行している環境は？ -m64 オプション付けてみ
268:デフォルトの名無しさん
12/01/06 13:53:42.06
# nvcc -L /usr/local/cuda/lib64/ -I /usr/local/cuda/include/ Drawtest.cu -lglut -m64
/usr/bin/ld: skipping incompatible /usr/local/cuda/lib64//libglut.a when searching for -lglut
/usr/bin/ld: skipping incompatible /usr/local/cuda/bin/../lib64/libglut.a when searching for -lglut
ううむ・・・・。
centOS 64bit
Device 0: "Quadro 4000"
Device 1: "GeForce 9500 GT"
パスは
LD_LIBRARY_PATH = /usr/local/cuda/lib64
です。
openGLのないコードはコンパイル及び実行まで問題ありません。
openGLを使うにあたって
/usr/local/cuda/include/にglut.h freeglut.h freeglut_ext.h freeglut_std.h
/usr/local/cuda/lib64/にlibglut.a
をコピペしましたが、まだ足りないのでしょうか？
269:デフォルトの名無しさん
12/01/06 13:59:40.26
libglutは64ビット？
32ビットなんじゃねか？
270:デフォルトの名無しさん
12/01/06 15:03:44.65
glut.h freeglut.h freeglut_ext.h freeglut_std.h libglut.aをccが見ているであろう先のものと置き換えてしてみました。
# nvcc -L /usr/local/cuda/include/ -I /usr/local/cuda/include/ Drawtest.cu -lglut -m64
/tmp/tmpxft_000030a4_00000000-12_Drawtest.o: In function `disp()':
tmpxft_000030a4_00000000-1_Drawtest.cudafe1.cpp:(.text+0x14e6): undefined reference to `glClearColor'
tmpxft_000030a4_00000000-1_Drawtest.cudafe1.cpp:(.text+0x14f0): undefined reference to `glClear'
/usr/local/cuda/bin/../lib64/libglut.a(libglut_la-freeglut_init.o): In function `glutInit':
…
(中略)
…
(.text+0x898): undefined reference to `glPopClientAttrib'
collect2: ld はステータス 1 で終了しました
openGL系の関数が読まれていないみたいです。
271:デフォルトの名無しさん
12/01/06 15:37:17.18
いや、だからfreeglutは64ビットを使用してるのかと聞いてるんだが
質問に対しきちんと回答できないようでは教えようがない罠
272:デフォルトの名無しさん
12/01/06 18:39:39.36
>>271
freeglutは64ビットのものを入れ直してやってみましたが、同じ結果になりました。
>>257
そこでCUDAのヘッダフォルダとライブラリフォルダに入れていたものを消し、-Lと-Iで直接freeglutのフォルダを指定したら
解決しました。
ライブラリが足りなかったのか、それともCUDAと一緒くたにしたのがダメだったのか解かりませんが、ちゃんとPATH通せってことみたいです。
お騒がせしました。
273:デフォルトの名無しさん
12/01/06 19:45:07.40
>>262
だよねぇ・・・
274:デフォルトの名無しさん
12/01/06 20:14:34.67
環境
Windows7 Professional 64bit
Microsoft Visual C++ 2010 Express Version 10.0.40219.1
Microsoft .NET Framwork Version 4.0.30319
GeForce GTX 580
CUDA Toolkit 4.0.17
SDK 4.0.19
devdriver_4.0_winvista-win7_64_270.81_general
このサイトを参考に環境を構築しました。
URLﾘﾝｸ(feather.cocolog-nifty.com)
そして以下のサイトのサンプルプログラムを実行してみました。
URLﾘﾝｸ(www.gdep.jp)
Hello,Worldと99 bottles of beerはcpu、gpu共に実行できました。
しかし、Matrixのプログラムはcpuの方は実行できるのですがgpuの方が実行できません。
以下のエラーを吐きます。
matrix_gpu.cu(5) : fatal error C1083: Cannot open include file: 'cutil_inline.h'
: No such file or directory
どうやらVisualStudioのパスがうまく通ってないということまでわかり、以下のサイトなどを参考にCUDA_INC_PATHなどを変えてみましたが、一向に変わりません。
URLﾘﾝｸ(d.hatena.ne.jp)
SDK内のcutil_inline.h自体をtoolkitのincフォルダにコピペすると、他の.hファイルもいくつか同じエラーが出たのでエラーになったものをすべてコピペしたところ、LNK2019"link.exe"というエラーで先に進めませんでした。
一度VisualStudioを再インストールしてみましたが、状況は変わりません。
Nvidia GPU computing SDK Browserではサンプルプログラムを実行できているので、CUDAの環境は整っていると思われます。
どうすれば解決できますでしょうか・・・。かれこれ１週間以上格闘しています。
VisualStudioは2010よりも2008にした方がいいでしょうか？
275:デフォルトの名無しさん
12/01/06 20:23:04.24
C初心者にはきついと思うんだが…
とりあえず 'cutil_inline.h'のある場所を見つけて
そこを-I /'cutil_inline.hのある場所'と指定する
意味わからなければCのコンパイルを勉強すること
276:デフォルトの名無しさん
12/01/06 20:39:53.25
274ではありませんが､「Cのコンパイル」の勉強にオススメの書籍とかありましたら紹介していただけるとうれしいです
277:デフォルトの名無しさん
12/01/06 20:55:31.81
275だけど、お薦めは他の人に任せるけど
一つ言えるのはCUDAはC/C++より数段難しい
C/C++の質問をしてるようでは先が思いやられる、と正直思うんだが…
超低レベルな質問もOKなくだすれだけど、その程度は自力で調べる能力が必要ではないだろうか？
278:デフォルトの名無しさん
12/01/06 21:22:02.03
>>275
>>274です。ありがとうございます。
今実行できる環境にないので後で試してみます。
また結果報告しにきます。
279:デフォルトの名無しさん
12/01/07 12:14:43.10
一時期は同価格帯のCPUとGPUを比較してGPUは５０倍の性能があるって言われてたけど
今１０～１５倍程度しか無くない？
3930k
５万円
158GFlops
GTX590
６．３万円
2480GFlops
GTX 580
３．８万円
1500GFlops
次に出ると言われてるGTX700系は500系の２倍の性能らしいけどそれでも５０倍には遠く及ばない。
この「同価格帯でGPUはCPUの１０～１５倍程度の性能しかない」という評価は正しいだろうか？
280:デフォルトの名無しさん
12/01/07 12:43:15.62
>>279
それは単にCPUは1コアとして比較しているからGPUを使えばうん十倍と言われるんだよ。
単純にメモリ帯域で3～4倍、演算器としては、単純にAVXの8並列x6コアxクロック2倍で比較すると5倍にしかならない。
CPUは大きなキャッシュがあるから局所性の高い処理の場合はメモリ帯域の差が縮まるので、10倍差がつくならいい方。
もちろんアプリによるけど、CPU向けのコードがカリカリに最適化されていた場合は思ったほど差がつかない印象。
それでも3～5倍位は早くなるからいいんだけどね。
問題はメンテナンスだな。。。CUDAのみと割り切れればいいんだけど、なかなかそうは行かないからなあ。
281:デフォルトの名無しさん
12/01/07 16:09:01.15
>>279
GTX580で500GFLOPS程度。そのデータはおかしい
CPUなんてE8500で25GFLOPSくらいだから20倍程度だな。
もちろん最近のコア多いやつならもっと差は縮まるだろうな。
それに理論値の話だから多数のスレッドでまわしたらコンテキストスイッチの
オーバーヘッドとかGPUプログラミングは難しくて効率上げるの難しいとか
そういうので実効は変わってくるんじゃね？
282:デフォルトの名無しさん
12/01/07 16:14:51.24
あ、同一価格帯か、E8500が2万しなかったことを考えると
580にそろえて2倍したとして10倍程度だから、まあおおざっぱには
「同価格帯でGPUはCPUの１０～１５倍程度の性能しかない」は正しいんじゃないかと思う。
580と同時期の4万弱のCPUならもっと差は縮まるだろうし。
283:デフォルトの名無しさん
12/01/07 16:22:20.40
32nmと40nmで比較とか
284:デフォルトの名無しさん
12/01/07 21:34:45.46
void *p;
CUDA_SAFE_CALL(cudaSetDevice(0));
CUDA_SAFE_CALL(cudaMalloc(&p, 16));
でcudaMallocからcudaError_enum例外が飛ぶのですが
原因はなにが考えられますか？
ちなみにSDKのサンプルは動きます。
285:284
12/01/07 22:23:52.14
異なるディレクトリに複数のCUDAのライブラリが入ってて
古いバージョンが参照されてるのが原因だった・・・
286:デフォルトの名無しさん
12/01/08 14:34:35.46
CUDA4.1でcomputeProfはnvvcになったんですか？
だいぶ代わってしまって驚きますた
287:デフォルトの名無しさん
12/01/08 14:35:50.30
nvvp
288:デフォルトの名無しさん
12/01/08 20:50:11.30
>>282
GPUの論理性能ってFMAとかMADとかの積和で倍になっているから、
GTX 580 1500GFlopsとかって、それらを使わなければ750GFlopsとかになるんだよな。
だから実効れべるではもっと縮まるんだろうな。それでも5倍以上差がでるとかなり有効なんだけど。
289:デフォルトの名無しさん
12/01/09 10:37:28.98
個人で買うやつも居るのかね？
URLﾘﾝｸ(page16.auctions.yahoo.co.jp)
290:デフォルトの名無しさん
12/01/09 12:04:44.39
いる。というか俺が買う。
291:デフォルトの名無しさん
12/01/09 13:09:16.68
こういう需要が限られていてかつ堅くて高い製品はamazonに出した方が高く売れるんじゃないかねえ。
292:デフォルトの名無しさん
12/01/09 18:01:50.03
>>275
274です。
-Iオプションを付けてヘッダファイルのある場所を指定してみましたが、相変わらず同じエラーです。
C:\cuda_practice\test>nvcc -o matrix_gpu.exe matrix_gpu.cu -I C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.0\C\common\inc
matrix_gpu.cu
matrix_gpu.cu(5) : fatal error C1083: Cannot open include file: 'cutil_inline.h': No such file or directory
自分でも試行錯誤してみますが周りにcudaがわかる人がいなくて行き詰ってます・・・。
293:デフォルトの名無しさん
12/01/09 18:20:34.55
もっとコンパイル早くなんねーかなぁ…
thrust::sortとか使うようにしただけでコンパイルに十数秒かかってイライラする
294:デフォルトの名無しさん
12/01/09 20:40:26.69
速いマシンでやる。
295:デフォルトの名無しさん
12/01/09 20:45:40.42
nvccコンパイラをGPGPUで実装すればいい！
CUDAがオープンになるどさくさにまぎれてそういうのを作る人が出るかもしれないこともないかも
296:デフォルトの名無しさん
12/01/09 21:57:43.85
コンパイラの動作は現状GPU向きとは言えないんじゃね？
GPU自体が単純な性能向上だけではない、思いもよぬ進化でもしない限り
297:デフォルトの名無しさん
12/01/09 22:54:12.52
コンパイルなんて依存関係ありまくりで、並列処理が苦手な最たるものだからまずやる意味がないだろう。
CUDAにすれば何でも速くなると思ってんじゃね？
298:デフォルトの名無しさん
12/01/10 17:15:02.61
sharedメモリって同一WARP内でのやり取りだったら同期なしで大丈夫なのでしょうか？
たとえば全部で32スレッドで次のkernelを実行した場合、
WARP内で一斉にsharedメモリに書きに行くので同期しないで大丈夫かと
思ったのですが、実際にはうまくいきません。
globalメモリに読みに行く段階でコアレッシングが発生していないので
それが原因なのでしょうか？
どなたか教えてください。
__global__ void kernel(float *g_v, float *g_x){
float x = 0.0f;
int i = blockDim.x * blockIdx.x + threadIdx.x;
__shared__ float s_v[32];
s_v[i] = g_v[i+i%3];
__syncthreads(); // これが必要かどうか？
x = s_v[(i+3)%32];
g_x[i] = x;
}
299:デフォルトの名無しさん
12/01/10 17:35:07.81
>>292
参照パスを ""で囲む♪
300:デフォルトの名無しさん
12/01/10 17:39:16.65
>>298
iが３２以上になることはないかい？？
ブロックの数は？？
301:274
12/01/10 17:52:39.67
>>299
書き込みしてからそのミスに気付きました・・・。
それを直してもうまくPathが通らなかったので、CUDA4.0、VisualStudio2008でやることにしました。
302:デフォルトの名無しさん
12/01/10 17:56:37.24
>>298
fence入れてないからじゃね。
303:298
12/01/10 17:59:16.44
>>300
ごめんなさい、書いたコードが一部誤ってました。
iが32以上になることがあります。
ブロック数は数100程度になります。
このときは、下のソースのようになると思うのですが、
やはりうまくいきません。
__global__ void kernel(float *g_v, float *g_x){
float x = 0.0f;
int i = blockDim.x * blockIdx.x + threadIdx.x;
__shared__ float s_v[blockDim.x];
s_v[threadIdx.x] = g_v[i+i%3];
__syncthreads(); // これが必要かどうか？
x = s_v[(threadIdx.x+3)%32];
g_x[i] = x;
}
304:デフォルトの名無しさん
12/01/10 18:12:32.43
>>303
コンパイル後のバイナリを見たわけじゃないから予測だが、
s_v[threadIdx.x] = g_v[i+i%3];
x = s_v[(threadIdx.x+3)%32];
g_x[i] = x;
は
1.1:tmp=g_v[i+i%3]; //グローバルアクセスなので遅い
1.2:s_v[threadIdx.x]=tmp; //
2.1:x = s_v[(threadIdx.x+3)%32];
3.1:g_x[i] = x;
みたいに解釈される。
で、単純な式だからコンパイラは
s_v[threadIdx.x]とs_v[(threadIdx.x+3)%32];
が必ず別のアドレスになる事を検知して
1.1 2.1 3.1 1.2のように命令を入れ替えてしまう。
だから__threadfence_block()がいるはず。
305:298
12/01/10 18:17:13.87
>>302 >>304
ありがとうございました。
試してみたいと思います。
306:デフォルトの名無しさん
12/01/10 19:45:54.97
>> 305
だいたい x = s_v[(threadIdx.x+3)%32]; の前にすべてのs_v[??}が定義されていないといけないだろ？？
スレッドが３２に設定されているのならば (threadIdx.x+3)%32 は０か１にしかならないが．．．
307:デフォルトの名無しさん
12/01/10 21:10:42.27
>>301
-lcutil32 とその参照パスを追加すれば Visual Studio 2010 でもコンパイル，実行できる♪
ちなみに私の場合は
nvcc matrix.cu -I "C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.0\CUDALibraries\common\inc" -L "C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.0\CUDALibraries\common\lib\Win32" -lcutil32
tmpxft_000024b8_00000000-3_matrix.cudafe1.gpu
tmpxft_000024b8_00000000-8_matrix.cudafe2.gpu
matrix.cu
tmpxft_000024b8_00000000-3_matrix.cudafe1.cpp
tmpxft_000024b8_00000000-14_matrix.ii
C:\cuda> a
Processing time: 1936.354858 (msec)
Press ENTER to exit...
308:デフォルトの名無しさん
12/01/10 21:12:58.21
ちなみに家のはGeForce GT 220なので非力．．．．
309:306
12/01/10 22:21:29.79
間違えた
↓
(threadIdx.x+3)%32 は０か１にしかならないが．．．
３２で割った余りであった．．．
310:298
12/01/11 09:16:02.50
>>306 >>309
いろいろ試してみたところ、とりあえず上記の問題は解決しました。
回答ありがとうございました。
311:デフォルトの名無しさん
12/01/11 11:25:42.40
すごい。
見た感じ、ちょっと不安だったから
参考になったみたいでよかった。
312:274
12/01/11 19:30:44.93
>>307
VS2008でもうまくできませんでした。
再度VS2010を入れてみるも手順通りいかず、当然実行もできませんでした。
リカバリしたのでこれからもう一度2010でやってみます。もう心が折れそうです・・・。
ちなみに2010expressとprofessionalの違いで実行できないこととかあり得ますか？
別のPCでprofessionalで試してみてくれた方がいて、うまくできたようなのですが・・・。
313:デフォルトの名無しさん
12/01/11 20:43:20.68
>>312
どんなエラーがでますか？
そもそもnvccでコンパイルするのだから，Visual Studioは関係ないでしょう？
VisualStudioのないLinux上でもコンパイル，実行できるはず．
もう一度やるのならばライブラリ lcutil32を付ければ良いかと（３２ビットならば）
314:274
12/01/12 01:51:21.33
>>313
VisualStudioであれば多少使ったことがあるので導入しました。
少なくともwindowsでは「nvccがVisualStudioのcl.exeを必要とする」と複数のサイトに書いてあるようです。
URLﾘﾝｸ(gpu.fixstars.com)
URLﾘﾝｸ(exth.net)
エラーはプロジェクトをビルドしたときに「CL.exeがないからビルドできない」というものでした。
そもそもcl.exeは32ビットのコンパイラらしいのですが、64ビットPCでも使うのでしょうか？
また以前、いい感じのところまでいったときのエラーは
C:\cuda_practice\test>nvcc -o matrix_gpu.exe matrix_gpu.cu
matrix_gpu.cu
tmpxft_00000bc0_00000000-3_matrix_gpu.cudafe1.gpu
（省略）
symbol __imp_cutCheckCmdLineFlag referenced in function "void __cdecl __cutilExi
t(int,char * *)" (?__cutilExit@@YAXHPEAPEAD@Z)
matrix_gpu.exe : fatal error LNK1120: 6 unresolved externals
でした。
どうやらリンカーの設定を見直せとのことだったのでURLﾘﾝｸ(www.scribd.com)を参考に
VisualStudioでプロジェクトのプロパティから「リンカ > 入力 > 追加の依存ファイル」にcutil32.lib, cutil64.lib を追加。
しかし「cutil64.libが見つからない」とのエラー。
実際に探してみてもcutil64.libがどこにも見つかりませんでした。
今まで3回もリカバリしてやってみましたが、どうもPathがうまく通せていない気がします。
もうすでにリカバリしてしまったので、明日また環境構築をはじめからやり直します。。
Linuxの環境はないんです、すみません・・・。
今の環境はWindows7 64bitです。
詳しくはこちら>>274に書きました。
315:デフォルトの名無しさん
12/01/12 07:17:27.30
>>314
はじめはVSのプロジェクトを構築するのではなく、VSの（32bit用の）コマンドプロンプトを立ち上げて、そこで
コマンドを打ち込み、コンパイル、実行した方がよいのでは。
これがうまく行ってから、プロジェクトの構築をした方がよいでしょう。
cutil64.libがないと言うことは64bit用のSDKがインストールされていないのでは？
ここにあるはず。
C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.0\CUDALibraries\common\lib\x64
プロジェクトを構築するときに32bit, 64bit、あるいは両方の実行ファイルを作るかどうか設定できますよ。
316:デフォルトの名無しさん
12/01/12 07:18:06.91
>>314
64bit版のcl.exeがある。
見つからないなら-ccbinオプションでcl.exeの場所を指定。
cutil32.libとcutil64.libを両方指定してはいけない。どちらか必要なものを。
で、cutil64は確か自分でビルドする必要がある。
SDKのディレクトリの中にプロジェクトファイルがあるからそれをビルド。
317:デフォルトの名無しさん
12/01/12 08:36:26.90
URLﾘﾝｸ(codepad.org)
上記のプログラムが期待と違う結果を出力するのだが原因がわかる方がいれば教えていただけないだろうか．
（CUDA 3.2 + GeForce 320M + Fedora 14 環境と CUDA 3.2 + Tesla C1060 + CentOS 5.5 環境でテスト）
-- 期待した出力 --
1.00
0.00
1.00
0.00
1.00
0.00
1.00
0.00
-- 実際の出力 --
0.00
0.00
0.00
0.00
0.00
0.00
0.00
0.00
ちなみに，11行目の右辺を 2.0f に変更すると期待通りの結果を出力してくれる．
-- 出力 --
1.00
2.00
1.00
2.00
1.00
2.00
1.00
2.00
318:317
12/01/12 08:37:35.44
すみません．
>>317のコードのURLを間違えました．
誤: URLﾘﾝｸ(codepad.org)
正: URLﾘﾝｸ(codepad.org)
319:デフォルトの名無しさん
12/01/12 09:55:39.89
やってみたけど、問題は再現されなかった
cuda_funcの中に
printf(" %d %d %f\n",threadIdx.x,threadIdx.x & 1,a[threadIdx.x]);
を入れて、チェックしてみれば？？
ちなみにこんな結果となる
0 0 1.000000 <- cuda_funcの中では問題はなさそう
1 1 0.000000
2 0 1.000000
3 1 0.000000
4 0 1.000000
5 1 0.000000
6 0 1.000000
7 1 0.000000
1.00 <-- 外でも問題はなかった
0.00
1.00
0.00
1.00
0.00
1.00
0.00
320:デフォルトの名無しさん
12/01/12 10:24:08.53
>>317
同じく問題再現されず（CUDA4.1 + GF560Ti）
そのデバイス両方とも compute capability 1.x だし
2.x 環境で試してみたらどうだ
321:317
12/01/12 10:56:51.02
>>319 >>320
コメントありがとうございます
現在の私の環境ではカーネル関数内で printf は使えないですが
変なことになるのは私の環境だけのようなので Compute capability 2.x 環境で後日試してみます
322:デフォルトの名無しさん
12/01/12 11:18:04.11
ちなみにcuPrintf.cuh とcuPrintf.cu をネット上からひらってくれば、printfと同等の関数（cuPrintf）が使えますよ♪
323:デフォルトの名無しさん
12/01/12 11:57:42.71
CUDA 3.2のコード出力に問題がある気がします。
カーネル部分をどうも
a[threadIdx.x] = (float)(FUNC(-(threadIdx.x & 1));
に変換しようとしている気配が。
FUNCはISET.S32命令による処理だけどそれがミスっているとしか思えない。
自分で以下のように書くか、あるいはFermi以降またはCUDA4.0以降を使うべし。
a[threadIdx.x] = (float)(1 - (1 & threadIdx.x));
ちなみに実行はしてないので同じ問題が起こるかどうかは知りません。
324:274
12/01/12 13:42:05.81
>>315 >>316
ありがとうございます。
とりあえず315さんの通りコマンドライン上で実行できるところを目指そうと思います。
cl.exeを指定するオプションがあるんですね。
自分はコンパイラオプションの勉強不足のようです・・・。
cutil64は自分でビルドするんですか、それは今までやってなかったと思います。
恐縮ですが、これから環境を整えるので何か参考になるサイトがあれば教えていただけますか？
いくつかサイトを見て設定したのですが、どれもうまくいかなかったので・・・。
何度もすみません。よろしくお願いします。
325:デフォルトの名無しさん
12/01/12 16:35:09.29
質問です。
同じkernel関数をfor文で何度も実行するのですが、
回数によって1回あたりのkernel関数の実行時間が異なるみたいです。
具体的には
for(i=0; i<10000; i++){
kernel_foo <<<grid_size, block_size >>>(...);
}
とすると300,000msかかるのが
for(i=0; i<100; i++){
kernel_foo <<<grid_size, block_size >>>(...);
}
では1,000msくらいになります。同じ関数を実行しているので
単純に考えれば10000回ループを回したときは100回の100倍時間がかかるはずですが、
実際には300倍程度になりました。
もう少し細かく見てみると、最初60回程度はkernel関数を1回実行するのに0.04ms程で済んでいたのですが
それ以降は1回あたり25ms程になっていました。
似た事例として
URLﾘﾝｸ(d.hatena.ne.jp)
という記事を見つけましたが、この記事内でも原因はよく分かっていません。
どなたか詳細をご存じないでしょうか？
326:デフォルトの名無しさん
12/01/12 17:29:28.93
>> 324
参考になるやら、ならないやら
URLﾘﾝｸ(www.fxfrog.com)
URLﾘﾝｸ(www.ademiller.com)
URLﾘﾝｸ(cudasample.net)
URLﾘﾝｸ(tech.ckme.co.jp)
URLﾘﾝｸ(cudasample.net)
URLﾘﾝｸ(others2.blog.so-net.ne.jp)
>> 325
1． カーネルの中の演算数がiによって異なる
2. 60回目で解が発散し NaNになっているが、お構いなしに計算は続けられている
3. カーネルからCPUにメモリーコピーしている時間が紛れ込んでいる
かな？
327:274
12/01/12 17:37:04.49
>>326
ありがとうございます。
全て拝見しましたが既読のものでした。やはり普通に設定すればできるようですね・・・。
今格闘しているので、また報告します。
328:325
12/01/12 17:50:39.27
>>326
ありがとうございます。
全項目について確認してみましたが、
1. 演算数はどれも一定
2. ループ回数によらず正しい値が得られている
3. 測定時間中にcudaMemcpy等の関数は呼ばれていない
という状態です。
329:デフォルトの名無しさん
12/01/12 18:22:29.85
>>328
1プロセスで2つのカーネル実行してるなら順番入れ替えて試してみた？
最初の実行は起動コストだかで時間かかるのが普通なんだが
330:デフォルトの名無しさん
12/01/12 19:29:37.96
時間の計測が間違っている。
cudaThreadSynchronize()を読んだあとに終了時刻を取得する必要がある。
1024個までのカーネル呼び出しはキューに積むだけですぐに次に進んでしまう。
30000回ループの方はキューが詰まるので(30000-1024)回くらいの実行完了を待つことになり、比較的妥当な計算時間に見えるはず。
もちろん正確ではないのだが。
331:325
12/01/12 19:35:41.01
>>329-330
回答ありがとうございました。
cudaThreadSynchronize()を入れるのを失念しておりました。
332:デフォルトの名無しさん
12/01/12 20:59:37.22
cudaDeviceSynchronize()
がナウいらしい
333:sage
12/01/13 08:22:49.98
>>330
便乗して質問ですが、
キューに一度に入るカーネル数が1024個で、
それ以上はカーネル呼び出しが待ちになるというのは、
CUDAのマニュアルかどこかに書いてありますか？
334:デフォルトの名無しさん
12/01/13 09:57:50.47
>>333
実験したらそうだったけど、文書化されているかどうかは知らない。
これだけ多ければ普通は問題ないし、バージョンアップで変わるかも。
335:デフォルトの名無しさん
12/01/13 11:09:53.09
>>334
了解しました。
どうも有難うございました。
336:デフォルトの名無しさん
12/01/13 16:58:47.32
それにしても きゅー にカーネルが溜まっているからと言って，GPUの処理能力が落ちるとは思えないぞ
337:デフォルトの名無しさん
12/01/13 17:17:25.93
きゅー にカーネル田丸？
それってstream使ってる時のこと？
文脈呼んでないからわからんけど
338:デフォルトの名無しさん
12/01/13 20:07:11.66
キューに空きがあれば積んで即制御が戻る。
空きが無ければ空くまで待ってる。
詰まってるから遅くなるわけではなく、遅い方が正しい実行時間なだけ。
339:デフォルトの名無しさん
12/01/18 22:19:57.72
CUDA4.0からデバイスエミュレーションがなくなったらしいけど
実機で実行するとOSごと固まることがある
みんなどうしているのでしょうか
340:デフォルトの名無しさん
12/01/19 00:43:55.87
デバイスを物理的に叩き割る
341:デフォルトの名無しさん
12/01/19 04:33:29.60
2.3くらいからなかったような気が
342:デフォルトの名無しさん
12/01/19 15:25:40.77
>>339
具体的には？？
343:デフォルトの名無しさん
12/01/20 18:49:16.04
>>342
メモリアクセスのバグがあるとWindowsが固まって動かなくなったり
ディスプレイにノイズが出てくる
OSが画面表示に使っている領域を壊しているようだけど
普通は起こらないことなんでしょうか？
344:デフォルトの名無しさん
12/01/21 10:42:26.35
私の経験では、CUDAのプログラムをチェックするために printf を使って変数を書き出したところ、
その量が多すぎて、画面が真っ暗になったり、PCがシャットダウンしたことがあった。
書き出し量を減らしたところ、問題はなくなった。
こんなことかな？？
345:344
12/01/21 11:00:22.89
追加です。
グラフィックカードを2枚刺して、一枚をCUDA用、もう一枚をディスプレイ用にすればよいとのことを聞いたことがあるが、
私のはGTX５９０でカードが2枚、入っているとのこと。
CUDAの計算に Device(0)にしても(1)にしても、前述の問題が起きたので、枚数には関係ないこともあるらしい？？
346:デフォルトの名無しさん
12/01/22 05:18:00.51
printf使えるとか便利になったわ
347:デフォルトの名無しさん
12/01/22 16:18:23.06
>>344
fermiではないのでprintfは使えないです。
GTX260なんですけど、それが悪いのかも。
計算用GPUを別に刺すってのは考えていました。
ただ、そうするべきという常識みたいなものがあるのかなっていう疑問があったので。
あとこういう状態でデバイスエミュレーション外したならひどい話だなと思ったのですが
みなさんあまり起きていないようで。
348:デフォルトの名無しさん
12/01/22 16:23:00.85
グローバルメモリの添え字計算計算が間違っていたりで確保した範囲外にアクセスすると
Windowsが固まったりディスプレイにノイズが出てくることがあるって状態です。
349:やんやん ◆yanyan72E.
12/01/22 16:48:55.32
グラフィックカード二枚差しにすると、
nSghtでハードウェアデバッグができるとかじゃなかったっけ？
350:デフォルトの名無しさん
12/01/22 17:02:34.79
YES
もしくはPC2台で片方をデバッグ用にしても良い。
実際に動かすGPUがCUDA対応なら、開発側は非nVIDIAでもOK
351:デフォルトの名無しさん
12/01/23 09:39:54.36
streamについて分からない点があるので
どなたか教えてください。
以下で「カーネル」はカーネル関数の呼び出し、
「転送」はcudaMemcpyAsync、()内はstreamだとします。
(質問1)以下の状態のとき、転送(1)は、カーネル(1)と同じstreamなので
実行できませんが、キュー内でそれより後にある転送(2)は実行できますか？
カーネル(1) 実行中
転送(1) キュー内の最初のタスク
転送(2) キュー内の2番目のタスク
(質問2)以下の状態のとき、転送(2)は実行できますか？
(キュー内で転送(2)の前にstream(0)のタスクがあります）
カーネル(1) 実行中
カーネル(0) キュー内の最初のタスク
転送(2) キュー内の2番目のタスク
(質問3)以下の状態のとき、転送(1)は実行できますか？
(実行中のタスクはstream(0)です)
カーネル(0) 実行中
転送(1) キュー内の最初のタスク
「CUDA C Programming Guide」を読んでもよく分かりませんでした。
宜しくお願いします。
352:デフォルトの名無しさん
12/01/23 10:54:29.21
1はFermiならOK。古い世代ではダメ。
2と3は絶対ダメ。stream0は前後をぶった切る役目がある。
353:デフォルトの名無しさん
12/01/24 08:17:05.28
>>352
1の動作がFermiとそれ以前で異なるのは知りませんでした。
回答どうも有難うございました。
354:デフォルトの名無しさん
12/01/24 10:48:47.33
Fermiでプロファイラのtimestampを使ってテストしたところ
カーネル(1) 実行中
転送(HostToDevice)(1) キュー内の最初のタスク
転送(HostToDevice)(2) キュー内の2番目のタスク
だと転送(2)はカーネル(1)と同時に動かず、
カーネル(1) 実行中
転送(HostToDevice)(1) キュー内の最初のタスク
転送(DeviceToHost)(2) キュー内の2番目のタスク
だと転送(2)はカーネル(1)と同時に動きました。
355:デフォルトの名無しさん
12/01/24 19:53:55.48
コンパイル時に以下のエラーメッセージが出ているのですが、
ptxas error : Entry function '...' uses too much local data (0x16054 bytes, 0x4000 max)
これはローカルメモリーと何か関係あるような気がするのですけど、
ちょっとわからないので教えて頂けないでしょうか。
宜しくお願いします。