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