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

[表示 : 全て最新50 1-99 101- 201- 301- 401- 501- 601- 701- 2chのread.cgiへ]
Update time : 05/18 03:52 / Filesize : 192 KB / Number-of Response : 722
[このスレッドの書き込みを削除する]
[＋板最近立ったスレ＆熱いスレ一覧 : ＋板最近立ったスレ／記者別一覧] [類似スレッド一覧]


↑キャッシュ検索、類似スレ動作を修正しました、ご迷惑をお掛けしました

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 名前：デフォルトの名無しさん mailto:sage [2011/08/23(火) 22:11:49.24 ]: 関連サイト
CUDA
www.nvidia.co.jp/object/cuda_home_new_jp.html

CUDAに触れてみる
chihara.naist.jp/people/STAFF/imura/computer/OpenGL/cuda1/disp_content

cudaさわってみた
gpgpu.jp/article/61432191.html

CUDA のインストール
blog.goo.ne.jp/sdpaninf/e/9533f75438b670a174af345f4a33bd51

NVIDIAの「GeForce 8800 GT(G92)」と次に控える64-bit GPUアーキテクチャ
pc.watch.impress.co.jp/docs/2007/1031/kaigai398.htm

CUDAを使う
tech.ckme.co.jp/cuda.shtml

NVIDIA CUDAを弄ってみたその2
dvd-r.sblo.jp/article/10422960.html

CUDAベンチ
wataco.air-nifty.com/syacho/2008/02/cuda_2044.html

KNOPPIX for CUDA
www.yasuoka.mech.keio.ac.jp/cuda/
3 名前：デフォルトの名無しさん mailto:sage [2011/08/23(火) 22:12:38.55 ]: 宣伝になっちゃうと難なので、一応分けて。いろいろ載っているサイト。

ttp://www.softek.co.jp/SPG/Pgi/TIPS/para_guide.html
4 名前：デフォルトの名無しさん mailto:sage [2011/08/23(火) 22:26:00.72 ]: おい、pertだろ
まちがえんな
5 名前：デフォルトの名無しさん mailto:sage [2011/08/23(火) 22:48:26.89 ]: >>2
個人blogのは情報古すぎなような…
「cudaさわってみた」はリンク切れだし。
6 名前：デフォルトの名無しさん mailto:sage [2011/08/25(木) 19:36:34.86 ]: >>4
やっぱりCUDAやってる人は荒んでるんだな
7 名前：デフォルトの名無しさん mailto:sage [2011/08/26(金) 07:04:46.69 ]: ほんとCUDAらないことで騒ぐよな
8 名前：忍法帖【Lv=40,xxxPT】 mailto:sage [2011/08/26(金) 14:04:14.48 ]: ケンカするのはやめてCUDAさい＞＜
9 名前：デフォルトの名無しさん mailto:sage [2011/08/26(金) 20:59:54.29 ]: そうだよ。喧嘩はやめ…GeForんｹﾞﾌｫﾝ
10 名前：デフォルトの名無しさん mailto:sage [2011/08/26(金) 21:57:35.41 ]: まだスレが始まっTeslaいないのに
11 名前：デフォルトの名無しさん mailto:sage [2011/08/27(土) 00:03:43.41 ]: ときに口論になるのもみんな自分の技術にそれだKeplerイドを持ってるってことだ。
12 名前：デフォルトの名無しさん mailto:sage [2011/08/27(土) 00:09:15.96 ]: amazon.comで検索して出てくるCUDA本ってことごとく出版がずいぶん先か延期になったりするな
なんでだろ
13 名前：デフォルトの名無しさん mailto:sage [2011/08/27(土) 21:26:38.44 ]: 無効とは情報の登録ポリシーが違う…ような気がする。
それがAmazonのポリシーなのか出版社のポリシーなのかは知らないけれど。

多分出版社のサイト直接見た方が正確。
14 名前：デフォルトの名無しさん mailto:sage [2011/08/27(土) 23:28:49.51 ]: cudaの本が必要、って訳でも無いが出版されるなら買ってしまうと思う
15 名前：デフォルトの名無しさん [2011/08/31(水) 20:00:08.40 ]: 4.0になってdeviceemuが使えんのだが、CUDA載せてないPCで動かす方法ない？
16 名前：デフォルトの名無しさん mailto:sage [2011/09/01(木) 07:46:58.96 ]: 3.2を使う。
17 名前：デフォルトの名無しさん mailto:sage [2011/09/01(木) 09:05:53.29 ]: >>16 アホしかおらんのか
18 名前：デフォルトの名無しさん mailto:sage [2011/09/01(木) 09:16:16.30 ]: >>17
現実的じゃないか。どうせemuで正確なエミュレーションができるわけじゃないし。
19 名前：デフォルトの名無しさん mailto:sage [2011/09/01(木) 09:25:26.25 ]: >>17
4.0でdeviceemu使う方法なんてねーよ、ってことだよ。
言わせんな恥ずかしい
20 名前：デフォルトの名無しさん mailto:sage [2011/09/01(木) 10:02:13.87 ]: >>15
Cuda x86
21 名前：デフォルトの名無しさん [2011/09/01(木) 15:33:21.69 ]: 理屈の上ではKnights Ferryでも動かせるのかな？
22 名前：デフォルトの名無しさん mailto:sage [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 名前：デフォルトの名無しさん mailto:sage [2011/09/01(木) 17:40:47.67 ]: >>22の続きです
assumed oldの値をcpu側に持ってきて表示してみたら、
oldの値(atomicCASの返す値)とassumedの値が末尾の数ビットが
異なってました。50回ほどループしてやっとビットレベルで値が一致して
ループを抜けているようでした。その間に過剰に加算が行われていたようです。
24 名前：デフォルトの名無しさん mailto:sage [2011/09/01(木) 18:47:40.63 ]: カードの計算が正しいかチェックするベンチマークってありませんか？

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

計算機が計算を間違うなんてありえないのに…
25 名前：デフォルトの名無しさん mailto:sage [2011/09/01(木) 19:11:55.88 ]: Teslaの存在意義全否定デスカ
26 名前：デフォルトの名無しさん mailto:sage [2011/09/01(木) 19:34:24.88 ]: >>24
過去スレを読むとわかってもらえると思うが、みんなそれを求めているんだ。
そのチェックにおける最先端はおそらく長崎大学。
27 名前：デフォルトの名無しさん mailto:sage [2011/09/01(木) 19:41:57.01 ]: 1スレッドの実行がどのコアかがわかる仕組みがどのような方法か分かればいいわけだよな
誰か調べてみろよ
28 名前：デフォルトの名無しさん mailto:sage [2011/09/01(木) 23:41:13.82 ]: >>22,23
もしかしてブロック辺りのスレッド数が32以下なら動いたりする？
29 名前：デフォルトの名無しさん mailto:sage [2011/09/01(木) 23:44:25.02 ]: >>27
何がしたいのか分からんがそれは簡単だ。
30 名前：デフォルトの名無しさん mailto:sage [2011/09/01(木) 23:56:07.27 ]: >>29
コアと計算結果の表が出来上がればコア間の相違が分かるじゃん
やり方知ってるなら教えてください
31 名前：デフォルトの名無しさん mailto:sage [2011/09/01(木) 23:58:10.21 ]: 煽っても教えてやらん
29じゃないけどｗ
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 名前：デフォルトの名無しさん mailto:sage [2011/09/02(金) 02:29:09.05 ]: ここはおまえの日記帳だから好きに使いなさい
34 名前：デフォルトの名無しさん mailto:sage [2011/09/02(金) 11:08:38.52 ]: >>28
ブロックあたりのスレッド数を32以下にして試してみましたが、
結果は変わらなかったです。
35 名前：デフォルトの名無しさん mailto:sage [2011/09/02(金) 18:20:08.65 ]: >>23
とりあえず、プログラミングガイドのコードそのままだし、
実際アルゴリズム自体には問題ないから、
意図どおり動かないとしたら、ドライバかSDKのバグか
ハードの故障じゃね。
まあ、ここで詳細なバージョンとか晒してみたり、
PTXやアセンブリを貼ってみるのもいいけど、
本家のフォーラムで聞いてみたほうが早いんじゃね。
36 名前：デフォルトの名無しさん mailto:sage [2011/09/05(月) 23:02:56.52 ]: 長崎大の言う演算間違えって、具体的にどういう演算を間違えたのかの情報ってないの？
間違えますって情報だけ垂れ流して具体例出さんって、どういうことなんだべ。
37 名前：デフォルトの名無しさん mailto:sage [2011/09/05(月) 23:24:45.84 ]: ×演算間違え
○演算間違い
38 名前：デフォルトの名無しさん mailto:sage [2011/09/07(水) 18:20:00.85 ]: CUDAを使用するプログラムで巨大なデータを扱いたいのですが、nvcc でのコンパイル時に
relocation truncated to fit: R_X86_64_32S against `.bss'
とエラーが出て完了しません。gcc ではデータ量制限を解除するのに-mcmodel=medium というオプションがあるようですが、nvcc でこれにあたるものはないでしょうか。
39 名前：デフォルトの名無しさん mailto:sage [2011/09/07(水) 19:43:00.72 ]: 巨大データってどれくらい？
CUDAは基本グローバルメモリ使用の制約があるから”巨大”は決して得意ではないのだけど
40 名前：38 mailto:sage [2011/09/07(水) 22:05:04.34 ]: >>39
今確認したところ、要素が数十万個の配列を使用していて、それらの合計は200MBほどです。
そのうち、実際にGPUに計算させるためGPUに渡すデータは50MB程度です。
上記のエラーメッセージでググったところ、データ量制限が問題だろうということだったのですが…それほど大きなデータでもないですよね。
なんなんでしょうこのエラー。

プログラミングに関してはまだ知識が浅いためチンプンカンプンなこと言ってたらすみません。
41 名前：デフォルトの名無しさん mailto:sage [2011/09/07(水) 22:36:52.16 ]: どうせデータをソース中に書き込んでるんだろw
42 名前：デフォルトの名無しさん mailto:sage [2011/09/07(水) 22:42:26.34 ]: >37
日本語の間違えより、長崎大の言ってる間違いの詳しい情報くれよぅ…。
43 名前：デフォルトの名無しさん mailto:sage [2011/09/07(水) 22:50:33.71 ]: 要素数が大きすぎるなら、mallocを使ってみたら？
44 名前：デフォルトの名無しさん mailto:sage [2011/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 名前：デフォルトの名無しさん mailto:sage [2011/09/07(水) 23:11:07.95 ]: ファイル読み込みで、cudaMallocHost(だったかな？？)でメモリ確保すればいいじゃん
46 名前：38 mailto:sage [2011/09/07(水) 23:11:33.99 ]: >>43 >>44
gccと同様にオプションでエラー回避できないかな…と思ったのですが、配列の取り方を変えるほうが良さそうですね。
特に>>44は知りませんでした。
ありがとうございました！
47 名前：デフォルトの名無しさん mailto:sage [2011/09/07(水) 23:23:20.08 ]: 良く分からんが、-Xcompiler -mcmodel=mediumじゃ駄目なん？
-Xlinkerも要るかもしれないが。
48 名前：デフォルトの名無しさん mailto:sage [2011/09/08(木) 08:01:27.80 ]: >>44
普通はありえない。そんなことしたら、ポインタアクセスが一回余計に要るから遅くなるし、
キャッシュアウトしやすくなるから更に遅くなる。
まぁ、GPUに転送するのは1スライス分だけだろうからGPUに任せてCPUは暇こいていられるならいいかも知れんが。
49 名前：デフォルトの名無しさん mailto:sage [2011/09/08(木) 17:24:44.67 ]: >>44
それは無い。
・メモリも無駄に消費
・アクセス性能がた落ち
・バグの温床になる

つーかC++ならvectorで初期化時に
サイズ指定したほうがいろいろ便利
50 名前：デフォルトの名無しさん mailto:sage [2011/09/08(木) 17:40:17.98 ]: 配列でやる場合、
配列の添え字がINT_MAXを越える大きさならそうするしかないと思うが・・
51 名前：デフォルトの名無しさん mailto:sage [2011/09/08(木) 20:47:04.41 ]: キャッシャがアクセスがとドヤ顔されても・・・。
要素のデカイ配列が確保できんって話なわけだし。
52 名前：デフォルトの名無しさん mailto:sage [2011/09/08(木) 22:19:09.58 ]: コンパイル時の話なのに
ヒープのフラグメンテーションの訳無いだろう。
nvccからgccに自由にオプション渡すやり方は
知らんが、最悪c(++)部分のコードを分割して
gccでコンパイルすりゃ良い。

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

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

特にGPGPUで巨大な配列扱うような用途なら
配列の配列なんて転送も困るだろうに。
53 名前：デフォルトの名無しさん mailto:sage [2011/09/08(木) 22:30:02.29 ]: >大きな配列確保する前に小さな物についてShrink-to-fitやればデフラグされるから
>その後巨大配列を確保すれば良い。

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

最終的に使うデータについて、大体処理した順番を守って
ひたすら同容量確保してはコピーを繰り返せば、
こういうデータは全て前方に圧縮されて配置されるようになる。
55 名前：デフォルトの名無しさん mailto:sage [2011/09/09(金) 21:48:41.98 ]: ？
まあスレ違いなんで他でやってよ。
56 名前：デフォルトの名無しさん mailto:sage [2011/09/09(金) 22:24:20.88 ]: kernelへの引数に構造体は使えんの？
今やってみたらコンパイルでエラーになった
57 名前：デフォルトの名無しさん mailto:sage [2011/09/09(金) 22:30:19.36 ]: あ、間違えてただ
構造体引数使えるべな
58 名前：デフォルトの名無しさん mailto:sage [2011/09/14(水) 11:12:56.96 ]: >>15
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 名前：デフォルトの名無しさん mailto:sage [2011/09/15(木) 17:52:31.49 ]: cuda4.0、gtx580で
Maximum dimension of a grid( 65535 , 65535 , 65535 )
ってなってるんですがz軸のブロックの取れる数は1じゃなくなったんですか？
60 名前：デフォルトの名無しさん mailto:sage [2011/09/19(月) 22:05:37.25 ]: デバイス上の処理が長時間になる場合、途中経過を取得する方法ってないかな？
例えば１００万回のモンテカルロを実行して、途中の１０万回、２０万回が済んだらその情報を取りたいんだけど…
61 名前：デフォルトの名無しさん mailto:sage [2011/09/19(月) 22:31:31.07 ]: >>60
streamingでできないかな?
62 名前：デフォルトの名無しさん mailto:sage [2011/09/19(月) 22:52:18.06 ]: streamingってなに？
SMのSじゃないよね？
63 名前：デフォルトの名無しさん mailto:sage [2011/09/20(火) 05:00:39.03 ]: カーネル呼び出し時の<<<>>>の第3引数がstreamingの何かじゃなかったっけ？
俺は>>61を見てそういうのがあったなと思い出したレベルなので使い方は知らないけど。
64 名前：デフォルトの名無しさん mailto:sage [2011/09/20(火) 08:29:52.67 ]: 第3パラメータはsharedMemoryのサイズだよ。そこで指定しなくても使えるけど。
65 名前：デフォルトの名無しさん mailto:sage [2011/09/21(水) 22:11:25.09 ]: ごめん第4引数だったみたい。それとAsync系関数がうんたらかんたら。
66 名前：デフォルトの名無しさん mailto:sage [2011/09/22(木) 00:50:42.07 ]: あー、その辺だね。
67 名前：デフォルトの名無しさん mailto:sage [2011/09/22(木) 18:05:24.11 ]: 【トリップ検索】CUDA SHA-1 Tripper【GeForce】
hibari.2ch.net/test/read.cgi/software/1311428038/

すげえな
68 名前：デフォルトの名無しさん mailto:sage [2011/09/23(金) 11:16:26.70 ]: 個人でCUDAプログラミングやってうれしいのはその程度なのか？
と逆に不安になるんだが…
69 名前：やんやん ◆yanyan72E. mailto:sage [2011/09/23(金) 11:35:33.50 ]: 世の中なんでも膨大なデータをいかに早く処理するかが勝負の世界に
なってきているから、CUDAプログラミングの旨味はいくらでもあるだろ。
70 名前：デフォルトの名無しさん mailto:sage [2011/09/23(金) 12:04:06.11 ]: いつかやってみようと思って半年くらいこのスレROMってるが
いまだに開発環境のセッティングの仕方すら調べていない・・・
71 名前：デフォルトの名無しさん mailto:sage [2011/09/24(土) 00:02:19.96 ]: >>68
>>67くらいのコード書ける？
72 名前：デフォルトの名無しさん mailto:sage [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. mailto:sage [2011/10/04(火) 21:13:10.02 ]: -archコマンドってなんだそら？
compute capabilityはnvccなりcl.exeのオプションに渡すものであって、
-archコマンドなんてものはないよ。
77 名前：デフォルトの名無しさん mailto:sage [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 名前：デフォルトの名無しさん mailto:sage [2011/10/04(火) 23:39:11.06 ]: CUDA以前の話だな。コンパイラオプションの意味がわかってない。

というか、VCなら、カスタムビルド規則でCUDA Runtime API Build Ruleを追加したほうが早い
81 名前：忍法帖【Lv=40,xxxPT】 mailto:sage [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 名前：デフォルトの名無しさん mailto:sage [2011/10/05(水) 12:14:48.16 ]: -arch=compute_20
にすればいいんじゃねえの？
85 名前：デフォルトの名無しさん [2011/10/05(水) 12:17:32.55 ]: >>84
ありがとうございます。
おっしゃるように>>75のように追加するとエラーを吐きます。
追加の仕方が違うのでしょうか。
86 名前：80 mailto:sage [2011/10/05(水) 13:36:15.15 ]: VS2010は色々変わってるんだな。
知らんかったすまん。
とりあえず「CUDA VS2010」でググって出てくる一番上のサイト参考にして、CUDA4.0のビルドカスタマイズファイルを導入した後、プロパティからCode Generationをcompute_20,sm_20にしたらいけた
87 名前：デフォルトの名無しさん mailto:sage [2011/10/05(水) 15:09:51.57 ]: GPUのリソースを全部CUDAに割り当てるため，グラボを2枚さしてディスプレイ出力用とCUDA用に分けるというのはありますが
同じようなことをsandy bridgeとグラボ1枚でも出来るのでしょうか？
88 名前：デフォルトの名無しさん mailto:sage [2011/10/05(水) 15:36:59.85 ]: >>87
Z68マザーなら、マザー側DVIにディスプレイつないでおけばいい話じゃ？
CUDAで計算したデータをintelで表示したいなら、
1.DirectXかOpenGLのリソースをCUDAにバインドしてデータ書き込み
2.メインメモリに転送
3.intel側に転送
って形になるかなぁ。多分。
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を起動しているなら兎も角も。
190 名前：デフォルトの名無しさん mailto:sage [2011/12/06(火) 12:18:58.00 ]: cuファイルとcppファイルに分けるのはどういう意図があるんでしょうか？thrustって何に使えるんですか？
191 名前：デフォルトの名無しさん mailto:sage [2011/12/06(火) 12:26:24.32 ]: >>190
・nvccに任せておきたくない部分は、.cppに書く。
・抽象的に取り敢えず書くのに便利。
192 名前：デフォルトの名無しさん mailto:sage [2011/12/06(火) 16:39:18.49 ]: cuにkernel_foo()を書いて
cppに書いたfoo()からkernel_foo()を呼ぶのが定石かと思ってた
そうすればopenclに移行したりしてもocl_kernel_foo()を呼び出すようにすれば変更すくないし。
193 名前：デフォルトの名無しさん mailto:sage [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コ,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 名前：デフォルトの名無しさん mailto:sage [2011/12/07(水) 00:45:36.69 ]: +=で同じアドレスに同時書き込みしてるから
196 名前：デフォルトの名無しさん [2011/12/07(水) 01:09:59.92 ]: >>195
+=で同じアドレスに同時書き込みすると、内部で何が起こるんですか？
197 名前：デフォルトの名無しさん mailto:sage [2011/12/07(水) 01:27:31.21 ]: 競合状態が発生してんじゃないの？
atomic演算とか同期が必要だと思うよ。
512スレッドで同一アドレスの変数の読み書きしてんだから。

まず５１２個の要素の配列作って、添え字にスレッド番号（上のi）を指定して確認してみたら？
書籍ならcuda exampleも買って読むといいかもね
198 名前：デフォルトの名無しさん [2011/12/07(水) 02:05:18.21 ]: >>197
>競合状態が発生してんじゃないの？
>>１９５のコメと合わせて考えるに、なんとなく予想はしてましたけど、取り合いになってるんですね...

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

>まず５１２個の要素の配列作って、添え字にスレッド番号（上のi）を指定して確認してみたら？
それは分岐条件がきちんと実行しているのか見るためのテストをしたとき確認しました。
199 名前：デフォルトの名無しさん mailto:sage [2011/12/07(水) 07:39:11.25 ]: 根本的にプログラミングの基礎が抜けている悪寒。
200 名前：デフォルトの名無しさん mailto:sage [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 名前：デフォルトの名無しさん mailto:sage [2011/12/07(水) 14:10:39.88 ]: >>202
同期とロック
204 名前：デフォルトの名無しさん mailto:sage [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 名前：デフォルトの名無しさん mailto:sage [2011/12/07(水) 23:52:25.72 ]: この本でも読んでみると良い。
日本語でわかりやすい。

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の本が出たようだ。
www.amazon.co.jp/dp/4906608000/ref=cm_sw_r_fa_dp_454Zob03VNZ25
211 名前：デフォルトの名無しさん mailto:sage [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: 洋書
www.amazon.co.jp/CUDA-Application-Design-Development-Farber/dp/0123884268/
212 名前：デフォルトの名無しさん mailto:sage [2011/12/09(金) 05:54:08.89 ]: CUDAというかVisualStudioが分からない。鬱だ死のう。
213 名前：デフォルトの名無しさん mailto:sage [2011/12/09(金) 09:40:45.47 ]: VSがわからないんじゃなく、C++がわからないんだろ
214 名前：デフォルトの名無しさん mailto:sage [2011/12/11(日) 23:39:15.74 ]: fortranもここでいいかな？
pgiのコンパイラ買って、三重ループの前後に指示行追加してやったけどまったく速くならない。憂鬱。。
215 名前：デフォルトの名無しさん mailto:sage [2011/12/11(日) 23:56:05.05 ]: fortran使う人って量子論とかやってる人？
どの世界の人が使うのかしら？
216 名前：デフォルトの名無しさん mailto:sage [2011/12/12(月) 08:31:48.58 ]: fortranを使うのは過去の遺産のためだろう。
fortranを使えば速くなるわけでもないしなあ。
217 名前：デフォルトの名無しさん mailto:sage [2011/12/12(月) 09:39:06.64 ]: fortran懐かし過ぎる
先日、二度と使うことは無いと思って本を捨てたばっかりだわ
218 名前：デフォルトの名無しさん mailto:sage [2011/12/12(月) 20:09:34.35 ]: 高レベル言語使うのは時代の趨勢じゃないかな。

LAPACKはfortranで書かれている。
fortranはnvccよりもアプリケーションに近い。
CUDAを使うのにnvccのレベルでなきゃダメということは無いと思う。
逆に、nvcc使わない人をバカにするヤツは、本物のバカだと思う。
219 名前：デフォルトの名無しさん [2011/12/12(月) 20:20:14.55 ]: 道具にこだわることが目的となって、肝心の結果を出せない人間を馬鹿というでは？
220 名前：デフォルトの名無しさん mailto:sage [2011/12/12(月) 21:18:03.01 ]: TSUBAME2.0でnvccを使うのと京でfortran使うのと比べるとどうなんだろう。
京はSPARCだけで計算して、TSUBAME2.0の10倍くらいの速度みたいだ。

東工大は教育機関なのに対して、理化学研究所は結果を求められる独立行政法人。
221 名前：デフォルトの名無しさん mailto:sage [2011/12/12(月) 22:25:59.20 ]: 既存のコードを使う限りでは京のほうが速いだろう。
ただ、富士通のコンパイラがダメダメだと聞いたけど、
それは解消されたのかな？
222 名前：デフォルトの名無しさん [2011/12/12(月) 22:29:03.43 ]: 富士通のコンパイラは、8コアCPUへの割り当ては自動らしい。
これ、すごいことだと思うんだけど、
このコンパイラは完成しているのだろうか。
223 名前：デフォルトの名無しさん mailto:sage [2011/12/12(月) 22:46:19.38 ]: >>222
openMPに対応してたら各CPUのコアへの割り当ては普通にコンパイラ側でやると思うけど？
224 名前：デフォルトの名無しさん mailto:sage [2011/12/13(火) 11:47:52.86 ]: >>223
OpenMPがそんな事やるかよ。やるのはOSだぞ。
>>222が言っているのはNUMAのことだ。
現時点でも普通のLinuxならnumactlを使えばかなりのケースでノードの割り当てができるようになる。
CUDAを使う場合でも、複数GPUを使う場合に、MPIを使う場合に有効だ。
225 名前：やんやん ◆yanyan72E. mailto:sage [2011/12/13(火) 14:16:05.38 ]: なんだ？OpenMPやNUMAを何のことだと思ってるんだ？
というか釣られたか？
226 名前：デフォルトの名無しさん mailto:sage [2011/12/13(火) 17:55:43.92 ]: >>222
> これ、すごいことだと思うんだけど、
スパコンの世界では、ずうううっと前からやってるんじゃない？
227 名前：デフォルトの名無しさん mailto:sage [2011/12/13(火) 20:56:29.90 ]: じゃあ、そろそろ俺らも本気出すか
228 名前：デフォルトの名無しさん mailto:sage [2011/12/14(水) 09:03:43.27 ]: >>227
本気だしたら、負けだと思うんだ。
229 名前：デフォルトの名無しさん [2011/12/14(水) 20:16:13.92 ]: 4Gamer.net ― NVIDIA，CUDA 4.1をリリース。CUDAコンパイラのソースコード公開も
www.4gamer.net/games/076/G007660/20111214033/
中国時間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 名前：デフォルトの名無しさん mailto:sage [2011/12/14(水) 22:38:03.27 ]: これって、個人でもソースもらえるのかな
231 名前：デフォルトの名無しさん mailto:sage [2011/12/14(水) 23:11:41.54 ]: へえ、ソースを公開したんだ。
でも思いっきりOpenCLとかぶるな。
OpenCLが死亡か。
232 名前：デフォルトの名無しさん mailto:sage [2011/12/15(木) 08:04:39.43 ]: なんで、OpneCLが死亡するんだ？
233 名前：デフォルトの名無しさん mailto:sage [2011/12/15(木) 11:54:57.42 ]: まだ来てなくない？
234 名前：デフォルトの名無しさん mailto:sage [2011/12/15(木) 11:59:38.42 ]: Radeon用CUDA作っても、特定の機種用になりそうだな
235 名前：デフォルトの名無しさん mailto:sage [2011/12/18(日) 02:35:32.30 ]: そもそもRADEON上でCUDAってマトモに動くのか？
一応HD3000世代のRV670からはFP64対応らしいが
236 名前：デフォルトの名無しさん mailto:sage [2011/12/19(月) 02:14:51.68 ]: 動かないよ
動かそうとする人も現時点ではいないだろ
まだまだGPGPU向けには問題ありそうだし
237 名前：デフォルトの名無しさん mailto:sage [2011/12/20(火) 01:48:20.74 ]: cudaスレだが、
研究室や自分の周りでけで高速化すればいいんだったらcudaでいいんだろうけど
製品に組みこむならintelの普及度やなるべく多くの人に使ってもらえることを
考えるとOpenCLのほうがいいんじゃないかと思ってる。

photoshopのGPU対応もOpenCLでやってるみたいだし。
238 名前：デフォルトの名無しさん mailto:sage [2011/12/20(火) 03:10:58.69 ]: Photoshopは現状CUDAとATi Streamです
高速化しようとすると結局はデバイス毎に最適化したプログラムを書くことになるので、
デバイスに特化したライブラリの方が融通が利きます
Parallel NsightでのデバッグはOpenCLではできません
239 名前：デフォルトの名無しさん [2011/12/20(火) 03:38:50.36 ]: NVIDIA Opens Up CUDA, Compiler Source Code is Out - Bright Side Of News*
www.brightsideofnews.com/news/2011/12/19/nvidia-opens-up-cuda2c-compiler-source-code-is-out.aspx
240 名前：デフォルトの名無しさん mailto:sage [2011/12/20(火) 18:56:37.44 ]: CUDAのハードウェア種別を問わない展開が可能になったという事は、
CUDAの上に乗っかる形で実装されているPhysXやOptiX等もまた
ハードウェア種別を問わず動作可能になる可能性があるという事だよな
期待したい
241 名前：デフォルトの名無しさん mailto:sage [2011/12/20(火) 22:23:45.55 ]: 別にソースコードなくたって、ある程度の知識を持った人なら
CUDAを移植できたと思うけど
242 名前：デフォルトの名無しさん mailto:sage [2011/12/20(火) 22:49:46.93 ]: RADEONでPhysXはマーケティングの理由で使えないんだろ？
243 名前：デフォルトの名無しさん mailto:sage [2011/12/20(火) 23:19:03.65 ]: AMDがCUDAを実装されるの嫌がってPhysX蹴ったんだっけ
244 名前：デフォルトの名無しさん [2011/12/21(水) 19:57:27.52 ]: 質問：
GPU側に2次元配列を送るには、
１.2次元を1次元に直してからcudaMallocでメモリ確保→cudaMemcpyで送信
２.cudaMallocPitchでメモリ確保→cudaMemcpy2Dで送信
という2つの方法がありますけど、
構造体のメンバーに2次元配列があるものをGPU側に送る場合はどのようにしたらいいんでしょうか？
245 名前：デフォルトの名無しさん mailto:sage [2011/12/21(水) 20:20:10.20 ]: ポインタじゃなく配列ならそのままでうまくいかんか？
246 名前：デフォルトの名無しさん [2011/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 名前：デフォルトの名無しさん mailto:sage [2011/12/21(水) 21:36:40.66 ]: そのまま転送すればいいじゃん
248 名前：デフォルトの名無しさん [2011/12/21(水) 21:49:19.70 ]: >>246です。
すいません。できました。お騒がせしました。
249 名前：デフォルトの名無しさん mailto:sage [2011/12/22(木) 08:21:40.56 ]: >>229
このコンパイラって、謎のCをPTXに変換する部分だけじゃないの？
今のフロントエンド(cudafe)は、他社製品のEDGベースだから、出てこないと思うが。
　cudafe -> PTXへのコンパイラ(今だとnvopencc) -> ptxas
のまんなかが出来るだけなので、この部分のソースがあったところで、CUDAコンパイラを自分で作れるとか、
CUDAを他のプロセッサで動かせるようになるとかいうものではない。
250 名前：デフォルトの名無しさん mailto:sage [2011/12/22(木) 13:07:41.69 ]: .cファイルがコンパイルできません
たすけて
251 名前：デフォルトの名無しさん mailto:sage [2011/12/22(木) 13:13:00.92 ]: んｖｃｃじゃできんだよ
252 名前：デフォルトの名無しさん [2011/12/26(月) 15:17:58.63 ]: あの、OpenCLってオワコンなの？IntelとAMDが支持してるのに全然盛り上がってないんだが。
253 名前：デフォルトの名無しさん mailto:sage [2011/12/26(月) 21:25:42.32 ]: 科学技術計算向けがtesla一択だから他を使う理由もないってっところなのかね。
わざわざcudaで組んだものを他に移植する理由もないだろうし。
254 名前：デフォルトの名無しさん [2011/12/26(月) 21:42:01.58 ]: OpenCLで書くと長くなってしまう。
それだけのこと
255 名前：デフォルトの名無しさん mailto:sage [2011/12/27(火) 02:04:51.35 ]: C++ Bindings使えよ。
256 名前：デフォルトの名無しさん mailto:sage [2011/12/27(火) 10:05:27.99 ]: ちょっとしたラッパーかくのも嫌なのか
257 名前：デフォルトの名無しさん [2011/12/27(火) 11:26:29.65 ]: 超素人の質問です。
コマンドラインでnvccを使いコンパイルを行っているのですが、OpenGLのサンプルコードをコンパイルできません。
freeglutを入れてます。
gcc -lglut program.cではコンパイル可能でした。
しかしnvcc -lglut program.cuとするとコンパイル不可となります。
-Lや-lによってライブラリやヘッダファイル先を指定したところで、コンパイルできませんでした。
阿呆な質問かもしれませんが、自己解決できそうにないのでよろしくおねがいします。
258 名前：デフォルトの名無しさん mailto:sage [2011/12/27(火) 11:33:58.46 ]: 凄いアフォだな
エラーメッセージすべてコピペし
259 名前：帝徒＝繪璃奈＝啓北商業の野島えり [2011/12/27(火) 14:53:52.03 ]: 主犯少頭劣一族＝蔗冽一族とは

中国華喃の山の梺の村八分の家。

二間位の横長で玄関の右側がお勝手。

大正に猿のままで生まれたのが

鈴木あゆみ(网(アミ) 范襤の子)

フィリピン人の范襤と同じくフィリピン人のモンゴルに逃げた『シバ』との間の男児。

日本名鈴木ひろしと聞いた。

その後、親戚の鈴木大樹を殺し戸籍を使用。

今は一文字雉。
260 名前：デフォルトの名無しさん mailto:sage [2011/12/29(木) 21:08:02.00 ]: 母体になるプログラムなしで、
CUDAのコード単体で動かせるツールって無いかねぇ。
外出先でCUDAのサンプル見せるのがメンドイ。
例えば ./cuda-run < source.cuda みたいな感じで実験できるといい感じ。
codepadみたいな感じで動かせればなおよしだけど。
261 名前：デフォルトの名無しさん [2011/12/29(木) 22:01:29.67 ]: そういうスクリプトかけばよい
262 名前：デフォルトの名無しさん mailto:sage [2012/01/05(木) 10:51:17.15 ]: >>260
＞母体になるプログラムなしで、
それってなんのこと? 外出先のPCで動かしたいってこと?
コンパイラのことなら、事前にコンパイルしておくかリモートでアクセスすればいい。
ドライバのことなら、外出先の端末にインストールさせて貰うしかない。
外出先のPCなんか頼らず、自前のPCを持っていけばいいじゃん。
263 名前：デフォルトの名無しさん [2012/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 名前：デフォルトの名無しさん mailto:sage [2012/01/05(木) 14:25:08.28 ]: インクルードパスの指定は -l じゃなくて -I な
265 名前：デフォルトの名無しさん mailto:sage [2012/01/05(木) 15:05:37.06 ]: /usr/bin/ld: skipping incompatible～って64/32ビット混在時のエラーメッセージじゃなかったかな
266 名前：デフォルトの名無しさん [2012/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 名前：デフォルトの名無しさん mailto:sage [2012/01/05(木) 15:37:13.27 ]: おれも64/32の違いかと思ったが、パス指定から指摘してみた
実行している環境は？ -m64 オプション付けてみ
268 名前：デフォルトの名無しさん [2012/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 名前：デフォルトの名無しさん mailto:sage [2012/01/06(金) 13:59:40.26 ]: libglutは64ビット？
32ビットなんじゃねか？
270 名前：デフォルトの名無しさん [2012/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 名前：デフォルトの名無しさん mailto:sage [2012/01/06(金) 15:37:17.18 ]: いや、だからfreeglutは64ビットを使用してるのかと聞いてるんだが
質問に対しきちんと回答できないようでは教えようがない罠
272 名前：デフォルトの名無しさん [2012/01/06(金) 18:39:39.36 ]: >>271
freeglutは64ビットのものを入れ直してやってみましたが、同じ結果になりました。

>>257
そこでCUDAのヘッダフォルダとライブラリフォルダに入れていたものを消し、-Lと-Iで直接freeglutのフォルダを指定したら
解決しました。
ライブラリが足りなかったのか、それともCUDAと一緒くたにしたのがダメだったのか解かりませんが、ちゃんとPATH通せってことみたいです。
お騒がせしました。
273 名前：デフォルトの名無しさん mailto:sage [2012/01/06(金) 19:45:07.40 ]: >>262
だよねぇ・・・
274 名前：デフォルトの名無しさん [2012/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

このサイトを参考に環境を構築しました。
feather.cocolog-nifty.com/weblog/2011/07/visual-studio-2.html
そして以下のサイトのサンプルプログラムを実行してみました。
www.gdep.jp/page/view/218
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などを変えてみましたが、一向に変わりません。
d.hatena.ne.jp/Levi/20090921/1253535802#c
SDK内のcutil_inline.h自体をtoolkitのincフォルダにコピペすると、他の.hファイルもいくつか同じエラーが出たのでエラーになったものをすべてコピペしたところ、LNK2019"link.exe"というエラーで先に進めませんでした。
一度VisualStudioを再インストールしてみましたが、状況は変わりません。
Nvidia GPU computing SDK Browserではサンプルプログラムを実行できているので、CUDAの環境は整っていると思われます。
どうすれば解決できますでしょうか・・・。かれこれ１週間以上格闘しています。
VisualStudioは2010よりも2008にした方がいいでしょうか？
275 名前：デフォルトの名無しさん mailto:sage [2012/01/06(金) 20:23:04.24 ]: C初心者にはきついと思うんだが…
とりあえず 'cutil_inline.h'のある場所を見つけて
そこを-I　/'cutil_inline.hのある場所'と指定する

意味わからなければCのコンパイルを勉強すること
276 名前：デフォルトの名無しさん mailto:sage [2012/01/06(金) 20:39:53.25 ]: 274ではありませんが､「Cのコンパイル」の勉強にオススメの書籍とかありましたら紹介していただけるとうれしいです
277 名前：デフォルトの名無しさん mailto:sage [2012/01/06(金) 20:55:31.81 ]: 275だけど、お薦めは他の人に任せるけど
一つ言えるのはCUDAはC/C++より数段難しい
C/C++の質問をしてるようでは先が思いやられる、と正直思うんだが…

超低レベルな質問もOKなくだすれだけど、その程度は自力で調べる能力が必要ではないだろうか？
278 名前：デフォルトの名無しさん mailto:sage [2012/01/06(金) 21:22:02.03 ]: >>275
>>274です。ありがとうございます。
今実行できる環境にないので後で試してみます。
また結果報告しにきます。
279 名前：デフォルトの名無しさん [2012/01/07(土) 12:14:43.10 ]: 一時期は同価格帯のCPUとGPUを比較してGPUは５０倍の性能があるって言われてたけど
今１０～１５倍程度しか無くない？

3930k
５万円
158GFlops

GTX590
６．３万円
2480GFlops

GTX 580
３．８万円
1500GFlops

次に出ると言われてるGTX700系は500系の２倍の性能らしいけどそれでも５０倍には遠く及ばない。
この「同価格帯でGPUはCPUの１０～１５倍程度の性能しかない」という評価は正しいだろうか？
280 名前：デフォルトの名無しさん mailto:sage [2012/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 名前：デフォルトの名無しさん mailto:sage [2012/01/07(土) 16:09:01.15 ]: >>279
GTX580で500GFLOPS程度。そのデータはおかしい
CPUなんてE8500で25GFLOPSくらいだから20倍程度だな。
もちろん最近のコア多いやつならもっと差は縮まるだろうな。

それに理論値の話だから多数のスレッドでまわしたらコンテキストスイッチの
オーバーヘッドとかGPUプログラミングは難しくて効率上げるの難しいとか
そういうので実効は変わってくるんじゃね？
282 名前：デフォルトの名無しさん mailto:sage [2012/01/07(土) 16:14:51.24 ]: あ、同一価格帯か、E8500が2万しなかったことを考えると
580にそろえて2倍したとして10倍程度だから、まあおおざっぱには
「同価格帯でGPUはCPUの１０～１５倍程度の性能しかない」は正しいんじゃないかと思う。

580と同時期の4万弱のCPUならもっと差は縮まるだろうし。
283 名前：デフォルトの名無しさん mailto:sage [2012/01/07(土) 16:22:20.40 ]: 32nmと40nmで比較とか
284 名前：デフォルトの名無しさん mailto:sage [2012/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 mailto:sage [2012/01/07(土) 22:23:52.14 ]: 異なるディレクトリに複数のCUDAのライブラリが入ってて
古いバージョンが参照されてるのが原因だった・・・
286 名前：デフォルトの名無しさん mailto:sage [2012/01/08(日) 14:34:35.46 ]: CUDA4.1でcomputeProfはnvvcになったんですか？
だいぶ代わってしまって驚きますた
287 名前：デフォルトの名無しさん mailto:sage [2012/01/08(日) 14:35:50.30 ]: nvvp
288 名前：デフォルトの名無しさん mailto:sage [2012/01/08(日) 20:50:11.30 ]: >>282
GPUの論理性能ってFMAとかMADとかの積和で倍になっているから、
GTX 580 1500GFlopsとかって、それらを使わなければ750GFlopsとかになるんだよな。
だから実効れべるではもっと縮まるんだろうな。それでも5倍以上差がでるとかなり有効なんだけど。
289 名前：デフォルトの名無しさん mailto:sage [2012/01/09(月) 10:37:28.98 ]: 個人で買うやつも居るのかね？
ttp://page16.auctions.yahoo.co.jp/jp/auction/u38857709
290 名前：デフォルトの名無しさん mailto:sage [2012/01/09(月) 12:04:44.39 ]: いる。というか俺が買う。
291 名前：デフォルトの名無しさん mailto:sage [2012/01/09(月) 13:09:16.68 ]: こういう需要が限られていてかつ堅くて高い製品はamazonに出した方が高く売れるんじゃないかねえ。
292 名前：デフォルトの名無しさん mailto:sage [2012/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 名前：デフォルトの名無しさん mailto:sage [2012/01/09(月) 18:20:34.55 ]: もっとコンパイル早くなんねーかなぁ…
thrust::sortとか使うようにしただけでコンパイルに十数秒かかってイライラする
294 名前：デフォルトの名無しさん mailto:sage [2012/01/09(月) 20:40:26.69 ]: 速いマシンでやる。
295 名前：デフォルトの名無しさん mailto:sage [2012/01/09(月) 20:45:40.42 ]: nvccコンパイラをGPGPUで実装すればいい！

CUDAがオープンになるどさくさにまぎれてそういうのを作る人が出るかもしれないこともないかも
296 名前：デフォルトの名無しさん mailto:sage [2012/01/09(月) 21:57:43.85 ]: コンパイラの動作は現状GPU向きとは言えないんじゃね？
GPU自体が単純な性能向上だけではない、思いもよぬ進化でもしない限り
297 名前：デフォルトの名無しさん mailto:sage [2012/01/09(月) 22:54:12.52 ]: コンパイルなんて依存関係ありまくりで、並列処理が苦手な最たるものだからまずやる意味がないだろう。
CUDAにすれば何でも速くなると思ってんじゃね？
298 名前：デフォルトの名無しさん mailto:sage [2012/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 名前：デフォルトの名無しさん mailto:sage [2012/01/10(火) 17:35:07.81 ]: >>292

参照パスを ""で囲む♪
300 名前：デフォルトの名無しさん mailto:sage [2012/01/10(火) 17:39:16.65 ]: >>298

iが３２以上になることはないかい？？

ブロックの数は？？
301 名前：274 mailto:sage [2012/01/10(火) 17:52:39.67 ]: >>299
書き込みしてからそのミスに気付きました・・・。
それを直してもうまくPathが通らなかったので、CUDA4.0、VisualStudio2008でやることにしました。
302 名前：デフォルトの名無しさん mailto:sage [2012/01/10(火) 17:56:37.24 ]: >>298
fence入れてないからじゃね。
303 名前：298 mailto:sage [2012/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 名前：デフォルトの名無しさん mailto:sage [2012/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 mailto:sage [2012/01/10(火) 18:17:13.87 ]: >>302 >>304
ありがとうございました。
試してみたいと思います。
306 名前：デフォルトの名無しさん mailto:sage [2012/01/10(火) 19:45:54.97 ]: >> 305

だいたい　　x = s_v[(threadIdx.x+3)%32]; 　の前にすべてのs_v[??}が定義されていないといけないだろ？？

スレッドが３２に設定されているのならば　(threadIdx.x+3)%32　は０か１にしかならないが．．．
307 名前：デフォルトの名無しさん mailto:sage [2012/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 名前：デフォルトの名無しさん mailto:sage [2012/01/10(火) 21:12:58.21 ]: ちなみに家のはGeForce GT 220なので非力．．．．
309 名前：306 mailto:sage [2012/01/10(火) 22:21:29.79 ]: 間違えた
　↓
(threadIdx.x+3)%32　は０か１にしかならないが．．．

３２で割った余りであった．．．
310 名前：298 mailto:sage [2012/01/11(水) 09:16:02.50 ]: >>306 >>309
いろいろ試してみたところ、とりあえず上記の問題は解決しました。
回答ありがとうございました。
311 名前：デフォルトの名無しさん mailto:sage [2012/01/11(水) 11:25:42.40 ]: すごい。
見た感じ、ちょっと不安だったから
参考になったみたいでよかった。
312 名前：274 mailto:sage [2012/01/11(水) 19:30:44.93 ]: >>307
VS2008でもうまくできませんでした。
再度VS2010を入れてみるも手順通りいかず、当然実行もできませんでした。
リカバリしたのでこれからもう一度2010でやってみます。もう心が折れそうです・・・。
ちなみに2010expressとprofessionalの違いで実行できないこととかあり得ますか？
別のPCでprofessionalで試してみてくれた方がいて、うまくできたようなのですが・・・。
313 名前：デフォルトの名無しさん mailto:sage [2012/01/11(水) 20:43:20.68 ]: >>312

どんなエラーがでますか？

そもそもnvccでコンパイルするのだから，Visual Studioは関係ないでしょう？

VisualStudioのないLinux上でもコンパイル，実行できるはず．

もう一度やるのならばライブラリ lcutil32を付ければ良いかと（３２ビットならば）
314 名前：274 mailto:sage [2012/01/12(木) 01:51:21.33 ]: >>313
VisualStudioであれば多少使ったことがあるので導入しました。
少なくともwindowsでは「nvccがVisualStudioのcl.exeを必要とする」と複数のサイトに書いてあるようです。
gpu.fixstars.com/index.php/Windows_Vista_%E3%81%ABCUDA%E9%96%8B%E7%99%BA%E7%92%B0%E5%A2%83%E3%82%92%E3%82%A4%E3%83%B3%E3%82%B9%E3%83%88%E3%83%BC%E3%83%AB%E3%81%99%E3%82%8B
exth.net/~ohshima/wordpress/category/windows/

エラーはプロジェクトをビルドしたときに「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

でした。
どうやらリンカーの設定を見直せとのことだったのでwww.scribd.com/doc/66757447/を参考に
VisualStudioでプロジェクトのプロパティから「リンカ > 入力 > 追加の依存ファイル」にcutil32.lib, cutil64.lib を追加。
しかし「cutil64.libが見つからない」とのエラー。
実際に探してみてもcutil64.libがどこにも見つかりませんでした。

今まで3回もリカバリしてやってみましたが、どうもPathがうまく通せていない気がします。

もうすでにリカバリしてしまったので、明日また環境構築をはじめからやり直します。。

Linuxの環境はないんです、すみません・・・。
今の環境はWindows7 64bitです。
詳しくはこちら>>274に書きました。
315 名前：デフォルトの名無しさん mailto:sage [2012/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 名前：デフォルトの名無しさん mailto:sage [2012/01/12(木) 07:18:06.91 ]: >>314
64bit版のcl.exeがある。
見つからないなら-ccbinオプションでcl.exeの場所を指定。

cutil32.libとcutil64.libを両方指定してはいけない。どちらか必要なものを。
で、cutil64は確か自分でビルドする必要がある。
SDKのディレクトリの中にプロジェクトファイルがあるからそれをビルド。
317 名前：デフォルトの名無しさん mailto:sage [2012/01/12(木) 08:36:26.90 ]: codepad.org/Kyns70af
上記のプログラムが期待と違う結果を出力するのだが原因がわかる方がいれば教えていただけないだろうか．
（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 mailto:sage [2012/01/12(木) 08:37:35.44 ]: すみません．
>>317のコードのURLを間違えました．
誤: codepad.org/Kyns70af
正: codepad.org/cfbTSsdF
319 名前：デフォルトの名無しさん mailto:sage [2012/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 名前：デフォルトの名無しさん mailto:sage [2012/01/12(木) 10:24:08.53 ]: >>317
同じく問題再現されず（CUDA4.1 + GF560Ti）
そのデバイス両方とも compute capability 1.x だし
2.x 環境で試してみたらどうだ
321 名前：317 mailto:sage [2012/01/12(木) 10:56:51.02 ]: >>319 >>320
コメントありがとうございます
現在の私の環境ではカーネル関数内で printf は使えないですが
変なことになるのは私の環境だけのようなので Compute capability 2.x 環境で後日試してみます
322 名前：デフォルトの名無しさん mailto:sage [2012/01/12(木) 11:18:04.11 ]: ちなみにcuPrintf.cuh　とcuPrintf.cu　をネット上からひらってくれば、printfと同等の関数（cuPrintf）が使えますよ♪
323 名前：デフォルトの名無しさん mailto:sage [2012/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 mailto:sage [2012/01/12(木) 13:42:05.81 ]: >>315 >>316
ありがとうございます。
とりあえず315さんの通りコマンドライン上で実行できるところを目指そうと思います。

cl.exeを指定するオプションがあるんですね。
自分はコンパイラオプションの勉強不足のようです・・・。
cutil64は自分でビルドするんですか、それは今までやってなかったと思います。

恐縮ですが、これから環境を整えるので何か参考になるサイトがあれば教えていただけますか？
いくつかサイトを見て設定したのですが、どれもうまくいかなかったので・・・。

何度もすみません。よろしくお願いします。
325 名前：デフォルトの名無しさん mailto:sage [2012/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程になっていました。
似た事例として
ttp://d.hatena.ne.jp/ymizushi/20091214/p1
という記事を見つけましたが、この記事内でも原因はよく分かっていません。
どなたか詳細をご存じないでしょうか？
326 名前：デフォルトの名無しさん mailto:sage [2012/01/12(木) 17:29:28.93 ]: >> 324

参考になるやら、ならないやら

www.fxfrog.com/?p=3660
www.ademiller.com/blogs/tech/2011/05/visual-studio-2010-and-cuda-easier-with-rc2/
cudasample.net/sample/1/001.html
tech.ckme.co.jp/cuda_first.shtml
cudasample.net/
others2.blog.so-net.ne.jp/2010-09-12

>> 325

1．カーネルの中の演算数がiによって異なる

2. 60回目で解が発散し NaNになっているが、お構いなしに計算は続けられている

3. カーネルからCPUにメモリーコピーしている時間が紛れ込んでいる

かな？
327 名前：274 mailto:sage [2012/01/12(木) 17:37:04.49 ]: >>326
ありがとうございます。
全て拝見しましたが既読のものでした。やはり普通に設定すればできるようですね・・・。
今格闘しているので、また報告します。
328 名前：325 mailto:sage [2012/01/12(木) 17:50:39.27 ]: >>326
ありがとうございます。
全項目について確認してみましたが、
1. 演算数はどれも一定
2. ループ回数によらず正しい値が得られている
3. 測定時間中にcudaMemcpy等の関数は呼ばれていない
という状態です。
329 名前：デフォルトの名無しさん mailto:sage [2012/01/12(木) 18:22:29.85 ]: >>328
1プロセスで2つのカーネル実行してるなら順番入れ替えて試してみた？
最初の実行は起動コストだかで時間かかるのが普通なんだが
330 名前：デフォルトの名無しさん mailto:sage [2012/01/12(木) 19:29:37.96 ]: 時間の計測が間違っている。
cudaThreadSynchronize()を読んだあとに終了時刻を取得する必要がある。
1024個までのカーネル呼び出しはキューに積むだけですぐに次に進んでしまう。
30000回ループの方はキューが詰まるので(30000-1024)回くらいの実行完了を待つことになり、比較的妥当な計算時間に見えるはず。
もちろん正確ではないのだが。
331 名前：325 mailto:sage [2012/01/12(木) 19:35:41.01 ]: >>329-330
回答ありがとうございました。
cudaThreadSynchronize()を入れるのを失念しておりました。
332 名前：デフォルトの名無しさん mailto:sage [2012/01/12(木) 20:59:37.22 ]: cudaDeviceSynchronize()
がナウいらしい
333 名前：sage [2012/01/13(金) 08:22:49.98 ]: >>330
便乗して質問ですが、
キューに一度に入るカーネル数が1024個で、
それ以上はカーネル呼び出しが待ちになるというのは、
CUDAのマニュアルかどこかに書いてありますか？
334 名前：デフォルトの名無しさん mailto:sage [2012/01/13(金) 09:57:50.47 ]: >>333
実験したらそうだったけど、文書化されているかどうかは知らない。
これだけ多ければ普通は問題ないし、バージョンアップで変わるかも。
335 名前：デフォルトの名無しさん mailto:sage [2012/01/13(金) 11:09:53.09 ]: >>334

了解しました。
どうも有難うございました。
336 名前：デフォルトの名無しさん mailto:sage [2012/01/13(金) 16:58:47.32 ]: それにしても　きゅー　にカーネルが溜まっているからと言って，GPUの処理能力が落ちるとは思えないぞ
337 名前：デフォルトの名無しさん mailto:sage [2012/01/13(金) 17:17:25.93 ]: きゅー　にカーネル田丸？
それってstream使ってる時のこと？
文脈呼んでないからわからんけど
338 名前：デフォルトの名無しさん mailto:sage [2012/01/13(金) 20:07:11.66 ]: キューに空きがあれば積んで即制御が戻る。
空きが無ければ空くまで待ってる。

詰まってるから遅くなるわけではなく、遅い方が正しい実行時間なだけ。
339 名前：デフォルトの名無しさん mailto:sage [2012/01/18(水) 22:19:57.72 ]: CUDA4.0からデバイスエミュレーションがなくなったらしいけど
実機で実行するとOSごと固まることがある
みんなどうしているのでしょうか
340 名前：デフォルトの名無しさん mailto:sage [2012/01/19(木) 00:43:55.87 ]: デバイスを物理的に叩き割る
341 名前：デフォルトの名無しさん mailto:sage [2012/01/19(木) 04:33:29.60 ]: 2.3くらいからなかったような気が
342 名前：デフォルトの名無しさん mailto:sage [2012/01/19(木) 15:25:40.77 ]: >>339

具体的には？？
343 名前：デフォルトの名無しさん mailto:sage [2012/01/20(金) 18:49:16.04 ]: >>342
メモリアクセスのバグがあるとWindowsが固まって動かなくなったり
ディスプレイにノイズが出てくる
OSが画面表示に使っている領域を壊しているようだけど
普通は起こらないことなんでしょうか？
344 名前：デフォルトの名無しさん mailto:sage [2012/01/21(土) 10:42:26.35 ]: 私の経験では、CUDAのプログラムをチェックするために printf を使って変数を書き出したところ、
その量が多すぎて、画面が真っ暗になったり、PCがシャットダウンしたことがあった。

書き出し量を減らしたところ、問題はなくなった。

こんなことかな？？
345 名前：344 mailto:sage [2012/01/21(土) 11:00:22.89 ]: 追加です。

グラフィックカードを2枚刺して、一枚をCUDA用、もう一枚をディスプレイ用にすればよいとのことを聞いたことがあるが、
私のはGTX５９０でカードが2枚、入っているとのこと。

CUDAの計算に Device(0)にしても(1)にしても、前述の問題が起きたので、枚数には関係ないこともあるらしい？？
346 名前：デフォルトの名無しさん mailto:sage [2012/01/22(日) 05:18:00.51 ]: printf使えるとか便利になったわ
347 名前：デフォルトの名無しさん mailto:sage [2012/01/22(日) 16:18:23.06 ]: >>344
fermiではないのでprintfは使えないです。
GTX260なんですけど、それが悪いのかも。
計算用GPUを別に刺すってのは考えていました。
ただ、そうするべきという常識みたいなものがあるのかなっていう疑問があったので。

あとこういう状態でデバイスエミュレーション外したならひどい話だなと思ったのですが
みなさんあまり起きていないようで。
348 名前：デフォルトの名無しさん mailto:sage [2012/01/22(日) 16:23:00.85 ]: グローバルメモリの添え字計算計算が間違っていたりで確保した範囲外にアクセスすると
Windowsが固まったりディスプレイにノイズが出てくることがあるって状態です。
349 名前：やんやん ◆yanyan72E. mailto:sage [2012/01/22(日) 16:48:55.32 ]: グラフィックカード二枚差しにすると、
nSghtでハードウェアデバッグができるとかじゃなかったっけ？
350 名前：デフォルトの名無しさん [2012/01/22(日) 17:02:34.79 ]: YES
もしくはPC2台で片方をデバッグ用にしても良い。

実際に動かすGPUがCUDA対応なら、開発側は非nVIDIAでもOK
351 名前：デフォルトの名無しさん mailto:sage [2012/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 名前：デフォルトの名無しさん mailto:sage [2012/01/23(月) 10:54:29.21 ]: 1はFermiならOK。古い世代ではダメ。
2と3は絶対ダメ。stream0は前後をぶった切る役目がある。
353 名前：デフォルトの名無しさん mailto:sage [2012/01/24(火) 08:17:05.28 ]: >>352
1の動作がFermiとそれ以前で異なるのは知りませんでした。
回答どうも有難うございました。
354 名前：デフォルトの名無しさん mailto:sage [2012/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 名前：デフォルトの名無しさん [2012/01/24(火) 19:53:55.48 ]: コンパイル時に以下のエラーメッセージが出ているのですが、
ptxas error : Entry function '...' uses too much local data (0x16054 bytes, 0x4000 max)
これはローカルメモリーと何か関係あるような気がするのですけど、
ちょっとわからないので教えて頂けないでしょうか。
宜しくお願いします。
356 名前：デフォルトの名無しさん [2012/01/24(火) 21:03:52.65 ]: それともう1つ。
externを使って、グローバル変数を共有させたいのですが、
ホスト・デバイスともに共有させたい場合はどうしたらいいのでしょうか?
357 名前：デフォルトの名無しさん mailto:sage [2012/01/24(火) 21:56:32.01 ]: >　ptxas error : Entry function '...' uses too much local data (0x16054 bytes, 0x4000 max)

単に配列の次元が大きすぎるのでは？

最大　0x4000バイトのところを0x16054 bytes使おうとしている？

> ホスト・デバイスともに共有させたい場合はどうしたらいいのでしょうか?

物理的に無理でないの？

ホスト←→デバイス間でコピーしないといけないのでは？？
358 名前：デフォルトの名無しさん mailto:sage [2012/01/24(火) 21:56:54.10 ]: >>355
メッセージそのままの意味だろ。
359 名前：デフォルトの名無しさん mailto:sage [2012/01/25(水) 01:32:34.89 ]: CUDAの思想はいいと思うけど、ターゲットデバイスを完全には
隠蔽しきれないのが残念だな。
360 名前：やんやん ◆yanyan72E. mailto:sage [2012/01/25(水) 01:41:55.08 ]: スピード勝負の世界で、しかもかなり複雑なアーキテクチャなのに、
ハードウェアを完全に隠蔽する訳にいかないだろ。
361 名前：357 mailto:sage [2012/01/25(水) 12:34:12.36 ]: > ホスト・デバイスともに共有させたい場合はどうしたらいいのでしょうか?

ヘッダーファイルに

#DEFINE

などで定義して、ホストとデバイスのプログラムにそのヘッダーファイルをinclude する手があった♪
362 名前：355,356 [2012/01/25(水) 15:21:00.98 ]: >>358
青木氏の「はじめてのCUDA」に似たようなエラーメッセージが載ってるのですが、
local dataがどの部分の話なのか、ちょっとわからない状態でして...

>>357
>単に配列の次元が大きすぎるのでは？
予定だと1×256スレッドしか使うので、足りないはずないような... 意味が違うか。
ちょっとわからないので詳細お願いします。

>ヘッダーファイルに
>#DEFINE
>などで定義して、ホストとデバイスのプログラムにそのヘッダーファイルをinclude する手があった♪

ちょっと試してみます。それができなかったら、コピーするしかないかもしれません。
__device__ extern TEST test;
のような宣言してて、実行ファイルができたときがあったのですが、何か違うよな、と引っかかってたので、
ここで質問しました。ありがとうございます。
363 名前：デフォルトの名無しさん mailto:sage [2012/01/25(水) 17:40:16.56 ]: 差し支えなければプログラムをアップしてくださいませ♪
364 名前：デフォルトの名無しさん mailto:sage [2012/01/25(水) 18:01:06.82 ]: >>362

このサイトの一番下にある資料の3-41ページにそのエラーメッセージが載っている模様。
ttp://accc.riken.jp/HPC/training.html
365 名前：355,356 [2012/01/25(水) 18:04:09.31 ]: >>362です。

>>363
申し訳ないのですが、プログラムのアップは出来ないです。すみません。

なんかもう分からなくなったので、ホストとデバイスの住み分けを行ったところ、
「ptxas error...」云々のメッセージが消え、コンパイルできました。
（さっきまでhost,device,global宣言関数がごちゃ混ぜな状態だった。）
何が原因だったのか分からず仕舞いです。

グローバル変数の共有はとりあえず、コンスタントメモリにデータコピーで様子見することにしました。
元コード(.c)をもう一度読み直したところ、デバイス側の方はReadOnlyで十分なようでしたから。

皆様回答ありがとうございました。またよろしくお願いします。
366 名前：365 [2012/01/25(水) 21:05:50.75 ]: >>364
資料ありがとうございます。
うまくいったと思った途端、また同じエラーが出てきてしまったので、確認します。

.cファイルに.cuファイルを組み込ませるようにしたら、
__host__ __device__で修飾した関数が定義されてないと.cファイル側に言われ、
.cファイルと.cuファイルそれぞれ単独で動かせるようにしたら、
（同じ内容の関数を.cと.cuそれぞれ別名で実体作った。.c:test .cu:ktestみたいな）
今度はさっきと同じエラー。

実行ファイル作るだけなのに難しい...
367 名前：デフォルトの名無しさん mailto:sage [2012/01/26(木) 08:42:18.26 ]: Host側とDevice側で型を別にすれば開発もっと楽になると思うんだけど
368 名前：デフォルトの名無しさん mailto:sage [2012/01/26(木) 08:53:13.45 ]: 取り敢えずCudaのサンプルは捨てて、インクルード関係とオブジェクトの生成手順を確認するんだ。
369 名前：デフォルトの名無しさん mailto:sage [2012/01/26(木) 10:14:19.95 ]: 見せられないところは削除して、問題の部分だけ残して、アップすれば？？
370 名前：366 [2012/01/26(木) 16:19:08.01 ]: >>368
>取り敢えずCudaのサンプルは捨てて、インクルード関係とオブジェクトの生成手順を確認するんだ。
のように手順を踏んでコンパイルしたところ、実行ファイルができました。
皆様ありがとうございました。
371 名前：370 [2012/01/26(木) 20:33:24.55 ]: お恥ずかしながら、また戻ってきました。
あれからセグメンテーション違反が出てきてしまったので、あれこれ探していた結果、
どうやらデバイスのメモリ確保＆送信に失敗していたようです。

しかし解せないことがあって、
構造体A a,構造体B b[100],構造体C　c[100],...(以下略　をデバイス側に送るのですが、
（1つ１つのサイズは結構大きめ。グローバルは1GBあって余裕で確保出きるハズ...）
void main(){
・・・
test(&a,b,c....);
}
void test(A *a,B *b,C *c...){
A *d_a; B *d_b; C *d_c;
CUT_SAFE_CALL(cudaMalloc((void**)&d_a,sizeof(A)));
CUT_SAFE_CALL(cudaMemcpy(d_a,a,sizeof(A),cudaMemcpyHtoD));

CUT_SAFE_CALL(cudaMalloc((void**)&d_b,sizeof(B)*100));
CUT_SAFE_CALL(cudaMemcpy(d_a,b,sizeof(B)*100,cudaMemcpyHtoD));
・・・
｝
はじめてcuda,cuda_by_exampleで確認したところ、
文法的ミスはないはずなのに確保ミスしてるらしく、中断してます。

この原因は一体全体なんなんでしょうか。
372 名前：デフォルトの名無しさん [2012/01/26(木) 21:28:56.32 ]: 4.1正式版
CUDA Toolkit 4.1 | NVIDIA Developer Zone
developer.nvidia.com/cuda-toolkit-41
373 名前：デフォルトの名無しさん mailto:sage [2012/01/26(木) 21:55:28.42 ]: 構造体の中身がよく分からんが，allocateのところで　*100は必要なのかな？

構造体Bのメモリーを１００個分用意しようとしている？？

構造体Bの中にすでに [100]個の配列を取っているのに？？？

数は合っておるのか？？？？
374 名前：デフォルトの名無しさん mailto:sage [2012/01/26(木) 22:17:44.08 ]: ｒｈｅｌ6.0のがリンクミスかでダウンロドできん
375 名前：デフォルトの名無しさん mailto:sage [2012/01/27(金) 00:03:32.49 ]: >>372
遂にRC取れたか！4.0からどう変わってるか知らんけど。
Kepler向けのプログラムを作れるのは5.0とかになるんだろうか。
そしてそれはいつ出るんだろう。

>>371
1GBのVRAM積んでるカードでも400MBぐらいの確保で失敗したことがあったような覚えがある。
確保するサイズを小さくしてもエラーが出るってんならこの話は忘れてくれ。
376 名前：デフォルトの名無しさん mailto:sage [2012/01/27(金) 01:34:46.65 ]: 100x100の行列の積を1万回くらい実行したいのですが、
CUBLASだと小さすぎてCPUより遅くなってしまいます。
この1万回は相互に依存関係は無いので、並列化が可能なのですが、
このプログラムは自分で書くしかないのでしょうか。
TESLA C1070があるので、丸ごとVRAMに載ります。
377 名前：デフォルトの名無しさん mailto:sage [2012/01/27(金) 03:17:55.02 ]: 4.1にcublas{S,D,C,Z}gemmBatchedってのが追加されてました。
especially for smaller matricesに効くってあったので、使ってみます。
378 名前：デフォルトの名無しさん mailto:sage [2012/01/27(金) 07:52:34.05 ]: 一つのブロック内で100x100の行列の積を１回行う？
379 名前：デフォルトの名無しさん mailto:sage [2012/01/27(金) 07:52:51.42 ]: 間違った
380 名前：デフォルトの名無しさん mailto:sage [2012/01/27(金) 07:53:00.44 ]: あれ？
381 名前：デフォルトの名無しさん mailto:sage [2012/01/27(金) 07:53:07.78 ]: 一つのブロック内で100x100の行列の積を１回行う？
382 名前：デフォルトの名無しさん mailto:sage [2012/01/27(金) 07:55:11.45 ]: たびたび，失礼

一つのブロック内で100x100の行列の積を１回行う？

そのブロックを１万個用意する？

それらを並列に計算する？

と言うコンセプトかな？？？

自分でプログラムを作った方が，よほど勉強になると思う
383 名前：371 [2012/01/27(金) 18:15:48.81 ]: >>373
CUT_SAFE_CALL(cudaMalloc((void**)&d_b,sizeof(B)*100));
CUT_SAFE_CALL(cudaMemcpy(d_a,b,sizeof(B)*100,cudaMemcpyHtoD));
これはミスですね....
正しくは
CUT_SAFE_CALL(cudaMalloc((void**)&d_b,sizeof(B)));
CUT_SAFE_CALL(cudaMemcpy(d_a,b,sizeof(B),cudaMemcpyHtoD));
です。すみません。

>>375
4000MBの領域確保でも失敗することがあったんですか...
ちょっと試してみます。
384 名前：デフォルトの名無しさん mailto:sage [2012/01/27(金) 20:31:39.53 ]: d_aに送信してるあたりが
385 名前：デフォルトの名無しさん mailto:sage [2012/01/27(金) 21:52:59.41 ]: d_b のサイズは構造体Bと同じ？？

d_aのサイズはbと同じ？すなわち構造体Bと同じ？？

違っていたらエラーが出て当たり前では？？？
386 名前：デフォルトの名無しさん mailto:sage [2012/01/28(土) 13:33:36.75 ]: それから，ホスト側で　も　メモリーは確保しているのか？？？？
387 名前：383 [2012/01/28(土) 17:54:35.37 ]: >>371
なんかサンプルがめちゃくちゃなんで書き直します。
388 名前：387 [2012/01/28(土) 20:55:21.49 ]: 確認なのですけど、カーネル関数の引数はポインタ限定ですか?
389 名前：デフォルトの名無しさん mailto:sage [2012/01/28(土) 22:58:27.22 ]: 配列はポインタで，１変数はそのままではなかったかい？

なぜか知らんけど
390 名前：デフォルトの名無しさん mailto:sage [2012/01/28(土) 23:07:13.34 ]: 夜の埼玉は
391 名前：デフォルトの名無しさん mailto:sage [2012/01/29(日) 01:49:05.91 ]: >388
ホスト側のポインタ以外なら何でもOKのはず。少なくともintは渡せる。
クラスを渡せるかどうかは知らないけど。
392 名前：デフォルトの名無しさん mailto:sage [2012/01/29(日) 06:58:36.12 ]: ホストのポインタ送って、GPU側のload/store時に
メインメモリから自動転送ってこともできるから、
その説明は正しくない。
393 名前：388 [2012/01/30(月) 15:13:56.68 ]: >>391
セグメンテーション違反してた部分の引数に（デバイスで使用する）double（実体）を指定してたのですけど、
（デバイスで使用する）ポインタに変えたら、問題なく実行できました。
dtempはcudaMalloc/cudaMemcpy使ってます。

そもそもな話cudaはポインタ宣言したdtempの実体は*dtempではないのかな。
394 名前：デフォルトの名無しさん mailto:sage [2012/01/30(月) 15:16:54.53 ]: >>393
何を言っているのか判らない。問題が再現する最低限のコードでも貼ってみていただきたい。
395 名前：デフォルトの名無しさん [2012/01/30(月) 16:19:41.78 ]: >>393
__global__ void kernel(double dtemp){
}

int main(){
double temp=0.0;
double *dtemp;
cudaMalloc((void**)&dtemp,sizeof(double));
cudaMemcpy(dtemp,&temp,sizeof(double),cudaMemcpyHostToDevice);
kernel<<<1,1>>>(*dtemp);//(←※1)
cudaMemcpy(&temp,dtemp,sizeof(double),cudaMemcpyDeviceToHost);
return 0;
}
※1の部分でsegmentation fault。
__global__ kernel(double *dtemp){
}として
kernel<<<1,1>>>(dtemp);
とすると問題無し。
dtempのアドレスを引数にするのはいいけど、dtempの中身（0.0）を引数にしようとするのはダメ？
そもそも※1の宣言自体間違ってる?どうなんだろう。
396 名前：デフォルトの名無しさん mailto:sage [2012/01/30(月) 16:29:17.80 ]: 393じゃないけど試してみた
__global__ void kernel(double dtemp){
}

int main(){
double temp=0.0;
double dtemp;
// cudaMalloc((void**)&dtemp,sizeof(double));
// cudaMemcpy(dtemp,&temp,sizeof(double),cudaMemcpyHostToDevice);
kernel<<<1,1>>>(dtemp);//(1)
// cudaMemcpy(&temp,dtemp,sizeof(double),cudaMemcpyDeviceToHost);
return 0;
}

これでコンパイルも実行もおkだったけど
397 名前：デフォルトの名無しさん mailto:sage [2012/01/30(月) 16:57:38.16 ]: >>395
dtempはdevicePointerなんだから、Hostコード中でデリファレンスしちゃダメ。
double値を渡したいだけなら引き数はdoubleで充分。
それと、※1は宣言じゃなくて呼び出しだから構文としてはあっている。

そもそも何がしたいのだろう。値を渡したいなら値を渡せばいいし、
ポインタを渡したいならポインタをそのまま渡せばいい。
devicePointerをHostでデリファレンスしちゃダメだし、
(その後どうせdevice側でデリファレンスできないから)HostPointerをdeviceにコピーしてもダメ。
398 名前：デフォルトの名無しさん [2012/01/30(月) 17:04:05.78 ]: >>396
>>397
色々ごちゃまぜだったようです。スッキリしました。
ありがとうございました。
399 名前：デフォルトの名無しさん mailto:sage [2012/01/30(月) 20:58:23.57 ]: これからCUDAの勉強を始めようと思っています。
GPGPUプログラミングはシェアードメモリを有効活用できるかがポイントであるように思います。
そこで、あえて最新のグラボに買い換えず、キャッシュのないTesla世代のGTX285を所持し続けています。
キャッシュがあると、シェアードメモリをうまく使えてなくても、それをカバーしてしまって、
うまくプログラムできているかどうかを把握し辛いと考えたからです。

ただ、心配なのが、Fermi世代以降に発売された書籍だと、内容がTesla世代のGPUにそぐわない
といった状況が起きてしまうことです。
GTX285のまま勉強するに際して気をつけること等あれば教えてください。
問題がなければ関連書籍（日本語）を全部購入して読もうと思います。

よろしくお願いします。
400 名前：デフォルトの名無しさん mailto:sage [2012/01/30(月) 21:08:43.55 ]: ＞＞GPGPUプログラミングはシェアードメモリを有効活用できるかがポイントであるように思います。
そんなの計算する内容によるわ
活用しようがないのもあるからな

計算（使用）目的は明確なの？
401 名前：399 mailto:sage [2012/01/30(月) 21:39:43.45 ]: >>400
信号処理に使う予定です。
FIRフィルタリングやFFTなどを大きな1次元、あるいは2次元配列に適用したいと思っています。
402 名前：デフォルトの名無しさん mailto:sage [2012/01/30(月) 21:45:58.41 ]: 俺は400じゃないが、気をつけることは、勉強がひと段落していざ使うときになって知識が陳腐化してても泣かないことかな。
ハードウェアはどんどん変わるし、TeslaやFermiで良しとされたことが数年後に通用するかはわからないからね。
せっかく今から始めるならFermi・Keplerを使ったほうがいいと俺は思うけど、
10年後にはどれも等しく陳腐化してるだろうから長期的に考えると確かにTeslaでもいいのかもしれない。
ただ長期的に考えること自体がGPGPUでは…と1行目にループする。
403 名前：デフォルトの名無しさん mailto:sage [2012/01/30(月) 22:12:50.10 ]: 古い方がいろいろ制限があって，勉強になるとは思う．
404 名前：デフォルトの名無しさん mailto:sage [2012/01/30(月) 22:17:32.23 ]: ｆｆｔはcpu側からはインターフェースあるけどディバイス内では自前で作るしかない
405 名前：デフォルトの名無しさん mailto:sage [2012/01/30(月) 22:18:26.89 ]: > キャッシュがあると、シェアードメモリをうまく使えてなくても、それをカバーしてしまって、
> うまくプログラムできているかどうかを把握し辛いと考えたからです。

シェアードメモリーを使うようにプログラミングすればよいかと

どのメモリーを使っているのか分からない，と言うことはないと思う

あれば，そこをしっかり勉強すべきかと
406 名前：デフォルトの名無しさん mailto:sage [2012/01/30(月) 22:19:12.36 ]: そう言えば古いのは倍精度の計算ができないのでは？？
407 名前：デフォルトの名無しさん mailto:sage [2012/01/30(月) 22:22:30.66 ]: まあ，デバイスメモリーだけで計算してみるとか，
シェアードメモリーも使って計算してみるとか，
いろんなプログラムを作って比較するのが良いかと
408 名前：デフォルトの名無しさん mailto:sage [2012/01/30(月) 22:26:01.98 ]: 最初はシェアドするスレッドはどれだってだけでも頭が痛くなった
409 名前：402 mailto:sage [2012/01/30(月) 23:28:19.90 ]: あとTeslaでやるなら、FermiとTeslaでバンクコンフリクトの発生条件が違ってるから
Teslaでのこれの回避には深入りしないほうがいいかと。

他にシェアードメモリの大きさなどの量的な違いがあるけど、そういった量的な
制約のために質的なというかアルゴリズムの種類を変えるようなことまで深入りするのも
避けたほうが。だってFermiやそれ以降で量が変わると一から見直しになるから。
これはFermi使ってる人もいずれ同じことなんだけど、勉強じゃなくて性能が欲しいんだからしょうがない。
410 名前：デフォルトの名無しさん mailto:sage [2012/01/30(月) 23:57:48.21 ]: sharedがregister並に速いなんて嘘だから。
occupancyなんて上げるな。
warpあたりのregisterを増やして
ILPを利用してレイテンシを隠蔽すべきだと。

www.cs.berkeley.edu/~volkov/volkov10-GTC.pdf
411 名前：399 mailto:sage [2012/01/31(火) 00:31:29.61 ]: >>403
やはり、そうですか！

>>404
参考になります。
FFTは重要になるのでしっかりと取り組みたいです。

>>405
>>407
>>408
なるほど、どのメモリかをよく意識してプログラミングします。
CPUの最適化に取り組んでいて思ったのですが、
キャッシュは便利な反面、こちらから意図的に挙動を制御できないことが
パフォーマンスの考察を難しくしてしまう側面もあると感じました。
キャッシュなしのTeslaで鍛えたいです。

>>406
単精度しか使わないので問題なしです。
もし半精度でパフォーマンスが上がるなら、そのほうがイイくらいです。

>>402
>>409
世代を跨ぐにあたって非常に参考になるご助言をありがとうございます。
各種メモリの特性や帯域を意識して取り組むことで、
固有のデバイスに限定されない定性的なところを理解したいと思います。
Keplerではこんな組み方がイイかもしれないな、と思索できたら楽しそうです。

>>410
興味深い資料をありがとうございます。
各部のレイテンシ、スループットを頭に叩き込みます。
412 名前：デフォルトの名無しさん mailto:sage [2012/01/31(火) 08:31:17.17 ]: >>411
最終的に動かす機器で最もよい性能が出るように最適化するべきだし
自動でやってくれることを手でやるってのは勉強にはなるかもしれないけど意味ないのではって思うけど。
コンパイラの最適化を切ってアセンブリ書くようなものでしょ。
413 名前：デフォルトの名無しさん mailto:sage [2012/01/31(火) 12:08:28.56 ]: 時間がいくらでもあるorそういう研究ならともかく
短期間で組めてそこそこのパフォーマンスが出せればいいなら
キャッシュのあるFermiで適当に書いてもいいと思うんだけどね
414 名前：399 mailto:sage [2012/01/31(火) 21:32:30.15 ]: >>412
おっしゃる通りです
そのあたりは目的をよく考えて判断したいです。
今回はどちらかというと、GPUのポテンシャルを見極めたいというのが主です。
コードの可読性を優先して、性能はコンパイラやキャッシュ等にできるだけ委ねるというスタイルも
今後検討していきたいと思っています。

>>413
今回は研究用途ですが、
将来的にはFermi以降のGPUで、ちゃっちゃと書いたコードが快適に動くというところを目指したいです。
ただ、あまりにも下手なプログラムにならないよう、最初に少し深くやっておきたいです。
415 名前：デフォルトの名無しさん mailto:sage [2012/02/02(木) 04:23:36.19 ]: レジスタといえば、GK104では強化されるのか、それともさらにピーキーになるのか気になる。
GF104はグラフィック特化で、レジスタ数やワープ数の上限に悩まされたし。

とりあえずアーキテクチャが一般公開されるまで待つしかないかな。
416 名前：デフォルトの名無しさん mailto:sage [2012/02/04(土) 00:16:58.91 ]: また出てるな、これで最後らしいが
ttp://page4.auctions.yahoo.co.jp/jp/auction/d125096624
417 名前：デフォルトの名無しさん mailto:sage [2012/02/08(水) 09:16:57.43 ]: 誰かいるかな…

includehoge.blog.fc2.com/blog-entry-71.html
このブログのソースをコンパイルして動かしてみようと思ったのですが

1>cudamatrix.obj : error LNK2019: 未解決の外部シンボル _cufftDestroy@4 が関数 _main で参照されました。
1>cudamatrix.obj : error LNK2019: 未解決の外部シンボル _cufftExecC2C@16 が関数 _main で参照されました。
1>cudamatrix.obj : error LNK2019: 未解決の外部シンボル _cufftPlan1d@16 が関数 _main で参照されました。
1>../../bin/win32/Debug/template.exe : fatal error LNK1120: 外部参照 3 が未解決です。

というエラーが出ます。
cufft.h自体は読み込んでいて、引数の過少があるときはそれに対してエラーを返すので
定義されていないという扱いになっているわけではないみたいなんですけど

解決法があったら教えて下さい
418 名前：デフォルトの名無しさん mailto:sage [2012/02/08(水) 09:17:57.58 ]: 誰かいるかな…

includehoge.blog.fc2.com/blog-entry-71.html
このブログのソースをコンパイルして動かしてみようと思ったのですが

1>cudamatrix.obj : error LNK2019: 未解決の外部シンボル _cufftDestroy@4 が関数 _main で参照されました。
1>cudamatrix.obj : error LNK2019: 未解決の外部シンボル _cufftExecC2C@16 が関数 _main で参照されました。
1>cudamatrix.obj : error LNK2019: 未解決の外部シンボル _cufftPlan1d@16 が関数 _main で参照されました。
1>../../bin/win32/Debug/template.exe : fatal error LNK1120: 外部参照 3 が未解決です。

というエラーが出ます。
cufft.h自体は読み込んでいて、引数の過少があるときはそれに対してエラーを返すので
定義されていないという扱いになっているわけではないみたいなんですけど

解決法があったら教えて下さい
419 名前：デフォルトの名無しさん mailto:sage [2012/02/08(水) 09:22:20.66 ]: 誰か居るかな・・・いたら助けてください

includehoge.blog.fc2.com/blog-entry-71.html
このブログのソースをコンパイルして動かしてみようと思ったのですが

1>cudamatrix.obj : error LNK2019: 未解決の外部シンボル _cufftDestroy@4 が関数 _main で参照されました。
1>cudamatrix.obj : error LNK2019: 未解決の外部シンボル _cufftExecC2C@16 が関数 _main で参照されました。
1>cudamatrix.obj : error LNK2019: 未解決の外部シンボル _cufftPlan1d@16 が関数 _main で参照されました。
1>../../bin/win32/Debug/template.exe : fatal error LNK1120: 外部参照 3 が未解決です。

以上のようなエラーが出てコンパイルが通りません
構造体の生成や関数の過少には反応するので、cufft.hが読み込めてないってことはないと思うんですが…

解決法があったら教えて下さい
420 名前：デフォルトの名無しさん mailto:sage [2012/02/08(水) 09:22:49.74 ]: 誰か居るかな・・・いたら助けてください

includehoge.blog.fc2.com/blog-entry-71.html
このブログのソースをコンパイルして動かしてみようと思ったのですが

1>cudamatrix.obj : error LNK2019: 未解決の外部シンボル _cufftDestroy@4 が関数 _main で参照されました。
1>cudamatrix.obj : error LNK2019: 未解決の外部シンボル _cufftExecC2C@16 が関数 _main で参照されました。
1>cudamatrix.obj : error LNK2019: 未解決の外部シンボル _cufftPlan1d@16 が関数 _main で参照されました。
1>../../bin/win32/Debug/template.exe : fatal error LNK1120: 外部参照 3 が未解決です。

以上のようなエラーが出てコンパイルが通りません
構造体の生成や関数の過少には反応するので、cufft.hが読み込めてないってことはないと思うんですが…

解決法があったら教えて下さい
421 名前：デフォルトの名無しさん mailto:sage [2012/02/08(水) 09:23:33.66 ]: うわなんか四回も書き込んじゃったごめん
422 名前：デフォルトの名無しさん mailto:sage [2012/02/08(水) 09:25:46.81 ]: コンパイル詳しくないけどlibcufft.so（linuxの場合）のリンクしてる？
windowsならlibcufft.dllだとおもう
423 名前：デフォルトの名無しさん mailto:sage [2012/02/08(水) 09:27:16.52 ]: コンパイルじゃなくてリンクの問題でしょ
424 名前：デフォルトの名無しさん mailto:sage [2012/02/08(水) 09:58:44.53 ]: nvcc -lcufft sample.cu
425 名前：デフォルトの名無しさん mailto:sage [2012/02/08(水) 10:20:43.62 ]: このスレでときどき思うのはC/C++をほとんど知らないのにもかかわらず
CUDAに挑戦しようとする勇壮な挑戦者が多いな、ということ

彼らはその後うまくいってるのだろうか？
426 名前：デフォルトの名無しさん mailto:sage [2012/02/08(水) 13:41:41.97 ]: CUDA.NETも3.0で止まってるしなぁ・・・
427 名前：デフォルトの名無しさん mailto:sage [2012/02/08(水) 21:44:42.85 ]: コンピュート・ケイパビリティによって何が変わったかを俯瞰できるような一覧を掲載したサイトってご存知ありませんか？？
428 名前：デフォルトの名無しさん mailto:sage [2012/02/08(水) 22:16:51.83 ]: >>417
追加のライブラリファイルにナントカ.libを追加すればいけそう。
ナントカは>>422さんのレスから察するにlibcufft.libかなぁ。
このへんは>>425さんの言うとおりcuda云々じゃなくてVisual Studio(C++)を使う上で頻繁に設定するところだねー。

>>427
俺も知りたいなぁそれ。その手の情報はSDKpdfのProgramming GuideとBest Practiceくらいしか知らない（片方だけかも）
429 名前：デフォルトの名無しさん mailto:sage [2012/02/08(水) 22:26:54.49 ]: cufft.libなのだが。
430 名前：デフォルトの名無しさん mailto:sage [2012/02/09(木) 07:17:56.91 ]: てすてす
431 名前：デフォルトの名無しさん mailto:sage [2012/02/09(木) 07:18:12.24 ]: コマンドなら
432 名前：デフォルトの名無しさん mailto:sage [2012/02/09(木) 07:18:18.71 ]: あれ？
433 名前：デフォルトの名無しさん mailto:sage [2012/02/09(木) 07:18:34.70 ]: 以下のよう
434 名前：デフォルトの名無しさん mailto:sage [2012/02/09(木) 07:18:39.54 ]: nvcc fft.cu --compiler-bindir="C:\Program Files\Microsoft Visual Studio 10.0\VC\bin" -arch=sm_10 -L"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.1\lib\Win32" -lcufft
435 名前：デフォルトの名無しさん mailto:sage [2012/02/09(木) 07:19:01.92 ]: ファイル名をfft.cuとしてみた
436 名前：デフォルトの名無しさん mailto:sage [2012/02/09(木) 07:20:23.21 ]: 結果の一部

1014 6.347414 0.587785
1015 6.318954 -0.000000
1016 6.293659 -0.587785
1017 6.271449 -0.951057
1018 6.252275 -0.951057
1019 6.236131 -0.587785
1020 6.222946 0.000000
1021 6.212736 0.587785
1022 6.205431 0.951057
1023 6.201046 0.951057
437 名前：デフォルトの名無しさん mailto:sage [2012/02/09(木) 07:21:26.93 ]: -arch=sm_10 は適宜，変更のこと
438 名前：デフォルトの名無しさん mailto:sage [2012/02/09(木) 07:25:00.70 ]: これでも行けた♪

nvcc fft.cu -lcufft
439 名前：デフォルトの名無しさん mailto:sage [2012/02/09(木) 07:26:45.21 ]: 既出でした．失礼 >> 424
440 名前：デフォルトの名無しさん mailto:sage [2012/02/09(木) 11:41:33.64 ]: GPUの勉強を始めたばかりなのですが、スパース行列の格納形式種類で
ELLの格納方法とメリットが分かりません

調べたらlisのマニュアルに配列方法はあるのですが、説明と図とが噛み合ない点があり
理解できずにいます。

だれか教えてください
441 名前：デフォルトの名無しさん mailto:sage [2012/02/09(木) 13:59:04.44 ]: >>426
CUDA#とかCUDAfy.NETとか代替のライブラリはあるっぽいよ
使ったことないから詳細は知らないけど
442 名前：デフォルトの名無しさん mailto:sage [2012/02/09(木) 14:23:54.18 ]: >>440
LisはGPU上で動くの？
443 名前：デフォルトの名無しさん mailto:sage [2012/02/09(木) 14:39:02.94 ]: sparse　matrixの格納形式はCUDAの問題じゃね～だろ
444 名前：デフォルトの名無しさん mailto:sage [2012/02/09(木) 15:45:25.65 ]: >>440
コアレッシング状態で転送できるようになる。
445 名前：デフォルトの名無しさん mailto:sage [2012/02/09(木) 17:19:58.52 ]: CUDA.NETは使えなくはないぞ♪
446 名前：デフォルトの名無しさん mailto:sage [2012/02/09(木) 18:44:06.72 ]: >>440
スパース行列の格納方法がのっているのがLisのオンラインマニュアルくらいだったので

>>443
ELLはGPUに適した格納方法と、ネットのどこかにあったので関連があるかと思い書き込みました
たしかにCUDA以前の問題ですよね

>>444
なるほど、見えてきました！ありがとうございます
447 名前：デフォルトの名無しさん mailto:sage [2012/02/09(木) 19:02:29.63 ]: >>422->>439
こんなにたくさん答えてくれるとは・・・
おかげで動きましたーありがとう
448 名前：デフォルトの名無しさん mailto:sage [2012/02/09(木) 19:08:48.94 ]: おお、よかった^^
449 名前：デフォルトの名無しさん mailto:sage [2012/02/10(金) 10:39:58.03 ]: detail.chiebukuro.yahoo.co.jp/qa/question_detail/q1481115004
450 名前：デフォルトの名無しさん mailto:sage [2012/02/10(金) 13:35:49.60 ]: >>441
CUDAfy.NETは知らんかった
最近出てきたのね、ありがとう
451 名前：デフォルトの名無しさん mailto:sage [2012/02/10(金) 21:23:36.65 ]: 2次元配列にアクセスするときのインデックス計算で

gy = blockDim.y * blockIdx.y + threadIdx.y;
gx = blockDim.x * blockIdx.x + threadIdx.x;

というのを頻繁にやると思うのですが、
この積和演算ってCUDAコアが使われるのでしょうか？
それとも、CUDAコアとは別にアドレス演算器（？）的なハードウェアがあって、
それで計算してくれるのでしょうか？
後者だとうれしいです。
452 名前：デフォルトの名無しさん [2012/02/10(金) 21:37:17.01 ]: アドレスもプログラムもデータ
453 名前：デフォルトの名無しさん mailto:sage [2012/02/10(金) 23:52:12.69 ]: >>451

CUDAコア上に展開された（定義された）ブロックの中で演算をしていると思う
454 名前：451 mailto:sage [2012/02/11(土) 01:18:08.05 ]: >>453
CUDAコアのサイクルを消費して計算しているということですね？
ありがとうございました。
455 名前：デフォルトの名無しさん mailto:sage [2012/02/11(土) 01:24:48.08 ]: インデックス計算かどうかをGPUがどう見分けるというんだ
456 名前：451 mailto:sage [2012/02/11(土) 10:05:14.94 ]: >>455
単にCUDAの予約語であるblockDimやthreadIdだけで構成されていて、
○○ * ○○ + ○○ の積和の形になっていれば、
専用の回路に計算させる、とかできそうなものですが。
CPUではメモリアドレス式のカタチを解析してアドレス演算器に割り当てています。
457 名前：デフォルトの名無しさん mailto:sage [2012/02/11(土) 11:29:23.45 ]: そもそもグリッド，ブロック，スレッドの概念を理解しているのかな？
458 名前：デフォルトの名無しさん mailto:sage [2012/02/11(土) 11:29:54.96 ]: ○○ * ○○ + ○○ だけではプログラムは作れないと思うのだが
459 名前：デフォルトの名無しさん [2012/02/11(土) 11:55:45.33 ]: >>456
それ、どのCPUでの話ですか？
460 名前：デフォルトの名無しさん [2012/02/11(土) 19:16:02.78 ]: 環境：　CUDA 4.1 + Visual Studio 2010 Pro. + Windows 7 Pro. 64bit

C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.1\C\common\cutil_vs2010.sln
をバッチビルドして、
cutil64.lib, cutil64.dll,
cutil64D.lib, cutil64D.dll
を作ろうとしたところエラーがでてしまいます。
エラーは出るのですが、cutil64.lib　以外は生成でき、
そのライブラリを使った行列積のサンプルプログラムもコンパイル＆実行できました。
＃ちょっと気持ち悪いですが。

しかし、cutil64.lib　はどうしても生成できません。
何が悪いのでしょうか。

cutil64.lib, cutil64.dll,
cutil64D.lib, cutil64D.dll
をビルドしたときのエラーメッセージは以下の通りです。
461 名前：デフォルトの名無しさん [2012/02/11(土) 19:16:34.66 ]: ------ すべてのリビルド開始: プロジェクト: cutil, 構成: Release x64 ------
C:\Program Files (x86)\MSBuild\Microsoft.Cpp\v4.0\Microsoft.CppClean.targets(74,5): warning : パス 'C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.1\C\common\lib\x64\cutil64.dll' へのアクセスが拒否されました。
bank_checker.cpp
cmd_arg_reader.cpp
cutil.cpp
stopwatch.cpp
stopwatch_win.cpp
Generating Code...
LINK : warning LNK4075: /INCREMENTAL は /OPT:ICF の指定によって無視されます。
LINK : fatal error LNK1104: ファイル 'lib\x64\\cutil64.dll' を開くことができません。
------ すべてのリビルド開始: プロジェクト: cutil, 構成: Debug x64 ------
bank_checker.cpp
cmd_arg_reader.cpp
cutil.cpp
stopwatch.cpp
stopwatch_win.cpp
Generating Code...
ライブラリ C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.1\C\common\lib\x64\\cutil64D.lib とオブジェクト C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.1\C\common\lib\x64\\cutil64D.exp を作成中
LINK : fatal error LNK1158: 'cvtres.exe' を実行できません。
========== すべてリビルド: 0 正常終了、2 失敗、0 スキップ ==========
462 名前：デフォルトの名無しさん [2012/02/11(土) 19:21:14.69 ]: 追記です。
cutil64.lib　がないので行列積サンプルプログラムは、Releace モードではビルドできていません。
ビルド＆実行でてきているのは、debug モードです。
463 名前：デフォルトの名無しさん mailto:sage [2012/02/11(土) 20:12:55.27 ]: Releace + x64モードでcutil64.dllはできるが，cutil64.libはできないと言うことですか？
464 名前：デフォルトの名無しさん mailto:sage [2012/02/11(土) 20:14:45.64 ]: エラーが出てもライブラリはできている，と言うことはないですか？

この中に↓

C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.1\C\common\lib\x64
465 名前：デフォルトの名無しさん mailto:sage [2012/02/11(土) 20:16:27.24 ]: 失礼　Releace -> Release
466 名前：デフォルトの名無しさん mailto:sage [2012/02/11(土) 20:26:05.18 ]: まあ、CUDAのインデックス計算ほど
アホ臭いものは無いと思うがな。
テクスチャアドレスではもっと複雑な計算しているのに。
467 名前：デフォルトの名無しさん mailto:sage [2012/02/11(土) 20:50:55.57 ]: CUDAを数値計算に使おうという話だよ
468 名前：デフォルトの名無しさん [2012/02/11(土) 22:18:37.30 ]: >>463
>>464
お返事ありがとうございます。

C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.1\C\common\cutil_vs2010.sln
のバッチビルドの前後で以下のようにライブラリが増えます。

ビルド前
common\lib\x64\　には　cutil64.dll　のみ存在

ビルド後
common\lib\x64\ には cutil64.dll cutil64D.dll cutil64D.lib　の３つになる

cutil64.libは生成されません。ビルド時に　>>461のエラーのためと思われます。
　
469 名前：デフォルトの名無しさん mailto:sage [2012/02/12(日) 00:49:32.40 ]: 変だね．こちらも同じようなエラーが出るが（xx.dllを開始できません），xx.libはできている．

試しに，ソルーションエクスプローラーのcutilを右クリックして，「追加」，「既存の項目」で，cutil64.dll　を追加してみる．．．．とか
470 名前：デフォルトの名無しさん mailto:sage [2012/02/12(日) 00:51:08.01 ]: で，何でビルド前に「cutil64.dll　のみ存在」するのかな？

一度，全部クリーンしてみる．．．．とか
471 名前：デフォルトの名無しさん [2012/02/12(日) 01:08:37.66 ]: >>469
お返事ありがとうございます。
エラーは出るんですね。貴重な情報です。
エラーを気にしないですみます。

>cutil64.dll　を追加してみる．．．．とか
無理を承知でやってみました。当然ですがダメでした。

>>470
そういうものだと思っていましたが、確かにそうですね。
クリーンインストールしてみます。
472 名前：デフォルトの名無しさん [2012/02/12(日) 02:45:15.54 ]: クリーンインストールしてみました。
C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.1\C\common\lib\x64
に最初から　cutil64.dll　が作られるみたいです。

なお、インストールたCUDAファイルは、
　devdriver_4.1_winvista-win7_64_286.16_notebook.exe
　cudatoolkit_4.1.28_win_64.msi
　gpucomputingsdk_4.1.28_win_64.exe
の３つです。
473 名前：デフォルトの名無しさん mailto:sage [2012/02/12(日) 07:49:26.70 ]: cutilのクリーンを実行すれば，すべてのdll，libが消えるはずなのに，もし， cutil64.dllが残っているとすれば，
書き込み，読み込みの権限がないと言うことでは？

cutil64.dll　はマニュアルで消せますか？
474 名前：デフォルトの名無しさん [2012/02/12(日) 15:39:25.13 ]: >>437
お返事ありがとうございます。
cutil64.dll　をマニュアルで削除して、再度
　　C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.1\C\common\cutil_vs2010.sln
でビルドしたら、すべてのライブラリが生成されました。
ご推察の通り、cutil64.dll　にアクセスできなかったのが原因だったようです。
本当にありがとうございました。

＜追記＞
adminユーザーで環境構築作業をしていたので、
cutil64.dll にアクセスできなかった原因は謎のままです。
自分で構築してしまえば、その後は、上書きビルドは問題なくできます。

同じく、CUDA-SDKの中に、simpleTemplate　というソースコードがあるのですが、
このソースコードを書き換えて、Releaseモードでビルドすると、
インストール初期からある、simpleTemplate.exe　にコンフリクトするらしく、
ビルドエラーが出ました。この場合はも、初期からある、simpleTemplate.exe　を削除すれば
問題なく繰り返しビルドができるようになりました。
同じ原因だと思います。
475 名前：デフォルトの名無しさん [2012/02/12(日) 15:40:34.77 ]: ↑　番号ミスです。ごめんなさい。
>>473　が正しいです。
476 名前：デフォルトの名無しさん mailto:sage [2012/02/12(日) 15:49:09.03 ]: よかった，よかった♪
477 名前：デフォルトの名無しさん [2012/02/12(日) 22:34:46.87 ]: OS：32bit　CUDA:32bit　VisualStudio:32bit
でCUDAプログラムをしている人が、速度を上げることを目的に
OS：64bit　CUDA:64bit　VisualStudio:64bit
に変更することって意味あります？
478 名前：デフォルトの名無しさん mailto:sage [2012/02/12(日) 23:58:34.00 ]: ないと思う。ホスト側で2GB以上のメモリを使いたいわけでもないのだろうし。

64bitビルドして動かすとポインタがレジスタを2つ使ってしまうという話があるので、逆に遅くなる可能性も。
本当に64bitビルド時にデバイス側でもポインタが64bitで扱われるかどうかを
自分の目で確かめたわけではないけど、少なくともここではそういうことを言っている。

高速演算記第3回「チューニング技法その１　ＣＵＤＡプログラミングガイドからピックアップ」 | G-DEP
www.gdep.jp/column/view/3
＞ここでレジスタ数は32bitレジスタの数で、long long, double値や64ビット環境を対象としたCUDAプログラムでは
＞ポインタなどが64bitとして扱われますのでレジスタを2つ使用することになります。
479 名前：デフォルトの名無しさん [2012/02/13(月) 10:45:03.24 ]: >>478

なるほど。ありがとうございます。

これと少し環境が違う場合ですが、
　　OS：64bit　CUDA:32bit　VisualStudio:32bit
と、
　　OS：64bit　CUDA:64bit　VisualStudio:64bit
では、速度差はどうなると思われますか？

ネット上でよく目にするmatrix積のサンプルコードで実験したところ、
同程度か、どうかすると前者の方が速い場合がありました。

前者は、WOW上で動いていると思われるので、
遅いだろうと予測したのですが、それに反する結果でした。
これをどう解釈すればよいか謎なんです。
480 名前：デフォルトの名無しさん [2012/02/13(月) 10:46:42.16 ]: 479です。OSは、Windows7 64bit Pro です。
481 名前：デフォルトの名無しさん mailto:sage [2012/02/13(月) 11:17:41.23 ]: プログラムのどの部分を計測したのであろうか？

CUDAの部分のみ？？
482 名前：デフォルトの名無しさん [2012/02/13(月) 12:26:28.94 ]: CUDAの部分のみです。
483 名前：デフォルトの名無しさん [2012/02/13(月) 12:35:43.06 ]: CUDAはOSに依存せずに動くとかで、OSのbit数には無関係ということですか？
そのため、
>>478で説明して頂いたように、同程度もしくは32bitが若干速くなる場合もあるということでしょうか。
484 名前：デフォルトの名無しさん mailto:sage [2012/02/13(月) 13:17:49.82 ]: ちなみに測定時間はいくらでしょうか？

演算量を増やすとどうなるでしょうか？
485 名前：デフォルトの名無しさん [2012/02/13(月) 16:09:51.62 ]: ソースは、以下のHPの最後の「Matrix_GPU」です。
www.gdep.jp/page/view/218

計算は、スレッドサイズ一杯で組まれているようですので、
　　#define MATRIX_SIZE 1024/*行列１辺の数*/
をこれ以上上げることができませんでした。

ちなみに、ノートパソコン搭載のGPUチップは、GT-540M。
Releaseでビルドして　650ms付近を中心に数10msばらつきます。
win32とx64で大差はなく、
win32の方が平均10ms程度速いような気がするという程度です。
486 名前：デフォルトの名無しさん mailto:sage [2012/02/13(月) 17:03:00.91 ]: CUDAの計算のところを１０００回程度，繰り返すとかはできるでしょうか？

そうすれば有意な差が出るかも知れません．
487 名前：デフォルトの名無しさん [2012/02/13(月) 19:18:58.94 ]: 200回繰り返してみました。

OS：64bit　CUDA:32bit　VisualStudio:32bit 　→　131,671 ms
OS：64bit　CUDA:64bit　VisualStudio:64bit 　→　132,146 ms

0.3%の差で、前者が速かったです。
100回繰り返しでも、0.3%差でした。
488 名前：デフォルトの名無しさん mailto:sage [2012/02/13(月) 19:39:39.79 ]: なるほど，有り難うございます．

有意な差が出たと言うことですね．

では，後は分かる人，よろしく♪

ε=ε=ε=ε=ε=┌(；　・＿・)┘
489 名前：デフォルトの名無しさん mailto:sage [2012/02/14(火) 10:35:58.19 ]: GPU側では64bitのメリットは使用可能なメモリが増えることだけで、その代わりにアドレス計算の命令数や使用レジスタ数が増えるデメリットがある。
前はGPU側だけ32bitモードで動かすという恐ろしい手法があったが、今使うのはなかなか難しい。
490 名前：デフォルトの名無しさん mailto:sage [2012/02/14(火) 16:21:21.98 ]: ptxオプションつけてptx出力を見れば見当がつくね。
64ビットアドレッシングだとアドレス計算の為にint32→int64の変換がしばしば行なわれる。
491 名前：デフォルトの名無しさん [2012/02/14(火) 18:19:16.13 ]: ざっくり捉えると、32bitと64bitでは、
　CUDA自体はさほど変わらない。
　CUDAと切り離されてコンパイルされるC++言語部分は、64bitの方がいくらか速い。
という感じでしょうか。
492 名前：デフォルトの名無しさん mailto:sage [2012/02/14(火) 18:24:53.01 ]: >>491
アドレス計算多用したり、レジスタ沢山使うようなカーネルだと
CUDA自体32bit版の方が無視できないほど
速くなる可能性はある。
493 名前：デフォルトの名無しさん [2012/02/14(火) 18:53:10.54 ]: doubleが速くなるっていう噂はどうなったの？
494 名前：デフォルトの名無しさん mailto:sage [2012/02/14(火) 22:53:20.23 ]: GPU上にあるコンスタントキャッシュメモリって、CUDAではどのようにして使われるのでしょうか？
計算の中にリテラルがあった場合（CPUでは命令の即値として埋め込まれますが）や、
constantを付けて宣言した定数が割り当てられたりするのでしょうか？？
495 名前：デフォルトの名無しさん mailto:sage [2012/02/16(木) 05:51:50.91 ]: llvmで使えるようになったのなら
boostも使える？
496 名前：デフォルトの名無しさん mailto:sage [2012/02/16(木) 08:10:03.40 ]: >>494

適当に数値を入れて，確かめてみるとか
497 名前：デフォルトの名無しさん mailto:sage [2012/02/16(木) 08:17:55.24 ]: >>495

２年前にパッチが出ているようです．boostが何か知らんけど

https://svn.boost.org/trac/boost/ticket/3919
498 名前：494 mailto:sage [2012/02/17(金) 00:11:19.70 ]: >>496
試してみました。
const変数もリテラルもレジスタにマッピングされました。

ひょっとして、DirectXAPIなどのグラフィックスAPIを利用していた場合に
使われるだけ（要は旧来通りの使われ方）かも･･･
プロファイラを見ると、テクスチャキャッシュのヒット/ミスヒットが見れるようになっていますが
これなんかはグラフィックスAPI専用のような気が。
しかし同じようにコンスタントキャッシュの項目もあって然りと思うが在らず。

う～ん、気にしないことにしますｗ
499 名前：デフォルトの名無しさん mailto:sage [2012/02/17(金) 06:31:17.52 ]: 貴重な報告、ありがとうございます♪
500 名前：デフォルトの名無しさん mailto:sage [2012/02/17(金) 07:54:24.51 ]: GPUを演算器としてフルに使いたい場合って
ビデオ出力をオンボードのものになるようBIOSで切り替えた方がいいのでしょうか
linuxだったらxconfigいじるだけでいけるとか？
501 名前：デフォルトの名無しさん mailto:sage [2012/02/17(金) 07:57:38.45 ]: フルの定義は？
502 名前：デフォルトの名無しさん mailto:sage [2012/02/17(金) 07:58:52.97 ]: >>494
CPUでも即値として埋め込めない配列や文字列リテラルは
.rdataセクションのデータとして書き込まれるのと同様に、
GPUでも即値として使えないconstデータはコンスタント領域に置かれる。

で、コンスタント領域の値を直接オペランド指定できない
命令を使う場合は当然レジスタを介して使用することになる。

逆に一部の変数にデフォルトで定数が入っているなんてのは
ABIで決まったもの以外は普通無理だから、コンスタント境域などの
メモリに置いて読み込む以外やり様が無い。
503 名前：デフォルトの名無しさん mailto:sage [2012/02/17(金) 08:33:23.95 ]: >>500
普通はGPUカードに何もモニタを繋がなければ(≒出力する設定をしなければ)
GVRAMをほぼ全部使い切れるよ。
Linuxなら、Xを立ち上げなければGPUカードから出力していてもかなり使えるよ。
ついでに、3秒だか5秒だからの制限もなくなるからGPUカーネルで長時間計算させられるよ。
504 名前：デフォルトの名無しさん [2012/02/17(金) 20:22:12.27 ]: VRAM 4GBのボードが六千円台
akiba-pc.watch.impress.co.jp/hotline/20120218/etc_inno.html
505 名前：494 mailto:sage [2012/02/18(土) 16:32:22.87 ]: >>502
ありがとうございます。
コンスタント変数の扱われ方、レジスタが必要な理由、理解しました。

書籍を読んでいたら、コンスタントメモリおよびテクスチャメモリについて書かれているのを見つけました。

コンスタントメモリは__constant__修飾子を付けて宣言することで割り当てられるようです。
物理的にはVRAMに置かれ、GPU内部のコンスタントキャッシュでキャッシュが効くものと思われます。
空間フィルタリングなどでフィルタ係数を格納しておくのに適しているのではと思いました。
（繰り返し参照されるのでキャッシュが効きそう）

テクスチャメモリはテンプレートクラス型のtexture<･･･>で宣言することで割り当てられるようです。
物理的にはVRAMに置かれ、GPU内部のテクスチャキャッシュでキャッシュが効くものと思われます。
さらに、グラフィック用途時の同様、テクスチャフィルタリングやサンプリング方式などのハードワイヤード回路機能を
利用できるよう、それらを指定するプロパティを保持しているようです。
506 名前：デフォルトの名無しさん mailto:sage [2012/02/18(土) 21:02:09.38 ]: 複数カーネルを呼び出すんだけどカーネルのソースはそれぞれ別ファイル
それぞれのカーネルから同じDevice関数を呼び出したいんだけどファイルまたがることできないようだし
ソースごとに同じ内容の別名関数作るしかないのかな？
507 名前：デフォルトの名無しさん mailto:sage [2012/02/18(土) 22:43:38.30 ]: ファイルが別でも，プログラムとしてつながっていれば問題はないと思うが．

関数のプロトタイプ宣言をしっかりしておけば良いかと．．．
508 名前：デフォルトの名無しさん mailto:sage [2012/02/18(土) 22:46:25.76 ]: 共通のdevice関数を別のヘッダに纏めて
カーネルソースからincludeして使う。

device関数はどうせinline展開されるのだから
通常のプログラムでのinline関数と同様に扱えばいい。
509 名前：デフォルトの名無しさん mailto:sage [2012/02/18(土) 22:54:19.36 ]: >>508
サンクス
でも今やってみたんだけど全ファイルで展開されるんで
multiple definition of 関数名
ってエラーになっちゃう

実際やってみました？
510 名前：デフォルトの名無しさん mailto:sage [2012/02/18(土) 23:22:56.53 ]: カーネル起動時、1スレッドの起動につき1クロックかかり、128スレッドなら128クロックかかると書籍に書いてあるのですが、
これはCPU側でもかかるのでしょうか？
それとも、CPU側はカーネル起動の合図を送ってすぐ制御を戻し、
GPU側の制御ブロックがクロックをかけてスレッドを起動するというかたちでしょうか？
511 名前：デフォルトの名無しさん [2012/02/19(日) 01:11:11.40 ]: 教えて欲しいことがあります．

CPU側で複数スレッドを立てて，互いに独立した処理を行っています．
そして，それぞれスレッドにおいて，CUDAを使って画像処理をさせたいのですが，（たとえば，グレースケール化とか）
この時，CUDAでの処理にテクスチャメモリを使いたい場合は，どのようなコードを書けばいいのでしょうか？
テクスチャメモリを使う場合，グローバルで宣言しなけらばならないですよね？

たとえば，5枚のRGB画像をグレースケール化するときに，
CPU側で5スレッド立てて，各スレッドでCUDAを使ってグレースケールへの変換処理をしたいのですが，
テクスチャメモリをグローバルで宣言するとおかしなことになるきがするんですが．

どなたか教えて頂けないでしょうか？
宜しくお願い致します．
512 名前：デフォルトの名無しさん mailto:sage [2012/02/19(日) 01:21:41.45 ]: CUDA対応のGPUを5枚挿す。
またはGTX590などを3枚挿す。
513 名前：デフォルトの名無しさん [2012/02/19(日) 01:26:39.83 ]: >>512
回答ありがとうございます．
テクスチャメモリを使った処理の場合，マルチGPUじゃないと出来ない
とういことでしょうか？
514 名前：デフォルトの名無しさん [2012/02/19(日) 01:44:29.68 ]: 公式マニュアル見れば分かること
試してみれば分かること

これを聞く人が多すぎる
515 名前：デフォルトの名無しさん mailto:sage [2012/02/19(日) 01:53:02.78 ]: 4-way SLI でも、ホスト上の4スレッドで利用する場合、
各々1GPUを独立して割り当てるため、4GPU間で通信して仕事をこなす本来的なSLIにはならない、って認識でイイんですよね？
（SLIコネクタがなくても動く？）
516 名前：デフォルトの名無しさん mailto:sage [2012/02/19(日) 02:32:56.76 ]: >>511
テクスチャを5個宣言じゃダメなの？
517 名前：デフォルトの名無しさん [2012/02/19(日) 02:47:03.16 ]: >>516
回答ありがとうございます．
CPU側でスレッドを立てる段階で決めたいので，
できることなら，可変にしたいと考えています．

514さんは「これを聞く人が多すぎる」と仰っていますが，
そもそも，このようなことをシングルGPUでやろうとするのが間違いなのでしょうか？
518 名前：デフォルトの名無しさん [2012/02/19(日) 02:59:57.92 ]: いやいや、まず試せば分かることじゃん。
その労力を2chに押しつけるのは、今この瞬間は楽では良いかもしれないけど、
自分の成長には繋がらないよ。
519 名前：デフォルトの名無しさん mailto:sage [2012/02/19(日) 05:31:26.29 ]: >>513
OS側のスレッド何本あろうとGPUが1つしかないんだから
順番に処理されるだけだろってことでしょ
520 名前：デフォルトの名無しさん [2012/02/19(日) 05:43:18.09 ]: concurrent kernel execution…
521 名前：509 mailto:sage [2012/02/19(日) 07:12:00.07 ]: >>508
失礼しました
インライン関数化で無事できました
ありがとうございました
522 名前：デフォルトの名無しさん [2012/02/19(日) 11:42:52.58 ]: >>518
心遣い感謝します．

>>519
>>520
回答ありがとうございます．
やっぱり，そうなんですね．
「concurrent kernel execution」ついて調べて確信しました．

回答・アドバイスして下さった皆様，ありがとうございました．
おかげで解決しました．
523 名前：デフォルトの名無しさん [2012/02/19(日) 11:49:26.29 ]: 次に同じ疑問を持たれた方のために，
参考資料のアドレスを貼っておきます．
ttp://www.nvidia.co.jp/docs/IO/81860/NVIDIA_Fermi_Architecture_Whitepaper_FINAL_J.pdf
上記アドレスにあるpdfの「コンカレントカーネル実行」に書かれていることが，参考になるかと思います．
524 名前：デフォルトの名無しさん mailto:sage [2012/02/19(日) 16:33:50.20 ]: 誰が質問者なのかわからないので謎な流れだが、とりあえずグレースケール化にテクスチャは不要だと思う今日この頃。
525 名前：デフォルトの名無しさん [2012/02/19(日) 17:00:58.89 ]: NPPで一発解決だよな
526 名前：デフォルトの名無しさん mailto:sage [2012/02/20(月) 22:07:59.44 ]: atomicCASのCASってどういう意味ですか？？
527 名前：デフォルトの名無しさん mailto:sage [2012/02/20(月) 22:24:51.56 ]: Compare And Swap
528 名前：526 mailto:sage [2012/02/20(月) 22:28:38.87 ]: >>527
ありがとうございました！
529 名前：デフォルトの名無しさん mailto:sage [2012/03/07(水) 03:16:50.68 ]: ｃｕｄａの本には2次元配列より1次元の方が速いってあったけど実際どのくらい違うんですか？
また、コピーとgpu上での演算の両方で遅くなるんですか？
530 名前：デフォルトの名無しさん [2012/03/07(水) 03:39:04.89 ]: なんでこのスレ突然止まってたの？
卒論修論シーズンが終わったから？
531 名前：デフォルトの名無しさん mailto:sage [2012/03/07(水) 06:34:38.99 ]: >>529
一行目: 速い
二行目: 遅くなる

馬鹿?
532 名前：デフォルトの名無しさん mailto:sage [2012/03/07(水) 16:56:25.77 ]: 学会はこれからだというのに
533 名前：デフォルトの名無しさん mailto:sage [2012/03/07(水) 17:15:29.75 ]: > ｃｕｄａの本には2次元配列より1次元の方が速いってあったけど実際どのくらい違うんですか？

メモリーのアクセス時間の早さのこと？

ちなみに２次元であっても，例えば f[i, j] で，先にｊを変化させるか，i　を変化させるのかで
アクセス時間が違ったと思う．

要はメモリーが並んでいる順番にアクセスするのが早いはず．どの位，早くなるかは知りませぬ♪

> また、コピーとgpu上での演算の両方で遅くなるんですか？

コピーとは？ CPUとGPU間のメモリーの転送のことかな？？
534 名前：デフォルトの名無しさん mailto:sage [2012/03/08(木) 10:13:29.60 ]: 人いなくなったな
春になってGPUネタで研究する学生が増えるのを待つしかないか
535 名前：デフォルトの名無しさん mailto:sage [2012/03/08(木) 11:44:54.65 ]: 二次元配列なんて存在しねぇ
って考えると楽なのに
536 名前：デフォルトの名無しさん mailto:sage [2012/03/08(木) 19:50:20.17 ]: 昔のオカルト本では「謎の四次元」「四次元失踪」とかあって、
「四次元」とか謎めいた雰囲気を感じた。

今なら四次元配列とか当たり前だｗ
537 名前：デフォルトの名無しさん mailto:sage [2012/03/09(金) 01:08:22.32 ]: >>536
そうそうｗ
四次元の神秘性が薄らいじゃうｗ
538 名前： ◆QZaw55cn4c mailto:sage [2012/03/09(金) 22:10:56.52 ]: ハミルトンの四元数というのがあって近年計算機工学に応用されるようになって云々かんぬん
539 名前：デフォルトの名無しさん mailto:sage [2012/03/09(金) 22:22:34.47 ]: 発見自体は1800年代だったっけ？
たしか橋の上かなんかで思いついたとかｗ
540 名前：デフォルトの名無しさん mailto:sage [2012/03/11(日) 11:39:47.33 ]: HLSLとどう違うの？
541 名前：デフォルトの名無しさん mailto:sage [2012/03/11(日) 11:49:19.59 ]: >>540
C言語っぽくなってる。
542 名前：デフォルトの名無しさん mailto:sage [2012/03/11(日) 14:43:17.16 ]: 四次元配列と四次元ベクトルは別物だろ
後者は要素数4の一次元配列
543 名前：デフォルトの名無しさん [2012/03/11(日) 17:04:08.60 ]: >>540
HLSLはShader Languageなんで整数とか扱えなかったと思う。
あとステップ数も制限があったような。
544 名前：デフォルトの名無しさん mailto:sage [2012/03/11(日) 17:26:35.76 ]: Compute Shaderを記述する言語も
HLSLじゃ無かったっけ？

リソースベースのHLSLと、ポインタ・配列ベースのCUDA
545 名前：デフォルトの名無しさん mailto:sage [2012/03/11(日) 17:38:23.47 ]: リソースベースって言い方、分かり易いね。
HLSLはまさにそんな感じだ。
546 名前：デフォルトの名無しさん mailto:sage [2012/03/12(月) 13:07:18.89 ]: >>541 >>543
サンクス
自由度がまして打ちやすくなってるんだな
547 名前：デフォルトの名無しさん mailto:sage [2012/03/12(月) 13:35:35.95 ]: >>544 >>545
サンクス
548 名前：デフォルトの名無しさん mailto:sage [2012/03/15(木) 21:09:26.09 ]: 誰かいますかね、三つほど質問が。

■__syncthreads()だけでは共有メモリへの書き込みを保証できない？
　１つのスレッドがグローバルメモリから共有メモリにデータを書き込み、
その後全てのスレッドがそのデータを使用して計算を行うような場合、
書き込み後に__syncthreads()だけではなく__threadfence_block()も必要なのでしょうか?
青木本には__threadfence_block()について特に言及ありませんでしたが・・・。

■ブロック内の全スレッドからの同一グローバルメモリへのアクセス
　ブロック内で共通で使用する構造体などをグローバル→共有メモリに移す場合
全スレッドで行うよりもやはり
　　 if(threadIdx.x==0)・・・
のようにした方が良いでしょうか?

■カーネル内でのreturn文の使用悪影響あるか
　スレッドごとに計算を行うか判定をする場合、if文で囲っている例をよく見ますが
これは
if(条件)return;
と書いてはいけないのでしょうか？
上のように書いてもとりあえず計算は流れたのですが何か悪影響はあるでしょうか？
549 名前：デフォルトの名無しさん mailto:sage [2012/03/15(木) 21:25:56.04 ]: >>548
・取り敢えずsyncしか使ってないけど問題になったことはない。
・全スレッドから共有メモリへの書き込みを行なうのは多分遅くなるんじゃないかな?
・どちらで書いても同じこと。普通のCPUのような分岐とは違うことを判っていればOK。
550 名前：548 mailto:sage [2012/03/15(木) 22:39:42.94 ]: >>549
ありがとうございます!!
一個目の_syncthreads()と__threadfence_block()の件ですが、
syncだけだと今日うまくいかなかったもので。
ただ他のバグの影響なども考えられるのでもうちょっと調べてみます。
551 名前：デフォルトの名無しさん mailto:sage [2012/03/16(金) 04:06:36.41 ]: >>548
・__syncthreads()は__threadfence_block()相当の処理を
含んでいた気がするけど気のせいかも。

・全スレッドで同じメモリにアクセスするのはたとえfermiでも遅くなるはず。

・カーネル内部で_syncthreads()使う必要があるなら
　returnは使っちゃ駄目だろう。
552 名前：デフォルトの名無しさん mailto:sage [2012/03/16(金) 20:10:45.37 ]: いまだにアプリ開発環境すらまともに構築できてない・・・
visual studio 2008でやろうと思って一応ビルドは通ったけど
実行するとまずcutil32.dllがありませんって出た。
次にcutil.dllをデバッグ.exeと同じフォルダに置き実行！！

CUDA version is insufficient for CUDART version.
ってなる・・・orz

まずなにからはじめるべきですか？
553 名前：デフォルトの名無しさん mailto:sage [2012/03/16(金) 20:17:42.17 ]: ちなみに.cuの中身は拾ってきたちょっと複雑なコード

#include <stdio.h>
#include <cutil.h>

int main( int argc, char** argv )
{
CUT_DEVICE_INIT(argc, argv);
CUT_EXIT(argc, argv);
return 0;
}

・・・・・orz
#include <stdio.h>
#include <cutil.h>

int main( int argc, char** argv )
{
return 0;
}

これに書き換えると
プログラムが完成し、エラーもなく実行もできる
554 名前：デフォルトの名無しさん mailto:sage [2012/03/16(金) 20:51:20.83 ]: GPUドライバのアップデート
555 名前：デフォルトの名無しさん mailto:sage [2012/03/17(土) 01:39:06.58 ]: >>554
ありがとうございました！！！！！！！！！！
動いた！！！！！！！！
556 名前：デフォルトの名無しさん mailto:sage [2012/03/17(土) 08:54:13.85 ]: おおｗ
よかったな！
557 名前：デフォルトの名無しさん mailto:sage [2012/03/17(土) 12:49:24.04 ]: >LINK : /LTCG が指定されましたが、コードの生成は必要ありません。リンクコマンドラインから /LTCG を削除し、リンカの性能を改善してください。
と表示されるのですがリンクコマンドラインは固定されて編集できません。
解決方法はありますか？
558 名前：デフォルトの名無しさん mailto:sage [2012/03/17(土) 13:37:04.85 ]: >>577
補足：
開発環境はVisualStudio2008
cuda ver 2.3
559 名前：デフォルトの名無しさん mailto:sage [2012/03/17(土) 14:59:10.69 ]: windowsを窓から投げ捨てろ
560 名前：509 mailto:sage [2012/03/17(土) 15:16:23.02 ]: そんなことして道歩いてる人の頭に当たっちゃったら大変ですよ
561 名前：デフォルトの名無しさん mailto:sage [2012/03/17(土) 15:44:28.84 ]: 角に当たったら痛そうだもんね・・・
562 名前：デフォルトの名無しさん mailto:sage [2012/03/17(土) 18:01:09.23 ]: >>557
リンカ -> 最適化 -> リンク時のコード生成 (/LTGG)
C/C++ -> 最適化 -> プログラム全体の最適化 (/GL)
Visual Studio 2008 の使い方なのでスレが違うかも。
563 名前：デフォルトの名無しさん mailto:sage [2012/03/18(日) 15:04:39.12 ]: >>562
ありがとうございます。

CUDA-Zの実行結果はどのように見たらいいですか？
日本のサイトが全然ないです。
564 名前：デフォルトの名無しさん mailto:sage [2012/03/18(日) 15:07:18.71 ]: >>563です
すみません。解決しました。
565 名前：デフォルトの名無しさん mailto:sage [2012/03/21(水) 21:06:39.97 ]: コンスタントメモリキャッシュへのアクセスはバンクコンフリクトとかないんでしょうか？？
566 名前：デフォルトの名無しさん mailto:sage [2012/03/21(水) 22:20:52.26 ]: >>565
そりゃキャッシュはバンクになってないからねー
567 名前：565 mailto:sage [2012/03/21(水) 22:44:39.32 ]: >>566
おお、やっぱり。
できるだけコンスタントメモリ使うようにしまつ。
568 名前：デフォルトの名無しさん mailto:sage [2012/03/21(水) 23:02:34.40 ]: アドレスが静的に解決できないというのが前提だけど
16ポートのSRAMなんてコスト的に不可能だからマルチバンク以外無いんじゃないの？
569 名前：デフォルトの名無しさん [2012/03/22(木) 00:27:08.40 ]: Fermi以前はコンスタントメモリ使う意味あったけど、
Fermi以降はL2キャッシュとあんまり変わらない印象
570 名前：デフォルトの名無しさん [2012/03/22(木) 22:48:01.36 ]: GTX680が発表されたけど、CUDA的には好ましくない方向の進化が多い。。
571 名前：デフォルトの名無しさん mailto:sage [2012/03/22(木) 23:03:25.80 ]: チップ名　　　　　GTX 680　　　GTX 580
GPC*1　　　　　　4　　　　　　　　4
SM*2　　　　　　　8　　　　　　　　16
CUDAコア　　　1536　　　　　　　512
テクスチャーユニット　128　　　　64
ROPユニット　　　32　　　　　48
ベースクロック*3　　1.006GHz　　　772M/1.544GHz
ブーストクロック　　　1.058GHz　　－
メモリー転送レート　6.008Gbps　4.008Gbps
メモリー容量　　　GDDR5 2048MB　　GDDR5 1536MB
メモリーバス幅　　　256ビット　　384ビット
メモリー転送速度　192.26GB／秒　192.4GB／秒
製造プロセス　　　　28nm　　　40nm
補助電源端子　　　　6ピン×2　　8ピン＋6ピン
推奨電源ユニット出力　550W　600W
TDP*4　　　　　195W　　244W
572 名前：デフォルトの名無しさん mailto:sage [2012/03/22(木) 23:06:16.47 ]: GK104はミドルレンジだからGK110は全体的に上回ってくるでしょ
573 名前：デフォルトの名無しさん [2012/03/22(木) 23:16:01.94 ]: 　　kepler誕生おめ！
　　　　　　　　　　.o゜*。o
　　　　　　　　　／⌒ヽ*゜*
　　　∧_∧　／ヽ　　）｡*o　　ｯﾊﾟ
　　 (・ω・)丿゛￣￣' ゜
.　ノ/　 /
　　ノ￣ゝ
574 名前：デフォルトの名無しさん mailto:sage [2012/03/22(木) 23:21:31.67 ]: Keplerｷﾀ━━━━━━(ﾟ∀ﾟ)━━━━━━ !!!!!
575 名前：デフォルトの名無しさん mailto:sage [2012/03/22(木) 23:22:01.12 ]: gen3じゃないんだっけ？
576 名前：デフォルトの名無しさん mailto:sage [2012/03/22(木) 23:42:00.14 ]: >>570
もともとはミドルレンジでグラフィック向けだったから仕方ない気もする。
予想以上にグラフィック方面に舵を切ったという感はあるけど。

このままグラフィック向けとGPGPU向けで大きく分かれていくのではないかという心配はあるかな。
577 名前：デフォルトの名無しさん mailto:sage [2012/03/22(木) 23:53:32.45 ]: 1SM = 192コアか。おっそろしいなあ。
578 名前：デフォルトの名無しさん mailto:sage [2012/03/23(金) 00:00:34.24 ]: nVidia始まったな。
579 名前：デフォルトの名無しさん mailto:sage [2012/03/23(金) 00:17:06.29 ]: >>577
warp の扱いどうなるんかな。。。
580 名前：デフォルトの名無しさん mailto:sage [2012/03/23(金) 00:54:42.91 ]: >>579
ttp://pc.watch.impress.co.jp/docs/column/kaigai/20120322_520640.html
> 32スレッドのWARPに同じ命令を実行する、この基本は、Keplerでも変わっていない。
らしいから、変わらないんじゃないかな。

GF104/114のSMには48コアと2ワープスケジューラ、4ワープディスパッチャで
GK104のSMXには192コアと4ワープスケジューラ、8ワープディスパッチャになっている。

その上レジスタ数は倍、L1キャッシュ/シェアードメモリはそのままってことは
GF104/114よりさらにピーキーになっているのかな？
581 名前：デフォルトの名無しさん mailto:sage [2012/03/23(金) 01:01:32.84 ]: >>580
あれ？
コンスタントキャッシュって無くなった？
L1/L2キャッシュがその役割を担ってる？
ということはFermiからか･･･
582 名前：デフォルトの名無しさん [2012/03/23(金) 01:06:15.11 ]: 48コアが192コアになったのに
レジスタは2倍、
共有メモリは据え置き。

どーすんだこれ。。
583 名前：デフォルトの名無しさん mailto:sage [2012/03/23(金) 01:54:37.82 ]: レジスタ足りんくなりそうな。
584 名前：デフォルトの名無しさん mailto:sage [2012/03/23(金) 04:09:37.63 ]: Keplerはクロック落としてパイプラインを浅くする設計

演算器のレイテンシが小さくなるならレジスタの消費量は変わらない
Fermiの18cycleは頭おかしすぎた
これが例えば6cycleにになればレイテンシ隠蔽に必要なスレッド数が1/3になるから問題ない
585 名前：デフォルトの名無しさん mailto:sage [2012/03/23(金) 04:14:42.78 ]: x86 CPUと同じ道を辿ってるのか
586 名前：デフォルトの名無しさん [2012/03/23(金) 15:07:43.87 ]: 誰か２６次元計算してくれ、1000ｺｱくらいじゃマジに足らんぞｗ
587 名前：デフォルトの名無しさん mailto:sage [2012/03/23(金) 15:59:10.59 ]: float a[1][1][1][1][1][1][1][1][1][1][1][1][1][1][1][1][1][1][1][1][1][1][1][1][1][1] = NULL;
*a += 1;
588 名前：デフォルトの名無しさん [2012/03/23(金) 17:47:27.55 ]: こんにちは
国際暗号学会のプレプリントサーバにこんな論文があがってました

Usable assembly language for GPUs: a success story
Daniel J. Bernstein, et. al.
eprint.iacr.org/2012/137

GPUのことはさっぱりわかりませんが、なにかこのスレの足しにでもなれば幸いです
それでは
589 名前：デフォルトの名無しさん mailto:sage [2012/03/23(金) 20:49:13.93 ]: >584
グローバルメモリアクセスのレイテンシ隠匿とか、ループが遅いとかの情報が頭にあったんで
今まで深く考えず1024スレッド突っ込んでたんだけど、
スレッド減らしてループ回すような構造にしたほうがいい、って解釈でいいんだろうか？
590 名前：デフォルトの名無しさん mailto:sage [2012/03/25(日) 02:00:59.95 ]: >>588
PTXよりもっとネイティブ寄りのアセンブラ言語qhasm-cudasmを使って
パフォーマンスクリティカルな場面で力を発揮(nvccの148％)するよ！って話かな？
暗号学会で発表されるんだね。
591 名前：デフォルトの名無しさん mailto:sage [2012/03/25(日) 20:09:00.66 ]: メモリアクセスに対する演算の比率を上げないと、性能をフルに発揮できないことは分かったんですが、
具体的にどれくらいの比まで高めるべきかの目標はどうやって決めればイイでしょうか？？
592 名前：デフォルトの名無しさん mailto:sage [2012/03/26(月) 10:43:34.08 ]: 理論性能（カタログ値）がでるまで頑張れば良いのでは？

それからグローバルメモリーのアクセス速度が，カタログ値の何％になっているのかも
チェックすべきだと思う．
593 名前：591 mailto:sage [2012/03/26(月) 18:59:54.70 ]: >>592
ありがとうございます！
やはり、

理論性能が出ない → ボトルネックを割り出して改善 → 先頭に戻る

のループで追い込んでいくやり方ですね。
594 名前：デフォルトの名無しさん mailto:sage [2012/03/27(火) 00:14:52.50 ]: 本を読んで、Visual Profilerを知ったのですが、
ひょっとして今はParallel Nsightで同じことができるでしょうか？
595 名前：デフォルトの名無しさん mailto:sage [2012/03/28(水) 05:31:51.00 ]: GTX 680駄目すぎるわ
死んだ
596 名前：デフォルトの名無しさん mailto:sage [2012/03/28(水) 06:40:03.41 ]: まだ未公開？
ttp://developer.download.nvidia.com/compute/cuda/4_2/rc/toolkit/cudatoolkit_4.2.6_win_64.msi
ttp://www.abload.de/img/desktop_2012_03_27_22wif3f.png
PTX ISA3.0
597 名前：デフォルトの名無しさん mailto:sage [2012/03/28(水) 08:00:17.14 ]: warp shuffle
598 名前：デフォルトの名無しさん mailto:sage [2012/03/28(水) 10:36:38.09 ]: VLIWの腐ったようなアーキテクチャになったくさいな
599 名前：デフォルトの名無しさん mailto:sage [2012/03/28(水) 10:51:03.77 ]: どのへんが？
600 名前：デフォルトの名無しさん mailto:sage [2012/03/28(水) 10:54:22.36 ]: 斜め読みしたうえで完全にESPだが
命令は各スロットごとに別という点でVLIWでデータパスはSIMDみたいに独立とみた
ソフトウェア的にはもちろん別スレッドとして書けるみたいな
全然違ったらごめんね
601 名前：デフォルトの名無しさん mailto:sage [2012/03/28(水) 13:35:57.98 ]: いやPTXレベルの命令だから全然関係ないね
>>598-600は忘れてくれ
602 名前：デフォルトの名無しさん mailto:sage [2012/03/28(水) 21:20:01.15 ]: やっぱりGCNと同じでshuffle入れてきたな。
603 名前：デフォルトの名無しさん mailto:sage [2012/03/28(水) 22:40:28.29 ]: CUDA C Programming Guide Version 4.2
74p.
Table 5-1. Throughput of Native Arithmetic Instructions (Operations per Clock Cycle per Multiprocessor)
いろいろやばすぎるな。
604 名前：営利利用に関するLR審議中＠詳細は自治スレへ mailto:sage [2012/03/29(木) 11:33:59.98 ]: >>589
いまさらだが
物理レジスタが足りてるなら同時に多数スレッドを保持しておけるが
複雑なカーネルだとレジスタは不足しがち
605 名前：営利利用に関するLR審議中＠詳細は自治スレへ mailto:sage [2012/03/29(木) 12:27:07.17 ]: >604があるから時として64ビットアドレッシングより32ビットアドレッシングの方が有利なんだよね。
606 名前：営利利用に関するLR審議中＠詳細は自治スレへ mailto:sage [2012/03/29(木) 18:48:17.20 ]: Visual studio2010、CUDA4.1、Windows7 64bitではじめようと思ったんだけど、ネットで拾ったプログラムとか動かすとcutil_inline.hが見つからないって出る。
これってどうすればいいの？
607 名前：営利利用に関するLR審議中＠詳細は自治スレへ mailto:sage [2012/03/29(木) 19:47:51.42 ]: >>606
GPU Computing SDKをインストールしてパスを通す
608 名前：営利利用に関するLR審議中＠詳細は自治スレへ mailto:sage [2012/03/29(木) 22:36:38.83 ]: スレタイがイイね。
くだすれくーだすれｗ
609 名前：営利利用に関するLR審議中＠詳細は自治スレへ mailto:sage [2012/04/06(金) 21:14:51.29 ]: なんでGPGPUはみんな同じようなアプリしかつくらないん？
想像力が欠如してるから他人の猿真似ばかりしてんの？

GTX680の性能を生かすアプリ教えろ
610 名前：営利利用に関するLR審議中＠詳細は自治スレへ mailto:sage [2012/04/06(金) 21:37:23.28 ]: >>609
邪魔
611 名前：営利利用に関するLR審議中＠詳細は自治スレへ mailto:sage [2012/04/06(金) 22:02:22.54 ]: わりぃなぁ。そんじょそこらにはないアプリ作っているんだが公表できないんだわ。
612 名前：営利利用に関するLR審議中＠詳細は自治スレへ mailto:sage [2012/04/06(金) 22:56:59.64 ]: 恥ずかしいやつが湧いたな
613 名前：営利利用に関するLR審議中＠詳細は自治スレへ mailto:sage [2012/04/07(土) 01:09:23.67 ]: >>609
BOINCでもやってな
なんぼぶん回しても次から次へ宿題出してくれるから
614 名前：営利利用に関するLR審議中＠詳細は自治スレへ mailto:sage [2012/04/07(土) 10:44:11.21 ]: 数値流体力学シンポジウムにでも参加すれば良いぞ♪
615 名前：営利利用に関するLR審議中＠詳細は自治スレへ mailto:sage [2012/04/07(土) 19:20:32.56 ]: 公表できないと言ってるけど、どうせCUDA ZONEに登録されてるようなものだろ？
外人の真似と負け惜しみしかできないの？

素人だからよくわからないけど、欧米に対して技術面で遅れているから
この分野で日本に有名な人がいないでしょ？
自称一流の教授に俺の書き込み見せてあげて
616 名前：営利利用に関するLR審議中＠詳細は自治スレへ mailto:sage [2012/04/07(土) 23:56:42.42 ]: だから数値流体力学シンポジウムにでも参加すれば？

この分野では例えば東京工大の青木先生が有名だが
617 名前：営利利用に関するLR審議中＠詳細は自治スレへ mailto:sage [2012/04/07(土) 23:57:49.78 ]: ちなみに 615
618 名前：営利利用に関するLR審議中＠詳細は自治スレへ mailto:sage [2012/04/07(土) 23:58:19.73 ]: あれ？

ちなみに６１５は大学生か？？
619 名前：営利利用に関するLR審議中＠詳細は自治スレへ mailto:sage [2012/04/08(日) 00:08:49.21 ]: 大学生なら，物理や力学関係の学会に参加すれば，GPGPUを使った
シミュレーションの研究結果が報告されていることがわかるはず♪

高卒なら縁はないが．．．
620 名前：営利利用に関するLR審議中＠詳細は自治スレへ mailto:sage [2012/04/08(日) 00:11:11.59 ]: 大学生がこんな幼稚な文章書いてたら日本終っちまうぞｗ
621 名前：営利利用に関するLR審議中＠詳細は自治スレへ mailto:sage [2012/04/08(日) 09:03:46.81 ]: GTX480である程度大きなサイズをホストからデバイスに転送するのに3回に2回ぐらいセグメンテーションエラーで落ちる。
うまく行くとなんの問題もなく実行できる。
デバイス側でメモリ確保ができてないみたいなんだが、こんなもんなのかね？
622 名前：営利利用に関するLR審議中＠詳細は自治スレへ mailto:sage [2012/04/08(日) 09:17:09.30 ]: そのカードでモニターを表示させているとか？
623 名前：営利利用に関するLR審議中＠詳細は自治スレへ mailto:sage [2012/04/08(日) 12:09:14.45 ]: ケプラーの倍精度計算は速くなったの？
それとも以前と同じ？
624 名前：営利利用に関するLR審議中＠詳細は自治スレへ mailto:sage [2012/04/08(日) 12:27:33.01 ]: HPC向けがでてみないと分からんけど
GTX680じゃSPに対して1/24だよ
625 名前：営利利用に関するLR審議中＠詳細は自治スレへ mailto:sage [2012/04/08(日) 15:32:40.75 ]: www.tml.tkk.fi/~timo/HPG2009/
レイトレは最適化するとスペックなりのパフォーマンスだな680
626 名前：営利利用に関するLR審議中＠詳細は自治スレへ mailto:sage [2012/04/08(日) 20:11:08.50 ]: >>622
表示させている。
それがダメなのかな？
627 名前：営利利用に関するLR審議中＠詳細は自治スレへ mailto:sage [2012/04/08(日) 21:17:10.19 ]: ２枚さして，一枚はディスプレイ用
もう一枚は演算用にしないと，一枚では負荷に耐えられないのでは？
628 名前：営利利用に関するLR審議中＠詳細は自治スレへ mailto:sage [2012/04/10(火) 00:35:51.46 ]: みんなすごいね
せいぜい準備して
九九の掛け算一括処理くらいしかできないよ
629 名前：営利利用に関するLR審議中＠詳細は自治スレへ [2012/04/10(火) 00:59:53.96 ]: CUDAを学ぼうとするからそうなるんだと思うよ。
何か問題があって、それをCUDAで解こう！って始めた方が早く習得できる。
630 名前：営利利用に関するLR審議中＠詳細は自治スレへ mailto:sage [2012/04/10(火) 16:05:07.35 ]: >>624
それって今までよりも遅いってこと？
どこかに倍精度のベンチマークの比較はありませんか？
631 名前：営利利用に関するLR審議中＠詳細は自治スレへ mailto:sage [2012/04/10(火) 16:47:44.33 ]: 今研究室でGTX680か580のどっちかを買おうって話になってるんだけど
CUDA的にはどっちがいいと思う？
一任されて困ってる・・・
632 名前：営利利用に関するLR審議中＠詳細は自治スレへ mailto:sage [2012/04/10(火) 17:43:14.82 ]: floatかdoubleか。
あと整数演算も遅くなったらしい。
テスラ売るためとはいえ、なんかいやーんな感じ
633 名前：営利利用に関するLR審議中＠詳細は自治スレへ mailto:sage [2012/04/10(火) 18:28:46.60 ]: > 631

両方買う．

が，テスラの方が安定していると聞いているよ（速度は若干GTXより落ちるが）
634 名前：営利利用に関するLR審議中＠詳細は自治スレへ mailto:sage [2012/04/10(火) 18:32:30.11 ]: >>631
迷わず580かと。

680はレジスタとか整数演算・論理演算のスループットとか色々と問題になりそう。
それに最新の正式版Toolkit 4.1のプログラミングガイド見てもKepler載っていないし・・・

>>632
グラフィック性能のワットパフォーマンスを上げるためというのが一番じゃないかな。
635 名前：営利利用に関するLR審議中＠詳細は自治スレへ mailto:sage [2012/04/10(火) 19:03:34.66 ]: >>631
勉強用に1枚2枚買うって話ならGTX580だろうね。
１．資料がぜんぜん揃ってないKeplerを今買ってもしょうがない。
２．GPGPUとしての性能がGTX580のほうが「上」
（完全上位互換というわけではないし処理にもよるしスペック上FLOPSでは負けてはいるが）

GTX680のグラフィック・GPGPU性能を調べる ≪ dokumaru
dokumaru.wordpress.com/2012/03/27/gtx680-spec/

10数枚買って研究室全体に大量導入…なら先生が決めるよね。
両方買ってもいい。でもそれならKeplerは対応するToolkitが出てからでも遅くはないかと。
あるいは一刻一秒を争うならなおさらKeplerは冒険かと。
636 名前：営利利用に関するLR審議中＠詳細は自治スレへ mailto:sage [2012/04/10(火) 20:08:49.43 ]: 今はまだ早いGK110をまて
5月にイベントあるから、そこでなんかあるかも
637 名前：営利利用に関するLR審議中＠詳細は自治スレへ mailto:sage [2012/04/11(水) 16:27:30.42 ]: Teslaはメモリ容量が良いよね
GTXだと計算領域が足りないよ・・・
638 名前：デフォルトの名無しさん [2012/04/11(水) 23:21:32.32 ]: FLOPS/MBで見ると、Teslaでも全然足りない。
639 名前：デフォルトの名無しさん mailto:sage [2012/04/12(木) 07:21:34.36 ]: 何の計算？？
640 名前：デフォルトの名無しさん [2012/04/13(金) 04:19:22.20 ]: 倍精度計算が主なのですが、Ivyと680と580、どれがコストパフォーマンス的にお薦めですか？
641 名前：デフォルトの名無しさん mailto:sage [2012/04/13(金) 09:48:41.62 ]: IvyはCUDA動かないよ。
642 名前：デフォルトの名無しさん mailto:sage [2012/04/13(金) 10:12:31.27 ]: 言語の問題じゃなくて、プログラムはこれから作るから倍精度計算をわんさかやろうと思うんだけどどれがいいかなあ？
程度の話じゃないかと。
643 名前：デフォルトの名無しさん mailto:sage [2012/04/13(金) 10:17:18.06 ]: そういうことか。
超並列に対応できるのであればGPUのほうがイイね。
644 名前：デフォルトの名無しさん mailto:sage [2012/04/13(金) 10:27:10.68 ]: えっ！？
ここCUDAのスレだよね？

てか、Ocelotとかどうなのかな。
645 名前：デフォルトの名無しさん mailto:sage [2012/04/13(金) 12:45:42.72 ]: 倍精度計算が中心なら、CPUで最適化するのが一番。
例えば近似計算のようにGPUの単精度で近づけてから、
CPUの倍精度で収束させるとかならありだけど。
646 名前：デフォルトの名無しさん mailto:sage [2012/04/13(金) 16:13:12.61 ]: cuda zoneがメンテナス中・・・
647 名前：デフォルトの名無しさん mailto:sage [2012/04/13(金) 17:11:45.77 ]: toolkit4.2が来るのかな
648 名前：デフォルトの名無しさん mailto:sage [2012/04/13(金) 18:28:30.79 ]: >>640
GPU用の倍精度プログラムを書く気があるならTeslaにしとけ。
コストが厳しいならRadeonの最上位にしとけ。
649 名前：デフォルトの名無しさん mailto:sage [2012/04/14(土) 01:50:49.98 ]: sdk 4.1とtoolkit 4.1インストールしたんだけど
アンインストールせずにそのまま
sdk 2.3とtoolkit２．３をインストールしたらコンパイルやリンクの挙動とかおかしくなりますか？
650 名前：デフォルトの名無しさん mailto:sage [2012/04/14(土) 12:46:31.14 ]: 自分でpathやMakefileなどを管理できるのなら無問題。
651 名前：デフォルトの名無しさん mailto:sage [2012/04/14(土) 22:18:52.58 ]: parallel nsightは甘え
652 名前：デフォルトの名無しさん mailto:sage [2012/04/14(土) 22:34:41.85 ]: >>651
甘えと言えるほどすごいのか。
今度使ってみようｗ
653 名前：デフォルトの名無しさん mailto:sage [2012/04/16(月) 19:24:07.51 ]: シングルGPUでもデバッグできるようになったのが凄くうれしい
654 名前：デフォルトの名無しさん [2012/04/17(火) 03:52:59.70 ]: 680を手に入れたんだけど、ガッカリ性能だった
ゲーム系は速くなってるんだけどね
655 名前：デフォルトの名無しさん mailto:sage [2012/04/17(火) 08:51:17.06 ]: 何をやったの
656 名前：デフォルトの名無しさん mailto:sage [2012/04/17(火) 15:39:00.54 ]: PTX直接書いてプログラミングする人とかいるの？
657 名前：デフォルトの名無しさん mailto:sage [2012/04/17(火) 19:51:20.08 ]: 一応いるけど。
658 名前：デフォルトの名無しさん mailto:sage [2012/04/17(火) 21:23:06.59 ]: GTX580が生産終了なんだとか
659 名前：デフォルトの名無しさん mailto:sage [2012/04/17(火) 23:38:03.55 ]: >658
在庫はけたからか。
しかしIntelのCPUはGPU載せてんのにGPGPUにはさっぱり対応しないからあんま意味ないな。
CUDA対応とかにしないのは戦略的判断なんだろうけど、なんとももったいない。
660 名前：デフォルトの名無しさん mailto:sage [2012/04/17(火) 23:51:03.54 ]: 大丈夫
Intelも来週からGPGPU対応する
661 名前：デフォルトの名無しさん [2012/04/18(水) 16:31:10.22 ]: 倍精度計算じゃまだインテルに分があるしね
662 名前：デフォルトの名無しさん mailto:sage [2012/04/18(水) 16:57:28.76 ]: ivyのeuどうなってのかね
663 名前：デフォルトの名無しさん mailto:sage [2012/04/20(金) 18:04:01.73 ]: > 654

今，それに触手を伸ばしているところだけど，どこがダメだった？？
664 名前：デフォルトの名無しさん mailto:sage [2012/04/21(土) 20:08:12.47 ]: >>660
ivyってGPGPUとして使えんの？
665 名前：デフォルトの名無しさん mailto:sage [2012/04/21(土) 20:18:04.15 ]: OpenCLならできなくもないのかな？
666 名前：デフォルトの名無しさん mailto:sage [2012/04/24(火) 13:30:34.56 ]: 誰かHMPP使ったことある人いる？
667 名前：デフォルトの名無しさん mailto:sage [2012/04/24(火) 20:58:18.61 ]: CUDA4.2きたな
ttp://developer.nvidia.com/cuda-downloads
668 名前：デフォルトの名無しさん mailto:sage [2012/04/25(水) 20:18:30.00 ]: Adobe Creative Suite 6: Bye bye CUDA, Hello OpenCL!
ttp://www.geeks3d.com/20120425/adobe-creative-suite-6-opencl-accelerated-mercury-graphics-engine-opengl/
669 名前：デフォルトの名無しさん mailto:sage [2012/04/25(水) 20:24:16.57 ]: ivyの影響だな
670 名前：デフォルトの名無しさん mailto:sage [2012/04/26(木) 23:10:43.77 ]: CPU、GPUを利用（プログラム）するには？
togetter.com/li/293863
671 名前：デフォルトの名無しさん mailto:sage [2012/04/29(日) 14:44:53.67 ]: GPUのデメリットは同じ変数計算を毎回糞真面目に超高速で行うところ
672 名前：デフォルトの名無しさん mailto:sage [2012/04/29(日) 18:55:51.68 ]: メモリ読むより速いからな
673 名前：デフォルトの名無しさん mailto:sage [2012/04/30(月) 07:47:23.73 ]: GTX690
pc.watch.impress.co.jp/docs/news/20120429_530569.html
674 名前：デフォルトの名無しさん mailto:sage [2012/04/30(月) 08:35:43.64 ]: >> 671
一つ一つの計算は超高速でもなんでもない
並列で行うので早くなるだけ
超高速になるか否かはプログラミングの問題

>>672
演算にはメモリーの読み書きを伴うので，演算が「メモリ読むより速い」とはならないのでは？
675 名前：デフォルトの名無しさん mailto:sage [2012/04/30(月) 21:06:45.57 ]: >>671
意味が分からん。
アーキの概念の理解ができていないじゃねーか？
676 名前：デフォルトの名無しさん mailto:sage [2012/05/01(火) 00:39:15.46 ]: >>675
明日短いわかりやすいソースアップするからコンパイルして実行してみて
言いたいことがわかると思う。
CPUにはあってGPUにはない機能を使うことになる、まぁホントしょうもないことだけど・・・
677 名前：デフォルトの名無しさん mailto:sage [2012/05/01(火) 00:53:32.34 ]: ？？？
678 名前：デフォルトの名無しさん mailto:sage [2012/05/01(火) 00:59:29.65 ]: 言ったら悪いかも知れんけど単にアルゴリズムが悪いんじゃないのか。
679 名前：デフォルトの名無しさん mailto:sage [2012/05/01(火) 08:09:04.59 ]: >>671
アドレス計算とかまさにそれだよね。
普通のループなら+4で済むところが、
ptr + threadIdx.x*4 + threadIdx.y*hoge
とかになっちゃう。
680 名前：デフォルトの名無しさん mailto:sage [2012/05/01(火) 08:24:43.67 ]: それはGPUのデメリットじゃないな。
GPU(nvcc)でもループなら普通に書いたら普通に最適化してくれる。
681 名前：デフォルトの名無しさん mailto:sage [2012/05/01(火) 10:38:58.63 ]: >>679
これはデメリットと違う。
CPUでマルチスレッドでやれば同じように明示的にアドレス計算を行う必要がある。
682 名前：デフォルトの名無しさん mailto:sage [2012/05/01(火) 13:31:24.76 ]: シングルスレッドでの最適化が、そのままマルチスレッドに使えると思ってるなら、並列で組むのに向いてないな。

ひとつの処理として見たとき無駄でも、それで大多数の演算を同時に走らせることができるなら、
並列処理においてはそれこそが効率的なんだよ。
683 名前：デフォルトの名無しさん mailto:sage [2012/05/02(水) 01:02:11.09 ]: >>669
残念、AMDとアドビのコラボでした
www.4gamer.net/games/133/G013372/20120426013/

Ivyは端から相手にされてません
684 名前：デフォルトの名無しさん [2012/05/02(水) 01:48:54.23 ]: ascii.jp/elem/000/000/672/672388/
このカード使ってる人居ませんか？

メモリがいっぱい欲しいけど、高いカードは買えないので
試しに買ってみようかと思うのですが。
685 名前：デフォルトの名無しさん mailto:sage [2012/05/02(水) 02:05:34.87 ]: 同じ世代のGPUでも生産地の違いで演算速度は全く違うからね
もっと言うと転送速度が全く違う
まあフラッシュメモリでも同じこと言えるけど
686 名前：デフォルトの名無しさん [2012/05/02(水) 04:31:00.56 ]: >683
ずいぶんニッチなところだな
687 名前：デフォルトの名無しさん mailto:sage [2012/05/02(水) 09:30:37.50 ]: >>684
マジレスすると、CUDAでやるメリットはない。
Sandyやivyの方がはるかに高速。
まあ、CUDA勉強するだけならいいが、もっと別のカードのほうがいいだろ。
688 名前：デフォルトの名無しさん mailto:sage [2012/05/02(水) 11:01:38.16 ]: >>684
メモリ転送が遅過ぎて4GBのメモリを活かしきれない悪寒。
689 名前：デフォルトの名無しさん mailto:sage [2012/05/02(水) 12:41:42.31 ]: SRAMを４GBつんでるカードはないのか？
690 名前：デフォルトの名無しさん mailto:sage [2012/05/02(水) 16:10:34.75 ]: >>667

>New Features
>Support for GK10x Kepler GPUs.

とりあえず、GK104対応にしました的か。
691 名前：デフォルトの名無しさん [2012/05/02(水) 18:28:26.54 ]: VRAM 4GB以上のカードって、ほとんどないんだね。
TeslaかQuadroしか見つからなかった。
お値段10万円越
692 名前：デフォルトの名無しさん mailto:sage [2012/05/02(水) 18:42:59.54 ]: 4年後にはVRAM16GBが普通にでまわるんだよ
693 名前：679 mailto:sage [2012/05/02(水) 22:18:41.81 ]: >>680 >>681
あれ、そういう話じゃないのか…
そうだとすると、 >>671 が何を言いたかったのか思い付かないな…
694 名前：デフォルトの名無しさん mailto:sage [2012/05/03(木) 15:04:28.21 ]: >671は皮肉だろ。
695 名前：デフォルトの名無しさん mailto:sage [2012/05/03(木) 16:07:25.13 ]: ということにしたいのですね。
696 名前：デフォルトの名無しさん mailto:sage [2012/05/04(金) 01:03:14.33 ]: デメリットに感じる境地まで辿りついたんだよ、きっと
俺にはまだメリットにしか思えないんだけど・・・
697 名前：デフォルトの名無しさん mailto:sage [2012/05/05(土) 21:18:13.51 ]: >>692
今から4年前の2008年頃はG80世代で大体1GBだったが、4年後のGTX680でまだ2GBだから、
4年後はせいぜい4GBなんじゃないの？Tesla系で16GBにはなっていそうだけど。
698 名前：デフォルトの名無しさん mailto:sage [2012/05/05(土) 21:24:38.22 ]: そういやPCのDRAM搭載量に比べて、あんまり伸びないよね＞ビデオカードのメモリ
699 名前：デフォルトの名無しさん mailto:sage [2012/05/05(土) 21:26:52.65 ]: GDDRは数が出ないからね。
DRAMメーカーがあんな状態だから尚更でしょう。
700 名前：デフォルトの名無しさん mailto:sage [2012/05/05(土) 21:33:02.38 ]: プロセスシュリンクが汎用DRAMと同じように進めば同じようにでかくなると思うんだけど。
だんだん引き離されてるってこと？
701 名前：デフォルトの名無しさん mailto:sage [2012/05/10(木) 01:35:01.75 ]: 日本のメモリの会社が潰れたのはかなり痛いな・・・
702 名前：デフォルトの名無しさん mailto:sage [2012/05/10(木) 15:36:00.11 ]: ptxコード読まなきゃいけなくなったんだけど、typeの.predって何なのかいまいちわかってない
703 名前：デフォルトの名無しさん mailto:sage [2012/05/10(木) 16:46:59.88 ]: 述部(predicate)だね。
ptxの場合は単に、比較などの結果を保持するだけのような希ガス。
で、そのレジスタの結果に依存してインストラクションの実行する、と。
例えば、
--
setp.gt.s32 %p1, %r5, %r7;
@%p1 bra $Lt_0_12802;
--
なら r5 > r7のときに分岐するし、
--
setp.lt.s32 %p2, %r9, %r11;
@%p2 sub.s32 %r14, %r11, %r14;
--
なら r9 < r11のときに引き算を行なう。
704 名前：デフォルトの名無しさん mailto:sage [2012/05/10(木) 21:33:54.24 ]: 分岐マスクのためのレジスタは何本あるんだろ
それとも汎用レジスタと共用なのか
705 名前：702 mailto:sage [2012/05/11(金) 16:21:50.31 ]: >>703
thx
そういう意味だったのか……
CUDAはC言語の延長だから大丈夫とか考えた三月の俺を叩きのめしたい

PTXコードの読み方って英語のやつしかないよねたぶん
706 名前：デフォルトの名無しさん mailto:sage [2012/05/11(金) 17:11:50.74 ]: >>705
私が書いたメモならあるよw

>>704
実験コードで見たところ、汎用レジスタと述語レジスタの合計で制限されてたかと。
述語レジスタだけでどこまで増やせるかは実験してない。
707 名前：702 mailto:sage [2012/05/11(金) 17:26:43.70 ]: >>706
恵んでください。
割と切実に。卒業したいので。
708 名前：デフォルトの名無しさん mailto:sage [2012/05/11(金) 18:02:12.84 ]: ISA的にはwarpあたり7本か6本じゃね。
3bitのどれかが常にalways扱いだったような。
709 名前：デフォルトの名無しさん [2012/05/14(月) 15:19:34.27 ]: 初心者質問です。
お願いします。
cufftってcuda3.2でも使えるのでしょうか？
cufftdestroyが未解決の外部シンボルだと言われてしまうのですが？
ただ単に、リンクできてないだけなのでしょうか？
710 名前：デフォルトの名無しさん mailto:sage [2012/05/14(月) 18:14:14.93 ]: >>709
使えたと思うよ。
つーか、cufftdestroy()が未解決って、あんたの間違いだろ。
711 名前：デフォルトの名無しさん [2012/05/15(火) 14:03:46.41 ]: >>710
返信ありがとうございます
他の関数はコンパイルが通る(通っているように見えるだけ？)のに
cufftdestroy()
cufftExecZ2Z()
cufftPlan1d()
だけが未解決となっているのですが、
この関数だけ、他のライブラリが必要だなんてことがあるのでしょうか？
712 名前：デフォルトの名無しさん mailto:sage [2012/05/15(火) 15:52:04.02 ]: destroyはDestroy。
z2zは未実装。
Plan1dはしらね。
警告レベル引き上げれば?
713 名前：デフォルトの名無しさん mailto:sage [2012/05/17(木) 11:36:49.94 ]: cufft.hはインクルードしているのかな？
714 名前：デフォルトの名無しさん mailto:sage [2012/05/17(木) 13:08:31.44 ]: ■後藤弘茂のWeekly海外ニュース■
NVIDIAが世界最多トランジスタ数のチップ「GK110」を公開

pc.watch.impress.co.jp/docs/column/kaigai/20120517_533500.html
715 名前：デフォルトの名無しさん [2012/05/17(木) 15:21:47.32 ]: 警告レベルって、デフォルトは最大なんですよね？
Destroyに関しては、タイプミスです。
z2zは未実装っていうのが、よくわからないんですけど。。。。
716 名前：デフォルトの名無しさん mailto:sage [2012/05/17(木) 17:10:10.69 ]: GPGPU上でソケット通信とかって出来るかな
717 名前：デフォルトの名無しさん mailto:sage [2012/05/17(木) 17:20:54.19 ]: GPGPUの仮想マシン同士のn対n通信をシミュレートとかそういうのをイメージした
718 名前：デフォルトの名無しさん mailto:sage [2012/05/17(木) 18:16:22.69 ]: >>715
未実装: 実装されていないこと。
cufftのライブラリの中にz2zの関数そのものが存在していないのよ。
で、あんたがどんな環境で開発しているか判らんのに警告レベルがどうなっているかなんか判るかい。
そんなことは自分で調べなさいよ。

>>716
cuda5でLAN接続されているGPU同士で連携させる機能がつくらしいよ。
719 名前：デフォルトの名無しさん mailto:sage [2012/05/17(木) 21:52:07.76 ]: 多次元配列を扱えないのは何でなんだろう．

ブロックとスレッドインデックスで一次元化するの面倒なんだけど．
720 名前：デフォルトの名無しさん mailto:sage [2012/05/17(木) 23:46:21.89 ]: ピンメモリを確保すると、スワップによる退避を防げるのは分かったのですが、
実際はスワップ以外にも、メモリフラグメンテーション解消のためのコンパクションでも
メモリアドレスの変化って起こり得ますよね？
それもないようにするのがピンメモリですよね？
721 名前：デフォルトの名無しさん mailto:sage [2012/05/18(金) 01:25:35.91 ]: >>719
別に扱えなくはないぞ。普通にdata[blockIdx.x][thiredIdx.x]ってできると思う。
スレッド数を定数にしなくちゃならなくなるから却って煩わしいと思うけど。

つーか、面倒ったってオフセット計算する関数を作るだけじゃん。

[ 新着レスの取得/表示 (agate) ] / [ 携帯版 ]

read.cgi ver5.27 [feat.BBS2 +1.6] / e.0.2 (02/09/03) / eucaly.net products.
担当:undef