12/01/05 10:51:17.15
>>260
＞母体になるプログラムなしで、
それってなんのこと? 外出先のPCで動かしたいってこと?
コンパイラのことなら、事前にコンパイルしておくかリモートでアクセスすればいい。
ドライバのことなら、外出先の端末にインストールさせて貰うしかない。
外出先のPCなんか頼らず、自前のPCを持っていけばいいじゃん。
263:デフォルトの名無しさん
12/01/05 13:36:51.62
>>257
手元に開発環境がなく、エラーメッセージをコピー忘れしてました。すみません。
# cc Drawtest.c -lglut
# ./a.out
（空のウィンドウが製作される）
（Drawtest.cをDrawtest.cuに中身は同じで拡張子だけ変えて）
# nvcc Drawtest.cu -lglut
/usr/bin/ld: skipping incompatible /usr/local/cuda/bin/../lib64/libglut.a when searching for -lglut
（-lglutを見付かられなかった？）
# nvcc -L /usr/local/cuda/lib64/ -l /usr/local/cuda/include/ Drawtest.cu -lglut
/usr/bin/ld: cannot find -l/usr/local/cuda/include/
collect2: ld はステータス 1 で終了しました
となりました。
264:デフォルトの名無しさん
12/01/05 14:25:08.28
インクルードパスの指定は -l じゃなくて -I な
265:デフォルトの名無しさん
12/01/05 15:05:37.06
/usr/bin/ld: skipping incompatible～って64/32ビット混在時のエラーメッセージじゃなかったかな
266:デフォルトの名無しさん
12/01/05 15:09:35.86
>>264
なんというミス。
目で見て違いが分からない。
L（小文字）→i（大文字）に変えて同じようにしてみました。
# nvcc -L /usr/local/cuda/lib64/ -I /usr/local/cuda/include/ Drawtest.cu -lglut
/usr/bin/ld: skipping incompatible /usr/local/cuda/lib64//libglut.a when searching for -lglut
/usr/bin/ld: skipping incompatible /usr/local/cuda/bin/../lib64/libglut.a when searching for -lglut
結局パス指定なしと同じ見たいですね。
267:デフォルトの名無しさん
12/01/05 15:37:13.27
おれも64/32の違いかと思ったが、パス指定から指摘してみた
実行している環境は？ -m64 オプション付けてみ
268:デフォルトの名無しさん
12/01/06 13:53:42.06
# nvcc -L /usr/local/cuda/lib64/ -I /usr/local/cuda/include/ Drawtest.cu -lglut -m64
/usr/bin/ld: skipping incompatible /usr/local/cuda/lib64//libglut.a when searching for -lglut
/usr/bin/ld: skipping incompatible /usr/local/cuda/bin/../lib64/libglut.a when searching for -lglut
ううむ・・・・。
centOS 64bit
Device 0: "Quadro 4000"
Device 1: "GeForce 9500 GT"
パスは
LD_LIBRARY_PATH = /usr/local/cuda/lib64
です。
openGLのないコードはコンパイル及び実行まで問題ありません。
openGLを使うにあたって
/usr/local/cuda/include/にglut.h freeglut.h freeglut_ext.h freeglut_std.h
/usr/local/cuda/lib64/にlibglut.a
をコピペしましたが、まだ足りないのでしょうか？
269:デフォルトの名無しさん
12/01/06 13:59:40.26
libglutは64ビット？
32ビットなんじゃねか？
270:デフォルトの名無しさん
12/01/06 15:03:44.65
glut.h freeglut.h freeglut_ext.h freeglut_std.h libglut.aをccが見ているであろう先のものと置き換えてしてみました。
# nvcc -L /usr/local/cuda/include/ -I /usr/local/cuda/include/ Drawtest.cu -lglut -m64
/tmp/tmpxft_000030a4_00000000-12_Drawtest.o: In function `disp()':
tmpxft_000030a4_00000000-1_Drawtest.cudafe1.cpp:(.text+0x14e6): undefined reference to `glClearColor'
tmpxft_000030a4_00000000-1_Drawtest.cudafe1.cpp:(.text+0x14f0): undefined reference to `glClear'
/usr/local/cuda/bin/../lib64/libglut.a(libglut_la-freeglut_init.o): In function `glutInit':
…
(中略)
…
(.text+0x898): undefined reference to `glPopClientAttrib'
collect2: ld はステータス 1 で終了しました
openGL系の関数が読まれていないみたいです。
271:デフォルトの名無しさん
12/01/06 15:37:17.18
いや、だからfreeglutは64ビットを使用してるのかと聞いてるんだが
質問に対しきちんと回答できないようでは教えようがない罠
272:デフォルトの名無しさん
12/01/06 18:39:39.36
>>271
freeglutは64ビットのものを入れ直してやってみましたが、同じ結果になりました。
>>257
そこでCUDAのヘッダフォルダとライブラリフォルダに入れていたものを消し、-Lと-Iで直接freeglutのフォルダを指定したら
解決しました。
ライブラリが足りなかったのか、それともCUDAと一緒くたにしたのがダメだったのか解かりませんが、ちゃんとPATH通せってことみたいです。
お騒がせしました。
273:デフォルトの名無しさん
12/01/06 19:45:07.40
>>262
だよねぇ・・・
274:デフォルトの名無しさん
12/01/06 20:14:34.67
環境
Windows7 Professional 64bit
Microsoft Visual C++ 2010 Express Version 10.0.40219.1
Microsoft .NET Framwork Version 4.0.30319
GeForce GTX 580
CUDA Toolkit 4.0.17
SDK 4.0.19
devdriver_4.0_winvista-win7_64_270.81_general
このサイトを参考に環境を構築しました。
URLﾘﾝｸ(feather.cocolog-nifty.com)
そして以下のサイトのサンプルプログラムを実行してみました。
URLﾘﾝｸ(www.gdep.jp)
Hello,Worldと99 bottles of beerはcpu、gpu共に実行できました。
しかし、Matrixのプログラムはcpuの方は実行できるのですがgpuの方が実行できません。
以下のエラーを吐きます。
matrix_gpu.cu(5) : fatal error C1083: Cannot open include file: 'cutil_inline.h'
: No such file or directory
どうやらVisualStudioのパスがうまく通ってないということまでわかり、以下のサイトなどを参考にCUDA_INC_PATHなどを変えてみましたが、一向に変わりません。
URLﾘﾝｸ(d.hatena.ne.jp)
SDK内のcutil_inline.h自体をtoolkitのincフォルダにコピペすると、他の.hファイルもいくつか同じエラーが出たのでエラーになったものをすべてコピペしたところ、LNK2019"link.exe"というエラーで先に進めませんでした。
一度VisualStudioを再インストールしてみましたが、状況は変わりません。
Nvidia GPU computing SDK Browserではサンプルプログラムを実行できているので、CUDAの環境は整っていると思われます。
どうすれば解決できますでしょうか・・・。かれこれ１週間以上格闘しています。
VisualStudioは2010よりも2008にした方がいいでしょうか？
275:デフォルトの名無しさん
12/01/06 20:23:04.24
C初心者にはきついと思うんだが…
とりあえず 'cutil_inline.h'のある場所を見つけて
そこを-I /'cutil_inline.hのある場所'と指定する
意味わからなければCのコンパイルを勉強すること
276:デフォルトの名無しさん
12/01/06 20:39:53.25
274ではありませんが､「Cのコンパイル」の勉強にオススメの書籍とかありましたら紹介していただけるとうれしいです
277:デフォルトの名無しさん
12/01/06 20:55:31.81
275だけど、お薦めは他の人に任せるけど
一つ言えるのはCUDAはC/C++より数段難しい
C/C++の質問をしてるようでは先が思いやられる、と正直思うんだが…
超低レベルな質問もOKなくだすれだけど、その程度は自力で調べる能力が必要ではないだろうか？
278:デフォルトの名無しさん
12/01/06 21:22:02.03
>>275
>>274です。ありがとうございます。
今実行できる環境にないので後で試してみます。
また結果報告しにきます。
279:デフォルトの名無しさん
12/01/07 12:14:43.10
一時期は同価格帯のCPUとGPUを比較してGPUは５０倍の性能があるって言われてたけど
今１０～１５倍程度しか無くない？
3930k
５万円
158GFlops
GTX590
６．３万円
2480GFlops
GTX 580
３．８万円
1500GFlops
次に出ると言われてるGTX700系は500系の２倍の性能らしいけどそれでも５０倍には遠く及ばない。
この「同価格帯でGPUはCPUの１０～１５倍程度の性能しかない」という評価は正しいだろうか？
280:デフォルトの名無しさん
12/01/07 12:43:15.62
>>279
それは単にCPUは1コアとして比較しているからGPUを使えばうん十倍と言われるんだよ。
単純にメモリ帯域で3～4倍、演算器としては、単純にAVXの8並列x6コアxクロック2倍で比較すると5倍にしかならない。
CPUは大きなキャッシュがあるから局所性の高い処理の場合はメモリ帯域の差が縮まるので、10倍差がつくならいい方。
もちろんアプリによるけど、CPU向けのコードがカリカリに最適化されていた場合は思ったほど差がつかない印象。
それでも3～5倍位は早くなるからいいんだけどね。
問題はメンテナンスだな。。。CUDAのみと割り切れればいいんだけど、なかなかそうは行かないからなあ。
281:デフォルトの名無しさん
12/01/07 16:09:01.15
>>279
GTX580で500GFLOPS程度。そのデータはおかしい
CPUなんてE8500で25GFLOPSくらいだから20倍程度だな。
もちろん最近のコア多いやつならもっと差は縮まるだろうな。
それに理論値の話だから多数のスレッドでまわしたらコンテキストスイッチの
オーバーヘッドとかGPUプログラミングは難しくて効率上げるの難しいとか
そういうので実効は変わってくるんじゃね？
282:デフォルトの名無しさん
12/01/07 16:14:51.24
あ、同一価格帯か、E8500が2万しなかったことを考えると
580にそろえて2倍したとして10倍程度だから、まあおおざっぱには
「同価格帯でGPUはCPUの１０～１５倍程度の性能しかない」は正しいんじゃないかと思う。
580と同時期の4万弱のCPUならもっと差は縮まるだろうし。
283:デフォルトの名無しさん
12/01/07 16:22:20.40
32nmと40nmで比較とか
284:デフォルトの名無しさん
12/01/07 21:34:45.46
void *p;
CUDA_SAFE_CALL(cudaSetDevice(0));
CUDA_SAFE_CALL(cudaMalloc(&p, 16));
でcudaMallocからcudaError_enum例外が飛ぶのですが
原因はなにが考えられますか？
ちなみにSDKのサンプルは動きます。
285:284
12/01/07 22:23:52.14
異なるディレクトリに複数のCUDAのライブラリが入ってて
古いバージョンが参照されてるのが原因だった・・・
286:デフォルトの名無しさん
12/01/08 14:34:35.46
CUDA4.1でcomputeProfはnvvcになったんですか？
だいぶ代わってしまって驚きますた
287:デフォルトの名無しさん
12/01/08 14:35:50.30
nvvp
288:デフォルトの名無しさん
12/01/08 20:50:11.30
>>282
GPUの論理性能ってFMAとかMADとかの積和で倍になっているから、
GTX 580 1500GFlopsとかって、それらを使わなければ750GFlopsとかになるんだよな。
だから実効れべるではもっと縮まるんだろうな。それでも5倍以上差がでるとかなり有効なんだけど。
289:デフォルトの名無しさん
12/01/09 10:37:28.98
個人で買うやつも居るのかね？
URLﾘﾝｸ(page16.auctions.yahoo.co.jp)
290:デフォルトの名無しさん
12/01/09 12:04:44.39
いる。というか俺が買う。
291:デフォルトの名無しさん
12/01/09 13:09:16.68
こういう需要が限られていてかつ堅くて高い製品はamazonに出した方が高く売れるんじゃないかねえ。
292:デフォルトの名無しさん
12/01/09 18:01:50.03
>>275
274です。
-Iオプションを付けてヘッダファイルのある場所を指定してみましたが、相変わらず同じエラーです。
C:\cuda_practice\test>nvcc -o matrix_gpu.exe matrix_gpu.cu -I C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.0\C\common\inc
matrix_gpu.cu
matrix_gpu.cu(5) : fatal error C1083: Cannot open include file: 'cutil_inline.h': No such file or directory
自分でも試行錯誤してみますが周りにcudaがわかる人がいなくて行き詰ってます・・・。
293:デフォルトの名無しさん
12/01/09 18:20:34.55
もっとコンパイル早くなんねーかなぁ…
thrust::sortとか使うようにしただけでコンパイルに十数秒かかってイライラする
294:デフォルトの名無しさん
12/01/09 20:40:26.69
速いマシンでやる。
295:デフォルトの名無しさん
12/01/09 20:45:40.42
nvccコンパイラをGPGPUで実装すればいい！
CUDAがオープンになるどさくさにまぎれてそういうのを作る人が出るかもしれないこともないかも
296:デフォルトの名無しさん
12/01/09 21:57:43.85
コンパイラの動作は現状GPU向きとは言えないんじゃね？
GPU自体が単純な性能向上だけではない、思いもよぬ進化でもしない限り
297:デフォルトの名無しさん
12/01/09 22:54:12.52
コンパイルなんて依存関係ありまくりで、並列処理が苦手な最たるものだからまずやる意味がないだろう。
CUDAにすれば何でも速くなると思ってんじゃね？
298:デフォルトの名無しさん
12/01/10 17:15:02.61
sharedメモリって同一WARP内でのやり取りだったら同期なしで大丈夫なのでしょうか？
たとえば全部で32スレッドで次のkernelを実行した場合、
WARP内で一斉にsharedメモリに書きに行くので同期しないで大丈夫かと
思ったのですが、実際にはうまくいきません。
globalメモリに読みに行く段階でコアレッシングが発生していないので
それが原因なのでしょうか？
どなたか教えてください。
__global__ void kernel(float *g_v, float *g_x){
float x = 0.0f;
int i = blockDim.x * blockIdx.x + threadIdx.x;
__shared__ float s_v[32];
s_v[i] = g_v[i+i%3];
__syncthreads(); // これが必要かどうか？
x = s_v[(i+3)%32];
g_x[i] = x;
}
299:デフォルトの名無しさん
12/01/10 17:35:07.81
>>292
参照パスを ""で囲む♪
300:デフォルトの名無しさん
12/01/10 17:39:16.65
>>298
iが３２以上になることはないかい？？
ブロックの数は？？
301:274
12/01/10 17:52:39.67
>>299
書き込みしてからそのミスに気付きました・・・。
それを直してもうまくPathが通らなかったので、CUDA4.0、VisualStudio2008でやることにしました。
302:デフォルトの名無しさん
12/01/10 17:56:37.24
>>298
fence入れてないからじゃね。
303:298
12/01/10 17:59:16.44
>>300
ごめんなさい、書いたコードが一部誤ってました。
iが32以上になることがあります。
ブロック数は数100程度になります。
このときは、下のソースのようになると思うのですが、
やはりうまくいきません。
__global__ void kernel(float *g_v, float *g_x){
float x = 0.0f;
int i = blockDim.x * blockIdx.x + threadIdx.x;
__shared__ float s_v[blockDim.x];
s_v[threadIdx.x] = g_v[i+i%3];
__syncthreads(); // これが必要かどうか？
x = s_v[(threadIdx.x+3)%32];
g_x[i] = x;
}
304:デフォルトの名無しさん
12/01/10 18:12:32.43
>>303
コンパイル後のバイナリを見たわけじゃないから予測だが、
s_v[threadIdx.x] = g_v[i+i%3];
x = s_v[(threadIdx.x+3)%32];
g_x[i] = x;
は
1.1:tmp=g_v[i+i%3]; //グローバルアクセスなので遅い
1.2:s_v[threadIdx.x]=tmp; //
2.1:x = s_v[(threadIdx.x+3)%32];
3.1:g_x[i] = x;
みたいに解釈される。
で、単純な式だからコンパイラは
s_v[threadIdx.x]とs_v[(threadIdx.x+3)%32];
が必ず別のアドレスになる事を検知して
1.1 2.1 3.1 1.2のように命令を入れ替えてしまう。
だから__threadfence_block()がいるはず。
305:298
12/01/10 18:17:13.87
>>302 >>304
ありがとうございました。
試してみたいと思います。
306:デフォルトの名無しさん
12/01/10 19:45:54.97
>> 305
だいたい x = s_v[(threadIdx.x+3)%32]; の前にすべてのs_v[??}が定義されていないといけないだろ？？
スレッドが３２に設定されているのならば (threadIdx.x+3)%32 は０か１にしかならないが．．．
307:デフォルトの名無しさん
12/01/10 21:10:42.27
>>301
-lcutil32 とその参照パスを追加すれば Visual Studio 2010 でもコンパイル，実行できる♪
ちなみに私の場合は
nvcc matrix.cu -I "C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.0\CUDALibraries\common\inc" -L "C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.0\CUDALibraries\common\lib\Win32" -lcutil32
tmpxft_000024b8_00000000-3_matrix.cudafe1.gpu
tmpxft_000024b8_00000000-8_matrix.cudafe2.gpu
matrix.cu
tmpxft_000024b8_00000000-3_matrix.cudafe1.cpp
tmpxft_000024b8_00000000-14_matrix.ii
C:\cuda> a
Processing time: 1936.354858 (msec)
Press ENTER to exit...
308:デフォルトの名無しさん
12/01/10 21:12:58.21
ちなみに家のはGeForce GT 220なので非力．．．．
309:306
12/01/10 22:21:29.79
間違えた
↓
(threadIdx.x+3)%32 は０か１にしかならないが．．．
３２で割った余りであった．．．
310:298
12/01/11 09:16:02.50
>>306 >>309
いろいろ試してみたところ、とりあえず上記の問題は解決しました。
回答ありがとうございました。
311:デフォルトの名無しさん
12/01/11 11:25:42.40
すごい。
見た感じ、ちょっと不安だったから
参考になったみたいでよかった。
312:274
12/01/11 19:30:44.93
>>307
VS2008でもうまくできませんでした。
再度VS2010を入れてみるも手順通りいかず、当然実行もできませんでした。
リカバリしたのでこれからもう一度2010でやってみます。もう心が折れそうです・・・。
ちなみに2010expressとprofessionalの違いで実行できないこととかあり得ますか？
別のPCでprofessionalで試してみてくれた方がいて、うまくできたようなのですが・・・。
313:デフォルトの名無しさん
12/01/11 20:43:20.68
>>312
どんなエラーがでますか？
そもそもnvccでコンパイルするのだから，Visual Studioは関係ないでしょう？
VisualStudioのないLinux上でもコンパイル，実行できるはず．
もう一度やるのならばライブラリ lcutil32を付ければ良いかと（３２ビットならば）
314:274
12/01/12 01:51:21.33
>>313
VisualStudioであれば多少使ったことがあるので導入しました。
少なくともwindowsでは「nvccがVisualStudioのcl.exeを必要とする」と複数のサイトに書いてあるようです。
URLﾘﾝｸ(gpu.fixstars.com)
URLﾘﾝｸ(exth.net)
エラーはプロジェクトをビルドしたときに「CL.exeがないからビルドできない」というものでした。
そもそもcl.exeは32ビットのコンパイラらしいのですが、64ビットPCでも使うのでしょうか？
また以前、いい感じのところまでいったときのエラーは
C:\cuda_practice\test>nvcc -o matrix_gpu.exe matrix_gpu.cu
matrix_gpu.cu
tmpxft_00000bc0_00000000-3_matrix_gpu.cudafe1.gpu
（省略）
symbol __imp_cutCheckCmdLineFlag referenced in function "void __cdecl __cutilExi
t(int,char * *)" (?__cutilExit@@YAXHPEAPEAD@Z)
matrix_gpu.exe : fatal error LNK1120: 6 unresolved externals
でした。
どうやらリンカーの設定を見直せとのことだったのでURLﾘﾝｸ(www.scribd.com)を参考に
VisualStudioでプロジェクトのプロパティから「リンカ > 入力 > 追加の依存ファイル」にcutil32.lib, cutil64.lib を追加。
しかし「cutil64.libが見つからない」とのエラー。
実際に探してみてもcutil64.libがどこにも見つかりませんでした。
今まで3回もリカバリしてやってみましたが、どうもPathがうまく通せていない気がします。
もうすでにリカバリしてしまったので、明日また環境構築をはじめからやり直します。。
Linuxの環境はないんです、すみません・・・。
今の環境はWindows7 64bitです。
詳しくはこちら>>274に書きました。
315:デフォルトの名無しさん
12/01/12 07:17:27.30
>>314
はじめはVSのプロジェクトを構築するのではなく、VSの（32bit用の）コマンドプロンプトを立ち上げて、そこで
コマンドを打ち込み、コンパイル、実行した方がよいのでは。
これがうまく行ってから、プロジェクトの構築をした方がよいでしょう。
cutil64.libがないと言うことは64bit用のSDKがインストールされていないのでは？
ここにあるはず。
C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.0\CUDALibraries\common\lib\x64
プロジェクトを構築するときに32bit, 64bit、あるいは両方の実行ファイルを作るかどうか設定できますよ。
316:デフォルトの名無しさん
12/01/12 07:18:06.91
>>314
64bit版のcl.exeがある。
見つからないなら-ccbinオプションでcl.exeの場所を指定。
cutil32.libとcutil64.libを両方指定してはいけない。どちらか必要なものを。
で、cutil64は確か自分でビルドする必要がある。
SDKのディレクトリの中にプロジェクトファイルがあるからそれをビルド。
317:デフォルトの名無しさん
12/01/12 08:36:26.90
URLﾘﾝｸ(codepad.org)
上記のプログラムが期待と違う結果を出力するのだが原因がわかる方がいれば教えていただけないだろうか．
（CUDA 3.2 + GeForce 320M + Fedora 14 環境と CUDA 3.2 + Tesla C1060 + CentOS 5.5 環境でテスト）
-- 期待した出力 --
1.00
0.00
1.00
0.00
1.00
0.00
1.00
0.00
-- 実際の出力 --
0.00
0.00
0.00
0.00
0.00
0.00
0.00
0.00
ちなみに，11行目の右辺を 2.0f に変更すると期待通りの結果を出力してくれる．
-- 出力 --
1.00
2.00
1.00
2.00
1.00
2.00
1.00
2.00
318:317
12/01/12 08:37:35.44
すみません．
>>317のコードのURLを間違えました．
誤: URLﾘﾝｸ(codepad.org)
正: URLﾘﾝｸ(codepad.org)
319:デフォルトの名無しさん
12/01/12 09:55:39.89
やってみたけど、問題は再現されなかった
cuda_funcの中に
printf(" %d %d %f\n",threadIdx.x,threadIdx.x & 1,a[threadIdx.x]);
を入れて、チェックしてみれば？？
ちなみにこんな結果となる
0 0 1.000000 <- cuda_funcの中では問題はなさそう
1 1 0.000000
2 0 1.000000
3 1 0.000000
4 0 1.000000
5 1 0.000000
6 0 1.000000
7 1 0.000000
1.00 <-- 外でも問題はなかった
0.00
1.00
0.00
1.00
0.00
1.00
0.00
320:デフォルトの名無しさん
12/01/12 10:24:08.53
>>317
同じく問題再現されず（CUDA4.1 + GF560Ti）
そのデバイス両方とも compute capability 1.x だし
2.x 環境で試してみたらどうだ
321:317
12/01/12 10:56:51.02
>>319 >>320
コメントありがとうございます
現在の私の環境ではカーネル関数内で printf は使えないですが
変なことになるのは私の環境だけのようなので Compute capability 2.x 環境で後日試してみます
322:デフォルトの名無しさん
12/01/12 11:18:04.11
ちなみにcuPrintf.cuh とcuPrintf.cu をネット上からひらってくれば、printfと同等の関数（cuPrintf）が使えますよ♪
323:デフォルトの名無しさん
12/01/12 11:57:42.71
CUDA 3.2のコード出力に問題がある気がします。
カーネル部分をどうも
a[threadIdx.x] = (float)(FUNC(-(threadIdx.x & 1));
に変換しようとしている気配が。
FUNCはISET.S32命令による処理だけどそれがミスっているとしか思えない。
自分で以下のように書くか、あるいはFermi以降またはCUDA4.0以降を使うべし。
a[threadIdx.x] = (float)(1 - (1 & threadIdx.x));
ちなみに実行はしてないので同じ問題が起こるかどうかは知りません。
324:274
12/01/12 13:42:05.81
>>315 >>316
ありがとうございます。
とりあえず315さんの通りコマンドライン上で実行できるところを目指そうと思います。
cl.exeを指定するオプションがあるんですね。
自分はコンパイラオプションの勉強不足のようです・・・。
cutil64は自分でビルドするんですか、それは今までやってなかったと思います。
恐縮ですが、これから環境を整えるので何か参考になるサイトがあれば教えていただけますか？
いくつかサイトを見て設定したのですが、どれもうまくいかなかったので・・・。
何度もすみません。よろしくお願いします。
325:デフォルトの名無しさん
12/01/12 16:35:09.29
質問です。
同じkernel関数をfor文で何度も実行するのですが、
回数によって1回あたりのkernel関数の実行時間が異なるみたいです。
具体的には
for(i=0; i<10000; i++){
kernel_foo <<<grid_size, block_size >>>(...);
}
とすると300,000msかかるのが
for(i=0; i<100; i++){
kernel_foo <<<grid_size, block_size >>>(...);
}
では1,000msくらいになります。同じ関数を実行しているので
単純に考えれば10000回ループを回したときは100回の100倍時間がかかるはずですが、
実際には300倍程度になりました。
もう少し細かく見てみると、最初60回程度はkernel関数を1回実行するのに0.04ms程で済んでいたのですが
それ以降は1回あたり25ms程になっていました。
似た事例として
URLﾘﾝｸ(d.hatena.ne.jp)
という記事を見つけましたが、この記事内でも原因はよく分かっていません。
どなたか詳細をご存じないでしょうか？
326:デフォルトの名無しさん
12/01/12 17:29:28.93
>> 324
参考になるやら、ならないやら
URLﾘﾝｸ(www.fxfrog.com)
URLﾘﾝｸ(www.ademiller.com)
URLﾘﾝｸ(cudasample.net)
URLﾘﾝｸ(tech.ckme.co.jp)
URLﾘﾝｸ(cudasample.net)
URLﾘﾝｸ(others2.blog.so-net.ne.jp)
>> 325
1． カーネルの中の演算数がiによって異なる
2. 60回目で解が発散し NaNになっているが、お構いなしに計算は続けられている
3. カーネルからCPUにメモリーコピーしている時間が紛れ込んでいる
かな？
327:274
12/01/12 17:37:04.49
>>326
ありがとうございます。
全て拝見しましたが既読のものでした。やはり普通に設定すればできるようですね・・・。
今格闘しているので、また報告します。
328:325
12/01/12 17:50:39.27
>>326
ありがとうございます。
全項目について確認してみましたが、
1. 演算数はどれも一定
2. ループ回数によらず正しい値が得られている
3. 測定時間中にcudaMemcpy等の関数は呼ばれていない
という状態です。
329:デフォルトの名無しさん
12/01/12 18:22:29.85
>>328
1プロセスで2つのカーネル実行してるなら順番入れ替えて試してみた？
最初の実行は起動コストだかで時間かかるのが普通なんだが
330:デフォルトの名無しさん
12/01/12 19:29:37.96
時間の計測が間違っている。
cudaThreadSynchronize()を読んだあとに終了時刻を取得する必要がある。
1024個までのカーネル呼び出しはキューに積むだけですぐに次に進んでしまう。
30000回ループの方はキューが詰まるので(30000-1024)回くらいの実行完了を待つことになり、比較的妥当な計算時間に見えるはず。
もちろん正確ではないのだが。
331:325
12/01/12 19:35:41.01
>>329-330
回答ありがとうございました。
cudaThreadSynchronize()を入れるのを失念しておりました。
332:デフォルトの名無しさん
12/01/12 20:59:37.22
cudaDeviceSynchronize()
がナウいらしい
333:sage
12/01/13 08:22:49.98
>>330
便乗して質問ですが、
キューに一度に入るカーネル数が1024個で、
それ以上はカーネル呼び出しが待ちになるというのは、
CUDAのマニュアルかどこかに書いてありますか？
334:デフォルトの名無しさん
12/01/13 09:57:50.47
>>333
実験したらそうだったけど、文書化されているかどうかは知らない。
これだけ多ければ普通は問題ないし、バージョンアップで変わるかも。
335:デフォルトの名無しさん
12/01/13 11:09:53.09
>>334
了解しました。
どうも有難うございました。
336:デフォルトの名無しさん
12/01/13 16:58:47.32
それにしても きゅー にカーネルが溜まっているからと言って，GPUの処理能力が落ちるとは思えないぞ
337:デフォルトの名無しさん
12/01/13 17:17:25.93
きゅー にカーネル田丸？
それってstream使ってる時のこと？
文脈呼んでないからわからんけど
338:デフォルトの名無しさん
12/01/13 20:07:11.66
キューに空きがあれば積んで即制御が戻る。
空きが無ければ空くまで待ってる。
詰まってるから遅くなるわけではなく、遅い方が正しい実行時間なだけ。
339:デフォルトの名無しさん
12/01/18 22:19:57.72
CUDA4.0からデバイスエミュレーションがなくなったらしいけど
実機で実行するとOSごと固まることがある
みんなどうしているのでしょうか
340:デフォルトの名無しさん
12/01/19 00:43:55.87
デバイスを物理的に叩き割る
341:デフォルトの名無しさん
12/01/19 04:33:29.60
2.3くらいからなかったような気が
342:デフォルトの名無しさん
12/01/19 15:25:40.77
>>339
具体的には？？
343:デフォルトの名無しさん
12/01/20 18:49:16.04
>>342
メモリアクセスのバグがあるとWindowsが固まって動かなくなったり
ディスプレイにノイズが出てくる
OSが画面表示に使っている領域を壊しているようだけど
普通は起こらないことなんでしょうか？
344:デフォルトの名無しさん
12/01/21 10:42:26.35
私の経験では、CUDAのプログラムをチェックするために printf を使って変数を書き出したところ、
その量が多すぎて、画面が真っ暗になったり、PCがシャットダウンしたことがあった。
書き出し量を減らしたところ、問題はなくなった。
こんなことかな？？
345:344
12/01/21 11:00:22.89
追加です。
グラフィックカードを2枚刺して、一枚をCUDA用、もう一枚をディスプレイ用にすればよいとのことを聞いたことがあるが、
私のはGTX５９０でカードが2枚、入っているとのこと。
CUDAの計算に Device(0)にしても(1)にしても、前述の問題が起きたので、枚数には関係ないこともあるらしい？？
346:デフォルトの名無しさん
12/01/22 05:18:00.51
printf使えるとか便利になったわ
347:デフォルトの名無しさん
12/01/22 16:18:23.06
>>344
fermiではないのでprintfは使えないです。
GTX260なんですけど、それが悪いのかも。
計算用GPUを別に刺すってのは考えていました。
ただ、そうするべきという常識みたいなものがあるのかなっていう疑問があったので。
あとこういう状態でデバイスエミュレーション外したならひどい話だなと思ったのですが
みなさんあまり起きていないようで。
348:デフォルトの名無しさん
12/01/22 16:23:00.85
グローバルメモリの添え字計算計算が間違っていたりで確保した範囲外にアクセスすると
Windowsが固まったりディスプレイにノイズが出てくることがあるって状態です。
349:やんやん ◆yanyan72E.
12/01/22 16:48:55.32
グラフィックカード二枚差しにすると、
nSghtでハードウェアデバッグができるとかじゃなかったっけ？
350:デフォルトの名無しさん
12/01/22 17:02:34.79
YES
もしくはPC2台で片方をデバッグ用にしても良い。
実際に動かすGPUがCUDA対応なら、開発側は非nVIDIAでもOK
351:デフォルトの名無しさん
12/01/23 09:39:54.36
streamについて分からない点があるので
どなたか教えてください。
以下で「カーネル」はカーネル関数の呼び出し、
「転送」はcudaMemcpyAsync、()内はstreamだとします。
(質問1)以下の状態のとき、転送(1)は、カーネル(1)と同じstreamなので
実行できませんが、キュー内でそれより後にある転送(2)は実行できますか？
カーネル(1) 実行中
転送(1) キュー内の最初のタスク
転送(2) キュー内の2番目のタスク
(質問2)以下の状態のとき、転送(2)は実行できますか？
(キュー内で転送(2)の前にstream(0)のタスクがあります）
カーネル(1) 実行中
カーネル(0) キュー内の最初のタスク
転送(2) キュー内の2番目のタスク
(質問3)以下の状態のとき、転送(1)は実行できますか？
(実行中のタスクはstream(0)です)
カーネル(0) 実行中
転送(1) キュー内の最初のタスク
「CUDA C Programming Guide」を読んでもよく分かりませんでした。
宜しくお願いします。
352:デフォルトの名無しさん
12/01/23 10:54:29.21
1はFermiならOK。古い世代ではダメ。
2と3は絶対ダメ。stream0は前後をぶった切る役目がある。
353:デフォルトの名無しさん
12/01/24 08:17:05.28
>>352
1の動作がFermiとそれ以前で異なるのは知りませんでした。
回答どうも有難うございました。
354:デフォルトの名無しさん
12/01/24 10:48:47.33
Fermiでプロファイラのtimestampを使ってテストしたところ
カーネル(1) 実行中
転送(HostToDevice)(1) キュー内の最初のタスク
転送(HostToDevice)(2) キュー内の2番目のタスク
だと転送(2)はカーネル(1)と同時に動かず、
カーネル(1) 実行中
転送(HostToDevice)(1) キュー内の最初のタスク
転送(DeviceToHost)(2) キュー内の2番目のタスク
だと転送(2)はカーネル(1)と同時に動きました。
355:デフォルトの名無しさん
12/01/24 19:53:55.48
コンパイル時に以下のエラーメッセージが出ているのですが、
ptxas error : Entry function '...' uses too much local data (0x16054 bytes, 0x4000 max)
これはローカルメモリーと何か関係あるような気がするのですけど、
ちょっとわからないので教えて頂けないでしょうか。
宜しくお願いします。
356:デフォルトの名無しさん
12/01/24 21:03:52.65
それともう1つ。
externを使って、グローバル変数を共有させたいのですが、
ホスト・デバイスともに共有させたい場合はどうしたらいいのでしょうか?
357:デフォルトの名無しさん
12/01/24 21:56:32.01
> ptxas error : Entry function '...' uses too much local data (0x16054 bytes, 0x4000 max)
単に配列の次元が大きすぎるのでは？
最大 0x4000バイトのところを0x16054 bytes使おうとしている？
> ホスト・デバイスともに共有させたい場合はどうしたらいいのでしょうか?
物理的に無理でないの？
ホスト←→デバイス間でコピーしないといけないのでは？？
358:デフォルトの名無しさん
12/01/24 21:56:54.10
>>355
メッセージそのままの意味だろ。
359:デフォルトの名無しさん
12/01/25 01:32:34.89
CUDAの思想はいいと思うけど、ターゲットデバイスを完全には
隠蔽しきれないのが残念だな。
360:やんやん ◆yanyan72E.
12/01/25 01:41:55.08
スピード勝負の世界で、しかもかなり複雑なアーキテクチャなのに、
ハードウェアを完全に隠蔽する訳にいかないだろ。
361:357
12/01/25 12:34:12.36
> ホスト・デバイスともに共有させたい場合はどうしたらいいのでしょうか?
ヘッダーファイルに
#DEFINE
などで定義して、ホストとデバイスのプログラムに そのヘッダーファイルをinclude する手があった♪
362:355,356
12/01/25 15:21:00.98
>>358
青木氏の「はじめてのCUDA」に似たようなエラーメッセージが載ってるのですが、
local dataがどの部分の話なのか、ちょっとわからない状態でして...
>>357
>単に配列の次元が大きすぎるのでは？
予定だと1×256スレッドしか使うので、足りないはずないような... 意味が違うか。
ちょっとわからないので詳細お願いします。
>ヘッダーファイルに
>#DEFINE
>などで定義して、ホストとデバイスのプログラムに そのヘッダーファイルをinclude する手があった♪
ちょっと試してみます。それができなかったら、コピーするしかないかもしれません。
__device__ extern TEST test;
のような宣言してて、実行ファイルができたときがあったのですが、何か違うよな、と引っかかってたので、
ここで質問しました。ありがとうございます。
363:デフォルトの名無しさん
12/01/25 17:40:16.56
差し支えなければプログラムをアップしてくださいませ♪
364:デフォルトの名無しさん
12/01/25 18:01:06.82
>>362
このサイトの一番下にある資料の3-41ページにそのエラーメッセージが載っている模様。
URLﾘﾝｸ(accc.riken.jp)
365:355,356
12/01/25 18:04:09.31
>>362です。
>>363
申し訳ないのですが、プログラムのアップは出来ないです。すみません。
なんかもう分からなくなったので、ホストとデバイスの住み分けを行ったところ、
「ptxas error...」云々のメッセージが消え、コンパイルできました。
（さっきまでhost,device,global宣言関数がごちゃ混ぜな状態だった。）
何が原因だったのか分からず仕舞いです。
グローバル変数の共有はとりあえず、コンスタントメモリにデータコピーで様子見することにしました。
元コード(.c)をもう一度読み直したところ、デバイス側の方はReadOnlyで十分なようでしたから。
皆様回答ありがとうございました。またよろしくお願いします。
366:365
12/01/25 21:05:50.75
>>364
資料ありがとうございます。
うまくいったと思った途端、また同じエラーが出てきてしまったので、確認します。
.cファイルに.cuファイルを組み込ませるようにしたら、
__host__ __device__で修飾した関数が定義されてないと.cファイル側に言われ、
.cファイルと.cuファイルそれぞれ単独で動かせるようにしたら、
（同じ内容の関数を.cと.cuそれぞれ別名で実体作った。.c:test .cu:ktestみたいな）
今度はさっきと同じエラー。
実行ファイル作るだけなのに難しい...
367:デフォルトの名無しさん
12/01/26 08:42:18.26
Host側とDevice側で型を別にすれば開発もっと楽になると思うんだけど
368:デフォルトの名無しさん
12/01/26 08:53:13.45
取り敢えずCudaのサンプルは捨てて、インクルード関係とオブジェクトの生成手順を確認するんだ。
369:デフォルトの名無しさん
12/01/26 10:14:19.95
見せられないところは削除して、問題の部分だけ残して、アップすれば？？
370:366
12/01/26 16:19:08.01
>>368
>取り敢えずCudaのサンプルは捨てて、インクルード関係とオブジェクトの生成手順を確認するんだ。
のように手順を踏んでコンパイルしたところ、実行ファイルができました。
皆様ありがとうございました。
371:370
12/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:デフォルトの名無しさん
12/01/26 21:28:56.32
4.1正式版
CUDA Toolkit 4.1 | NVIDIA Developer Zone
URLﾘﾝｸ(developer.nvidia.com)
373:デフォルトの名無しさん
12/01/26 21:55:28.42
構造体の中身がよく分からんが，allocateのところで *100は必要なのかな？
構造体Bのメモリーを１００個分用意しようとしている？？
構造体Bの中にすでに [100]個の配列を取っているのに？？？
数は合っておるのか？？？？
374:デフォルトの名無しさん
12/01/26 22:17:44.08
ｒｈｅｌ6.0のがリンクミスかでダウンロドできん
375:デフォルトの名無しさん
12/01/27 00:03:32.49
>>372
遂にRC取れたか！4.0からどう変わってるか知らんけど。
Kepler向けのプログラムを作れるのは5.0とかになるんだろうか。
そしてそれはいつ出るんだろう。
>>371
1GBのVRAM積んでるカードでも400MBぐらいの確保で失敗したことがあったような覚えがある。
確保するサイズを小さくしてもエラーが出るってんならこの話は忘れてくれ。
376:デフォルトの名無しさん
12/01/27 01:34:46.65
100x100の行列の積を1万回くらい実行したいのですが、
CUBLASだと小さすぎてCPUより遅くなってしまいます。
この1万回は相互に依存関係は無いので、並列化が可能なのですが、
このプログラムは自分で書くしかないのでしょうか。
TESLA C1070があるので、丸ごとVRAMに載ります。
377:デフォルトの名無しさん
12/01/27 03:17:55.02
4.1にcublas{S,D,C,Z}gemmBatchedってのが追加されてました。
especially for smaller matricesに効くってあったので、使ってみます。
378:デフォルトの名無しさん
12/01/27 07:52:34.05
一つのブロック内で100x100の行列の積を１回行う？
379:デフォルトの名無しさん
12/01/27 07:52:51.42
間違った
380:デフォルトの名無しさん
12/01/27 07:53:00.44
あれ？
381:デフォルトの名無しさん
12/01/27 07:53:07.78
一つのブロック内で100x100の行列の積を１回行う？
382:デフォルトの名無しさん
12/01/27 07:55:11.45
たびたび，失礼
一つのブロック内で100x100の行列の積を１回行う？
そのブロックを１万個用意する？
それらを並列に計算する？
と言うコンセプトかな？？？
自分でプログラムを作った方が，よほど勉強になると思う
383:371
12/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:デフォルトの名無しさん
12/01/27 20:31:39.53
d_aに送信してるあたりが
385:デフォルトの名無しさん
12/01/27 21:52:59.41
d_b のサイズは構造体Bと同じ？？
d_aのサイズはbと同じ？すなわち構造体Bと同じ？？
違っていたらエラーが出て当たり前では？？？
386:デフォルトの名無しさん
12/01/28 13:33:36.75
それから，ホスト側で も メモリーは確保しているのか？？？？
387:383
12/01/28 17:54:35.37
>>371
なんかサンプルがめちゃくちゃなんで書き直します。
388:387
12/01/28 20:55:21.49
確認なのですけど、カーネル関数の引数はポインタ限定ですか?
389:デフォルトの名無しさん
12/01/28 22:58:27.22
配列はポインタで，１変数はそのままではなかったかい？
なぜか知らんけど
390:デフォルトの名無しさん
12/01/28 23:07:13.34
夜の埼玉は
391:デフォルトの名無しさん
12/01/29 01:49:05.91
>388
ホスト側のポインタ以外なら何でもOKのはず。少なくともintは渡せる。
クラスを渡せるかどうかは知らないけど。
392:デフォルトの名無しさん
12/01/29 06:58:36.12
ホストのポインタ送って、GPU側のload/store時に
メインメモリから自動転送ってこともできるから、
その説明は正しくない。
393:388
12/01/30 15:13:56.68
>>391
セグメンテーション違反してた部分の引数に（デバイスで使用する）double（実体）を指定してたのですけど、
（デバイスで使用する）ポインタに変えたら、問題なく実行できました。
dtempはcudaMalloc/cudaMemcpy使ってます。
そもそもな話cudaはポインタ宣言したdtempの実体は*dtempではないのかな。
394:デフォルトの名無しさん
12/01/30 15:16:54.53
>>393
何を言っているのか判らない。問題が再現する最低限のコードでも貼ってみていただきたい。
395:デフォルトの名無しさん
12/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:デフォルトの名無しさん
12/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:デフォルトの名無しさん
12/01/30 16:57:38.16
>>395
dtempはdevicePointerなんだから、Hostコード中でデリファレンスしちゃダメ。
double値を渡したいだけなら引き数はdoubleで充分。
それと、※1は宣言じゃなくて呼び出しだから構文としてはあっている。
そもそも何がしたいのだろう。値を渡したいなら値を渡せばいいし、
ポインタを渡したいならポインタをそのまま渡せばいい。
devicePointerをHostでデリファレンスしちゃダメだし、
(その後どうせdevice側でデリファレンスできないから)HostPointerをdeviceにコピーしてもダメ。
398:デフォルトの名無しさん
12/01/30 17:04:05.78
>>396
>>397
色々ごちゃまぜだったようです。スッキリしました。
ありがとうございました。
399:デフォルトの名無しさん
12/01/30 20:58:23.57
これからCUDAの勉強を始めようと思っています。
GPGPUプログラミングはシェアードメモリを有効活用できるかがポイントであるように思います。
そこで、あえて最新のグラボに買い換えず、キャッシュのないTesla世代のGTX285を所持し続けています。
キャッシュがあると、シェアードメモリをうまく使えてなくても、それをカバーしてしまって、
うまくプログラムできているかどうかを把握し辛いと考えたからです。
ただ、心配なのが、Fermi世代以降に発売された書籍だと、内容がTesla世代のGPUにそぐわない
といった状況が起きてしまうことです。
GTX285のまま勉強するに際して気をつけること等あれば教えてください。
問題がなければ関連書籍（日本語）を全部購入して読もうと思います。
よろしくお願いします。
400:デフォルトの名無しさん
12/01/30 21:08:43.55
＞＞GPGPUプログラミングはシェアードメモリを有効活用できるかがポイントであるように思います。
そんなの計算する内容によるわ
活用しようがないのもあるからな
計算（使用）目的は明確なの？
401:399
12/01/30 21:39:43.45
>>400
信号処理に使う予定です。
FIRフィルタリングやFFTなどを大きな1次元、あるいは2次元配列に適用したいと思っています。
402:デフォルトの名無しさん
12/01/30 21:45:58.41
俺は400じゃないが、気をつけることは、勉強がひと段落していざ使うときになって知識が陳腐化してても泣かないことかな。
ハードウェアはどんどん変わるし、TeslaやFermiで良しとされたことが数年後に通用するかはわからないからね。
せっかく今から始めるならFermi・Keplerを使ったほうがいいと俺は思うけど、
10年後にはどれも等しく陳腐化してるだろうから長期的に考えると確かにTeslaでもいいのかもしれない。
ただ長期的に考えること自体がGPGPUでは…と1行目にループする。
403:デフォルトの名無しさん
12/01/30 22:12:50.10
古い方がいろいろ制限があって，勉強になるとは思う．
404:デフォルトの名無しさん
12/01/30 22:17:32.23
ｆｆｔはcpu側からはインターフェースあるけどディバイス内では自前で作るしかない
405:デフォルトの名無しさん
12/01/30 22:18:26.89
> キャッシュがあると、シェアードメモリをうまく使えてなくても、それをカバーしてしまって、
> うまくプログラムできているかどうかを把握し辛いと考えたからです。
シェアードメモリーを使うようにプログラミングすればよいかと
どのメモリーを使っているのか分からない，と言うことはないと思う
あれば，そこをしっかり勉強すべきかと
406:デフォルトの名無しさん
12/01/30 22:19:12.36
そう言えば古いのは倍精度の計算ができないのでは？？
407:デフォルトの名無しさん
12/01/30 22:22:30.66
まあ，デバイスメモリーだけで計算してみるとか，
シェアードメモリーも使って計算してみるとか，
いろんなプログラムを作って比較するのが良いかと
408:デフォルトの名無しさん
12/01/30 22:26:01.98
最初はシェアドするスレッドはどれだってだけでも頭が痛くなった
409:402
12/01/30 23:28:19.90
あとTeslaでやるなら、FermiとTeslaでバンクコンフリクトの発生条件が違ってるから
Teslaでのこれの回避には深入りしないほうがいいかと。
他にシェアードメモリの大きさなどの量的な違いがあるけど、そういった量的な
制約のために質的なというかアルゴリズムの種類を変えるようなことまで深入りするのも
避けたほうが。だってFermiやそれ以降で量が変わると一から見直しになるから。
これはFermi使ってる人もいずれ同じことなんだけど、勉強じゃなくて性能が欲しいんだからしょうがない。
410:デフォルトの名無しさん
12/01/30 23:57:48.21
sharedがregister並に速いなんて嘘だから。
occupancyなんて上げるな。
warpあたりのregisterを増やして
ILPを利用してレイテンシを隠蔽すべきだと。
URLﾘﾝｸ(www.cs.berkeley.edu)
411:399
12/01/31 00:31:29.61
>>403
やはり、そうですか！
>>404
参考になります。
FFTは重要になるのでしっかりと取り組みたいです。
>>405
>>407
>>408
なるほど、どのメモリかをよく意識してプログラミングします。
CPUの最適化に取り組んでいて思ったのですが、
キャッシュは便利な反面、こちらから意図的に挙動を制御できないことが
パフォーマンスの考察を難しくしてしまう側面もあると感じました。
キャッシュなしのTeslaで鍛えたいです。
>>406
単精度しか使わないので問題なしです。
もし半精度でパフォーマンスが上がるなら、そのほうがイイくらいです。
>>402
>>409
世代を跨ぐにあたって非常に参考になるご助言をありがとうございます。
各種メモリの特性や帯域を意識して取り組むことで、
固有のデバイスに限定されない定性的なところを理解したいと思います。
Keplerではこんな組み方がイイかもしれないな、と思索できたら楽しそうです。
>>410
興味深い資料をありがとうございます。
各部のレイテンシ、スループットを頭に叩き込みます。
412:デフォルトの名無しさん
12/01/31 08:31:17.17
>>411
最終的に動かす機器で最もよい性能が出るように最適化するべきだし
自動でやってくれることを手でやるってのは勉強にはなるかもしれないけど意味ないのではって思うけど。
コンパイラの最適化を切ってアセンブリ書くようなものでしょ。
413:デフォルトの名無しさん
12/01/31 12:08:28.56
時間がいくらでもあるorそういう研究ならともかく
短期間で組めてそこそこのパフォーマンスが出せればいいなら
キャッシュのあるFermiで適当に書いてもいいと思うんだけどね
414:399
12/01/31 21:32:30.15
>>412
おっしゃる通りです
そのあたりは目的をよく考えて判断したいです。
今回はどちらかというと、GPUのポテンシャルを見極めたいというのが主です。
コードの可読性を優先して、性能はコンパイラやキャッシュ等にできるだけ委ねるというスタイルも
今後検討していきたいと思っています。
>>413
今回は研究用途ですが、
将来的にはFermi以降のGPUで、ちゃっちゃと書いたコードが快適に動くというところを目指したいです。
ただ、あまりにも下手なプログラムにならないよう、最初に少し深くやっておきたいです。
415:デフォルトの名無しさん
12/02/02 04:23:36.19
レジスタといえば、GK104では強化されるのか、それともさらにピーキーになるのか気になる。
GF104はグラフィック特化で、レジスタ数やワープ数の上限に悩まされたし。
とりあえずアーキテクチャが一般公開されるまで待つしかないかな。
416:デフォルトの名無しさん
12/02/04 00:16:58.91
また出てるな、これで最後らしいが
URLﾘﾝｸ(page4.auctions.yahoo.co.jp)
417:デフォルトの名無しさん
12/02/08 09:16:57.43
誰かいるかな…
URLﾘﾝｸ(includehoge.blog.fc2.com)
このブログのソースをコンパイルして動かしてみようと思ったのですが
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:デフォルトの名無しさん
12/02/08 09:17:57.58
誰かいるかな…
URLﾘﾝｸ(includehoge.blog.fc2.com)
このブログのソースをコンパイルして動かしてみようと思ったのですが
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:デフォルトの名無しさん
12/02/08 09:22:20.66
誰か居るかな・・・いたら助けてください
URLﾘﾝｸ(includehoge.blog.fc2.com)
このブログのソースをコンパイルして動かしてみようと思ったのですが
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:デフォルトの名無しさん
12/02/08 09:22:49.74
誰か居るかな・・・いたら助けてください
URLﾘﾝｸ(includehoge.blog.fc2.com)
このブログのソースをコンパイルして動かしてみようと思ったのですが
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:デフォルトの名無しさん
12/02/08 09:23:33.66
うわなんか四回も書き込んじゃったごめん
422:デフォルトの名無しさん
12/02/08 09:25:46.81
コンパイル詳しくないけどlibcufft.so（linuxの場合）のリンクしてる？
windowsならlibcufft.dllだとおもう
423:デフォルトの名無しさん
12/02/08 09:27:16.52
コンパイルじゃなくてリンクの問題でしょ
424:デフォルトの名無しさん
12/02/08 09:58:44.53
nvcc -lcufft sample.cu
425:デフォルトの名無しさん
12/02/08 10:20:43.62
このスレでときどき思うのはC/C++をほとんど知らないのにもかかわらず
CUDAに挑戦しようとする勇壮な挑戦者が多いな、ということ
彼らはその後うまくいってるのだろうか？
426:デフォルトの名無しさん
12/02/08 13:41:41.97
CUDA.NETも3.0で止まってるしなぁ・・・
427:デフォルトの名無しさん
12/02/08 21:44:42.85
コンピュート・ケイパビリティによって何が変わったかを俯瞰できるような一覧を掲載したサイトってご存知ありませんか？？
428:デフォルトの名無しさん
12/02/08 22:16:51.83
>>417
追加のライブラリファイルにナントカ.libを追加すればいけそう。
ナントカは>>422さんのレスから察するにlibcufft.libかなぁ。
このへんは>>425さんの言うとおりcuda云々じゃなくてVisual Studio(C++)を使う上で頻繁に設定するところだねー。
>>427
俺も知りたいなぁそれ。その手の情報はSDKpdfのProgramming GuideとBest Practiceくらいしか知らない（片方だけかも）
429:デフォルトの名無しさん
12/02/08 22:26:54.49
cufft.libなのだが。
430:デフォルトの名無しさん
12/02/09 07:17:56.91
てすてす
431:デフォルトの名無しさん
12/02/09 07:18:12.24
コマンドなら
432:デフォルトの名無しさん
12/02/09 07:18:18.71
あれ？
433:デフォルトの名無しさん
12/02/09 07:18:34.70
以下のよう
434:デフォルトの名無しさん
12/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:デフォルトの名無しさん
12/02/09 07:19:01.92
ファイル名をfft.cuとしてみた
436:デフォルトの名無しさん
12/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:デフォルトの名無しさん
12/02/09 07:21:26.93
-arch=sm_10 は適宜，変更のこと
438:デフォルトの名無しさん
12/02/09 07:25:00.70
これでも行けた♪
nvcc fft.cu -lcufft
439:デフォルトの名無しさん
12/02/09 07:26:45.21
既出でした．失礼 >> 424
440:デフォルトの名無しさん
12/02/09 11:41:33.64
GPUの勉強を始めたばかりなのですが、スパース行列の格納形式種類で
ELLの格納方法とメリットが分かりません
調べたらlisのマニュアルに配列方法はあるのですが、説明と図とが噛み合ない点があり
理解できずにいます。
だれか教えてください
441:デフォルトの名無しさん
12/02/09 13:59:04.44
>>426
CUDA#とかCUDAfy.NETとか代替のライブラリはあるっぽいよ
使ったことないから詳細は知らないけど
442:デフォルトの名無しさん
12/02/09 14:23:54.18
>>440
LisはGPU上で動くの？
443:デフォルトの名無しさん
12/02/09 14:39:02.94
sparse matrixの格納形式はCUDAの問題じゃね～だろ
444:デフォルトの名無しさん
12/02/09 15:45:25.65
>>440
コアレッシング状態で転送できるようになる。
445:デフォルトの名無しさん
12/02/09 17:19:58.52
CUDA.NETは使えなくはないぞ♪
446:デフォルトの名無しさん
12/02/09 18:44:06.72
>>440
スパース行列の格納方法がのっているのがLisのオンラインマニュアルくらいだったので
>>443
ELLはGPUに適した格納方法と、ネットのどこかにあったので関連があるかと思い書き込みました
たしかにCUDA以前の問題ですよね
>>444
なるほど、見えてきました！ありがとうございます
447:デフォルトの名無しさん
12/02/09 19:02:29.63
>>422->>439
こんなにたくさん答えてくれるとは・・・
おかげで動きましたーありがとう
448:デフォルトの名無しさん
12/02/09 19:08:48.94
おお、よかった^^
449:デフォルトの名無しさん
12/02/10 10:39:58.03
URLﾘﾝｸ(detail.chiebukuro.yahoo.co.jp)
450:デフォルトの名無しさん
12/02/10 13:35:49.60
>>441
CUDAfy.NETは知らんかった
最近出てきたのね、ありがとう
451:デフォルトの名無しさん
12/02/10 21:23:36.65
2次元配列にアクセスするときのインデックス計算で
gy = blockDim.y * blockIdx.y + threadIdx.y;
gx = blockDim.x * blockIdx.x + threadIdx.x;
というのを頻繁にやると思うのですが、
この積和演算ってCUDAコアが使われるのでしょうか？
それとも、CUDAコアとは別にアドレス演算器（？）的なハードウェアがあって、
それで計算してくれるのでしょうか？
後者だとうれしいです。
452:デフォルトの名無しさん
12/02/10 21:37:17.01
アドレスもプログラムもデータ
453:デフォルトの名無しさん
12/02/10 23:52:12.69
>>451
CUDAコア上に展開された（定義された）ブロックの中で演算をしていると思う
454:451
12/02/11 01:18:08.05
>>453
CUDAコアのサイクルを消費して計算しているということですね？
ありがとうございました。
455:デフォルトの名無しさん
12/02/11 01:24:48.08
インデックス計算かどうかをGPUがどう見分けるというんだ
456:451
12/02/11 10:05:14.94
>>455
単にCUDAの予約語であるblockDimやthreadIdだけで構成されていて、
○○ * ○○ + ○○ の積和の形になっていれば、
専用の回路に計算させる、とかできそうなものですが。
CPUではメモリアドレス式のカタチを解析してアドレス演算器に割り当てています。
457:デフォルトの名無しさん
12/02/11 11:29:23.45
そもそもグリッド，ブロック，スレッドの概念を理解しているのかな？
458:デフォルトの名無しさん
12/02/11 11:29:54.96
○○ * ○○ + ○○ だけではプログラムは作れないと思うのだが
459:デフォルトの名無しさん
12/02/11 11:55:45.33
>>456
それ、どのCPUでの話ですか？
460:デフォルトの名無しさん
12/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:デフォルトの名無しさん
12/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:デフォルトの名無しさん
12/02/11 19:21:14.69
追記です。
cutil64.lib がないので 行列積サンプルプログラムは、Releace モードではビルドできていません。
ビルド＆実行でてきているのは、debug モードです。
463:デフォルトの名無しさん
12/02/11 20:12:55.27
Releace + x64モードでcutil64.dllはできるが，cutil64.libはできないと言うことですか？
464:デフォルトの名無しさん
12/02/11 20:14:45.64
エラーが出てもライブラリはできている，と言うことはないですか？
この中に↓
C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.1\C\common\lib\x64
465:デフォルトの名無しさん
12/02/11 20:16:27.24
失礼 Releace -> Release
466:デフォルトの名無しさん
12/02/11 20:26:05.18
まあ、CUDAのインデックス計算ほど
アホ臭いものは無いと思うがな。
テクスチャアドレスではもっと複雑な計算しているのに。
467:デフォルトの名無しさん
12/02/11 20:50:55.57
CUDAを数値計算に使おうという話だよ
468:デフォルトの名無しさん
12/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:デフォルトの名無しさん
12/02/12 00:49:32.40
変だね．こちらも同じようなエラーが出るが（xx.dllを開始できません），xx.libはできている．
試しに，ソルーションエクスプローラーのcutilを右クリックして，「追加」，「既存の項目」で，cutil64.dll を追加してみる．．．．とか
470:デフォルトの名無しさん
12/02/12 00:51:08.01
で，何でビルド前に「cutil64.dll のみ存在 」するのかな？
一度，全部クリーンしてみる．．．．とか
471:デフォルトの名無しさん
12/02/12 01:08:37.66
>>469
お返事ありがとうございます。
エラーは出るんですね。貴重な情報です。
エラーを気にしないですみます。
>cutil64.dll を追加してみる．．．．とか
無理を承知でやってみました。当然ですがダメでした。
>>470
そういうものだと思っていましたが、確かにそうですね。
クリーンインストールしてみます。
472:デフォルトの名無しさん
12/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:デフォルトの名無しさん
12/02/12 07:49:26.70
cutilのクリーンを実行すれば，すべてのdll，libが消えるはずなのに，もし， cutil64.dllが残っているとすれば，
書き込み，読み込みの権限がないと言うことでは？
cutil64.dll はマニュアルで消せますか？
474:デフォルトの名無しさん
12/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:デフォルトの名無しさん
12/02/12 15:40:34.77
↑ 番号ミスです。ごめんなさい。
>>473 が正しいです。
476:デフォルトの名無しさん
12/02/12 15:49:09.03
よかった，よかった♪
477:デフォルトの名無しさん
12/02/12 22:34:46.87
OS：32bit CUDA:32bit VisualStudio:32bit
でCUDAプログラムをしている人が、速度を上げることを目的に
OS：64bit CUDA:64bit VisualStudio:64bit
に変更することって意味あります？
478:デフォルトの名無しさん
12/02/12 23:58:34.00
ないと思う。ホスト側で2GB以上のメモリを使いたいわけでもないのだろうし。
64bitビルドして動かすとポインタがレジスタを2つ使ってしまうという話があるので、逆に遅くなる可能性も。
本当に64bitビルド時にデバイス側でもポインタが64bitで扱われるかどうかを
自分の目で確かめたわけではないけど、少なくともここではそういうことを言っている。
高速演算記 第3回 「チューニング技法その１ ＣＵＤＡプログラミングガイドからピックアップ」 | G-DEP
URLﾘﾝｸ(www.gdep.jp)
＞ここでレジスタ数は32bitレジスタの数で、long long, double値や64ビット環境を対象としたCUDAプログラムでは
＞ポインタなどが64bitとして扱われますのでレジスタを2つ使用することになります。
479:デフォルトの名無しさん
12/02/13 10:45:03.24
>>478
なるほど。ありがとうございます。
これと少し環境が違う場合ですが、
OS：64bit CUDA:32bit VisualStudio:32bit
と、
OS：64bit CUDA:64bit VisualStudio:64bit
では、速度差はどうなると思われますか？
ネット上でよく目にするmatrix積のサンプルコードで実験したところ、
同程度か、どうかすると前者の方が速い場合がありました。
前者は、WOW上で動いていると思われるので、
遅いだろうと予測したのですが、それに反する結果でした。
これをどう解釈すればよいか謎なんです。
480:デフォルトの名無しさん
12/02/13 10:46:42.16
479です。OSは、Windows7 64bit Pro です。
481:デフォルトの名無しさん
12/02/13 11:17:41.23
プログラムのどの部分を計測したのであろうか？
CUDAの部分のみ？？
482:デフォルトの名無しさん
12/02/13 12:26:28.94
CUDAの部分のみです。
483:デフォルトの名無しさん
12/02/13 12:35:43.06
CUDAはOSに依存せずに動くとかで、OSのbit数には無関係ということですか？
そのため、
>>478で説明して頂いたように、同程度もしくは32bitが若干速くなる場合もあるということでしょうか。
484:デフォルトの名無しさん
12/02/13 13:17:49.82
ちなみに測定時間はいくらでしょうか？
演算量を増やすとどうなるでしょうか？
485:デフォルトの名無しさん
12/02/13 16:09:51.62
ソースは、以下のHPの最後の「Matrix_GPU」です。
URLﾘﾝｸ(www.gdep.jp)
計算は、スレッドサイズ一杯で組まれているようですので、
#define MATRIX_SIZE 1024/*行列１辺の数*/
をこれ以上上げることができませんでした。
ちなみに、ノートパソコン搭載のGPUチップは、GT-540M。
Releaseでビルドして 650ms付近を中心に数10msばらつきます。
win32とx64で大差はなく、
win32の方が平均10ms程度速いような気がするという程度です。
486:デフォルトの名無しさん
12/02/13 17:03:00.91
CUDAの計算のところを１０００回程度，繰り返すとかはできるでしょうか？
そうすれば有意な差が出るかも知れません．
487:デフォルトの名無しさん
12/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:デフォルトの名無しさん
12/02/13 19:39:39.79
なるほど，有り難うございます．
有意な差が出たと言うことですね．
では，後は分かる人，よろしく♪
ε=ε=ε=ε=ε=┌(； ・＿・)┘
489:デフォルトの名無しさん
12/02/14 10:35:58.19
GPU側では64bitのメリットは使用可能なメモリが増えることだけで、その代わりにアドレス計算の命令数や使用レジスタ数が増えるデメリットがある。
前はGPU側だけ32bitモードで動かすという恐ろしい手法があったが、今使うのはなかなか難しい。
490:デフォルトの名無しさん
12/02/14 16:21:21.98
ptxオプションつけてptx出力を見れば見当がつくね。
64ビットアドレッシングだとアドレス計算の為にint32→int64の変換がしばしば行なわれる。
491:デフォルトの名無しさん
12/02/14 18:19:16.13
ざっくり捉えると、32bitと64bitでは、
CUDA自体はさほど変わらない。
CUDAと切り離されてコンパイルされるC++言語部分は、64bitの方がいくらか速い。
という感じでしょうか。
492:デフォルトの名無しさん
12/02/14 18:24:53.01
>>491
アドレス計算多用したり、レジスタ沢山使うようなカーネルだと
CUDA自体32bit版の方が無視できないほど
速くなる可能性はある。
493:デフォルトの名無しさん
12/02/14 18:53:10.54
doubleが速くなるっていう噂はどうなったの？
494:デフォルトの名無しさん
12/02/14 22:53:20.23
GPU上にあるコンスタントキャッシュメモリって、CUDAではどのようにして使われるのでしょうか？
計算の中にリテラルがあった場合（CPUでは命令の即値として埋め込まれますが）や、
constantを付けて宣言した定数が割り当てられたりするのでしょうか？？
495:デフォルトの名無しさん
12/02/16 05:51:50.91
llvmで使えるようになったのなら
boostも使える？
496:デフォルトの名無しさん
12/02/16 08:10:03.40
>>494
適当に数値を入れて，確かめてみるとか
497:デフォルトの名無しさん
12/02/16 08:17:55.24
>>495
２年前にパッチが出ているようです．boostが何か知らんけど
URLﾘﾝｸ(svn.boost.org)
498:494
12/02/17 00:11:19.70
>>496
試してみました。
const変数もリテラルもレジスタにマッピングされました。
ひょっとして、DirectXAPIなどのグラフィックスAPIを利用していた場合に
使われるだけ（要は旧来通りの使われ方）かも･･･
プロファイラを見ると、テクスチャキャッシュのヒット/ミスヒットが見れるようになっていますが
これなんかはグラフィックスAPI専用のような気が。
しかし同じようにコンスタントキャッシュの項目もあって然りと思うが在らず。
う～ん、気にしないことにしますｗ
499:デフォルトの名無しさん
12/02/17 06:31:17.52
貴重な報告、ありがとうございます♪
500:デフォルトの名無しさん
12/02/17 07:54:24.51
GPUを演算器としてフルに使いたい場合って
ビデオ出力をオンボードのものになるようBIOSで切り替えた方がいいのでしょうか
linuxだったらxconfigいじるだけでいけるとか？
501:デフォルトの名無しさん
12/02/17 07:57:38.45
フルの定義は？
502:デフォルトの名無しさん
12/02/17 07:58:52.97
>>494
CPUでも即値として埋め込めない配列や文字列リテラルは
.rdataセクションのデータとして書き込まれるのと同様に、
GPUでも即値として使えないconstデータはコンスタント領域に置かれる。
で、コンスタント領域の値を直接オペランド指定できない
命令を使う場合は当然レジスタを介して使用することになる。
逆に一部の変数にデフォルトで定数が入っているなんてのは
ABIで決まったもの以外は普通無理だから、コンスタント境域などの
メモリに置いて読み込む以外やり様が無い。
503:デフォルトの名無しさん
12/02/17 08:33:23.95
>>500
普通はGPUカードに何もモニタを繋がなければ(≒出力する設定をしなければ)
GVRAMをほぼ全部使い切れるよ。
Linuxなら、Xを立ち上げなければGPUカードから出力していてもかなり使えるよ。
ついでに、3秒だか5秒だからの制限もなくなるからGPUカーネルで長時間計算させられるよ。
504:デフォルトの名無しさん
12/02/17 20:22:12.27
VRAM 4GBのボードが六千円台
URLﾘﾝｸ(akiba-pc.watch.impress.co.jp)
505:494
12/02/18 16:32:22.87
>>502
ありがとうございます。
コンスタント変数の扱われ方、レジスタが必要な理由、理解しました。
書籍を読んでいたら、コンスタントメモリおよびテクスチャメモリについて書かれているのを見つけました。
コンスタントメモリは__constant__修飾子を付けて宣言することで割り当てられるようです。
物理的にはVRAMに置かれ、GPU内部のコンスタントキャッシュでキャッシュが効くものと思われます。
空間フィルタリングなどでフィルタ係数を格納しておくのに適しているのではと思いました。
（繰り返し参照されるのでキャッシュが効きそう）
テクスチャメモリはテンプレートクラス型のtexture<･･･>で宣言することで割り当てられるようです。
物理的にはVRAMに置かれ、GPU内部のテクスチャキャッシュでキャッシュが効くものと思われます。
さらに、グラフィック用途時の同様、テクスチャフィルタリングやサンプリング方式などのハードワイヤード回路機能を
利用できるよう、それらを指定するプロパティを保持しているようです。
506:デフォルトの名無しさん
12/02/18 21:02:09.38
複数カーネルを呼び出すんだけどカーネルのソースはそれぞれ別ファイル
それぞれのカーネルから同じDevice関数を呼び出したいんだけどファイルまたがることできないようだし
ソースごとに同じ内容の別名関数作るしかないのかな？
507:デフォルトの名無しさん
12/02/18 22:43:38.30
ファイルが別でも，プログラムとしてつながっていれば問題はないと思うが．
関数のプロトタイプ宣言をしっかりしておけば良いかと．．．
508:デフォルトの名無しさん
12/02/18 22:46:25.76
共通のdevice関数を別のヘッダに纏めて
カーネルソースからincludeして使う。
device関数はどうせinline展開されるのだから
通常のプログラムでのinline関数と同様に扱えばいい。
509:デフォルトの名無しさん
12/02/18 22:54:19.36
>>508
サンクス
でも今やってみたんだけど全ファイルで展開されるんで
multiple definition of 関数名
ってエラーになっちゃう
実際やってみました？
510:デフォルトの名無しさん
12/02/18 23:22:56.53
カーネル起動時、1スレッドの起動につき1クロックかかり、128スレッドなら128クロックかかると書籍に書いてあるのですが、
これはCPU側でもかかるのでしょうか？
それとも、CPU側はカーネル起動の合図を送ってすぐ制御を戻し、
GPU側の制御ブロックがクロックをかけてスレッドを起動するというかたちでしょうか？
511:デフォルトの名無しさん
12/02/19 01:11:11.40
教えて欲しいことがあります．
CPU側で複数スレッドを立てて，互いに独立した処理を行っています．
そして，それぞれスレッドにおいて，CUDAを使って画像処理をさせたいのですが，（たとえば，グレースケール化とか）
この時，CUDAでの処理にテクスチャメモリを使いたい場合は，どのようなコードを書けばいいのでしょうか？
テクスチャメモリを使う場合，グローバルで宣言しなけらばならないですよね？
たとえば，5枚のRGB画像をグレースケール化するときに，
CPU側で5スレッド立てて，各スレッドでCUDAを使ってグレースケールへの変換処理をしたいのですが，
テクスチャメモリをグローバルで宣言するとおかしなことになるきがするんですが．
どなたか教えて頂けないでしょうか？
宜しくお願い致します．
512:デフォルトの名無しさん
12/02/19 01:21:41.45
CUDA対応のGPUを5枚挿す。
またはGTX590などを3枚挿す。
513:デフォルトの名無しさん
12/02/19 01:26:39.83
>>512
回答ありがとうございます．
テクスチャメモリを使った処理の場合，マルチGPUじゃないと出来ない
とういことでしょうか？
514:デフォルトの名無しさん
12/02/19 01:44:29.68
公式マニュアル見れば分かること
試してみれば分かること
これを聞く人が多すぎる
515:デフォルトの名無しさん
12/02/19 01:53:02.78
4-way SLI でも、ホスト上の4スレッドで利用する場合、
各々1GPUを独立して割り当てるため、4GPU間で通信して仕事をこなす本来的なSLIにはならない、って認識でイイんですよね？
（SLIコネクタがなくても動く？）
516:デフォルトの名無しさん
12/02/19 02:32:56.76
>>511
テクスチャを5個宣言じゃダメなの？
517:デフォルトの名無しさん
12/02/19 02:47:03.16
>>516
回答ありがとうございます．
CPU側でスレッドを立てる段階で決めたいので，
できることなら，可変にしたいと考えています．
514さんは「これを聞く人が多すぎる 」と仰っていますが，
そもそも，このようなことをシングルGPUでやろうとするのが間違いなのでしょうか？
518:デフォルトの名無しさん
12/02/19 02:59:57.92
いやいや、まず試せば分かることじゃん。
その労力を2chに押しつけるのは、今この瞬間は楽では良いかもしれないけど、
自分の成長には繋がらないよ。
519:デフォルトの名無しさん
12/02/19 05:31:26.29
>>513
OS側のスレッド何本あろうとGPUが1つしかないんだから
順番に処理されるだけだろってことでしょ
520:デフォルトの名無しさん
12/02/19 05:43:18.09
concurrent kernel execution…
521:509
12/02/19 07:12:00.07
>>508
失礼しました
インライン関数化で無事できました
ありがとうございました
522:デフォルトの名無しさん
12/02/19 11:42:52.58
>>518
心遣い感謝します．
>>519
>>520
回答ありがとうございます．
やっぱり，そうなんですね．
「concurrent kernel execution」ついて調べて確信しました．
回答・アドバイスして下さった皆様，ありがとうございました．
おかげで解決しました．
523:デフォルトの名無しさん
12/02/19 11:49:26.29
次に同じ疑問を持たれた方のために，
参考資料のアドレスを貼っておきます．
URLﾘﾝｸ(www.nvidia.co.jp)
上記アドレスにあるpdfの「コンカレントカーネル実行」に書かれていることが，参考になるかと思います．
524:デフォルトの名無しさん
12/02/19 16:33:50.20
誰が質問者なのかわからないので謎な流れだが、とりあえずグレースケール化にテクスチャは不要だと思う今日この頃。
525:デフォルトの名無しさん
12/02/19 17:00:58.89
NPPで一発解決だよな
526:デフォルトの名無しさん
12/02/20 22:07:59.44
atomicCASのCASってどういう意味ですか？？
527:デフォルトの名無しさん
12/02/20 22:24:51.56
Compare And Swap
528:526
12/02/20 22:28:38.87
>>527
ありがとうございました！
529:デフォルトの名無しさん
12/03/07 03:16:50.68
ｃｕｄａの本には2次元配列より1次元の方が速いってあったけど実際どのくらい違うんですか？
また、コピーとgpu上での演算の両方で遅くなるんですか？
530:デフォルトの名無しさん
12/03/07 03:39:04.89
なんでこのスレ突然止まってたの？
卒論修論シーズンが終わったから？
531:デフォルトの名無しさん
12/03/07 06:34:38.99
>>529
一行目: 速い
二行目: 遅くなる
馬鹿?
532:デフォルトの名無しさん
12/03/07 16:56:25.77
学会はこれからだというのに
533:デフォルトの名無しさん
12/03/07 17:15:29.75
> ｃｕｄａの本には2次元配列より1次元の方が速いってあったけど実際どのくらい違うんですか？
メモリーのアクセス時間の早さのこと？
ちなみに２次元であっても，例えば f[i, j] で，先にｊを変化させるか，i を変化させるのかで
アクセス時間が違ったと思う．
要はメモリーが並んでいる順番にアクセスするのが早いはず．どの位，早くなるかは知りませぬ♪
> また、コピーとgpu上での演算の両方で遅くなるんですか？
コピーとは？ CPUとGPU間のメモリーの転送のことかな？？
534:デフォルトの名無しさん
12/03/08 10:13:29.60
人いなくなったな
春になってGPUネタで研究する学生が増えるのを待つしかないか
535:デフォルトの名無しさん
12/03/08 11:44:54.65
二次元配列なんて存在しねぇ
って考えると楽なのに
536:デフォルトの名無しさん
12/03/08 19:50:20.17
昔のオカルト本では「謎の四次元」「四次元失踪」とかあって、
「四次元」とか謎めいた雰囲気を感じた。
今なら四次元配列とか当たり前だｗ
537:デフォルトの名無しさん
12/03/09 01:08:22.32
>>536
そうそうｗ
四次元の神秘性が薄らいじゃうｗ
538: ◆QZaw55cn4c
12/03/09 22:10:56.52
ハミルトンの四元数というのがあって近年計算機工学に応用されるようになって云々かんぬん
539:デフォルトの名無しさん
12/03/09 22:22:34.47
発見自体は1800年代だったっけ？
たしか橋の上かなんかで思いついたとかｗ
540:デフォルトの名無しさん
12/03/11 11:39:47.33
HLSLとどう違うの？
541:デフォルトの名無しさん
12/03/11 11:49:19.59
>>540
C言語っぽくなってる。
542:デフォルトの名無しさん
12/03/11 14:43:17.16
四次元配列と四次元ベクトルは別物だろ
後者は要素数4の一次元配列
543:デフォルトの名無しさん
12/03/11 17:04:08.60
>>540
HLSLはShader Languageなんで整数とか扱えなかったと思う。
あとステップ数も制限があったような。
544:デフォルトの名無しさん
12/03/11 17:26:35.76
Compute Shaderを記述する言語も
HLSLじゃ無かったっけ？
リソースベースのHLSLと、ポインタ・配列ベースのCUDA
545:デフォルトの名無しさん
12/03/11 17:38:23.47
リソースベースって言い方、分かり易いね。
HLSLはまさにそんな感じだ。
546:デフォルトの名無しさん
12/03/12 13:07:18.89
>>541>>543
サンクス
自由度がまして打ちやすくなってるんだな
547:デフォルトの名無しさん
12/03/12 13:35:35.95
>>544>>545
サンクス
548:デフォルトの名無しさん
12/03/15 21:09:26.09
誰かいますかね、三つほど質問が。
■__syncthreads()だけでは共有メモリへの書き込みを保証できない？
１つのスレッドがグローバルメモリから共有メモリにデータを書き込み、
その後全てのスレッドがそのデータを使用して計算を行うような場合、
書き込み後に__syncthreads()だけではなく__threadfence_block()も必要なのでしょうか?
青木本には__threadfence_block()について特に言及ありませんでしたが・・・。
■ブロック内の全スレッドからの同一グローバルメモリへのアクセス
ブロック内で共通で使用する構造体などをグローバル→共有メモリに移す場合
全スレッドで行うよりもやはり
if(threadIdx.x==0)・・・
のようにした方が良いでしょうか?
■カーネル内でのreturn文の使用悪影響あるか
スレッドごとに計算を行うか判定をする場合、if文で囲っている例をよく見ますが
これは
if(条件)return;
と書いてはいけないのでしょうか？
上のように書いてもとりあえず計算は流れたのですが何か悪影響はあるでしょうか？
549:デフォルトの名無しさん
12/03/15 21:25:56.04
>>548
・取り敢えずsyncしか使ってないけど問題になったことはない。
・全スレッドから共有メモリへの書き込みを行なうのは多分遅くなるんじゃないかな?
・どちらで書いても同じこと。普通のCPUのような分岐とは違うことを判っていればOK。
550:548
12/03/15 22:39:42.94
>>549
ありがとうございます!!
一個目の_syncthreads()と__threadfence_block()の件ですが、
syncだけだと今日うまくいかなかったもので。
ただ他のバグの影響なども考えられるのでもうちょっと調べてみます。
551:デフォルトの名無しさん
12/03/16 04:06:36.41
>>548
・__syncthreads()は__threadfence_block()相当の処理を
含んでいた気がするけど気のせいかも。
・全スレッドで同じメモリにアクセスするのはたとえfermiでも遅くなるはず。
・カーネル内部で_syncthreads()使う必要があるなら
returnは使っちゃ駄目だろう。
552:デフォルトの名無しさん
12/03/16 20:10:45.37
いまだにアプリ開発環境すらまともに構築できてない・・・
visual studio 2008でやろうと思って一応ビルドは通ったけど
実行するとまずcutil32.dllがありませんって出た。
次にcutil.dllをデバッグ.exeと同じフォルダに置き実行！！
CUDA version is insufficient for CUDART version.
ってなる・・・orz
まずなにからはじめるべきですか？
553:デフォルトの名無しさん
12/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:デフォルトの名無しさん
12/03/16 20:51:20.83
GPUドライバのアップデート
555:デフォルトの名無しさん
12/03/17 01:39:06.58
>>554
ありがとうございました！！！！！！！！！！
動いた！！！！！！！！
556:デフォルトの名無しさん
12/03/17 08:54:13.85
おおｗ
よかったな！
557:デフォルトの名無しさん
12/03/17 12:49:24.04
>LINK : /LTCG が指定されましたが、コードの生成は必要ありません。リンク コマンド ラインから /LTCG を削除し、リンカの性能を改善してください。
と表示されるのですがリンク コマンド ラインは固定されて編集できません。
解決方法はありますか？
558:デフォルトの名無しさん
12/03/17 13:37:04.85
>>577
補足：
開発環境はVisualStudio2008
cuda ver 2.3
559:デフォルトの名無しさん
12/03/17 14:59:10.69
windowsを窓から投げ捨てろ
560:509
12/03/17 15:16:23.02
そんなことして道歩いてる人の頭に当たっちゃったら大変ですよ
561:デフォルトの名無しさん
12/03/17 15:44:28.84
角に当たったら痛そうだもんね・・・
562:デフォルトの名無しさん
12/03/17 18:01:09.23
>>557
リンカ -> 最適化 -> リンク時のコード生成 (/LTGG)
C/C++ -> 最適化 -> プログラム全体の最適化 (/GL)
Visual Studio 2008 の使い方なのでスレが違うかも。
563:デフォルトの名無しさん
12/03/18 15:04:39.12
>>562
ありがとうございます。
CUDA-Zの実行結果はどのように見たらいいですか？
日本のサイトが全然ないです。
564:デフォルトの名無しさん
12/03/18 15:07:18.71
>>563です
すみません。解決しました。
565:デフォルトの名無しさん
12/03/21 21:06:39.97
コンスタントメモリキャッシュへのアクセスはバンクコンフリクトとかないんでしょうか？？
566:デフォルトの名無しさん
12/03/21 22:20:52.26
>>565
そりゃキャッシュはバンクになってないからねー
567:565
12/03/21 22:44:39.32
>>566
おお、やっぱり。
できるだけコンスタントメモリ使うようにしまつ。
568:デフォルトの名無しさん
12/03/21 23:02:34.40
アドレスが静的に解決できないというのが前提だけど
16ポートのSRAMなんてコスト的に不可能だからマルチバンク以外無いんじゃないの？
569:デフォルトの名無しさん
12/03/22 00:27:08.40
Fermi以前はコンスタントメモリ使う意味あったけど、
Fermi以降はL2キャッシュとあんまり変わらない印象
570:デフォルトの名無しさん
12/03/22 22:48:01.36
GTX680が発表されたけど、CUDA的には好ましくない方向の進化が多い。。
571:デフォルトの名無しさん
12/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:デフォルトの名無しさん
12/03/22 23:06:16.47
GK104はミドルレンジだからGK110は全体的に上回ってくるでしょ
573:デフォルトの名無しさん
12/03/22 23:16:01.94
kepler誕生おめ！
.o゜*。o
／⌒ヽ*゜*
∧_∧ ／ヽ ）｡*o ｯﾊﾟ
(・ω・)丿゛￣￣' ゜
. ノ/ /
ノ￣ゝ
574:デフォルトの名無しさん
12/03/22 23:21:31.67
Keplerｷﾀ━━━(ﾟ∀ﾟ)━━━ !!!!!
575:デフォルトの名無しさん
12/03/22 23:22:01.12
gen3じゃないんだっけ？
576:デフォルトの名無しさん
12/03/22 23:42:00.14
>>570
もともとはミドルレンジでグラフィック向けだったから仕方ない気もする。
予想以上にグラフィック方面に舵を切ったという感はあるけど。
このままグラフィック向けとGPGPU向けで大きく分かれていくのではないかという心配はあるかな。
577:デフォルトの名無しさん
12/03/22 23:53:32.45
1SM = 192コアか。おっそろしいなあ。
578:デフォルトの名無しさん
12/03/23 00:00:34.24
nVidia始まったな。
579:デフォルトの名無しさん
12/03/23 00:17:06.29
>>577
warp の扱いどうなるんかな。。。
580:デフォルトの名無しさん
12/03/23 00:54:42.91
>>579
URLﾘﾝｸ(pc.watch.impress.co.jp)
> 32スレッドのWARPに同じ命令を実行する、この基本は、Keplerでも変わっていない。
らしいから、変わらないんじゃないかな。
GF104/114のSMには48コアと2ワープスケジューラ、4ワープディスパッチャで
GK104のSMXには192コアと4ワープスケジューラ、8ワープディスパッチャになっている。
その上レジスタ数は倍、L1キャッシュ/シェアードメモリはそのままってことは
GF104/114よりさらにピーキーになっているのかな？
581:デフォルトの名無しさん
12/03/23 01:01:32.84
>>580
あれ？
コンスタントキャッシュって無くなった？
L1/L2キャッシュがその役割を担ってる？
ということはFermiからか･･･
582:デフォルトの名無しさん
12/03/23 01:06:15.11
48コアが192コアになったのに
レジスタは2倍、
共有メモリは据え置き。
どーすんだこれ。。
583:デフォルトの名無しさん
12/03/23 01:54:37.82
レジスタ足りんくなりそうな。
584:デフォルトの名無しさん
12/03/23 04:09:37.63
Keplerはクロック落としてパイプラインを浅くする設計
演算器のレイテンシが小さくなるならレジスタの消費量は変わらない
Fermiの18cycleは頭おかしすぎた
これが例えば6cycleにになればレイテンシ隠蔽に必要なスレッド数が1/3になるから問題ない
585:デフォルトの名無しさん
12/03/23 04:14:42.78
x86 CPUと同じ道を辿ってるのか
586:デフォルトの名無しさん
12/03/23 15:07:43.87
誰か２６次元計算してくれ、1000ｺｱくらいじゃマジに足らんぞｗ
587:デフォルトの名無しさん
12/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:デフォルトの名無しさん
12/03/23 17:47:27.55
こんにちは
国際暗号学会のプレプリントサーバにこんな論文があがってました
Usable assembly language for GPUs: a success story
Daniel J. Bernstein, et. al.
URLﾘﾝｸ(eprint.iacr.org)
GPUのことはさっぱりわかりませんが、なにかこのスレの足しにでもなれば幸いです
それでは