- 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/
- 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を使ってグレースケールへの変換処理をしたいのですが，
テクスチャメモリをグローバルで宣言するとおかしなことになるきがするんですが．
どなたか教えて頂けないでしょうか？
宜しくお願い致します．
|
 |
