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

このスレッドは、他のスレッドでは書き込めない超低レベル、
もしくは質問者自身何が何だが分からない質問を勇気を持って書き込むスレッドです。
CUDA使いが優しくコメントを返しますが、
お礼はCUDAの布教と初心者の救済をお願いします。

CUDA・HomePage
http://www.nvidia.com/cuda

関連スレ
GPUで汎用コンピューティングを行うスレ
http://pc11.2ch.net/test/read.cgi/tech/1167989627/
GPGPU#3
http://pc12.2ch.net/test/read.cgi/tech/1237630694/

関連サイト
CUDA Zone
http://www.nvidia.co.jp/object/cuda_home_jp.html
フォーラム
http://forum.nvidia.co.jp/EokpControl?&event=TE0001

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

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

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

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

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

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

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

KNOPPIX for CUDA
http://www.yasuoka.mech.keio.ac.jp/cuda/

前スレ

【GPGPU】くだすれCUDAスレ【NVIDIA】
http://pc12.2ch.net/test/read.cgi/tech/1206152032/

忘れてた・・・。

pert2という時点で>>1の頭が超低レベルであることがわかっちゃうけどいちおういっておこう乙

>>1さんの「超低レベル」というお言葉に甘えて
インストール以前の質問なんですが、

普通にPCを使ってるときのグラフィックボード用のドライバと
CUDAドライバって別物ですよね？

http://www.a2art.com/?p=399
これの「ドライバは特に入れなくてok」ってどゆこと？

>>5
リンク先を読まずに当て推量で回答。

リビジョンの違いこそあれ、原則同じもの。
勿論、対応さえしていればどちらを使っても同じ。
逆に、NVIDIAのドライバが全く入っていないのなら入れる必要あり。

>>5
取り合えず最新のドライバいれて、CUDAツールキットを入れればおｋ
リンク読んでねえけど

レスサンクスです

具体的に言うとですね。
GeForce 9500M GT の入ったモニタ一体型VAIOがありまして、
それでCUDAってみようと思ったのですが、
CUDAドライバがうまくインストールできなくて、
とりあえず、ツールキットとSDK入れて、
bandwidthtestしてみたら、
認識できるGPUがないとかなんとか（今手元にないのでスマソ）
言われたんですよね。
やっぱCUDAドライバ入れなきゃいけないのかな。

>>8
モバイル版のドライバじゃないとけないんじゃね？

Mってのがモバイルみたいですね
ノートじゃないから、きづかなかった

ちょっといろいろやってみます
ありがと

画像処理なんかのサンプルを見ると
メモリ転送を行っていないのにカーネル呼び出しの引数になり
GPU側で使われている変数があるのですが
どのメモリに格納されているのでしょうか

filter_kernel<<<grid, block>>>(d_result, d_data, width);
こんな感じでwidthのメモリ転送を明示的に行っていないのに使用されています

言葉じゃよく判らん。 もちっと例くれ。
あるいは誰でもダウンロードできるサンプルで説明してくれ

まぁ，
__device__
__constant__
__shared__
のいずれかの修飾子の変数が予め確保されてるんだろなー･･･

>>11
配列でない引数はレジスタに入るんじゃないの

>>12
http://developer.download.nvidia.com/compute/cuda/sdk/website/samples.html
ここのboxfilerというサンプルに載っています
boxfilter_x、boxfilter_yというカーネルを呼び出しているのですが、width,height,radiusが
メモリ転送されていないにもかかわらず引数として使われています

>>14
レジスタでしたか
希少なのでコンスタントメモリあたりで代用したほうがよさそうですね

引数は勝手に転送される
格納場所はLocal Memory

>>16
ありがとうございます
助かりました

相変わらずCUDAれません(;o;

ツールキットとSDKはちゃんとインストールできているようです。

CUDA2.2を入れようとしているのですが、
http://www.nvidia.co.jp/object/cuda_get_jp.html
ここの
Windows Vista 64ビット 185.85用 NVIDIA ドライバ
も
Windows Vista 64ビット 185.85用 NVIDIA ノートブックドライバ
もはいりません。
一見インストールうまくいくのですが、再起動してデスクトップ右クリックから
NVIDIAコントロールパネル見ても、相変わらずForceWareのバージョンは176.76のままです。

nvcc -Vはうごきますが、bandwidthTestをすると
no CUDA-capable device is availableと言われます

http://www.nvidia.co.jp/content/DriverDownload-March2009/confirmation.php?url=/Windows/185.85/185.85_notebook_winvista_64bit_whql.exe&lang=jp&type=GeForce%20M
ここの一番下にVAIOは対象外って書いてあるから、VAIOじゃだめなのかな、、、

それかもね。ご愁傷様。

どんどんダウングレードして、2.0用のCUDAドライバ入れたら入った。
そしたら、今度はなぜか2.2のツールキットとSDKがアンインストールできないヽ(`Д´)ﾉ
エラー6003でぐぐっていろいろやったけど、できない。
アンインスコしないで、2.0インストールしたらまずいですよねぇ

連投スマソ

アンインストールしなくても医院で内科医。

>>21
あなたを信じてインスコしたら、
2.0のインストーラーが自動的に2.2をアンインスコしてくれました。
無事2.0環境整いました。
どうもでした(^^

多倍長整数使いたいときは自分で作るしかないん？

cudaに多倍長整数ライブラリは付属していないので、そういうことになりますね。

for文のネストにどこまで耐えられるのかがわからない＼(^o^)／ﾑﾘﾎﾟ

CUDAのスタックって案外少ない？
ハードウェアによってまちまち？

束胃腸やるのはいいが、そもそもキャリーアップフラグとかあったっけ

>>25
スタックはないです

>>25
template使った再帰もどきとかやらない限り
心配することはないかと思われる。

CUDAカーネルのルーチンはインラインとして扱われるんで、スタックなし。

んでもって、レジスタが使用できないデータタイプだったり、レジスタが枯渇してたりすると、ローカルメモリ(デバイスメモリ)を使用しだすので、パフォーマンスが一気に低下するんで心配してください。

25ですが，みなさんありがとうございます．
最近になってサンプルコードをじっくり読み始めましたが，
自分のやろうとしてるくらいのことは素直に出来そうで安心です．
んでもって，もう少しアーキテクチャを知らないといかんな．．．

CPU側のプログラミングにインテルのiccを使いたいのだけど、どこかにやり方は載っていませんでしょうか？
OSはWindowsXPです。

>>31
ソース分けてオブジェクトリンクすれば？

>>31
つか勝手に使えばいいじゃん

CUDA2.0入れてVS2005でtemplateビルドしたら
LINK : fatal error LNK1181: 入力ファイル 'cutil32D.lib' を開けません。
って言われます
VISTA64版なので、cutil64D.libはあっても32はないのですが、、
どこを修正すればいいのかもわかりません。
ちなみにVSも使ったことない初心者です
どうしたらよいでしょう？

>>34
上のツールバーの真ん中の方にWin32となっているところをx64に変更。
そもそも64bit用のツールをインストールしているかどうか・・・・。

>>35
レスサンクス

x64に変更したら、ビルドがスキップされました、、、
プログラマでもないのに、突然仕事でCUDAやってる俺涙目
聞ける人もほとんどいない

ちなみにVSのバージョンは
Microsoft Visual Studio 2005
Version 8.0.50727.867 (vsvista.050727-8600)
Microsoft .NET Framework
Version 2.0.50727 SP2
でつ

スレ違い、と言うか人生の道を踏み外したな…。

>>37
せめて正しいスレに導いておくんなまし
とりあえずコンパイルできないことには段ボーラーです。
踏み外す以前に舗装道路なんて歩いたことございません。

メモ、とりあえず報告しておきます。
OS：Windows Vista Home Premium 64bit
IDE：Visual C++ 2008 Express Edition
ドライバ：cudadriver_2.3_winvista_64_190.38_general
ツールキット：cudatoolkit_2.3_win_64
SDK：cudasdk_2.3_win_64
ぐらぼ：GeForce 9600GT
で、64bitビルドを行う。
まずはttp://www.sharkpp.net/blog/2009/04/26/visual-c-2008-express-edition-enable-64-program.html
で64bitビルド出来るようにする。
このまま進めてくと、
nvcc fatal : Visual Studio configuration file '(null)' could not be found
みたいなエラーが出てくると思う。
これを解決する方法が→ttp://forums.nvidia.com/index.php?showtopic=98319
vcvarsamd64.batは必ず、
C:\Program Files (x86)\Microsoft Visual Studio 9.0\VC\bin\amd64\vcvarsamd64.bat
か、または
nvcc fatal : Visual Studio configuration file '(null)' could not be found....のところで実行。
後は、nvccのコマンドラインとかちゃんと設定しておけばいける！はず・・・

NVIDIAが目指す究極のプロセッサへと進む「Fermi」アーキテクチャ
http://pc.watch.impress.co.jp/docs/column/kaigai/20091023_323529.html

>>39　の意味はわからないけど、64ビットで使うときは、コンパイルとリンクのオプションに
-ccbin "C:\Program Files (x86)\Microsoft Visual Studio 9.0\VC\bin"
をつければたいていうまくいく。

>>36
その仕事、うちの会社にくれ。あんたがやるより速いものが早くできるぞ。

> VSも使ったことない初心者です
でいきなりCUDAか…　なんというかどういう会社だと…

「まず32bitで試させてくださいお願いします」
で後から64bitのことを考えたほうがいいんじゃね？

>>36　今、
Windows 7 Professional 英語版
VisualStudio 2008 Pro(C++の64bitツール入れる指定)
CUDA Driver　(190.38)
CUDA Toolkit
CUDA SDK
入れて、
OceanFFTのoceanFFT_vc90.slnを開いてx64のDebugでリビルドして
さくっと動いたなぁ、と思ったところですよ。
buildのログには確かに-ccbin "C:\Program Files (x86)\Microsoft Visual Studio 9.0\VC\bin"　と入ってる。

次はCUDA.NETと、PyCUDA＋VisualPythonで遊ぼうかなと

>>44
追記...　Win7は64bitです。
ToolKitとSDKは2.3です。

ボリュームテクスチャをルックアップテーブルとして使いたいのだが
cuda kernel上で、tex3D関数を使うと異常に処理時間が増加してしまいます。

case 1
ret=tex3D(tex,yy[0]/256.0-1, yy[1]/256.0-1, yy[2]/10.0-1);

case 2
ret=tex3D(0.0,0.0,0.0);

で実行時間に１０倍以上の差がついてしまうんだがなぜだろう。
メモリ上の配置でかなりパフォーマンスが落ちるということはありますか？

>>46
テクスチャキャッシュにヒットするかどうかでパフォーマンスは大きく変わります。

win7x64にCUDA_VS_Wizard_W64.2.0.1入れてもテンプレに表示されるようになる
けどエラーでプロジェクトが生成できない・・・
Vistax64の時はこれで一発だったんだが・・・
7でも32bitのSDK入れてCUDA_VS_Wizard_W32使ったらすんなり入ったが

>>48
使えると便利だから使えるようになってほしいな。
フォーラム見に行ったら10/31までmergeでdownだよと書いてあったorz
復活したらエラーメッセージをポストしてください

誰かtesla使ってる人いる？
tesla一枚挿しのマシンに、リモートからwinデフォのtelnet使って動かそうとすると、
tesla自体認識しないんだが・・これどうやって使うの？

認識してすらいない、っと。

>>50
TeslaがささっているマシンがWindowsなのかどうか読み取れませんが、
リモートからWindowsマシンのTesla使うのであればVNCを使うことが多いですね。
リモートからの利用中心であればLinuxにしてしまった方が・・・。

>>47
詳しく！ お願いします 助かります

逆に、デバイスから高速かつランダムな位置に参照したいときはどのメモリ使えばいい？

>>54
サイズも示さないで高速かつランダムって、馬鹿なの? 間抜けなの?

レジスタに乗るなら1クロックでアクセスできる。
共有メモリに乗るなら4クロックでアクセスできる。
グローバルメモリに乗るならランダムと言う前提からcoalscedであることは期待できないから最長数百クロック掛かる。

>>50
んじゃ私はLinuxサーバ機であるとして。
一部のメーカのサーバ機(BIOS依存かマザボ依存かは不明)ではTeslaをGPUとして認識する。
その場合は(オンボのGPUが切り離されて涙目になることはあっても)なんとかなる。
処が、一部のサーバ機ではGPUとして認識できないので自前でデバイスを叩かないと認識しない。

確かにそうだな サイズは20メガバイトくらいだ

8800GTXがあるんだが四倍率早く処理するにはなにがある？

ちょっと日本語でＯＫ

8800GTXを4本さす

クロックを４倍に

>>52
レスありがとう。
winでGTX280とTeslaさした状態だと、両方のデバイス認識してて、GTX280外してTesla単体にすると、
Teslaを認識しなくなるんだけど、原因わかる人いる？ドライバではないと思うけど・・
ちなみに、winデフォのリモートデスクトップで確認しました。

DeviceQueryでどうなるのかはっきりして欲しい

>>62
Tesla単体のときって画面は何に表示されるのでしょうか？
TeslaはGPUを搭載していますが、ビデオカードではないことに注意してください。

Windowsの制約で複数ベンダーのビデオドライバを同時にロードできないので
オンボードにATIやIntel製のビデオチップとかがあると駄目です。
どうしてもその構成にしたいのであればLinuxに。

というかWindowsのバージョンによるけどまず無理です。

馬鹿みたいに高いのに、すごい仕様だな

>64
Win7からはWDDM1.1対応ドライバであれば普通に混載できる

なんかメモリアロケートのアルゴリズムいじったら五倍以上速くなった、、

あやうく三枚発注するとこだったよ。よかった

CPUの100倍速い、になるまで満足しちゃいかん

将来性は銅なのよ?

GeForce 8600 GT使って、3日目の初心者です。

素朴な疑問なのですが、GeForce 8600 GTには512MBのメモリが実装されているのですが、
cudaMalloc を行って、alloc領域が512MBを超えた場合には、どの様になるのでしょうか？
ご存知の方がいらっしゃれば、事象や回避策を教えてください。m(_ _)m

>>70
素朴な通り取れない。
エラーが返ってくるはず。

>>67
これはCUDAの話？

>>71

ありがとうございます。
まだ、某サイトにあった手順でHellow…が出せた段階です。

もしや実装メモリ量を超えると、HDDとの間でガリガリswapでもするのかなぁ?っと思った次第です。

うっ…
w 綴り違うし orz...

>>73
大丈夫、swapなんてするはずもない。そもそも512MiB全部使えるわけでさえない。
その世代で512MiB搭載のボードなら、実際に使えるのは高々511.69MiBだけだ。
しかも、そのボードを実際に描画にも使っているならそこから更に画面解像度に依存した分使える量が減る。

逆に、目一杯cudaMalloc()で確保した状態で画面解像度を上げたりすると面白いことになる。

>>73
つかQueryつかってデバイスメモリの上限チェックしないと

>>63,64
レスありがとう。どうやら、winではディスプレイ出力できるグラボを最低ひとつ積んでないと、
cuda使えないみたいです。つまりTesla単体では動かない。。deviceQueryしても、
there is no device supporting cudaしか返ってきませんでした。倍精度の数値計算したくて
Teslaシングルで動かしたいときは、nvidiaの公式が推奨してるようにLinuxでやるのがベスト
みたいです。いろいろとありがとうございました〜

snow leopardじゃつかえないの？

>>74
皮肉のつもりじゃなかったのか。

いまだにtemplateのコンパイルもできません。
32bitコンパイルしようとするとlibがないといわれ、
64bitコンパイルしようとするとスキップされる。

人に聞いたら、ソースが32bit用なんじゃないかって言われた。
ちゃんと64bit用SDKインスコしたつもりなんだけど、
64bit用のソースでなかったりするのでしょうか？

>>80
リンカが言うようにlibのファイルがないんでしょ？
パスが通ってるか、ファイルがあるかどうかくらい自分でなんとかしなよ。

>>81
レスサンクスです。
でも、>>34に書いたとおり、32bit用のlibはもともとないんです。

まったく初心者なので、勘違いしてたらすんません。

>>82
それなら、リンカに指定するライブラリをcutil64D.libに変えたらいいじゃん。

>>80
その書き込みから推測するに、出力先にtemplate.exeが存在しているため、スキップされると思われます。
SDKのtemplateをそのまま開いて64bitコンパイルしたいのなら、
出力先のtemplate.exeを削除してからビルドすればスキップされなくなるはずです。
C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK\C\bin\win64\Release\template.exe
私の環境ではこの場所。SDKにはすでにコンパイルされた.exeが入っているはずです。確認してみて下さい。

64bitSDKには32bitのlibが入っていないんじゃないかな？詳しく確認はしていないけれどね。

>>83,84
ありがとっす〜
試してみます！

>>75､76

どうもです。お返事遅れました。

先週はXPでやってましたが、今日からはFedoraに入れ直し実施してます。


結果ですが、画面が点滅したり、真っ黒になったり、となりました。
復旧には、sshで入ってプロセス消してもだめで、結局rebootで…。
ちゃんとsizeみて使う様にします。

m(_ _)m

済みません。連投です。

>>79

はぃ、
天然ってより、単に英語は苦手です。

exe消してもだめでした、、、

つか、なんかソリューションエクスプローラーのところで
template.cuを右クリック→プロパティにした時点で、
「操作を完了できませんでした。エラーを特定できません。」
ってなる(´･ω･`)だめぽ

リンカへの指定はちょっと勉強してみます。

どうにもうまくいかないので質問なのですが

short intでGPUにmemcpyしたデータをFFTしたいのですが、
short intのデータをfloatにキャストするには
どうしたらいいのでしょうか？

int→floatやfloat→intはプログラミングガイドに載っていた通りできたのですが、
short intだとうまくいかないのは何故だろう…？

>>89
何を試したのかよく分かりませんが、short2型をfloat2型に変換する感じがよさそうな。

プログラミングガイド読んでたらﾊｧﾊｧ(´д`*)してきた
なにこの気持ち（*´д`*）ﾊｧﾊｧ

>>90
おかげさまで出来ました。ほんとにありがとう

ちなみに試したのは
__int_as_float(int)のような形式のキャストと
__int2float_rn(int)のような形式のものです。
どちらもshortをキャストしたり、shortにキャストすることが出来ませんでした。

GPUに型変換するのは遅いと思うのだが。

その通りだな、そこで相当なロスが出ると思う

voidでコピーして、floatのポインタで計算すれば良いんじゃねーの?
何か勘違いしてる?

>>91
興奮してるなぁｗ

int -> floatのキャストなら1命令だからむしろ速いんじゃないのか
アライメントは考慮したほうがいいけど

いやだから、gpuでshort -> intはダメだろ。
intで渡しておいて、int -> floatはイントリンシックで。

別にshortで渡しても問題ないかと。
GPU内でshort -> intの変換をしてから int -> floatの変換が起こりますが、
そんなに頻繁に変換をするわけでなければメモリ転送量の削減効果の方が大きいでしょう。
＃32-bitアクセスにするためにshort2型をお薦め。

GPGPUのハードのいいベンダーだか機種だかオススメある？
予算は潤沢にあるとして

NVIDIAのTesla。つーか、Teslaのラック筐体マジお勧め。
漏れなくNVIDIAのサポートがついてくるから。

>>101
CUDAしたいからtesla　c1060を搭載した機種を買うのは
当然なんだけど、teslaを組み込んだ
ハード全体としてのオススメを知りたいのでした。
ようするにフラクティカだとかELSAだとか爆速だとか、
どこがいいのよっって話です。
teslaのラック筐体ってなに？
タワー型じゃなくてユニット型がいいってこと？
それともteslaの名を冠したサーバーがあるの？

すまんあんま詳しくないんでわからん

HPC用にTeslaC1060相当を4枚入れた、TeslaS1070っていうラック筐体があるですよ。
当然、PC筐体は別途必要。
PC本体なら、QuadroPlex2200S4ってーのがTeslaS1070にマザボをつけたような仕様だったかと。
ラック筐体じゃなければ、QuadroPlex2200D2がタワー型でGPU2枚挿しのPCになる。
NVIDIAの営業曰く、「GeForceなんてアキバ的発想はやめましょう」ということなので。
# GeForceでいいならELSA辺りがリファレンスボードをそのまま使うから安定しているけどね。
## つーか、QuadroPlex使うような予算があるならソフト開発受注したいぞw

>>103
ありがと〜
調べてみる

http://www.supermicro.com/products/system/4U/7046/SYS-7046GT-TRF.cfm?GPU=TC4
これ、自分はアメリカで買った。
国内の代理店でも売ってるとこあるよ。

>105のPCすげぇ。TeslaC1060が4台も載っている。なのにVGAはMatrox。

おまえらどんな仕事してんだよ

自宅の治安を守る仕事

>>106
だってサーバーだから。
すんごくうるさいよこのマシン。標準でリモート管理機能が付いているから、
リモートから電源のON/OFFやキーボード、マウス、VGA、IDEポートの
ネットワークリダイレクトができる。
だから管理用IPアドレスだけ設定してサーバールームに入れっぱなし。

>>105
国内のページ見つけたけど
http://www.able.across.or.jp/catalog/product_info.php?products_id=2919
スペックが書いてあるようで書いてないような、、、
この値段は一体なんの値段なんだろう

GeForceとQuadroってどう違うんですか？？

シールが違う。

>>93-99
色々参考になります。ありがとう

とりあえずそれぞれのやり方で試して処理時間比較してみます。

>>110
書いてあるじゃん。
7046GT-TRF-TC4はサーバーのベアボーン。
箱、電源、マザーボード、ドライブエンクロージャ、そしてTesla C1060 * 4だけ。
あとはCPU、メモリ、ハードディスクを買ってきて刺してやれば動くってこと。

>>112
それだけですかｗ
ありがとうございました m(_ _)m

>>109
自宅にサーバールームか。アニメみたいww

やはりワレワレはコストパフォーマンスの良いものをと考えて…
i7-920、12GB
ASUS　P6T7 WS　nForce200が二個乗り
とりあえず　GTX275あたりを二枚位かな
1000W電源
これでざっと20万コース？

randってつかえないのですか？

>>114
そっか単純にCPUもメモリもついてないのか
安いわけだ

サーバー用のメモリも安くなってきたから、サーバー用でいいかも。
Kingstonの4GB×3本で$350とかだ。

CUDA_SAFE_CALLって必ず必要なのですか？

> nForce200が二個乗り

おれの場合、O(n^2)だと、使うメモリ帯域<<計算量で、
x16で有る必要は無いなと思ったよ。

カーネルを実行する際に、引数で、変数を渡すことができますが、多数の数値を配列で渡したい場合、
配列をデバイスにコピーしてから、配列へのポインタを渡すしかないのでしょうか？
オーバーヘッドの少ない、数値の渡し方を差がしているのですが・・・

引数で配列のポインタを渡したって、どっちみちホストメモリ→デバイスメモリの
コピーは要るし？

すみません、CULAについて質問です。

Red Hat Enterprise Linux 5.3(64bit)に
まず、CUDA3.2(ドライバ、ツールキット、SDK)をインストール。
(~/CUDA/ 以下。ツールキットは /CUDA/cuda/bin/nvccのようになっている)
その後、CUDA3.2上でCUBLASが動作することを確認しました。（Dgemmを利用）

後に、CULAがあることを知り昨日CULA Basic 1.0をダウンロードし
~/cula/　以下にインストールしました。
（~/cula/lib64/libcula.so）

インストール後表示される
export CULA_ROOT="/home/nakata/cula"
export CULA_INC_PATH="$CULA_ROOT/include"
export CULA_BIN_PATH_32="$CULA_ROOT/bin"
export CULA_BIN_PATH_64="$CULA_ROOT/bin64"
export CULA_LIB_PATH_32="$CULA_ROOT/lib"
export CULA_LIB_PATH_64="$CULA_ROOT/lib64"

以上を実行後、example/geqrf にある、Makefileを用いて
make build64 を実行すると、
/home/nakata/cula/lib64/libcula.so: undefined reference to `cublasDtrmv'
/home/nakata/cula/lib64/libcula.so: undefined reference to `cublasZswap'
/home/nakata/cula/lib64/libcula.so: undefined reference to `cublasZaxpy'
/home/nakata/cula/lib64/libcula.so: undefined reference to `cublasCtrmv'
/home/nakata/cula/lib64/libcula.so: undefined reference to `cublasZtrmv'
/home/nakata/cula/lib64/libcula.so: undefined reference to `cublasZcopy'
/home/nakata/cula/lib64/libcula.so: undefined reference to `cublasCgemv'
/home/nakata/cula/lib64/libcula.so: undefined reference to `cublasCtrmm'
/home/nakata/cula/lib64/libcula.so: undefined reference to `cublasCtrsm'
/home/nakata/cula/lib64/libcula.so: undefined reference to `cublasZtrmm'

リファレンスを調べてみると、例えば、cublasにDtrmvが無いように思えます。
これはどのようにすればいいのでしょうか。

en_curr_regionがないってエラー出るんで、調べてたらbreakやcontinueのところでエラー出てるっぽいんだけど、CUDAってbreakやcontinueとの相性悪いのだろうか？

分岐は罪

質問です

CUDAを使ったプログラムをCUDA toolkitをインストールせずに実行する方法ってありますか？

試した事はないけど、driverと実行ファイルがあれば出来るだろうな

よ〜し、パパ、CUDAを覚えて会社で活躍しちゃうぞ〜！

VCランタイムがあるとして、cudart.dllと、cutil32.dll　だけ必要(Win32)。

>>132
それって再配布できるの？

cudatoolkitのEULA、cudart.dll　はredistributableと書いてあるな
CUTILは、cutil.cppの頭に「何の保証もしねぇよ」って書いてある。CUDAとは別で単なるサポート用なのかな。
ていうかcutil32.dllは無くても動くように書けると思う。ソースみたらくだらないぞこれ
nVidiaのドライバ入れればCUDAも動く、ようにはまだなってないのかな。

>>134
cudart.dll はCUDAの上位互換を保つために、配布するんだそうだ。
ちなみに、cutil32 or 64 .lib は使わないほうが。。。

ともかく、

cudart's version <= cuda driver's version
の条件が必須

cutilはクソすぐる

いつの間にか団子が２ちゃんのスポンサーになってる・・・

【２ちゃんねる10週年】神秘の「乳団子」の秘密とは？
http://pc12.2ch.net/test/read.cgi/tech/9240911011/

Tesla C2050　$2,499　発表になったじゃない。倍精度500-630GFLOPS。
倍精度なら10倍速い、ってどかすか買うお金持ちな人居るんだろうなぁ。

今まで散々売られてもいないLarrabeeを盾にして現存のプロセッサと
それを擁護する人たちを馬鹿にしていた癖して
いざ完成が近づいてくると

ディスクリート版には期待していないｷﾘｯ

とか馬鹿にしてるよな。

おまえは誰と戦って(ry

書き込むスレ間違えたが分かっている人はいると思うから問題ない。

>>138
倍精度性能だけ見るとHD5870と同じ性能だというのは本当なのか？
単精度の1/2の性能でいいって言ってるけど
HD5870の単精度性能が数倍早いだけだという噂が

>>142
AMDはなにかを犠牲にしていると思うんだけど何を犠牲にしているの？

>>143
たとえばプロセッサエレメントはFP演算とデータの読み書きは排他実行だったりとか
それでなくとも各プロセッサエレメントにぶら下がってる5つの演算ユニットが平均2〜3程度しか稼動してないとか

ららびーは倍精度で１TFLOPSとか行くのかな。
Core i7で60GFLOPSくらいだっけ。

>>144
データ読み書きは別スレッドの分が並列に動くだろ。
メモリアクセス性能低いからALU命令の比率が余程高くないと
読み書き速度がネックになるけど。

初歩的な疑問なのですが
ttp://tech.ckme.co.jp/cuda_const.shtml
のコードのように
__constant__ float g_idata[10000];
とか書いた時，このconstantメモリ領域は
いつ確保されるのでしょうか？

宣言だけ書かれてあっても
host,device共にその領域を使うコードを実行しない限り，
領域確保は行われないと思っていいのでしょうか？

>>147
GPUにはそもそも領域確保と言う概念がありません。
定数メモリ領域も共有メモリ領域もただそこにあるだけなので、宣言はあくまでもコード上のものでしかありません。

一応タイムシェアリングできてるみたいだけどGPUのメモリ領域の退避はしないの？

CUDAでCPU、GPUを並列に動作させられますか？
CPU GPU 並列 CUDA　あたりで検索しても出てきません。
原理的にできそうな気もするのですが、
GPUで操作させている間はドライバを働かせているので無理なのでしょうか？

スレッドを使えばできますよ。GPUのプログラムは殆んどCPUに負荷をかけません。

>>150
GPUを起動したスレッドを待機状態のまま放置しておけば、他のスレッドでCPU資源を遣り繰りできます。
但し、SPの個数を超えるようなグリッドを食わせると途中で転送処理が入るので要注意。

CentOSで動きますか？

FermiでC言語が使えるってどーいうこと？

>>153
動きます。

fermi用のCコンパイラを用意しましたってことじゃねーの？

スレッド内部でレジスタがどのように使われているかわかりません
変数の数＝レジスタの数という認識をしていたのですが、そうではないようです
どなたかご教授ください

すいません、テクスチャメモリで疑問があります。
よろしければどなたか教えて頂けませんか？
コードは
http://gpu.fixstars.com/index.php/%E3%83%86%E3%82%AF%E3%82%B9%E3%83%81%E3%83%A3%E3%83%A6%E3%83%8B%E3%83%83%E3%83%88%E3%82%92%E4%BD%BF%E3%81%86
のテクスチャメモリを参照するていう項目のソースをとりあえず勉強がてらコンパイルしてみました。
配列に値を3ずつ入れていって、バインドして、それにtex1Dでアクセスするというだけのものです

13行目にインデクスに3.14fを用いている事から線形補間なしの場合、小数切捨てで3になり
配列[3]の値を読んで、9が返ってくるのは理解できたんですが、線形補間を有りにすると
9と12の間を補間してるはずなのに7.921875なんていう値が帰ってきています。不思議に思ってv2.3のPrograming guideを読んだところ、P.139に記述があり、どうも与えた値が-0.5されているようです。
ので、はじめから+0.5シフトしてやれば正しい値になりそうですが、そもそも何故-0.5されているのかがわかりません。

検索してもぜんぜん出てこないので当たり前な事なのかもしれませんがドツボにハマってしまってわかりません。よろしければどなたか教えて頂けませんか？長文失礼しました

158です。
上記質問で訂正があります。programming guideのP.139ではなくP.137でした。すいません。

>>158
テクスチャだからとしか言い様が無いような。
配列の[0]の値は0 〜 1が守備範囲で中心は0.5
...
配列の[i]の値はi 〜 i+1が守備範囲で中心はi+0.5
...
配列の[N-1]の値はN-1 〜 Nが守備範囲で中心はN-0.5

N個の値を0〜Nの範囲に均等にマッピングするためにこうなっている。

>>160
なるほど、こんな感じか。
buffer[2.5] = 6.0
buffer[3.0] = (buffer[3.5]-buffer[2.5])*0.5 + buffer[2.5] = 7.5
buffer[3.5] = 9.0
buffer[4.0] = (buffer[4.5]-buffer[3.5])*0.5 + buffer[3.5] = 10.5
buffer[4.5] = 12.0

buffer[3.141592653589] = (buffer[3.5]-buffer[2.5])*0.641592653589 + buffer[2.5] = 7.924777960767
あれ？ちょっと違う値だね

>>159,160
ありがとうございます。
なんとなく理由が理解できた気がします。
が、こちらでも計算しましたが
buffer[3.141592653589]=7.924777960767
となって件のサイトの値7.921875とは合わないですね・・・
buffer[9.80665]=27.91995となって27.925781ではないですね・・・
低精度の線形補間というのは有効数字3桁って事なんですかね？

リニアフィルタつきのテクスチャでは、
座標(0,0)てのは左上端のドットが持つ四角い領域の左上端を意味するからね。
ドットの真ん中は座標(0.5, 0.5)になる。
真ん中にそのドットの本来の色（値）がくる。

精度はシラネ。リファレンスマニュアルに何か書いてあると思うけど。

CUDA3.2上で動くTesla1070Sを使っているんだが
doubleでは、sqrtとか動作しない？

//#define DOUBLE double
#define DOUBLE float
__global__ void sqrtTest(DOUBLE *A)
{
int x=threadIdx.x;
(*(A+x))=(DOUBLE)sqrt(*(A+x));
}

みたいにして、実験したんだが、doubleにすると
値がそのまんま帰ってくるんだが　orz

>164
コンパイル時に　-arch=sm_13　はつけてる？

>>165
それだっ！動作しました。
cublasは倍精度で動いたんで、いろいろ調べたんですが。
ありがとうございます。

>>164
CUDAって3．2までいっているの？
つい最近2．3が出たような気がする。
ひょっとして2．3の間違い？

おそらく。
今は3.0のβだね。

newでメモリを確保するのは反則ですか？

>>169
どうやってnewで確保された領域をGPUに転送するの？

>>169
CPU側のメモリでしたら反則ではありません。但し、VCで使う場合は*.cuでnewしてもmsvcrtにリンクできません。

>>169
パフォーマンス求めるなら、論外です。

newでメモリ確保するのがだめならどうやってCPU側のメモリ確保するの?
全部静的確保?
それともcudaHostAllocを使えって話?

メモリ転送ってＳＳＥで高速化されたりするんかな？
だとしたらnewより専用で用意されたものを使ったほうがいいかもね

>>174
メモリ転送ってCPUメモリ間？それともHOSTｰGPU間？
前者だったら高速化はされるけど、後者はDMAでPCIEにダイレクトに転送されるだろうから、
SSEは関係ないんじゃない？

これは・・・

108 名前：デフォルトの名無しさん[sage] 投稿日：2008/07/14(月) 19:24:11
>>106
GPUはWarp単位で同じインストラクションが走るから、要は16人17脚みたいに考えればいい。
メモリアクセスを16人17脚によるパン喰い競争みたいに考えると、自分のパンが目の前にある状態がcoalesced。
そのとき、2,3人パンを喰う必要がなくてもスルーするだけだから問題ない。
処が、二人のパンが入れ違っていたらそこで入れ替える間、みんなが待たされることになるって感じ。
# 判っている人には判るけど、判っていない人には判らない説明だなw

>>107
共有メモリを使うかどうか違うだけだと思うけど。ptx出力させて較べてみたら?

109 名前：デフォルトの名無しさん[sage] 投稿日：2008/07/15(火) 00:26:16
>>108
その説明、いただいてもいいですか？

110 名前：デフォルトの名無しさん[sage] 投稿日：2008/07/15(火) 01:56:27
>>109
本にするならもっと書かせてくれw
Vipのwikiに載せるなら是非やってくれ
金取って講習するのに使うのなら分け前よこせw

対称行列になるものをGPUに送信したいのだが、うまい方法はない？
一般の場合は、GPU上では対称ではないとみたいのだが
送る際対称になる場合が結構あって、転送時間無駄だなあと。

>>175
これじゃないの？
http://gpu.fixstars.com/index.php/WriteCombine%E3%83%A1%E3%83%A2%E3%83%AA%E3%82%A2%E3%82%AF%E3%82%BB%E3%82%B9%E3%82%92%E9%AB%98%E9%80%9F%E5%8C%96%E3%81%99%E3%82%8B

このコードはいろいろ酷いからそのまま使えると思っちゃ駄目よ。
srcのアドレスはこの場合64バイトでアラインメントされてないといけない。
destのほうも最低限16バイトアライン

>>176
なかなか良い説明だなｗ

>>176
あと、競争するコースに分岐があったら、
その分岐に用のある人が２，３人だったとしても全員一応付き合わされた後
本流に戻るというか、そんな感じだな。

>>180,181
あれ、みんな読んでないのかな。
最近ようやくCUDA本がでたわけだけど、
まんまこの文章書いてあるんだよね。
青木先生乙

>>182
絵もついてるしなｗ

地球シミュレータを蹴散らし一躍脚光を浴びたというのに
中身スカスカなスレだなｗ

CUDAスパコンってそのときだけのものなんじゃない？
研究機関のスパコンは定期的に更新されるもので、更新したら前の代の
ソフトウェア資産はさっぱり使えなくなりましたじゃとても困るだろう。

まあG80〜GT200世代のコードも一応次のFermiでは動くし
NVIDIAが父さんしない限りはずっと続くんじゃない？

>>182
もうある程度把握しちゃってるからいまさら入門書買ってもとも
思ったけど、どうなんだろ。

>>185
IBMがHPC向けにはCELLやめてOpenCLに舵を切った的なことが書いてある。
ま、TheInquirerの記事は眉唾ではあるけど。
http://www.theinquirer.net/inquirer/news/1563659/cell-hpc-material

あ、ここCUDAスレだけどまあ似たようなもんということで。

×CellやめてOpenCL
○Cellベッタリのコード書くのやめてOpenCL

段階的には切り捨てることも考えられるが
いきなりOpenCLで他のデバイスとかいっても、資産が無いじゃん

はじめてのCUDAプログラミング
買った人いる？
どうだった？

図書館に頼んだら陳列は来月からと言われてしまった

本なんか読まなくても分かるだろ

日販は使えない会社だ。

>>190
思ったより安かったのでぽちってきた。
ASCII.techの特集も買ったけどアクセスの最適化あたりで苦戦中なレベルなので、
どうだった？とか言われても答えられないかもしれない。

>>184
これ、とんでもない誤報
浮動小数点演算を理解していないバカコミの馬鹿記事

科学技術立国　日本の恥を世界に晒したもの
「ふざけたやつがペンもつな、馬鹿野郎」だ。

浮動小数点数

GTX295と電源買った〜
さて、何に使おうw

っBOINC

>>176
おにいさん、もっと！

smokeParticles.exe の動作が 8600GTS より遅く感じるんだがなぜだー

>>200
２個同時使用に対応してないんだろ

>>201
それでも8600GTSよりは早くなるはずじゃない？
あまり詳しくないんで間違ってたらすまん。

>>202
言われて見ればそうだな、約２倍に性能アップしてるはずだし
いつからか分からんけど古いＣＵＤＡと最近のＣＵＤＡのサンプルプロジェクトが入れ替わってるからな
パーティクル関係の数字が増えてるけど同じプログラムでやってみた？

>>203
同じプログラムでやってる。
ベンチマークでもやってみるかな。

文字列処理をさせてはみたものの
遅すぎて使い物にならねーぞこんちくしょー

>>205
単精度な数値計算に変換すれば良いんでないかい

>>206
UTF-8をどうやって数値計算にすればいいぉか？

UTFの全領域を使う分けじゃなければ、必要な部分だけを数値にマップするとか。

文字列処理って言ってもいろいろあるだろ

ベンチマークやったら電源が落ちた・・・
780Wじゃ足りないのかな

>>210
よっぽどの詐欺電源でも買ってない限りは、さすがに足りないってことは無いと思うけど
初期不良じゃないか？

熱落ちじゃないの？

Geforceは80 PLUS シルバー以上の
電源じゃないとまともに高負荷に耐えられないぞ

80 Plusって書いてないわ・・・
CORAZON ってやつです。

>>214
http://www.keian.co.jp/products/products_info/kt_780as_sli/kt_780as_sli.html
12Vが25AだなHDDとかいっぱい着いてたら微妙だな
最小構成なら問題ない範囲だと思うけど
てかこれ安いな、これ買えばよかったｗ

25Aが2つ書いてあるのはどういう意味？

定員25人のエレベーター2基と、定員50人のエレベーター1基は違うというのは分かるな？

ＨＤＤとかPCIE用の電源が２本電源の中にあるってことだ
ちゃんと分割して接続すれば12x25=300WあるからGTX295のTDP300Wは支えられる
HDDとかつけるともう足りない
ちなみにうちの700W表記の電源は一本36Aある

IDE電源にハードディスクは接続しているけれども、
PCIEの電源は2本ともGTX295に挿してるぜ。
そういう意味じゃない？
ごめん詳しくないんだ。

大人しくこのあたり買っておくかな。
http://www.scythe.co.jp/images/energia/energia-label0800.jpg

>>219
検索すりゃわかるが今の電源でGTX295を動かしてる人が居るから動くって
ＨＤＤ１個にしてみて周辺機器もはずしてＯＣしてるならデフォルトにして
配線を入れ替えたりしてみてダメなら初期不良だろう

つまり、安いもんじゃないし保障が切れたらもともこもないから
電源が原因だとしても保障があるうちに買った店に持っていって確認してもらったほうがいい
お店の人が電源がだめだと言うなら電源を買えばいいし
お店にあるちゃんとした電源でもダメだったら初期不良で交換してもらえるからね

そうしてみる。
電源も一緒に買ったものだし、持って行ってみるわ。
めっちゃ勉強になった。ありがとう。

追記
HDを1つにしたら、落ちるまでの時間が長くなりました

排熱もやばいんだろw
きちんとしたケースとFAN買えw

多次元配列の領域確保、コピーってcudaMallocとcudaMemcpyでできる？

できる。
てか普通のCみたいにコピーできる。
ちなみにCudaだと多次元配列だと面倒だから1次元配列として扱うことがおおい。

cudaMemcpyAsyncってのもある

面倒っていうか１次元しか扱えないし

1次元だけだったか

自分の技量が足りなくてできないのかと思ってた。

今日からCUDA触ってみたのですが、全然速くない…

device側で1MB×2(dest, src)をアロケートして、hostからデータをコピー

for (int n = 0; n < 1024*1024; n += 512)
{
CUDA_Func<<<1, 512, 0>>>(dest, src, nPos);
}

hostからdeviceへコピー

=====

__global__ static void CUDA_Func(BYTE* pDest, BYTE* pSrc, int nPos)
{
int i = blockIdx.x * blockDim.x + threadIdx.x + nPos;
pDest[i] = ((int)pDest[i] + (int)pSrc[i]) >> 1;
}

=====

なんて事をやっているのですが、CPUの方が速いです
Visual Profilerを見ると、各CUDA_FuncのGPU Time は 8〜9us で終わってますが、CPU Timeが80〜150us になってます
こんなものでしょうか？アドバイス頂けると嬉しいです

Win7/GF8800/SDK 2.3

>>229
CUDA_Func<<<1, 512, 0>>>(dest, src, nPos);を2ｋ回も呼んでるのがまず悪いんじゃね？
あとは詳しい人に任せた。

俺も勉強中。共有メモリのバンクコンフリクトがわけわからねえ。

>>229
そりゃカーネルをキックするコストはかなりでかいから、ループで何回も呼んだらCPUにまけるだろ。。。。

>>230
shared memoryのbank conflictは
ものすごーーーーく大雑把にいうと、thread_id順でshared memoryのアドレスにアクセスすると各バンクのチャネルがぶつからなくて、パラレルに出来るよってお話

レスありがとうございます。
リニアに1次元配列を処理するような事は意味がないと言う事でしょうか？
例えば、ある程度の長さの、サンプリング単位のPCMの演算や、ピクセル単位の画像の演算とか…

連投スマソ

なんかレスを書いてて、やっとピンと来たんですが、
例えば各スレッドでさらにループで回して、CUDA_Funcを減らせば良い的な話だったりします？
>>229 のコードで言うと

CUDA_Func内で1KB分ループさせて、各スレッドへは1KBのオフセットを渡す。
その分、CUDA_Funcの起動回数を減らす。

違う…？

>>299
通常CPUなら、forで何回もやるような処理を
CUDAのカーネルを一発たたくことによって処理させるっていうのが基本的な考え方じゃないの？
あと、メモリは一度になるべく大きくとって転送したほうが効率がいいらしいよん

>>233
そのコードを見た感じ、1024*1024*512スレッドつかってることになってるけど、
何をしてるの？

いや>>229のやり方は正しいよ
画面描画とＣＵＤＡは同期処理だから大きい単位でやると画面がタイムアウト起こす
これ以上の最適化はＣＵＤＡでは不可能
これで遅いというならそれがそのカードの性能限界だと考えるしかないな

ちなみに>>229のCPUとカードの具体的な名前と周波数と
PCIEの速度とx16 gen2とかね
遅いって実際にどれくらい遅かったのは知りたいね
上位のクアッドCPUと8400GSなんかじゃ勝負にならないのは当たり前だから

>>229
ごめん元のソースが何したいのかよく分からないんだけど、こんなことしたいわけ？
1024ピクセル×1024ピクセルの二つの領域の明るさ平均を取るみたいな？
CPUなら、縦座標が外側ループで1024回×横座標が内側ループで1024回、回ると。

//device側で1MB×3(dest, src1, src2)をアロケートして、hostからsrc1,src2にデータをコピー
//512スレッドを起動するなら、外側ループの回る回数はCPUだと1024、GPUだと1024÷スレッド数512
for (int n = 0; n < 1024 / 512; n ++)
{
　CUDA_Func<<<1, 512, 0>>>(dest, src1, src2, nPos);
}
//deviceからHostへdestをコピー

=====
__global__ static void CUDA_Func(BYTE* pDest, BYTE* pSrc1, BYTE* pSrc2, int nPos)
{
　int i = nPos*512 + threadIdx.x * 1024;　//各GPUスレッドが動き出す起点、縦にずれてるわけ
　for (k = 0; k < 1024; k++)　//各GPUスレッドは呼ばれると横向きに1024回ループする。内側ループをGPUスレッドで処理。
　{
　　pDest[i + k] = ((int)pSrc1[i + k] + (int)pSrc2[i + k]) >> 1;
　}
}
=====

ちなみにこれだといちいちグローバルメモリへのアクセスになるんで、

　//Sharedにスクラッチコピー
　//スクラッチコピー分だけループ処理
　//SharedからGlobalに書き出し

するともっと速くなる

>>238
早けりゃいいってもんじゃないぞ
そんなもん低クラスのカードで動かしたら一発で画面真っ暗だわ

あとsharedメモリはそんな使い方するもんじゃないだろう
毎回コピーしてたらそのコストの方がデカイわな

sharedメモリってあれだ
１スレッドでやる計算が複雑な時に頻繁に変数の値を更新するだろ
そういう時にグローバルメモリよりもアクセスが早い一時領域として利用するもんだ
こういう計算自体が単純なケースでは効果はない

>>239
別人だけど8400GSくらいだとそうなの？経験上何msを超えるとハングする？

ググルとOSにより2秒や5秒でタイムアウトとあるがギリギリまでやるのはまずそうな気はする。

あ、失礼　coalescedになってなかった。こうかな??

//4回CUDA_Funcを呼び出す方向で。
for (int n = 0; n < 1024 / 256; n ++)
{
　CUDA_Func<<<1, 512, 0>>>(dest, src1, src2, n);
}

//nは「縦」の分割数
//512スレッドが連続した512バイトを取り込む。二回動くと、1ピクセル×横に1024ピクセルを処理。
//上に向かって縦256回回る（k)
__global__ static void CUDA_Func(BYTE* pDest, BYTE* pSrc1, BYTE* pSrc2, int n)
{
　for (k=0; k < 256; k++)
　{
　　//動き出す起点は各スレッドで1バイトずつずれてる。
　　int address = n*1024*256 + k*1024 + threadIDx.x;
　　//1024バイトを512スレッドで処理するので、二回。
　　pDest[address] = ((int)pSrc1[address]+(int)pSrc2[address])>>1;
　　pDest[address+512] = ((int)pSrc1[address+512]+(int)pSrc2[address+512])>>1;
　}
}

実際書いて動かさないと良く分からないすな。グレーのビットマップ二つ用意してやってみる形かな。

>>242
タイムアウトが何秒とか議論することですらない
マウスすら動かない状態が２、３秒も続くようなアプリはアプリ失格だろ

CPUより遅いくらいの8400GSで動かそうとしたなら2,3秒のフリーズ程度なら止む無し。
クラッシュしてデータを失わせるかもしれないリスクを犯すよりは
起動時に簡単なベンチ走らせて遅いGPUはハネちゃうのもありかな。
8400GSを考慮したせいでミドル以上のグラボの足を引っ張るとか馬鹿すぎる。

プログレスバー表示したら遅くなるから表示しないで画面を固まらせるなんて
そんなものは個人で使うだけにするんだなｗ

ちょっと待て、みんな一台のGPUであれこれやろうとしているのか?
それじゃ出るスピードも出ないぞ。

皆様おはようございます。そして、レス感謝です。
朝イチで打ち合わせがあるので、結果だけ取り急ぎ報告します。


前のコード
0 memcpyHtoD 332.928 2155.51
1867.52 memcpyHtoD 332.512 1848.49
3403.26 CUDA_Func 10.624 1158.18
3588.86 CUDA_Func 8.864 119.289
(略)
767008 memcpyDtoH 289.504 997.333

CUDA_Funcでループ
0 memcpyHtoD 332.864 2149.65
1815.04 memcpyHtoD 332.512 1792.27
3264.26 CUDA_Func 11235.1 12351.3
28136.2 memcpyDtoH 286.368 1402.62

満足行く結果ではありませんが、速くはなりました。CPUでリニアに処理した方が速いです。AthlonX2 @1GHz〜3GHz
あと、気づいたのですが、当方の環境ではRDP経由でCUDAが動きませんでした。ちょっとヤバイかも…

詳細は追ってフォローさせて下さい。

cudaMemcpyは同期を取ってから転送するから、結果の利用のタイミングぎりぎりまで実行を遅らせられれば
見掛け上の処理時間を短縮できるよ。

>>248
RDP経由でCUDAが動かないのは仕様です。

俺もやってみたらこうなりましたが。
Using device 0: GeForce 9500 GT
GPU threads : 512
Processing time GPU: 5.406832 (ms)
Processing time CPU: 18.742046 (ms)
Test PASSED
GPU: 78.0000 87.0000 177.0000 1077.0000
CPU: 78.0000 87.0000 177.0000 1077.0000
Press ENTER to exit...

カーネルはこう。
__global__ void testKernel( float* g_idata1, float* g_idata2, float* g_odata, int n)
{
　// access thread id
　const unsigned int tid = threadIdx.x;
　// access number of threads in this block
　const unsigned int num_threads = blockDim.x;
　__syncthreads();
　unsigned int startaddress = n * 1024 * num_threads;
　for (int j = 0; j < num_threads; j++)
　{
　　for (int k = 0; k < 1024; k = k + num_threads)
　　{
　　　unsigned int accessAddress = startaddress + k + tid;
　　　g_odata[accessAddress] = (g_idata1[accessAddress] + g_idata2[accessAddress]) / 2.0;
　　}
　　__syncthreads();
　}
}

//ホストでこうして、
for( unsigned int i = 0; i < 1024 * 1024; ++i)
{
h_idata1[i] = (float) (66.0 + i);
h_idata2[i] = (float) (88.0 + i);
}
// こう実行。
for (int i = 0; i < (int)(1024 / num_threads) ; i++)
{
testKernel<<< grid, threads >>>( d_idata1, d_idata2, d_odata, i);
}

//CPUはこう。Athlon 2.3GHz。
computeCPU(float* idata1, float* idata2, float* reference)
{
for( unsigned int i = 0; i < 1024 * 1024; i++)
{
reference[i] = (idata1[i] + idata2[i]) / 2.0;
}
}

これじゃサイズが小さすぎてあんまり比較にならないと思うっすよ。
景気よく4096×4096でやってみるといいかな？

4096x4096でやってみたらだいぶ差が出てきましたよ。

GPU threads : 2048
Processing time GPU: 33.380219 (ms)
Processing time CPU: 260.214355 (ms)

accessAddressの計算おかしくありません？

コードちゃんと見て無いけど、floatならGPUでパラで動かした方が速いに決まってる
athlon系の浮動小数点演算のコストは、整数演算のざっと平均70倍
intel系は知らないにゃぁ

computeCPU(float* idata1, float* idata2, float* reference)
{
for( unsigned int i = 0; i < 1024 * 1024; i+=4)
{
static const __m128 div2 = { 0.5f, 0.5f, 0.5f, 0.5f };
__m128 tmp = _mm_load_ps(&idata1[i]);
tmp = _mm_add_ps(tmp, _mm_load_ps(&idata2[i]));
tmp = _mm_mul_ps(tmp, div2);
_mm_store_ps(&reference[i], tmp);
}
}

あとどっかでprefetchnta噛ませるといいかも。


CPU側は最低限SSE使おうや。
大学関係者も含めて比較用のCPU側コードが酷いのが多すぎる。

ぶたぎりですが、
geforce gtx285 と　tesla c1060 って
どのくらい違います？
c2050がどうせでるから、c1060と285が
対してかわらないなら、285がいいかと思うのですけど。

ところで
http://www.nvidia.co.jp/object/personal_supercomputing_jp.html
いつのまにかnvidiaのページに
「Tesla GPU コンピューティングプロセッサは発売されています:
- Tesla C2050/C2070
- Tesla C1060」
こんなことが書いてあるけど、まだ出てないですよね。

>>257
そんなこと言うと、NVIDIAの営業に「そんなアキバ的発想はダメですよ」って言われちゃいますよw

CPU側は、やっぱりCore i7で8スレッド並列とか動かしてあげないとだめじゃね？
そんでGTX285と勝負するみたいな。

率直に言って、なんで値段が5倍のTeslaが売れるのか良く分からないのです＜アキバ的発想ですが
Fermiアーキテクチャも、GeForceが先に出るんじゃないですか？　来年2月くらいでしょうかね。

>>260
そりゃぁ、GeForceが叩き台だモノ。NVIDIA曰く、GeForceよりも信頼性が10倍高いから5倍の価格差はペイできるって発想でしょ。

>>259
HPC用によく最適化されたコードだとHT切ったほうが速い

>>260
サポートが手厚い

Engadgetに、米軍が旧型PS3を2000台買ったとか書いてあったよな
$200ドルで256MB、100GFlops、て安いのかな??
10倍して$2000ドル、2.5GB、1TFLOPS、て。どうなんだろ。GPUのほうが安そうな。びみょう？

GTX285とか、240SPに、512とか1024とか2^nで放り込んでもうまく処理してくれるのですか？
半端な部分の処理がどうなるのか心配していると脳みそかゆくなります。

>>264
cudaライブラリが「適当に」割り振ります。グラフィック兼用だと、そもそも全部使う保証さえありません。

信頼性とか保証とかサポートとかの違いか

いや、メモリ搭載量全然違うし。

そうだね、GeForceの方が安いメモリを使っているしね。

はじめてのCUDAプログラミング買ったが
テクスチャには触れてないも同然だな。

>>263
ゲーム機としても使うつもりだから

>>251
実コード動かしたのは乙だけど、なんでデータがfloatになったの？

整数演算だったら、結局ちゃんとコードが書かれていればボトルネックは転送速度であって、
DDR2-CPUとDDR2-(PCI-Ex16)-GPUの比較になりそうな気がする。

CUDAの処理速度を計りたいならまずプロファイラにかけろ。話はそれからだ

>>263
コードの資産の問題なんでしょ。
Cell用に書いていたので、GPUにポーティングする手間を考えたら、
PS3を大量に買った方が安いのでしょう。
もちろんアプリによるけど、PS3は実効性能で100Gflopsはでるけど、
GPUは実効性能で1TFLOPSも出ないでしょ。
自分の経験では、Cｅｌｌに対して、G200は大体2倍〜3弱倍という感じだった。
Cellはカリカリにチューンしたけどね。

アキバ的発想ではやっぱり、GTX295を二枚というのが現時点ではさいきょうCUDA環境ですか

>>273
GPU内部だけで完結させる処理なら１Tいくかもしれないが
PCの場合はデータ転送がボトルネックになって１Gｆlopsとかになるが

>>256
Type Zでやってみた。P9500(C2D　2.53GHz)、GeForce9300M GS(8SP)

GPU Threads: 256, memsize = 16777216
Processing time GPU: 81.911575 (ms)
Processing time CPU: 218.790421 (ms)
Processing time SSE: 84.257759 (ms)
Test PASSED

Press ENTER to exit...

ほう

結局さぁ、「ゲーム」だったらCPUで全部演算しようなんて考えるやついないんだし、
別にGPUでやるほどのもんじゃねぇ、ってだけなんじゃないの＞229の例

>>274
アキバ的にはGTX295を４枚が最強だけど、
そこまでするならtesla買った方が幸せだわな。

Larrabeeがなんかよさそう

>>271
4096x4096でやってみたら(float)

GPU threads : 2048
Processing time GPU: 33.380219 (ms)
Processing time CPU: 260.214355 (ms)

BYTEにしてみた。4096*4096で
GPU threads : 2048
Processing time GPU: 27.527241 (ms)
Processing time CPU: 345.093292 (ms)

になった。BYTEのほうがCPU遅い…のはこれはこれでよいのかな

>>278
電源がそれだけで1200Wくらい要るなぁ。
1000W電源二台って感じか。うちのブレーカーは間違いなく飛ぶだろう

電子レンジとかドライヤーつけなければいけるだろう
たぶん

夏は地獄を見るだろう

CUDAで倍精度演算をしたいと思っているのですが、
・GeForce GTX 285
・GeForce GTX 295
なら、どっちがいいですか？

また、もっといいものがあれば教えてください。
Teslaは高いらしいのでそれ以外で、お願いします。

>>284
Fermiが出るまで待つ
でなければGTX295でしょうねぇ。(メモリ・帯域幅は小さくても倍精度演算リソースが二倍)

>>285
ありがとうございます。
Fermiがいつ出るかわからない状況なので、GTX295の方向で考えます。
いつ出るか発表されてませんよね？

cudaMallocHostってmallocのラッパーではない？
ライブラリつかってて、なぜかcudaMallocHostを指定していて
mallocで大丈夫だろうと思ったら動作せず
cudaMallocHostなら動作したので。
どう違うの？

>>287
つ　page lock

>>285
１個づつ付いてるのを２倍って呼ぶのは詐欺くさいですな

>>287
メモリコピーのハードウェアアクセラレートって決まった区切りで始まる決まったサイズのメモリ単位でしか扱えなくて
そういうのを考慮してメモリを確保するんでしょう

>>280
byteこそSSE使わないと。
最大16倍程度は速くなる

1byte単位のストアは部分書き込みになるからあまり性能的によろしくない。
アラインメント境界にあわせて128ビット単位のストアが一番パフォーマンスよろしい

#include <emmintrin.h>
computeCPU_SSE2(BYTE* idata1, BYTE* idata2, BYTE* reference)
{
for( unsigned int i = 0; i < 1024 * 1024; i+=16)
{
__m128 tmp = _mm_load_si128(&idata1[i]);
tmp = _mm_avg_epu8(tmp, _mm_load_si128(&idata2[i]));
_mm_store_si128(&reference[i], tmp);
}
}

>>279
最終的にはメモリ帯域ネックなんで同等のメモリ帯域を持つGTXと変わらんと思うよ。
ただ2つの配列の平均取るだけならほぼ帯域ベンチ。
PCIeかGPUのL/S帯域で律速。

キャッシュメモリが256KB/coreほどあって16コアあれば、1024x1024ならキャッシュにほぼ収まるので
それを有効に使って処理を連結していくなら、そこではじめてLarrabeeの旨味が出てくる。

個人的にはN-Bodyとかが良い勝負しそう。

>>293
だんごさんてきにはGrape-DRが
お勧めということですね

多体問題だけやりたいんならいいんじゃね？

おいおまいら、こんなのがあるぞ。
http://www.acceleware.com/index.cfm/cuda-training/2010sunnyvale/

リッチな会社にいたら旅行がてらいってみるといいぞ。
これ見るとＦｉｘｓｔａｒｓがいかに良心的かわかるぞ。

>>295
Larrabeeは机上計算だと1コアあたり16並列で15サイクル（＋α）くらいで回るんで、かなり理想的なデバイスなんだけどね。
（ただしvrcppsの精度補完ありならもう少しかかる）

GTX260 で決まりすよ。18000円なんだから4枚位買えばいい。

ららびーはもう終わったみたいな話が出てますけど
http://venturebeat.com/2009/12/04/intel-cancels-larrabee-consumer-graphics-chip/

読んで字のごとくコンシューマ向けディスクリートはね。
何のためにスパコン向けカンファレンスで発表したと思う？

まあなんにせよ「パソコン」に刺さるカードは無くなったな。
PTX2.0の資料請求しとかなきゃ。

>>289
っていうことは、285の方がよさそうでしょうか？
メモリの帯域が大きく、こちらの方が使いやすそうな気がします。
デュアルGPUのコーディングも大変そうなので……

>>301
待てるならFermiを待った方がいいよ

倍精度性能はGT200世代はおまけで１個演算機がついてるだけだから期待するだけ無駄だよ
Fermi世代から各コアに演算機付けるって言ってるけど、
一般向けには倍精度削ってGF100の名前で出すとか言ってるからどうなるか分からん
CUDAに拘らないならRADEON HD5xxxって選択もある
GT200世代より倍精度演算能力は圧倒的にHD5xxxの方が高いし
ただRADEONはドライバが糞だしCUDAも動かないしいろいろ中途半端
Larrabeeがどうなるかってところか

現状実用的なものは無いから、実験的な目的以外では買わない方がいいし
実験的な目的なら安いやつでいいじゃないかという話になる

Larrabeeはコンシューマ向けディスクリートGPUとしてはキャンセル。
HPC向けには続投しそうだから100万くらい出せば手に入るようになるかもよ？

RADEONはFFTを出せないところを見ると行列積が精一杯のようだよ。
GTX280は512bitのバンド幅がどうもよろしくないのでGTX260を奨めます。

もしRADEONで遊ぶなら、現時点ではDXCSでSM5.0のスレッド制御を
使ってどこまでできるかだろうなぁ。
DirectX 11 SDKにFFTのサンプルコードなんかもあるから、持ってる人は色々
ベンチマーク取ってみて欲しいな。

>>304
Tesla by Fermiの値段にぴったり張り付いて売るのがIntel式な気が
しないでもない。そのためにもnVidiaには頑張ってもらわないと！

295だろ、普通に考えて。
スレッド数を多く使える分だけ、高速化が容易
メモリのバンド幅とかよりも重要だと思うが？

迷っている段階なら、とりあえずGTX260を買って試すのがオススメかな

>>307
流石にIntelシンパの俺でも30万は出せない。
Fermi出た時点で一番コストパフォーマンスいいの選ぶわ。

試行錯誤するのにcompute capability 1.3に対応した
ファンレスカードとかあれば良いんだけどな。
260や295を付けっぱなしってのは、なんか精神衛生に良くない。

CUDA勉強中の者ですが共有メモリの利用で躓いてるところです。アドバイス頂けたら幸いです。
下記の二つのカーネルでCUDA_karnel_sの方が5倍時間がかかってしまうのですが原因がわかりません。
違いは読み込んだデータをグローバルメモリに保存するか共有メモリに保存するかです。
__global__ void CUDA_karnel_g(uchar4 *vram, int sx, int sy, uint1 *vram2)
{
　　int i = threadIdx.x;
　　vram += blockIdx.y*sx;
　　vram2[i] = ((uint1 *)vram)[i];
　　__syncthreads();
　　uchar4 px;
　　*((uint1 *)&px) = vram2[i];
　　unsigned char Y = (unsigned char)(0.2126f*px.z+0.7152f*px.y+0.0722f*px.x);
　　px.z = px.y = px.x = Y;
　　vram[i] = px;
}
__global__ void CUDA_karnel_s(uchar4 *vram, int sx, int sy, uint1 *vram2)
{
　　int i = threadIdx.x;
　　vram += blockIdx.y*sx;
　　__shared__ uint1 shared[125*32];
　　shared[i] = ((uint1 *)vram)[i];
　　__syncthreads();
　　uchar4 px;
　　*((uint1 *)&px) = shared[i];
　　unsigned char Y = (unsigned char)(0.2126f*px.z+0.7152f*px.y+0.0722f*px.x);
　　px.z = px.y = px.x = Y;
　　vram[i] = px;
}
まずバンクコンフリクトを疑ったのですがCUDA Visual Profilerでみるとwarp serializeは0で発生していませんでした。
意図的にバンクコンフリクトを発生させると更に１０％程遅くなるのでバンクコンフリクトは原因ではなさそうです。
共有メモリは速いはずなのにグローバルメモリよりなぜ遅くなるのか悩んでいます。初歩的なミスだろうとは思うのですが。

>>312
sharedが一定以上多いとOccupancyが下がるから、そこらへんじゃない？
Occupancyは実行効率にダイレクトに効いてくる。
Visual Profilerの実行ログにも出てくるし、SDKのtools/CUDA_Occupancy_calcurator.xlsで試算可能。
Shared Memory Per Block (bytes)のところに16000って入れると良い。(125*32*sizeof(uint1))
ちょっと計算してみると、
スレッドブロックのサイズが512ぐらいならまだマシ(67%)だけど、
64とかだと壊滅的に遅くなる(8%)。

あとpxがアドレス参照でローカルメモリ（VRAM上)に行ってないか心配だ。

vram2[i] がレジスタのってたりしないかな。
親から vram3 として渡してみるとか、参照を i+1 にするとか。

>>313
仰る通りだったようです。shared[125*32]をshared[32]としたら劇的に速くなりました。
バンクコンフリクトを疑うあまりブロックあたりのスレッド数を32にしてたのも不味かったようですね。まだ未確認。
>あとpxがアドレス参照でローカルメモリ（VRAM上)に行ってないか心配だ。
*((uint1 *)&px) = shared[i]　を　px.x = shared[i].x;
と書き換えたところ速度アップしたのでその可能性大です。最適化で同じコードになるんじゃね？と勝手に思ってました。

>>314
>vram2[i] がレジスタのってたりしないかな。
このコードは質問のために単純化したものだったのですが、単純化前にでptxファイルを出力して比較したときは
ld.global.u32のところがld.shared.u32に変化しただけでしたのでレジスタにのったり最適化で消えたりはしてなさそうです。
とはいえptxの書式の資料を見つけられなかったので自信ありません。
参照を i+1の意味はちょっとわかりませんでした。すみません。

>>313-314
おかげさまでグローバルメモリより共有メモリが遅いという現象は解決しました。ありがとうございます。
ブロック専用なのにextern __shared__ の構文がなぜあるのか不思議に思っていたのですがこういう理由だったのかと。
しかし大きな共有メモリで遅くなるのは厄介な仕様ですね。

>>315
おつ。
勘が鈍ってなくて安心したｗ
共有メモリは多数のコアで共有される有限リソースだから、そこも一定以下に抑えないといけないって話ね。

あと本当にsharedを使う必要があるのか、一時変数（レジスタ）で済ませられないか？
sharedを使うべきなのは、同時に動くコア同士のやりとりがある場合と、
どう処理を分配してもcoarescedアクセスにできなくて、でもsharedを使えばできるという場合ぐらいかと思う。

あ、そういや、逆にレジスタが多くてOccupancyが上がらない場合の退避用とかにも使えるだろうか？

訂正　*((uint1 *)&px) = shared[i]　を　px.x = shared[i].x;の部分は勘違いでした。
１バイトしか書き換えてないので等価ではありません。無視してくださいです…
>>315
実はどういう理屈で遅くなるのかよくわかってないのですがお蔭様で対処の方法がわかってひと安心です。
最終的にはいろいろな画像処理に使う予定なのでキャッシュ的な使い方をsharedでさせてみるテストでした。
比較のためテクスチャ版とshared版作ろうとしてはまってしまいました。

自宅でGPU4枚とかって人はなにに使うの？

エンコ

ベンチ見てほくそえむ

>>317
ptx出力を眺めれば判るけど、普通のCならできる最適化もGPU向けにはできないことが多いよ。
敢えて言えば、ポインタ変数に尽くvolatileがついているかのように振る舞うみたい。
例えば、
int function(int * foo) {int bar = 3; foo[0] = bar; return foo[0];}
みたいなコードはCPU向けには
{foo[0] = 3; return 3;}
のように最適化されるのにGPU向けには
{foo[0] = 3; return foo[0];}
のように律儀に解釈される。
なので、ptx出力を読めるかどうかは割りと重要かも知れず。
# つーか、メモリアクセスの個数を数えるくらいのことは普通にやってる。

すみません、かなり初心者です。
行ごとに要素数の違う配列をデバイス側に渡したいんですけど、

a=1;
//ホスト側
float** mat=(float**)malloc(size);

for(i=0;i<num;i++){
mat[i]=(float*)malloc(size / a);
a*=2;
}

//デバイス側
float** mat_d;

CUDA_SAFE_CALL(cudaMalloc((void**)&mat_d,size));

for(i=0;i<num;i++){
CUDA_SAFE_CALL(cudaMalloc((void**)&mat_d[i],size/a));
a*=2;
}


CUDA_SAFE_CALL(cudaMemcpy(mat_d,mat,size,cudaMemcpyHostToDevice));
for(i=0;i<M;i++){
CUDA_SAFE_CALL(cudaMemcpy(mat_d[i],mat[i],size/a,cudaMemcpyHostToDevice));
a*=2;
}
というようにはできないんでしょうか？

>>322
>CUDA_SAFE_CALL(cudaMemcpy(mat_d,mat,size,cudaMemcpyHostToDevice));

これがいらないんじゃない？
ホスト側のポインタ列をデバイスに渡しても使いようがない。メモリ空間が違う。

開発環境の入っていないマシンで動かすには、どのファイルを持って行けばよいのでしょうか？

>>291
個人的な興味なんですけど、CUDA使ってるcodecってなんで速いんでしょうね？
BYTEアクセスじゃ無いとしても、どの部分をGPUでやらせれば爆速になるのか不思議ﾁｬﾝﾃﾞｽ

SDKのドキュメントの通りに，VS2005で
サンプルプロジェクト template を元に使っているのですが
このプロジェクトに新しい.cuファイルを追加した場合に
そのcuファイルへの更新がビルド時に察知されません
毎回リビルドすれば一応反映できるのですが
通常のビルドで反映させるにはどうすればよいのでしょうか？

>>323
すみません。
抜けてましたが、
>CUDA_SAFE_CALL(cudaMemcpy(mat_d,mat,size,cudaMemcpyHostToDevice));

は、
ホスト側であらかじめ計算を行い、
その結果を初期値としてポインタ列、matに与えた上で、
その値をデバイス側のポインタ列、mat_dにコピーし、
その値を使ってデバイス側で計算をする。
というつもりで書きました。

このようにホスト側で計算した値をデバイス側に渡す時は
どのように記述するのがよいのでしょうか？

{1個、2個、4個、8個、16個、…}　みたいなデータを渡したいのかな??
固めてまとめて渡しちゃったほうがオーバーヘッドが少ないと思います。
トータル何列あるよ、は別にパラメータで渡す。
(実際にCUDAのルーチン書く前に、コピー/戻しの時間を色々計ってみるといいです)

cudaMallocしたデータにはホストからは触れないので、
ホストでmallocしたデータ(mat)に計算結果格納
→同じサイズでcudaMalloc(mat_d)
→cudaMemcpyHostToDeviceで渡す
なのでそれでいいです

>>327
>CUDA_SAFE_CALL(cudaMalloc((void**)&mat_d,size));
>
>for(i=0;i<num;i++){
>CUDA_SAFE_CALL(cudaMalloc((void**)&mat_d[i],size/a));
>a*=2;
>}
考えてみるとここも問題があって、cudaMallocということはデバイス側でポインタ列を確保しているんだけど、
そうすると&mat_d[i]というアドレスは、デバイス側にはバッファがあるが、
ホスト側には存在しないから、ここで例外になりそう。
やるならこんな感じかな？↓（未検証）

float** mat_d; // GPU側に確保する（ポインタ列用）バッファ
CUDA_SAFE_CALL(cudaMalloc((void**)&mat_d, count * sizeof(float*));
float** mat_d_tmp=(float**)malloc(count * sizeof(float*)); // ホスト側に確保する。内容はGPU側ポインタ列
for(i=0;i<num;i++){
// GPU側データバッファのポインタを格納
CUDA_SAFE_CALL(cudaMalloc((void**)&mat_d_tmp[i],size/a));
a*=2;
}
// GPU側ポインタ列をGPUに転送
CUDA_SAFE_CALL(cudaMemcpy(mat_d,mat_d_tmp,count * sizeof(float*),cudaMemcpyHostToDevice));
free(mat_d_tmp);
てやっておいて、mat_dをカーネル呼び出しの引数にしてやるとか。

ここまで、バッファ作りしかしてないので、データ転送は別途必要。
たぶん>>322の最後４行のとおりで構わない。

ただデータ転送回数多いと多少とも時間かかるから、>>328の言うように
固めたほうがベターではある。

>>329
ちなみに、ポインタサイズが32bitか64bitかはホスト側に依存してることを確認した。
http://forums.nvidia.com/index.php?s=&showtopic=149501&view=findpost&p=943899

>>325
あいまいな質問だけど答えてみる。

CPUよりGPUのほうが計算能力の総量はでかいから、本来GPUが速いほうが当然なんだけど、
以下のようなものは遅くなる場合もある。
・計算負荷よりもメモリ転送がボトルネックになるもの
・処理を細かく分解して並列化することが難しいもの

codecなどは、メモリ転送がボトルネックになりやすい傾向はあるものの、
だいたい画面の領域ごとに分解できる処理が多いから、多少とも速くはできるでしょう。
あと速くならない処理はCPU側でやればいい。

>>328>>329
参考になりました。ありがとうございます！
まだ不慣れなので色々試行錯誤しながらやってみます！

>>325
CUDAだけ使って早いと思ってる？ｗ

動き検出なんかは当然CPUにやらせたほうがいいぞ
SpursEngineなりCellなりがあるならそれでもいいが

>>333-334
流石にGPUだけとは思ってないけど、atomでh264なんかエンコすると日が暮れる勢いだけど、ionプラットフォームだと実用範囲になる。
エンコ中でも完全にCPUを使い切ってないところを見ると、やっぱGPUをかなり使ってるんだなと勝手な想像

IONはメインメモリをグローバルメモリとして使う、つまりマッピングしておけば
CPUからもGPUからもフルスピードでアクセスできるメモリ空間を作ることができる。

ION2は買おうかな・・・。

>>336
実際はそんなに甘くはなく、GPUからフルスピードでアクセスするためにはWriteCombinedモードに
する必要があって、それを付けるとCPUからのアクセスが死ぬほど遅くなります。
通常は素直に転送（実際にはメモリコピーですが）したほうがマシです。

低価格で4GB近いデバイスメモリが使えるというメリットはありますが。

そのためのMOVNTDQA/MOVTDQだろう。
まあ前者はSSE4.1でないと使えないがな。Atomは仕様不可。あしからず。

それらを使っても十分に遅いと思うのですが。

WCにしつつ、CPU->GPUはmovnt*でcopyすれば良さそうだね。
CUDAとか知らないけど:b

>>339
だからCPU専用のワーク領域と分けて読み書きを最小限にする。

>>337
そか、CPUのキャッシュが使えなかったね。
死ぬほどまで遅くなるとは思ってなかった。

GTX295 Vista x64 CUDA2.3 の環境で、CUDA VS WIZARDで作成したプロジェクトを使っています。
質問1. ブロックあたりのスレッド数を512にすると計算結果が正しくなりません。256なら正しいです。なぜでしょう？
shared memory はグローバル変数として宣言して、各スレッドで共用（リードオンリー）しています。スレッド内で__shared__の宣言はありません。

質問2. 完成したので、SDKのsimpleMultiGPUを参考にGPU 2個使おうとしています。
とりあえずマルチスレッド用の空の関数を書いたところ、コンパイルはできるのですが、リンクできません。
cutStartThreadなどが外部で定義されてないと言われます。
simpleMultiGPUはビルドできるので、プロジェクトのプロパティを見比べてみましたが特に違いは見あたりません。
何が悪いんでしょうか？

どなたかお助けください。

>>343
・sharedにはCPUから書けませんが、各スレッドがリードオンリーならどこで書き込んでいるのでしょう。
・cutで始まる名前の関数はcutilライブラリ内にあります。cutilライブラリもビルドし、あなたのプロジェクトからリンクできるようにしないといけません。

>>344
上については、言葉足らずでした。カーネルの最初に
if( threadIdx.x == 0 ) { /* デバイスメモリからコピー */ }
__syncthreads();
でコピーしています。
将来的にはコピーも並列化する予定です。

下についてはありがとうございます。既存のlib(dll?)だけでは駄目だってことですね。

threadIdx.xはスレッドの識別子で、実行順番とは無関係では？

共有メモリのつくりに関する知識が思いっきり欠如していると思われ。
threadIdxが0の人だけが書こうとしたら、他の人はみんな何もできなくて時間が無駄。
おまけに、スレッド数が一定数を超えると(GTX295の場合は256か)、別の実行単位になってしまうから(表に見えない)
共有メモリ間転送が行なわれてしまってその点でも時間の無駄。
尤も、>345の様にthreadIdx0でしかコピーしない場合は転送は発生しない代わりに結果がご覧の通りなわけで。

どうせthreadIdx0がコピーしている間他はなにもできないんだから、一斉に同じものを書いてしまえばいいんでない?

>>347
> threadIdxが0の人だけが書こうとしたら、他の人はみんな何もできなくて時間が無駄。
開発途中なので、とりあえずこうしているだけです。
将来的にはコアレスにします。

> スレッド数が一定数を超えると(GTX295の場合は256か)、別の実行単位になってしまうから(表に見えない)
512で上手くいかないのはこのあたりが原因のようですね。

>>347
たびたびすみません。
「別の実行単位」の意味がよく分からないのですが、
要は1つのブロックが2つのSMに割り当てられるって意味ですか？

>>345
Sharedは文字通り、共有できるメモリなので、
各スレッドが１ワードずつ協力して作成→どこでも読み書き可能、
ですよ。0番スレッドだけに任せなくてもいい

自分は>>343じゃないけど256スレッドでOKで512スレッドで計算結果がおかしくなる理由がわからない。
>おまけに、スレッド数が一定数を超えると(GTX295の場合は256か)、別の実行単位になってしまうから(表に見えない)
>共有メモリ間転送が行なわれてしまってその点でも時間の無駄。
というのは時間の無駄を指摘してるだけで結果がおかしくなる理由ではないという認識なのですがあってますか？
この時の共有メモリ間転送というのは別の実行単位に転送しているということ？

>>343
256スレッドでやればOKなのでもう興味はないかもしれないが、シェアードメモリの代わりにグローバルメモリで
共用して計算した場合は512スレッドの時も結果は正しくなるのですか？

>質問1. ブロックあたりのスレッド数を512にすると計算結果が正しくなりません。256なら正しいです。なぜでしょう？

カーネルコールの後でエラーキャッチのコード入れてる？エラー出てない？

質問の傾向をみると、
CUDAって、普通のCのように自由度高く書けるけど、
実際は、サンプルソースからは読み取れない
あれは駄目、こうしたら駄目、的な制約が多すぎるのかな？
DXAPI+HLSLのような、自由度の少ない環境の方が
むしろ、良質なソースが書けるのかも。

参考資料

Device 0: "GeForce GTX 295"
CUDA Driver Version: 3.0
CUDA Runtime Version: 2.30
CUDA Capability Major revision number: 1
CUDA Capability Minor revision number: 3
Total amount of global memory: 939524096 bytes
Number of multiprocessors: 30
Number of cores: 240
Total amount of constant memory: 65536 bytes
Total amount of shared memory per block: 16384 bytes
Total number of registers available per block: 16384
Warp size: 32
Maximum number of threads per block: 512
Maximum sizes of each dimension of a block: 512 x 512 x 64
Maximum sizes of each dimension of a grid: 65535 x 65535 x 1
Maximum memory pitch: 262144 bytes
Texture alignment: 256 bytes
Clock rate: 1.24 GHz
Concurrent copy and execution: Yes
Run time limit on kernels: No
Integrated: No
Support host page-locked memory mapping: Yes
Compute mode: Default (multiple host threads can use this device simultaneously)

>>353
CUDAをただの言語だと思っていてGPUのアーキテクチャを理解していないからかと。

後半については同意。限られた記述で実装可能なアプリしか書かないからというのがより正確かも。

>>351
スレッド数512、デバイスメモリ使用でやってみましたが、やはり正しくありませんでした。
Emuではちゃんと動くので、原因はよく分かりません。
コンスタントも使って多のでそっちが原因かも知れません。
分からん。もういい。

>>343 の質問2については、multithreading.cppをプロジェクトに追加してないだけでした。
俺死ね。

>353
テクスチャはまじで鬼門。
凝った使い方をしようとすると、コンパイル通っても動かない・動いても結果間違ってるがざらにある。

>>357
あるある。
deviceemuともまた動作違うこと多いしね。
NEXUS出ればデバッグもうちょっと楽になりそう。

>>343
歩留まり向上のため、重い処理をさせると計算が狂うことがあるらしい
どこかのPDFで見た

>>359
それちょっと表現が微妙に誤解を招くような・・・。
出荷されている状態で、計算がまともじゃないコアがたまにまぎれこんでいる。

グラフィック用途と違って、メーカーの想定より負荷が偏った計算をさせることがあるから、
やや耐性の低いコアの耐性がちょっと足りない部分がエラーを起こす。

高負荷プログラムなんかを使ってそういうコアというか製品をはじいてやれば、あとはほぼ安定稼働する、という話だと思う。

はじめてのCUDAがいつまでたっても届かない
一体どこに発注したんだよ上の人

>>360
CUDAのスパコンで1/6が不良品だったと言ってたが

えっ　Teslaとかでもそうなのか？？

Teslaは選別品

低価格帯で一番安定してるGPUってなに？

http://slidesha.re/5FtABc のP.26
長崎大の人に言って選別プログラムを貰うことだなｗ

CUDAは保障外だからCUDAやりたい奴は動くのが当たるまで買えってことなんだろう

仮に2/3超の確率で完動しないとしても、自前で問題無く選別できるなら
Tesla買うよりコンシューマGPU買うほうがまだ安い位だから、たくさん
買う人は自前で選別できるように頑張るのが正解だよなぁ。

多少クロック落とせば動くんじゃないのかな。CPUのオーバークロックと一緒で。

>>361
自分は11/27楽天でぽちって 12/01に到着。
ぼちったあとすぐ納期の表示が延びたから焦ったけど滑り込みセーフだったらしい。
書いてあることはわかり易い。しかし意外なことが触れられていなかったりする。

>>360
その状態が続く限り、コンシューマ用パッケージソフトじゃGPGPUを使えないね。
「動かない」サポートの爆発でたいへんなことになる。

てか>>356の作ってるソフトが選別ソフト代わりになるんじゃね？

おれらまだまだ、「主メモリ側が主役」って固定観念なくね？
いかんな。どかどかデバイスメモリにロード、大量のブロック×スレッドを駆使、
CPUはたまにお手伝いをさせていただく、位に思わないとな。俺まだまだ

メモリ転送と計算を非同期で出来るのかね？

>>374
マルチスレッドで、CPU側は「GPU様が計算し終わるまで寝とけや」くらいの扱いにするわけよ

>>374
できる。
基本GPUとCPUでは非同期。
ストリームを設定すれば、GPUで演算中にデータ転送させることも可能。

ちゃんと同期させるための関数もある。

> NEXUS
いつ出るんだっけ
VisualStudio2008でいいんだっけ

カーネルで計算した結果をCPU側に返すにはどうしたらいいの？
return aのように簡単にはできないのですか？

>>378
ちょｗ
カーネルfunctionの引数で　float * outdata
とかやってカーネルからこれに書き込む　outdata[i] = result;
ホスト側でcudaMemCopyのdeviceToHostで、ホスト側でcudamallocしたfloat * に戻す
みたいになる。ややこしいぞ？

thrust便利だなしかし

thrust、自分で書いたカーネル関数はどう使うのっすか

__shared__ int sh[num*2]; /* numはスレッド数 */
という共有メモリの配列をソートし、最大or最小のデータのみグローバルメモリに書き込みたいのですが
これをカーネルの中でやるいい方法を教えてください。

ソートは不要で、reductionで半分を繰り返せばいいのでは

>>383
言われてみればそうですね。
解決しました。thx

どこかにfortranでの使用例って
ありませんか

そういえば見かけたことねえな

CUDAはそもそも実行速度がシビアで
高級言語向けではないと思うけどな

共有メモリは16個のバンクに分割されて、半ワープのスレッドが同じバンクにアクセスしなければ競合は起こらない。
みたいな事の意味がイマイチ分かりません。

例えばブロックの共有メモリが8KBなら0.5KBづつに分けられて、半ワープのスレッドが0.5KBの中ならどこにアクセスしても良いって意味なの？
それとも、先頭から4ﾊﾞｲﾄとか8ﾊﾞｲﾄとかの固定の領域に連続して分けられてるの？
でもそれだと分けられる領域が4ﾊﾞｲﾄなのか8ﾊﾞｲﾄなのか分からないじゃまいかヽ(`Д´)ﾉ

カーネルの引数にvoid*とデータサイズを渡して任意のデータ長で処理させたいのですが、
こういう事をしようとすると元からバンク競合は避けられないのですか？教えて、エロイ人(´・ω・`)

今あるMPIのプログラムをGPUに対応させようと考えているんだけど、
CUDAでMPIを使うとき、データの転送は一端ホスト側のメモリに転送してから
他のノードへ転送することになるの？
だとするとレイテンシがすごく大きそうだね。
それとも専用のライブラリなんかあるのかな？
GPU側のグローバルメモリがホスト側のメモリにマップされていれば、
GPUを意識しないで転送ができそうなんだが。

>>388
おれの解釈だと、共有メモリには16個のアドレス入力と16個のデータ入出力が有る
と思ってるんだけど　違うかも

>>388
連続する64バイトの領域（先頭アドレスが４の倍数）について、4バイト(32bit)のバンク１６個に分かれています。

例えばバンク０に属するアドレスは０、６４、１２８、１９２、２５６、・・・・・から始まる４バイトの領域。

任意のデータ長が４バイトより大きいのか小さいのか分かりませんが、テンプレート等を使う場合には
型に合わせて場合分けをする必要があるでしょう。

>>390
「１個のアドレス入力と１個のデータ入出力ができるバンク」が１６個あるというのが適切な表現かと。

>>389
device memory on node 0 -> host memory on node 0 -> host memory on node 1 -> device memory on node 1
となるのでレイテンシは大きくなります。今のところCUDAではデバイスメモリをメモリマップする手段はありません。
ただし十分大きなデータを転送する場合にはパイプラン化すれば問題なくなると思います。

むしろpinnedメモリとMPIライブラリの干渉が一番の問題・・・・。

>>392
レスサンクス

やはりレイテンシは大きいのだね。
目的のアプリは数10kbyteから数100kbyteの転送を頻繁に行うから、
せっかくGPUで速くできても、転送ボトルネックで詰まりそう。
転送するサイズも小さいので、page lockさせない方がよいのかも。

GPUの購入で悩んでいるのだが、
Tesla C1060, GTX295,GTX285のうち結局どれが
速いんですか？ちなみに流体に使います。
GTX295ってGPU2基積んでるけど並列プログラミング
しないと2基機能しないとか？
素人質問で申し訳ない。
Teslaの保証とサポートも魅力的だが。

>>394
メモリ量でC1060になっちゃう　なんてことないかな？
295はちゃんとマルチスレッドでプログラムしないと二基使えないです。

>>390 >>391
なるほど…よく分かりました。ﾚｽ感謝。
処理させたい任意のデータ長は5ﾊﾞｲﾄだったり、11ﾊﾞｲﾄだったり半端な数もきます。
任意のデータ長に対してコピーやビット演算を行なうんですが、データ長が4ﾊﾞｲﾄより大きいと
もうバンク競合は避けられない感じなんですね。

プログラムも汎用にさせたかったけどここは我慢するしかないのか…(´・ω・`)

VCでプロジェクト作るのが面倒なんだけど、
なんかウィザード的なものはないのかな

http://forums.nvidia.com/index.php?showtopic=65111&pid=374520&st=0&#entry374520

HTML+JavaScriptだから好きに書き換えられるだろ

おいらも流体で使う予定だが
GTX260を二つ買ってCUDAのお勉強しつつFermi待ち
C1060とか中途半端に高いからねえ

Nvidia寄りのとある企業のエンジニア曰く、
スペックはまったく同じでGTX285のチップの選別品だといっていた。
クロックがGTX285のほうが高いし、GTX285のほうが若干早いかも。

でもこのスレでもあるように、なんか計算結果がずれる可能性があるし、
メモリの多いTeslaが使えるなら使いたいよね。

DodでさえIBM CellじゃなくてPS3を購入してるんだからGTX285で十分でしょう。

>>398
まだ完成してないんじゃん

>>396
ASCII本の絵がわかりやすいけど、
共有メモリーは[16][17]にすれば問題ないと思う。
共有メモリーは無駄飯食っている用心棒みたいなもんで、
冗長にしてでも使えるときに使わない手はない。

皆さん、レス感謝です。
なるほど、GTX285とC1060差はさほどないんですね。
メモリと耐久性とサポート,計算結果
に関してはTeslaが有利というわけですな。
とりあえずGTX295のマルチスレッドはちょっと遠慮しようかなと思います。
もうちょい1人で悩んでみます。

個人的には285の2GBメモリ版がお薦めなのです。

【SIGGRAPH Asia 2009レポート】
東工大、スクウェアエニックスがCUDA実装事例を紹介
ttp://pc.watch.impress.co.jp/docs/news/event/20091221_338290.html

【SIGGRAPH Asia 2009レポート】
NVIDIAがGPUによるグラフィックスワークスタイルの変革をアピール
ttp://pc.watch.impress.co.jp/docs/news/event/20091218_336837.html

>>405
どっか、285の4GB版とか安く出してくれないもんかね。

　　　　 ...|￣￣ |　＜ Fermiはまだかね？
　　　／:::|　　___|　　　　　　 ∧∧　 　 ∧∧
　　/::::＿|＿＿_|＿　　　　（　｡_｡）.　 （　｡_｡）
　　||:::::::（　・∀・）　　 　　／<▽>　 ／<▽>
　　||::／ <ヽ∞/>＼　　 |::::::;;;;::/ 　|::::::;;;;::/
　　||::|　　 <ヽ/>.-　|　　|:と),__｣ 　 |:と),__｣
＿..||::|　　 o　　o　...|_ξ|:::::::::| 　 　.|::::::::|
＼　　＼＿_(久)__／＿＼::::::|　　　 |:::::::|
.||.i＼　　　　　　　　､__ﾉﾌ ＼|　　 　|:::::::|
.||ヽ .i＼　_ __ ＿＿＿＿ __ _.＼　　 |::::::|
.||　ﾞヽ i　　　 ﾊ i　ﾊ i　ﾊ i　ﾊ　|　　し'_つ
.||　　 ﾞ|i〜^~^〜^~^〜^~^〜

めるせんぬついすた、GPU版の使い方＠Windowsがやっとわかったわー

　　　　　|┃三　　　　　　　　　 /::::::::ハ､＼､::::::::＼＼::::::::::::',
　　　　　|┃　　　　　　　　　 　 ｉ:::::::ｲ　　｀> ー─--ミ::::::::::::|
　　　　　|┃　　　　　 　　　　　 {::::::::|　 　 ::＼:::／:::: 　＼:::ﾘ-}
　ｶﾞﾗｯ.　|┃　　　　　　　　　 　 ',::r､:|　 <●>　　<●>　　!> ｲ
　　　　　|┃　 ﾉ//　　　　　　　　|:､`{　 ｀> .::　 ､　　　 　 __ﾉ
　　　　　|┃三　　　　　　　　　　|::∧ﾍ　　／､__r)＼　　　|:::::|
　　　　　|┃ 　　　　　　　　　　　|::::::`~',　〈　,＿ｨｪｧ 〉　　l::::::》 ＜フェニミストはまだかね　辛抱たまらん
　　　　　|┃　　　　　　 　　　　　|:::::::::::::'､　 `=＝'´　　,,ｲ::ﾉﾉ从
　　　　　|┃三　　　　　　　　　ノ从､:::::::::`ｉ､,,　...　..,,／ |::::://:从

おれの物欲も辛抱たまらん、Core i7＋GTX260Mノートが欲しくて。
でもFermiのモバイル版が乗ったノートが出るのを松

故あってCUDAを使うことになったのですが自宅にはRadeonを積んだものしかありません。
コンパイルオプションでCPUでのエミュレーションができると何かで読んだのですが
これを利用すればRadeonの環境でも一応の動作はするのでしょうか？

CPUでのエミュレーションなので
radeon関係なくCPU上で動くよ

即レスどうもです。学校には設備があるので問題ないのですが自宅でもやりたかったので

あともう一つ質問よろしいでしょうか。VineLinuxでの動作は可能でしょうか?
気になったのですがどうもサポートしてないようなので

>>412
悪魔：「単精度なら大して変わらないんだから、買っちゃえよ！Fermi出たらまた買えよ！」

>>415
Linuxでサポート期待しちゃいかんだろう
あとエミュレーションはテクスチャ周りなど微妙に？おかしい模様なんで、あんまり期待しないほうがいい。
5000円くらいで適当なカード買うのが一番無難かもよ。

なんか、GeForceの在庫がどんどん無くなっていってない？
まもなくFermi出るな！これは！

C2050/2070よりも先にFermi搭載のGeForceが出ると発表しているしね

NVIDIA、“Fermi”採用第1弾GPU「Tesla 20」シリーズ発表 - ITmedia +D PC USER
ttp://plusd.itmedia.co.jp/pcuser/articles/0911/17/news039.html

まあ、実際に2010Q1に出るかは怪しいわけだが

C2050が2499ドルってのは自宅用で用意するには結構きつい値段なので、
Fermi搭載のGeForceがどの程度の値段なのかが今から気になっている

>>415
とりあえず玄人思考のGTX260でも買ってみよう
でかいからPCケースに入るか確認してからね

Device 0: "GeForce 8600 GTS"
Total number of registers available per block: 8192

CUDAやろうと思ってますが、レジスタの領域が少なすぎませんか？
__device__ void swap(float *d1, float *d2);
例えばこのような関数呼び出すのに引数とtempで計12byte、他にもthIDや作業用でローカル変数使うから、
最大のパフォーマンス求めようとすると実質スレッドは300個くらいになるんだけど…

こんなんだと何万、何千のスレッドとか無理じゃね？
みんなカーネル以外の関数は作らずにカーネルの中に処理を直書きしてるのですか？

>>421
レジスタとスタックとかバイトとレジスタ個数とかごっちゃになってないか色々と

4Gamer.net ― Fermi時代の製品展開が分かってきたNVIDIAのGPUロードマップ。DX11世代が出揃うのは2010年Q2以降に（Fermi（開発コードネーム））
http://www.4gamer.net/games/099/G009929/20091012001/

Fermi搭載のGeForceは3月くらいかなあ…？
この記事自体が2ヶ月前だからなんともだけどね

あんどうさん更新
ttp://www.geocities.jp/andosprocinfo/wadai09/20091226.htm

>>421
>Device 0: "GeForce 8600 GTS"
>Total number of registers available per block: 8192
直訳すると、「ブロックあたり使用可能なレジスター数：8192」

ブロックあたりのスレッド数は数千・数万もいらない。
128〜256程度でだいたいパフォーマンスは出る。（それを複数コアで何セットも同時に動かす）

あとカーネル１個でなんでもかんでもするもんでもない。（直列的な）処理ごとに分割すればいい。

>>422 >>425
うん…見直すと自分とても恥ずかしい事言いましたね(´・ω・`)
当初の疑問は理解して解決しました。

で、また別の話のGPGPUに関して質問です。
今DirectX使ってゲームを作っていて処理の重い部分(当たり判定とか)をGPUにやらせようと思っているんだけど、
ゲームって元々グラフィック描画でGPU使うので、それに加えてCUDAやプログラマブルシェーダーでGPGPUしたら相当パフォーマンス落ちますか？

目的はCPU使用率を抑えたいんだけど、ゲームのような1フレーム毎にGPGPUをするのは実用で使えうるものなんでしょうか。

当たり前の話だけど、GPUをフルに使うようなレンダリングをしていたら遅くなる
フレーム枚にGPGPUの処理を入れるのはそんなに問題ない
ただし処理の重い部分ってのが上手く並列化できるアルゴリズムであることや、そこそこの量あることが求められる
分岐ばっかだったり、処理よりメインメモリ-VRAM間の転送のほうが時間かかると逆に遅くなる
ようは、本当にGPGPUを使ったほうがいい処理なのかを吟味する必要がある

質問です。

C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK\C\src\template\template_vc90.vcproj
を開き、ビルドした所、「入力ファイル 'cutil32D.lib' を開けません。」というエラーが出ました。「x64」では、スキップされてしまいます。
>>34 とほぼ同じ状況なので、このスレに書かれている事をしてみたのですが、変化ありません。

また、リンカの設定でcutil32D.libをcutil64D.libに変えた所、x64でビルド出来るようになったのですが、
「モジュールのコンピュータの種類 'x64' は対象コンピュータの種類 'X86' と競合しています。」と別のエラーが発生しました。

環境は
windows7 64bit
Visual C++ 2008 Express Edition
GTS250
CUDAのドライバ、toolkit、sdkは2.3でwin7 64bitの物を入れました。

改善策がありましたら、ぜひお願いします。

>>428
Expressは64bitでコンパイルできないかも？
http://www.microsoft.com/japan/msdn/vstudio/2008/product/compare.aspx

>>429
みたいです。やはりproを買うしかないのでしょうか…
レスありがとうございました。

>>428 >>430
答えるのに十分な情報量を持った質問を久しぶりに見たような気がしたｗ

proが良いけどねぇ。
今は最適化コンパイラもついてるし、standardで良い気もする。
（standard使ったことはないので参考まで）

PlatformSDKにx64コンパイラあったような

落ち着いて、なぜこういうエラーがでるのか？
と、考えるしかないよ。

CUDAが使えるGPUが載ってるか否かの判断をプログラム上で行うには
どうするといいのでしょう？

OpenCL入門 - マルチコアCPU・GPUのための並列プログラミング

が出版されるそうな
アスキードットテクノロジーズの記事書いた人たちらしいが

GPGPUでもっともメジャーなのはCUDAだと思うけど
OpenCLがこの本の売り文句通りスタンダードになるんだろうか…？

よくわからないけど、無関心でいるわけにも行かないのでとりあえずポチってみる

感想あったらアマゾンのレビューに書く

>>434
プログラム上でその判断をやる必要というのがよくわからない
コマンドラインでやるのでは駄目なのか

>>434
それはやはり
cudaGetDeviceCountして、CUDAデバイス個数調べ
cudaGetDevicePropertyをDeviceCountまわして、
.major、.minorでバージョンチェック
.multiProcessorCountでコアの個数調べ
して使うしかないんじゃ。

>>436
どうも

>>435
CUDA依存部分を共有ライブラリの形でプログラム本体から切り出しておいて、
実行時に動的にリンク出来たらとか考えてます。

…CUDA版バイナリと、CUDA無し版バイナリを用意して、インストール時に
選ばせればいいのか。

>>435
スタンダードはどう考えてもDirectComputeだろ。
いまのWindowsのシェアとDirectXの普及率から考えてCUDAがスタンダードになるにはポテンシャルが違いすぎる。

CUDAはスタンダードになれる可能性はあるけど、OpenCLはたぶん廃れる。
OpenCLって所詮はCUDAのラッパでしょ？
ラッパライブラリがスタンダードになった事例ってあんまり知らないんだけどさぁ。

Appleはラッパ好きだよな。
ObjCもラッパではないにしろマクロ駆使してるだけで、骨はCに変わりはないでしょ？
Appleは声は大きいけど技術が無いから、OpenCLもあんまり期待してないよ。
所詮ラッパライブラリだから、コンパイルしたらどのGPUでも使えるってわけではないし。
OpenCLのDLLロードしたらどのGPUでも同じバイナリでおｋみたいになったら使いやすいけどね。

てか俺が作っちゃう？ｗｗｗ

>>437
436でCUDA使えるかどうか確認して、
プログラムとしてはCPU用ルーチンとGPU用ルーチンを両方持っておく
でいいんじゃね？(やはり、動作確認用にCPU/GPU双方で試すのは必要と思うし)

>>438
おれ、
ttp://openvidia.sourceforge.net/index.php/DirectCompute
見て「うざっ！！」と思って挫折した。　ダメな俺

pyCUDAってどうかな
DriverAPIを使ってモジュールを実行時に作成、ロードするみたいですね。

3次元のデータをデバイス側に送って計算したいんですけど、
>>328の
＞固めてまとめて渡しちゃったほうがオーバーヘッドが少ないと思います。
＞トータル何列あるよ、は別にパラメータで渡す。

というのはどこにどのように記述すればいいのでしょう？

>>445
もちろん、少ないならデバイス関数の引数として渡す。パラメータ複数個まとめた配列で渡しても良いし。
計算の対象、計算の結果、に加えて、計算のパラメータ、も同様にホスト側で確保→cudaMemCopyするわけ。

>>446
ありがとうございます。

物分かりが悪くて申し訳ないのですが、
cudamemcpyは確保した領域にデータをコピーする使い方しか思いつかないので、

＞計算の対象、計算の結果、に加えて、計算のパラメータ、も同様にホスト側で確保→cudaMemCopyするわけ。
を簡単に例で説明していただけないでしょうか？

えっ　なぜ分からないんだ？　たとえばこんな感じだよ

__device__ calculate(float * input, float * output, int * params){

:
if (params[i] == 0)
　　　output[i] = func_A(input[i]);
if (params[i] == 1)
　　　output[i] = func_B(input[i]);

　__syncthreads();

>>448
ありがとうございます。
なんだか変な方向に考えすぎていました。

CUDAってGPUカーネルの中で関数読んで値をreturnさせることできたっけ？

あと__global__にしないとホストから関数呼べなくね？

calculateが__global__で、func_A()、func_B()が__device__ですたorz

Fermi搭載GeForceの発売がまた不透明になったとか

⇒を何らかの演算として　0⇒256 256⇒512 512⇒0　といった感じにしたいのですが

a = (a+256)%513;

とする以外思いつきません。けど剰余は遅いので違う計算にしたいんです。
いっそのことテーブル作ったりif文にしたほうが早いのでしょうか？
馬鹿すぎてわかりません。誰か助けて

>>453
Cudaだと四則演算はほとんど処理時間には影響しないし、その式をつかうといいんじゃない？
っていおうとおもったけど、a=512のとき0にならなくね？

257⇒0　258⇒1　…　でいいの？
でもそうすると512⇒0がヘンだけどな。

int val;
val = a + 256;
if (val > 512) val = val - 513;

でいいんじゃないの？
というかC言語の疑問をCUDAスレで聞くのはまちがっとる、ここは人口が少ないぞ

じゃあCUDAっぽく

const int TABLE[512] = { 256, ... , 512, ... , 0 }

__device__ int f(int a){
return TABLE[a];
}

レスありがとうございます
わけわかんないこと書いてすみません
テーブルとif文試して出直してきます

>>453
b⇒aとして、
a = (b+256)%768;
でいいのかな？
0⇒256
256⇒512
511⇒767
512⇒0

釣りだよな？

申し訳ないっす

ランダム関数の作成で悩んでます。
・CPUで作成した乱数リストをGPUに転送しない
・1スレッドの中で何度も呼ばれる
・グローバルメモリは使わない。スレッドIDやclock()をseedとして使う

この制約で良い乱数生成アルゴリズムを教えて下さい。
MTGPっていうのを参考にしようとしましたがムズ過ぎて挫折…まずコンパイルが通らず(´・ω・`)

おれMTGPコンパイルできたよ。GPU側で作成してそのままGPUで使えるし。
（シンボルコピー）

使ったファイルはこれだけ
mtgp32-fast.h
mtgp32-cuda.cu
mtgp32-fast.c
mtgp32-param-fast.c
make_single_random()とかの末尾を適当に書き換える。

誰かＯｐｔｉＸやってる人いませんかー？

CES 2010: NVIDIA Shows GeForce GF100 Fermi Video Card - Legit Reviews
http://www.legitreviews.com/news/7122/

3/2に本当に出るのかしらん

気にするな、出たら買う、出るまではCUDAの修行をする。
単純でｲｲじゃないか。

>>461
Perlinノイズやってみてよ
ttp://mrl.nyu.edu/~perlin/doc/oscar.html
Cのソースが有る

Fermiもダメそうってどこかのブログで見た
NVIDIAまじやばい

ようやくCUDAに慣れてきたところなのに、StreamSDKとかOpenCLとか
DirectComputeとかまた勉強しないといけないの？　
辛いのうｗｗプログラマは辛いのうｗｗｗ　　　orz

windows7 64bit、Visual C++ 2008 Pro環境に、CUDA　toolkit、sdk 2.3をインストールして、
C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK\C\srcにある、
_vc90.slnファイルを開いて、問題なくビルドできます。
exeファイルの出力先は、デフォルトでC:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK\C\bin\win64\ですが、
これ以外の出力先に変更する方法を教えてください。

プロジェクトのプロパティで変更できるよ
$(OutDir)\$(ProjectName).exe
これを丸ごとフルパスで指定してしまいなさい。

XP,32bitでコンパイルしたプログラムを7,64bitで動かそうとしたら動かなかった･･･

こういうものなのか

32/64ビットのcudart.dllが必要

>>471
thanks

現在CUDA SDKを導入しVisual C++ 2005 Express Editionを使用しCPUでのエミュレーションを行っています。
いくつかのサンプルは動かせたのですがparticlesなど複数のサンプルでコンパイルはできているようですが実行すると以下のようなエラーが出ます

device emulation mode and device execution mode cannot be mixed.

エラーが出ている行は

cutilCheckMsg("cudaMalloc failed");
cutilSafeCall(cudaGLRegisterBufferObject(vbo));
cutilSafeCall( cudaMalloc3DArray(&noiseArray, &channelDesc, size) );

のように「cutil〜」になってます。オプションから設定しているのでデバイスエミュレーション自体はうまく行ってるようなのですが
エミュレーションを行う際はソースを書き換える必要があるのでしょうか?

メッセージのとおりエミュONOFFまぜてる
またはリビルドしてないんじゃ？

プロジェクト設定じゃなくてビルド構成のほうでエミュレーションにかえられないっけ？

メッセージの内容は把握できています。
リビルドも多分できてると思います。
エミュレーションの設定は
プロジェクトのプロパティ→CUDA Build Rule→Emulation Mode はい
にしてます。ほかの方法はよくわからないのですがほかの方法があるのでしょうか?

ビルド構成のほうからエミュレーションの設定ができました。問題なくビルド&実行できました。ありがとうございます。

ただFPSが0.1なんでほぼ無意味でした。

エミュレーションなんてHello world以上のことをやるもんじゃないな

処理する画面領域を１０００分の１ぐらいに限定して、バグを調査するくらいのことはできる。

なんだっけ、エミュじゃなくて実機で実デバッグできるやつ
早く来てほしいな

CPU版アプリの移植の際に、エミュはデータ検証に使える。

-G　オプションって使える？WinXPで
「'assembler' は、内部コマンドまたは外部コマンド、
操作可能なプログラムまたはバッチ ファイルとして認識されていません。」
と出てビルドできないけど。

--deviceemuもつけないとだめなの？

CUDAのバージョンを書き出しておきたいのですが、
CUDAのバージョンを調べる関数か何かありませんか？

>>484
>>436　でどう？

こっちを要求している気がする。
cudaRuntimeGetVersion()
cudaDriverGetVersion()

WindowsでCUDAを使っています。
ループが一定回数超える度に、coutで経過を表示しています。
Linuxに比べて遅かったので、いろいろ試してみたところ、
DOSウインドウの前で、マウスを動かすと、速くなります。
垂直同期か何かで引っかかっているのかと思い、設定で切ってみましたが
状況が変わりません。
どなたか似たような状況になって解決された方はいないでしょうか？

追加報告です。

cudaMemcpyでcudaMemcpyDeviceToHostでコピーするところで異常に時間がかかっているようです。
この1行を消すだけで、100秒ほどかかっていたのが2秒にまでなりました。
逆方向のcudaMemcpyHostToDeviceは問題ありません。

CUDAプログラムがGPUで計算等をしている間は画面表示の更新が遅くなります。
１回のループにどれくらいの時間がかかるのか、一定回数がどれくらいかによりますが、
経過の表示される間隔が短いとその影響を受ける可能性があります。

あるいは、マウスを動かすことでCPUに負荷をかけないとパワーセーブモードに入ってしまうとか・・・・。

>>488
転送に１００秒かかるのは現実的にありえないので、恐らくcudaThreadSynchronize()等でカーネルの実行完了待ちを
していないのではないかと思います。カーネル呼び出しそのものは一瞬で終わりますが、
後続のcudaMemcpy(.....,...,...., cudaMemcpyDeviceToHost)時に暗黙的にそれまでに呼び出したカーネルの実行完了を待っています。

・・・・・ということかどうかは分かりませんが。

>>489,490さん

早速の返信ありがとうございます。

cudaMemcpyでcudaMemcpyDeviceToHostを消したプログラムで、
プログラムの一番最後にだけ結果を取り出してみたところ、
正しく計算されていました。

次に、カーネルの実行部分のみをコメントアウトした場合、
やはり非常に時間がかかってしまいました。

やはり、GPUからCPUへのデータ転送に（というよりなにか同期の部分の様な気もしますが）
時間がかかっているようです。

計算用マシンなので省電力の設定はしていません。

>>491

> cudaMemcpyでcudaMemcpyDeviceToHostを消したプログラムで、
> プログラムの一番最後にだけ結果を取り出してみたところ、
> 正しく計算されていました。

これを読むと、そもそも何度もDeviceToHostの転送をする必要がない処理という解釈で
よいのでしょうか？

DeviceToHostが遅くなる理由はハードウェア的なものから色々あります。
マザーボードを交換したという人もいました。

SDKに含まれているbandwidthTestの結果はどうなっていますか？
"--memory=pinned" を付けた場合と付けなかった場合をそれぞれ調べてみてください。

Device 0: GeForce GTX 285
Quick Mode

Host to Device Bandwidth, 1 Device(s), Paged memory
Transfer Size (Bytes) Bandwidth(MB/s)
33554432 4340.2
Device to Host Bandwidth, 1 Device(s), Paged memory
Transfer Size (Bytes) Bandwidth(MB/s)
33554432 4193.7
Device to Device Bandwidth, 1 Device(s)
Transfer Size (Bytes) Bandwidth(MB/s)
33554432 119962.3

--memory=pinned

Device 0: GeForce GTX 285
Quick Mode

Host to Device Bandwidth, 1 Device(s), Pinned memory, Write-Combined Memory Enabled
Transfer Size (Bytes) Bandwidth(MB/s)
33554432 5758.1
Device to Host Bandwidth, 1 Device(s), Pinned memory, Write-Combined Memory Enabled
Transfer Size (Bytes) Bandwidth(MB/s)
33554432 5471.9
Device to Device Bandwidth, 1 Device(s)
Transfer Size (Bytes) Bandwidth(MB/s)
33554432 119852.3

結果はこんな感じです。
特に片方だけ遅いというわけでもないようです。
Linuxで起動した場合、遅くなるという問題は起きませんでした。

DeviceToHostが必要なのは結果を取り出すときだけで、
計算自体には必要ありません。

似たような話をnvidiaのforumで見つけました

http://forums.nvidia.com/lofiversion/index.php?t68705.html
http://forums.nvidia.com/lofiversion/index.php?t73047.html

返信は結構あるのですが、解決策はよくわかりません。

・CUDA内での計算処理、ループ数とループ内のステップ数、スレッド数ブロック数、だいたいどんなもんすか
・HostToDevice、DeviceToHostで転送するメモリのサイズはどのくらいですか

とりあえず

タイマー1,2作成
cudaMemCopy(HostToDevice)
タイマー1,2起動
<<<>>>カーネル呼ぶ
cudaThreadSyncronize()
タイマー1停止
cudaMemCopy(DeviceToHost)
タイマー2停止

してほんとにMemCopyなのかカーネルなのか確かめてみる必要は。

コピーが細切れでforループでやたら回数呼んで転送してたりしない？
１発で全部転送するのと1/1000ずつ1000回コピーするのとでは、
かかる時間に雲泥の差があるけど。

PCを再起動したら、上記の問題は出なくなりました。
お騒がせしました。

おいｗｗ　終わりかよｗｗ

cutil64.dllってどこに置けばいいの？

GPU側で配列の合計を求める方法を教えてください

サンプルのreduction見れ

syncthreadsって異なるブロック間のスレッドも同期されるのでありましょうか？

>>504
programing guide よめ

Best Practices Guide
もいいこと盛りだくさん書いてあるよ

cudaでは多次元配列の使用は推奨されてないの？

pthreadとの組み合わせでうまく動作せずに困っています。

引数として渡す構造体に、float * device_Xを用意して
for(int k=0;k<NUM_THREADS;k++){
cudaMalloc((void**)&targ[k].device_X,
sizeof(float)*(hostX.nRow)*(hostX.nCol));
cudaMemcpy((void*)targ[k].device_X,
host_X,
sizeof(float)*(hostX.nRow)*(hostX.nCol),
cudaMemcpyHostToDevice);

pthread_create(&handle[k], NULL,
compute_thread_func,
(void *)&targ[k]);
}
と渡すと、pthread内で
cudaMemcpy(temp,argument->device_X, //argumentは渡されたvoid* argをキャストしたもの
sizeof(float)*(nRow)*(nCol),
cudaMemcpyDeviceToHost);
としてもうまく取り出せません。（中身がぐちゃぐちゃになる）

同様の処理をpthreadで呼び出す関数内で行うとうまく動作します。
スレッドごとにメモリを取るのだから、中でやってもいいのですが気になります。

何か原因になっていそうなことは無いでしょうか。

>>508
＃1個のGPUに複数のスレッドからアクセスしたいアプリケーションなのかよく分かりませんが・・・・。

CUDAはスレッド毎にGPUのリソース（コンテキスト）を管理しています。
従ってこの例では
子スレッドのcompute_thread_func内の処理は
親スレッドがcudaMallocした領域にアクセスできません。

>>509
ありがとうございます。
複数のスレッドから一台のGPUへのアクセスです。
（後々、スレッドごとに一個のGPUにしたいのですが）

スレッドごとに管理しているのははじめて知りました。
解決しました。

「Picture Motion Browser VAIO Edition」を追加。
超解像HDアップコンバートやノイズ除去、手ぶれ補正などの編集や動画作成を、NVIDIA CUDAによるGPU演算を使って行なえる。

　　　　　　　　　　　　　　　　　GT200　　　　　　　　　Fermi
トランジスタ数　　　　　　　　　14億　　　　　　　　　　30億
倍精度浮動小数点　　30FMA Oｐｓ/clock　　　　　　256FMA Oｐｓ/clock
単精度浮動小数点　　　　　　 240　　　　　　　　 　　512
シェアドメモリ(SM)　　　　　　　16KB　　　　　 　　 　64KB
L1メモリ　　　　　　　　　　　　　無し　　　　　　　　48KB or 16KB
L2メモリ　　　　　　　　　　　　　無し　　　　　　　　　　768KB
アドレス幅　　　　　　　　　　　　32bit　　　　　　　　　64bit
複数のカーネル実行
　　　 　　無し　　　　　　　　16個まで（GigaThredエンジン)

＊L1キャッシュ搭載
GT200では16KBのSMが32個あり、それぞれCUDAコアを8個割り当てられ、L1キャッシュが無し。
Fermiでは64KBのSMが16個。それぞれにCUDAコアが32個割り当てられ、SMから16KBか48KBのどちらかをL1キャッシュとして使用可能。
GT200に対して、3倍のSMと16KBのL1が使用可もしくは同じサイズのSMと48KBのL1が使用できるようになった。これにより、今までCUDA化できなかったプログラムの対応を増やし、さらに高速化もできる。
各CUDAコアに含まれるFPユニットは倍精度浮動少数演算を強化し、GT200に対し8倍の能力。

＊L2メモリの搭載　グローバルメモリに対する高速化。
＊C++をサポート
＊複数のカーネルの動作をサポート
SM内部のパイプラインを強化。SFUが複数に分けられたのでタスクとデータをより効率化。スレッドスケジューラを2個。

＊双方向PCIE通信
GT200ではPCIEバスの送受信をどちらか片方しか一度に実行できず、理論値8GB/s・実測4〜5GB/s程度だが
Fermiでは双方向通信が可能になり12GB/sを実現

＊新しいメモリコントローラ　FermiよりGDDR5まで対応し、ECCにも対応する。

＊コア内部を各部でモジュール化
設定された価格帯や用途ごとにコアを設計しなおさず、機能をカットオフしたり追加したりできる。
SM単位でCUDAのコアを減らしたり、D3DやOpenGLなどの固定ハードウェアが不要なTeslaなどでオフになったりする可能性もある。

　　　　　　　　　　　　　　　　　GT200　　　　　　　　　理想のFermi　　　　　　　　現実のFermi
トランジスタ数　　　　　　　　　14億　　　　　　　　　　30億　　　　　　　　　　　　　　30億 　
倍精度浮動小数点　　30FMA Oｐｓ/clock　　　　　256FMA Oｐｓ/clock　　　　224FMA Oｐｓ/clock
単精度浮動小数点　　　　　　 240　　　　　　　　 　　512　　 　　　　　　　　　　　448FMA Oｐｓ/clock
シェアドメモリ(SM)　　　　　　　16KB　　　　　 　　 　64KB　　　　　　　　　　　　　　〃
L1メモリ　　　　　　　　　　　　　無し　　　　　　　　48KB or 16KB　　　　　　　　　　〃
L2メモリ　　　　　　　　　　　　　無し　　　　　　　　　　768KB　　　　　　　　　　　　　〃
アドレス幅　　　　　　　　　　　　32bit　　　　　　　　　64bit　　　　　　　　　　　　　　　〃
複数のカーネル実行　　　 　　無し　　　　　　　　16個まで（GigaThredエンジン)　14個まで（GigaThredエンジン)

OpenCL待ちしてたけど、GPUの対応はCUDAだけっていう開発ばっかな
Fermiで本気出すから

GT100ってなんかメモリアクセスのレイテンシがすごくでかくない？
それを隠蔽するために、L2キャッシュが128Kになっているけど、
コヒーレンシとかどうなのかな？
やっぱりGPUはデータ並列なアプリしか向いていないのかね。

Fermiのグローバルメモリのレイテンシが遅いってのは何処から来た情報？GDDR5に対応するんだから、帯域は大きくなりそうだけど。

>>512
SFUは今のGT200でも複数ある。
ただ、SFUが外部のパイプラインに分かれたからSFUを使っているときに違う種類の演算ができるようになったってことじゃね・。

>>515
最適化しなければいけない

>>516
これをみて思った。
http://pc.watch.impress.co.jp/img/pcw/docs/343/352/html/kaigai-01.jpg.html
でもGT200と一緒なんだね。
http://pc.watch.impress.co.jp/docs/2008/0617/kaigai_10l.gif

いずれにしても、何でGPUがあんなにレイテンシが大きいのかがわかったよ。
だから、coalesced accessが必要なのね。

>>518
>いずれにしても、何でGPUがあんなにレイテンシが大きいのかがわかったよ。

俺のために説明してくれ！

GPUのメモリのレイテンシが大きいというよりも、CPUの場合と同じぐらいはかかるといった方がいい

CPUだとメモリチャンネルが精々2〜3チャネル位しか無いが、GT100だと12チャネルもあるから、DRAMのチャネルのスイッチングに
それ以上かかると思っていい。あと、仮にGPC1?がDRAM0チャネル上にあるデータと、
DRAM11チャネル上にあるデータを同時にアクセスする場合、当然レイテンシがかかる。
もし、GPC1がDRAM0、DRAM1だけのデータアクセスですむなら、CPUと変わらない。
最適化が必要というのはそういったところだと思う。

MTGPがコンパイルできません。
>>463の言う、4つのファイル使って試したところ100個近くコンパイルエラーが出てしまいます。
あと「末尾を適当に書き換える。」の意味がよく分からんどす…　
エラー内容は
1>\略)\mtgp32-fast.h(117) : error C2054: 'inline' の後に '(' が必要です。
1>\略)\mtgp32-fast.h(120) : error C2057: 定数式が必要です。
1>\略)\mtgp32-fast.h(120) : error C2466: サイズが 0 の配列を割り当てまたは宣言しようとしました。
1>\略)\mtgp32-fast.h(120) : error C2085: 'mtgp32_do_recursion' : 仮パラメータ リスト内にありません。
1>\略)\mtgp32-fast.h(121) : error C2061: 構文エラー : 識別子 'inline'
1>\略)\mtgp32-fast.h(122) : error C2054: 'inline' の後に '(' が必要です。
...

WinXP32bit、VC++2005、CUDA SDK2.3
該当箇所のコード見てもどこが悪いのか分からない…助けて…。

単にパイプラインステージが何百もあるからレイテンシがでかいのだと思ってました

>>522
あ。おれか。
元のソースと自分のソースでdiffしてみた。
・inline　を全部削除してみて。
・末尾を適当に書き換える。は、元だとプリントして捨ててしまっているので、
　利用したいように書き換えてねと。
おれはmake_uint32_random(d_status, num_data, hostbuff);　として
hostbuffを自分で使えるようにしました。

>>523
正解
メモリチャンネル云々は殆ど関係ない

CPUはメモリとのクロック差のため
GPUは長大なパイプラインのためレイテンシがデカイ

別にパイプラインが深いわけではなくてバッファが大きいだけなんだけど。

200-400段は十分に・・

>>524
レスサンクス。
inlineを削除したら見事にエラーが亡くなりました(`・ω・´)
これから自分のソースに組み込もうか…と状況に入れます。

重ねて質問申し訳無いですが、hostbuffの名前からしてMTGPの乱数は一度バッファにぶち込んでから使うという事になるのですか？
>>461ようにカーネルの中で特に意識せず使いたいのだけれども…
MTGPがglobalﾒﾓﾘやﾃｸｽﾁｬを使ってるのなら、Cのrand()のように連続で呼び出して使えたら良いなと思ってるんですが無理ですかね？

>>528
おれはホスト側で使いたいので、make_uint32_randomの末尾でホスト側のバッファにコピーして使ってる。
デバイス側で使いたいなら、make_xxxxx_randomにd_buffをmain()側でとって渡し、
make_xxxx_random()内ではfreeせずにそのまま。すると次にmain()内でデバイス側関数を呼び出すときに
そのまま渡して使える。　※スレッドが変わってはいかんので注意

CUDAを使って大量のデータのごく簡単な統計値（最小、最大、平均、標準偏差）を計算したいんですが、何かいいサンプルとかありますか？

CUBLASにあるんじゃ。

>>530
そういう、「最後が一個」な演算は、reductionで探すと有る気がする
並列性を有用に使うのがなかなか面倒らしい

OpenGL glutで時系列で変化する数値計算の結果を表示しようとしています。
GPU 2個で計算する場合、コールバック関数内でスレッドを生成することになると思うのですが、
この度（毎フレーム）にホストからグローバルメモリへデータ（時間による変化なし）を
コピーしなくてはいけません。
glutMainLoop(); 呼び出し前に両GPUにコピーした場合は、
コールバック関数内で生成したスレッドにグローバルメモリのポインタを渡しても参照できません。
データのコピーを一度だけですます方法はないでしょうか？ご教示ください。

スレッド生成は ttp://tech.ckme.co.jp/cuda_multi.shtml を参考にしました。

>530
それぐらいだったらたぶんCPUで実行した方が速いよ

>>533
GPU0用、GPU1用のスレッドをglutMainLoop()呼び出し前に生成するべきかと。

thread数　=　N
Mat[N];
id = threadIdx.x;

if(id%2==0){
演算Ａ
Mat[id]にコピー
}else if(id%2!=0){
演算Ｂ
Mat[id]にコピー
}

のようなプログラムを組んでいるのですが、結果をみると最後の2つの要素と同じ計算結果が全体に入ってしまいます。

N=16なら
14の結果が0,2,4…12に
15の結果が1,3,5…13に入ってしまいます。

どこに問題があるのでしょうか

>>536
とりあえずC言語から始めようか

>>530

俺もこれ知りたい。
N社のSDKを見ても、画像処理のサンプルとかたくさんあっても、単純な総和計算とかないもんだね。やはり向いてないからか・。。。

ご丁寧にCのほうが早いとか教えてくださる人もいるがｗ
マルチパスの画像フィルターとか、
デバイスの中にある中間結果を作って、
その統計値を、次のパスの計算で使ったりするのが常套手段だから
ここでいったんホストのほうにコピーするとボトルネックになってしまう。
デバイスの中のデータの統計値を出すライブラリとか作ってくれると本にありがたいんだが

>>530, 538
SDKのreductionサンプルが参考になると思う。 確かpdfのわかりやすいスライドが一緒に入っているはず。

CUDAのアーキテクチャって、総和とかのreduction系演算は
不向きとまでは言わないまでもちょっと頭をひねらなきゃいけない分野だよね。

演算途中で計算領域が被ってただけでした。失礼しました。

>>539
そうだね。
沢山のデータがあるなら、最終はCPUにやらせても良さそうだけど、
転送がボトルネックになるなあ。
それだったらGPU内でやらした方が速いか。

あとI/O系があるとだめだよね。
ダブルバッファやトリプルバッファが使えるようなアプリならいいんだけど。
そうなると画像、映像系ばかりになってしまうなあ。

処理に割り当てるmultiprocessorの数を指定とか出来ませんか？
出来ればgridサイズ変えずに
そもそもOSもGPU使う場合割り当てがどうなってるのか分からないんですが

>>542
CUDAプログラムが実行中はOSはGPUを使えないので全てのSMを使用しても問題ありません。

5秒以上使ったらエラーが出たよ。

SMを1個しか使用していなくても5秒以上１つのCUDAカーネルが実行されるとタイムアウト。

へー、なるほど。
あとgridサイズ小さくする以外に使うSM数を制限できますか？

できませんし、その必要もないでしょう。

>>547
そうですか、ありがとうございます
並列計算の研究の一環でCUDA使ってるんで
並列数の変化でデータが取れたら嬉しいな、という理由です

>>548
素直に描画用とは別にGPUを用意しましょう。

総和を取る処理は私も書いているけど、仮想的に二次元にデータを配置して、縦に集計するところまでをGPUでやっている。
残りはCPU側に吐き出してからCPUではSSEを使って集計。

>>548
1番最初に実行が始まるCTAの配置は予測可能なので、
その中で使わないSMに割り当てられたCTAはコア内でダミーの演算を長時間繰り返すことでそのSMを占有し続ける。
こうすれば本来の計算は残りのSMでのみ行われるようになる。

通常時間計測できるのは全CTAが完了するまでの時間なので以下のどちらかを。
１）最後のCTAがdevice memoryにマップされたhost pinned memoryに何か書き込む。
２）ダミー演算の繰り返し回数を段々少なくしていき、計測される時間の減少率が変わるところを探す。

なお、全SMを使わないとメモリアクセス効率も落ちるのであまり面白くない。

> 仮想的に二次元にデータを配置して、縦に集計するところまで
なるほど。ふつくしい。ｎ次元でCPUでは最後の一次元だけやらせれば、対称になるな

>>549
別に用意してこの場合メリットってありますか？
あと総和ならReductionで組んだけど今回はGPU単体の話だったんで最後までGPUでやりました
デバイスホスト間の転送時間って馬鹿にならないイメージあるんですが、CPUにやらせたほうが速いんですかね？
まあ最後のちょっとだから誤差の範囲な気がしないでもないｗ

>>550
結構シビアですね、
直接的な方法が無ければgridサイズ縮めてIdx周り弄ろうと思います
↑の方法で弊害無いですよね？；

少し前にMTGPについて質問した者ですが、どうやら自分の要求とズレた感じでした。
thread 256、block 1でmake_uint32_random()するとバッファに768個のデータが生成されるが、でっていう…。

これはﾒﾙｾﾝﾇﾂｲｽﾀの高周期な乱数列の１部分って事で、本当にあの高周期な乱数を使いたいならその分のメモリが必要だということなのかな。
独自の乱数ジェネレータを作って、その768個の中から１つを取り出して…みたいな事をやり始めるとまた性質が変わってしまうし、本末転倒に。

結局、カーネルの中でシミュレーション目的の使用方法としては微妙だったので絶望。。。
スレッドID毎に使える線形合同法(遷移型)の乱数として使えるようになんとかできないものか…

え、おれ100万個単位で作って使えてるよ。
int num_data,　にちゃんとでかい数与えてるかな

>>554
いや、num_dataの数を変えて生成される個数の事はあまり重要じゃないのよ…
問題はカーネルの中で使おうとした時、バッファに作成された乱数が並んでいる形態では微妙なのです。

例えば、100万個作ったとして256のスレ数で使うなら使用部分を分割する事になりますよね（thID==0は、バッファのindex0〜約4000、という感じ）
いや、各スレッドは100万個のうち初めのindexだけseedとして決めて、あとは順次indexを増やして使っていく感じでもいいけど、
両者とも乱数列の周期はバッファのサイズに依存してしまいます。

一方、よくある{x=x*1103515245+12345; return x&2147483647;}このような方法は、アルゴリズムが優秀だとxが4byteなら最大で2^32の周期を持ちますよね。
今の状態のﾒﾙｾﾝﾇﾂｲｽﾀで2^32の周期を得ようとしたら、どんだけ大きいバッファが必要かっていう…
精度の良い乱数という点では利点ありますが、globalﾒﾓﾘを物凄く使う割にはなんだかなぁ…という複雑な気持ち。

MTGPの形態を知らなかった自分が悪いんだけど、要はこれ乱数生成を並列化して高速にしたもので、
実際にパラレルな計算で使う用にはなりえない事が分かりました。
自分の要求としては、GPU Gems3の「CUDAによる効率的な乱数の生成と応用」がﾁﾗｯと見た感じでは合致してたので、今からこっちを参考にしてみます。
長文ｽﾏｿ。>>554にはとても感謝している。こんな結末で申し訳ない。

＞乱数列の周期はバッファのサイズに依存してしまいます

横から失礼しますが、
for( ; ; )
{
make_uint32_random(, , d_buff); //デバイスメモリに作らせて残す
my_kernel_func<<<>>>(d_buff, , , , ); //MTGPで作った乱数を消費
}
こんな感じとして、256スレッドが一回に8192個の乱数を使う
→make_uint32_randomは2097152個の乱数を生成する

で良いのでは?　make_uint32_random()　は複数回呼び出すごとに
前回のmtgp32_params_fast_tの続きから処理するわけで、周期は
2^11213-1　でしょう。÷256しても2^11205とかだ

stupid questionですみませんが、VC++ 9.0で
1. .cppと.hのように、.cuファイル内でインクルードしているファイルが更新されたら.cuを再コンパイル対象にしたい。
2. ptxを出力したい。nvcc -ptxでは無理でした。
以上について教えてください。お願いします。

>>557
コマンドラインから　nvcc -ptx　*.cu とやっても駄目？

-keepでいいんじゃね
1.についてはSource Dependenciesを個別に設定すれば一応できる

>>558
XP Pro 32bitとVista Ultimate 64bitの両環境で、
コマンドラインからnvcc と打つとcl.exeがないと怒られます。
プロジェクトのプロパティを参考にパスとincludeを指定してやっても
エラーは控えていませんがコンパイルできません。
VC使ってる人はどうしてるんでしょう？

-ccbin で指定しても駄目ですか？

もちろんVSコマンドプロンプトから打ってるよな

>>562
ふつうのコマンドプロンプトを使っていました。
どうもお騒がせしました。

普通にC++のコードを書けるようになってからじゃないと、学習効率が悪くてどうしようもないぞ。

>>559
俺は556じゃないが-keep知らなかった。ありがとう。

GPU Challenge 2010
ttp://www.hpcc.jp/sacsis/2010/gpu/

自由課題もあるそうな

まあ俺は学生だから規定課題でるけどな
しかしCellとかに比べてあんまり最適化するとこないな

どのくらい参加するんだろう？

俺もとりあえずエントリーしようかな。

自由課題の方で、パターン形成する発展方程式とかの数値計算すると、絵的にも面白そうなの出来そうじゃない？

jpegをデコードするライブラリもしくはCUDAのコードはどこかにありませんか？

GPU Challengeの課題が増えた
メルセンヌツイスタと言われるとHack the Cellを思い出すな

SLI環境で、プログラムを2つ実行した場合、それぞれ別のGPUを
利用する方法を教えてくれ

> SLI環境で
えっ、ひとつにしか見えないんじゃないの??　出来るの？

>>572
cudaSetDevice()でそれぞれ0と1を指定する。

>>573
東工大青木先生はGeForce4つ並べてた
OpenMPで並列化していたと思う

>>575
CUDA版OpenMPてあるの？

CUDA版と言えるOpenMPはない
OpenMPのスレッド指定とCUDAのdevice指定を組み合わせただけのこと

そういうことか、今度挑戦してみようかな

> GeForce4つ並べ
と
>　SLI
は違うんじゃね？　ケーブルで繋ぐのがSLI･･･かな

OpenMPを使えば複数のGPUを簡単に使えるのですか？
やりかたをおしえてくれろ

CUDAで顧客向けプログラムを作成しています。
CUDAプログラムの配布先には本体以外に何を配布すればよろしいのでしょうか？

>>581
GeForce人数分

TeslaかQuadroを配っておけば良いと思うよ

cutil使わなければcudart.dllだけでよろしよ

>>579
内部でケーブルでつながってても、
デバイスメモリが共有されるわけじゃないから
CUDA的には関係ない。

>>580
SDKにサンプルがある

CUDA FORTRANのセミナーが青木先生のとこで開催されるらしいが、おまいら行く？

>>587
青木先生か、Cだったら行くんだけどな

デバイスエミュレーション時の速度って、実際のCPUとの目安で考えたら
どのくらいスケールして考えればいいですか？

初めてエミュレーションモードを使ってみたんですが、3000倍以上の差が付いて明らかにおかしいと思うんです…
CPU: Core i7 [email protected]、 GPU:Tesla C1060

Q720って720QMのこと？
ノートPCにTeslaが搭載されているとか、聞いたことがないんだけど

追記。
grid(2,1,1)、block(256,1,1)でのカーネルで、3000倍になります。
これからgridを増やすと、さらに差が広がっていきます。

カーネルで実装した内容を、CPU版で実装したくないけど速度比較はしたい。
･･･無理でつか？

>>590
Teslaは別のデスクトップPCので、エミュを動かしたのはノートPCでの方です。
紛らわしくて申し訳ない。
どちらもPCもCUDA使えるんですが、CPU自体はノートの方が性能良かったのでこちらを使いました。

ノートPCのGPU： GeForce GTX 260M

そもそも、エミューレーションモードって非CUDA環境でも
CUDAカーネルのデバッグが出来るようにしたものでしょ

あくまでテスト用のもの

誤)エミューレーションモード
訂)エミュレーションモード

3000倍？そんなもんでしょ

効率の良いブロック分けの仕方？がわからず困っています．

実行時にＮ個のデータ系列が与えられて，
それぞれの系列へ施す処理内容自体は同一なのですが，
その処理に用いるパラメタ値が異なります．
例えばＮ＝３の場合，パラメタもp[0]〜p[2]の３個があって，
　データ系列0の全データ { D[0,0], D[0,1], D[0,2], ..., D[0,m0] } にはp[0]を加算，
　データ系列1の全データ { D[1,0], .... , D[1,m1] } にはp[1]を加算，
　データ系列2の全データ { D[2,0], .... , D[2,m2] } にはp[2]を加算
という具合です．

全系列のデータ数が同じ(m0=m1=m2)ならば
グリッドの次元の１方向を系列(0〜N-1)に割り当てれば良いかと思うのですが，
系列毎にデータ個数がかなり異なる場合はどうすればいいのでしょうか？
データ個数は系列ごとに少なくとも数千個くらいになります．
同じような割り振り方だと何もしないブロックが大量にできてしまいそうです．

CにD[0,0]というものはないのでよくわからないけど、
いったん長い配列にまとめて処理して、あとでCPUでばらせばいいのでは。
D[i,j]のjについて、自分はどのpに属するのか覚えさせて。

>>592
Nvidia Nexus使えば？

追記
Nexus使うとネットワーク経由で、
コード書く用のPCとデバッグするPCを分けられるよ、ってことね。

>>598
うーん、デバッグというよりは単にCPUとGPUで速度比較をしたいだけなんです。
うまく並列化して普通は、1〜50倍くらいの成果になると思うんですが･･･

目安でいいからエミュレーションモードから大体の速度が分からないのかなと。

いまいち意味がわかんないけどCPUコードとGPUコードをデバイスエミュで実行したらCPUコードのが3000倍早いって事？
それだったらそんなもんかと。デバイスエミュは重いし。
違うんだったらごめん。
CPUとGPUで速度比較したいなら普通にCPUとGPUそれぞれ向けのコード書いて実行したらいいんじゃない？？

>>601
あ、あれ？　自分のデバイスエミュの認識自体が間違ってたかな…？
言いたかったのは、実行するハードの方での両者の比較です。
CPUコードと言うのはありません。
カーネルや、その内部で呼ぶ__device__の関数らがGPUコードだよね？それを普通に「GPU」が実行した時の速度と、
デバイスエミュを使ってCPUが実行した場合（内部では逐次計算？）の速度では、普通に「GPU」で実行した方が3000倍速いということです。

>>591の通り、GPUコードが多くなりすぎて、同じ事をするCPUコードを実装するのが面倒なのです。
デバイスエミュはCPUが実行するとの事で淡い期待を抱いてましたが無理そうな感じなんですな…

>>602
比較する目的はなんでしょうか？
研究目的であれば面倒であろうがCPU用も実装しなければなりません。
そうでないなら比較なぞしなくてよいのではないかと。

ちなみにGPU：エミュが3000：１程度であればCPU用に普通に実装した方がGPUより速い可能性が十分にあります。

面倒でも計算結果の比較しろよｗ

nvidiaのサンプルコードでも結果の比較してるだろ

>>602
emulation modeは，名前の通りGPUでの動作を模擬しているだけで，
その計算速度とGPUの計算速度を比較することに意味はない．
emulation modeがある理由は，カーネル内にprintfとかのコードを書いて
debugしたり出来るようにするため．
従って，CPUとGPUの計算速度の比較を行いたいなら，それぞれに最適化した
２つのコードを書いて比較する必要がある．

関係ないけど，
CPUとの比較しているときにCPUのコア一つとの比較が多い気がするけど，
それフェアな比較じゃないよね．せめてOpenMPくらい使おうよと．
まぁ使うとGPUの存在感が薄れるのだけれども…．

>>603
研究なんて言えないようなものです。目的としては自己満足になりますね。
ただ、目安程度であれど比較できないとGPUとCUDAを使う意義に関わってきます。
早ければSUGEEEE!!ってなって満足し、遅ければそこできっぱり止めるという選択ができる。
そして3000：1ならまだCPUの方が早そうだというのは参考になりました。

>>605
おっしゃる通りですが、厳密に比較するまでは求めていないんです。
今自分がやってることは無意味なのか？（つまりCPUの方が普通に早い）が分かればいいんです。
grid(2,1,1)でフェアじゃないのは、ｺｰﾃﾞｨﾝｸﾞが糞なので2以上じゃないとGPUで動作しないんです('A`)･･･(ﾒﾓﾘ周りの関係で)

我侭な要求でスマン。

逆に考えるんだ
CPU側で動作をきっちり検証したプログラムを、
GPUに移植して、速度を比べる。

GPUに本当に適した問題なら、数十倍出る場合もあるし。

>>605
GPUの優位性をうたうような文脈で
比較対象のCPUコードが1スレッドだったら
それを評価する人間は目が余程の節穴でない限り、
CPUのコア数倍した結果と比べるでしょ。

物凄く差がある場合はそれでも十分優位性をうたえるから。

>>605
１コアの40倍とあれば4コアの10倍と読み替えればいい訳で。

４コアと比較したらどうなるかということではなく、
１コアと比較してる人がCPU版もまともにチューニングしてるとは思えない、ということかと。

まぁ、皆GPU版ばっかりチューニングしてますから。

たぶんCPU版はSSEすら使っていないんだろうね。
メモリ帯域がものを言うコードでなくて、CPUがNehalemだったら、
安いGPUなんかじゃ全く優位ないからね。

俺はNvidiaちゃんを信じるよ

http://twitter.com/NVIDIAGeForce/status/9265680539

誤爆った／(^o^)＼

うちは理論で「***手法より*%高速化して最速!」とかやってないってのもあるけど
GPUで組んだ手法と既存の手法を比べる場合、既存のほうはベーシックにしろと指導された。
複数CPUだとかSSEを使ってガチガチに最適化した手法と比べちゃうと基準が分からなくなるからだと。
他の高速化との差を知りたければその論文と比較しろということだと思う。

CPU最適化して無いなんて糞というのも分かるけど、こういうところもあるということで。

コードの比較もいろいろだよな。
同じアルゴリズムを採用しても、CPUでも書き手によってGPUでも明らかに差が出てくる。
でもGPUを使う場合、多くの場合はCPUよりも速くなりました。というのが目的な訳で、
CPUの方が速いならあえてGPUを使う必要はないからね。
基準が曖昧になるのもわかるけど、そもそも基準が曖昧な気がするなあ。
場合によってはかなり恣意的になることもあるし・・・・。

Femiやばいまた延期確定かも

一般人が入手できるのは１年後になる可能性もあるらしいね

なんでそんな度々延期になるの

>>618
ペーパーロンチで実際開発が
行われていないからだよ

今回のケースは大きな欠陥があることを知りながら、小手先の改良でなんとかしようとして
「完成品」を大量生産をして、まとにチップが取れなかったのが原因だろ

１％程度とされる歩留まり率で、１チップ当たり5000ドルの原価
これでは商売にならないね

3/19に東工大青木先生がCUDA　Fortranのセミナやるんだって

関係者の宣伝おつ

青木先生に集客されたくねえなあ正直

nexusをリモートで動かそうとしたが、ブレークポイントでとまらねぇ・・・。
色々試したがどうにも解決しないので教えてください。

状況としては、nexusのユーザーズガイドに沿って設定。
ためしにnexusサンプル動かそうとしたら、
ランタイムAPIプロジェクトはGetDeviceで引数に０が。
で、次の行で落ちる。
DriveAPIは落ちない。んでホスト側の画面右下に青いポップアップ出て、
ターゲットマシンにコンソール画面出てるのでプログラムは正常に動いてるっぽい。
でもカーネル関数内にブレークポイント置いても止まらず。
自分でSDKサンプルのプロジェクトの設定変えて試しても同じ。

マシン環境はこんな感じ。
ホストマシン
Vista 64bit SP2
.Net3.5 SP1
Host nexus1.0(jan 64bit)
GPU FX570
VC++ 2008 SP1
DirectX10 August 09
CUDA SDK2.3 32bit
CUDA ToolKit2.3 32bit

ターゲットマシン
Vista 64bit SP2
.Net3.5 SP1
Target nexus1.0(jan 64bit)
GTX285
VC++ 2008
DirectX 10 August 09
CUDA SDK2.3 32bit

他に設定としてはEnable Security Serverをfalse
nexus→option→Enable secure connectionをfalseにしてます。
使い方は、ホストマシンでVC起動→プロジェクト読み込み→nexusデバッグ
ターゲットマシンはデバッグモニタ起動のみ。
どこがおかしいのだろう？

Fermi終了したらこのスレも終了するんかなあ

the launch date for GeForce GTX 480 and GTX 470 is March 26

>>627
ペーパーリリースで全世界で１万枚以下の出荷といううわさだけどね
一般人が手に入れられるようになるのは、下手をすると来年
良くて年末という予想がある

やっぱ、シリコン丸ごと改良しないとだめなんか。
半分ダメで256コアでもいいんだけどｗｗ

俺はNexusインストールすらできなかった

消費電力280Wだっけ。。。GTX480

アム虫キモ

もうよくわからんからnexusのエミュモードでやることにした。
カーネル関数内でブレークするし、値もちゃんと表示されてそう。

>>630
OS対応してないとか？Vistaか7しかできない。
あとはOSとnexusのbitが違うとか。

>>633
インストーラがVS2008 SP1入れてあるのに読み取ってくれなくて、
インストールができない状態だった。

Microsoftからダウンロードしたばっかりのイメージを使ってインストールしたから
当然SP1はあたっているものかと思っていたらあたっていなかった、っていう初歩的なミスだった。

ちょっくら遊んでくる

　　　　　　　　 ,. -‐'''''""¨¨¨ヽ
　　　　 　 　 (.＿＿_,,,... -ｧァﾌ|　　　　　　　　　　あ…ありのまま 今　起こった事を話すぜ！
　 　 　 　 　 |i i|　 　 }!　}} /／|
　　　　 　 　 |l､{　 　j}　/,,ｨ//｜　　　　　　　『おれはNvidiaにARE YOU READY?と言われて
　　　　　　　 i|:!ヾ､_ﾉ／ u {:}//ﾍ　　　　　　　　　　準備していたら準備しているのはNvidiaの方だった』
　　　　　　　 |ﾘ u' }　 ,ﾉ　_,!V,ﾊ |
　　 　 　 ／´fト､_{ﾙ{,ィ'ｅﾗ　, ﾀ人　　　　　　　　な…　何を言ってるのか　わからねーと思うが
　　　　 /' 　 ヾ|宀| {´,)⌒`/ |<ヽﾄiゝ　　　　　　　　おれも何をされたのかわからなかった
　　　　,ﾞ　 ／ )ヽ iLﾚ 　u' |　| ヾｌﾄﾊ〉
　　 　 |／_／　 ﾊ !ニ⊇　'／:} 　V:::::ヽ　　　　　　　　頭がどうにかなりそうだった…
　　　 /／ 二二二7'T'' ／u'　__ /:::::::/｀ヽ
　　　/'´r　-―一ｧ‐ﾞＴ´　'"´ ／::::／-‐ 　＼　　　　128bitメモリバスだとかリネームテクノロジーだとか
　　 / // 　 广¨´ 　/'　　 ／:::::／´￣｀ヽ ⌒ヽ　　　　そんなチャチなもんじゃあ　断じてねえ
　　ﾉ ' /　 ノ:::::`ー-､___／:::::／/ 　 　 　 ヽ　　}
_／｀丶　/::::::::::::::::::::::::::￣`ー-{:::...　　　 　　　ｲ　 もっと恐ろしいものの片鱗を味わったぜ…

GTX480がそこそこ出回るらしい（core数は当初想定よりも少なくなる可能性があるが）
Fermiアーキテクチャを思っていたよりも早い時期に体験出来る可能性が出てきた

Fermi火事出すだろうな
電気食いすぎだ

GF100（GTX480, 470）は高電力でお値段も高めになるだろうから、
Fermiアーキテクチャをとにかく早く試したい人以外はその次のGF104がいいと思う

半分だけの460とか出ないかな？

この辺りの情報がそれかも知れん

The AMD's Cafe:ローコストなFermiやQuadroについて - livedoor Blog（ブログ）
http://blog.livedoor.jp/amd646464/archives/51525107.html

まだまだ先のようだ

売り物にならない奴はQuadro逝きか。
OpenGL市場ってほんと舐められてるな。
ゲイツに縛られない自由なAPIだったはずなのに。

http://www8.plala.or.jp/b4zabeat/

このレビューは何が言いたいのかよくわからんかった
http://pc.watch.impress.co.jp/docs/column/nishikawa/20090518_168541.html

CUDA上の命令がどれくらいのクロックで動くかまとめられていませんか？
整数の乗算やら三角関数はプログラミングガイドに載っていたのですが,ほかの命令も知りたいです

「どのくらい」でいいなら実測すればいいかと

エミュレーションモードでやるとうごくのですが、GPUをつかうと動きません。

const int c = border + (blockDim.x * blockIdx.x + threadIdx.x) * step * 2;
const int r = border + (blockDim.y * blockIdx.y + threadIdx.y) * step * 2;
const int i = 1;
// atomicAdd(&count[0], 1);//ここでは動くのでatomicAddの問題ではない
if(c >= i_width - border || r >= i_height - border) return;
int i_max = -1, r_max = -1, c_max = -1;
float max_val = 0;
for (int ii = i; ii < min(i+2, intervals-1); ii += 1)
for (int rr = r; rr < min(r+2*step, i_height - border); rr += step)
for (int cc = c; cc < min(c+2*step, i_width - border); cc += step) {
float val = getVal(d_m_det, o, ii, cc, rr, i_width, i_height);
if (val > max_val) {
max_val = val, i_max = ii, r_max = rr, c_max = cc;
}
} // Check the block extremum is an extremum across boundaries.
/***********ここでd_iptsにiptを加えてもうごく*********/
// float4 ipt ;
// d_ipts[atomicAdd(&counter, 1)] = ipt;
if (max_val > 0.0004f && i_max != -1 && isExtremum(d_m_det,o, i_max, c_max, r_max, i_width, i_height, intervals)) {
float4 ipt = interpolateExtremum(d_m_det, o, i_max, r_max, c_max, i_width, i_height);
if(ipt.x >= 1)
{
d_ipts[atomicAdd(&counter, 1)] = ipt;//ここの行をコメントアウトすると動く
ipt.x += 1;
}}
最後のif文の中でd_iptsやcountにアクセスするのがだめっぽいのですが・・・
なにかif文を書いたときに同じような症状になった方や、これを見ただけでわかる方いらっしゃったら、教えてください。
よろしくおねがいします。

とりあえず問題になってるatomicAdd(&counter, 1)の戻り値調べようか

counterの宣言にちゃんと__device__は付いているのだろうか・・・

>>647
エミュレーションで確認したところ戻り値はcounterと同じ値になっていました。
実際にGPUでうごかすと、
counterをデバイスからホストへ送るところでunspecific launchとなるか、
永久ループに入るか何かで画面が固まり、ブルースクリーンになって落ちます。

>>648
_device_をつけてグローバル変数(っていうのかわからないけど)として定義しています。

その反応だとアクセス違反のときが多いのだがそんなことないよな？

CUDAでつかうlong doubleってWindowsとLinuxでサイズは違うの？

>>650
アクセス違反ってcounterの位置がおかしいってことですか？
このプログラムをいれずに単にcounterをデバイスからホストに送ることは可能でした(初期値として０をおくっていたので０が帰ってきていました)。

つまり、このプログラムのようにatomicAddをif文とか分岐が多くなるような文章内でつかうと、
今回でいえばcounterのアドレスがかわるということですか？？

パスワードクラックでTeslaがHD5970に完敗

http://www.geeks3d.com/20100316/radeon-hd-5970-the-ultimate-password-cracking-hardware/
http://www.geeks3d.com/public/jegx/201003/elcomsoft-wpa-psk-password-test.jpg

Dual-GPUとSingle-GPUを比べてる時点でアウト

１スロットどうしの比較だから問題ない

>>655
それを言い出すと、TeslaDが出てくるぞ。ブリッジ自体は1スロットだからなw

はいはい、負け惜しみ
単純な計算性能では圧倒的にラデオンのほうが優れているんだよ

比較対象にHD5870が入っていない時点で
なんか違和感があるんだが

Tesla C1070ってなに？Fermi？

http://www.elsa-jp.co.jp/products/hpc/tesla_s1070/index.html
これ

GT200じゃなかった?

ttp://developer.nvidia.com/object/gpucomputing.html

CUDA Programming Guideがリンク切れしているんですが
誰か正しいアドレスを知りませんか

高い・遅い・熱い

S1070ってGT200×4の奴だろ。

CUDA 3.0 Downloads
http://developer.nvidia.com/object/cuda_3_0_downloads.html

CPUでの計算にインテルコンパイラを使いたいのですが、どのようにすればいいのでありまするか？

>>665
ありがとうございます

>>666
nvcc -cで*.cuをコンパイルしたら、できた*.oをiccでリンクすればいい。
Windowsの場合はnvccもiccも別々にオブジェクトを作ることになるからそれをVCでリンク。

利用者はCUDAなんて独占的技術を求めてませんよね

まして永久β版のSDKなんて論外

まあ利用者だけが決めるわけでもないのも悲しいけど現実なのよね

なんか面白いことに使えないかな
もったいない

「何かには使える」って言ってるうちは何にも使えないまま終わるんですけどね。
ターゲットアプリケーションがあってはじめて、そのニーズに合わせてハードの機能・性能の拡充が行われうるわけで
今までであればゲームがそうだった。

ウン十並列のデータを同時処理するような用途のニーズが仮に高まってるとしても
それはCPUのSIMD拡張によってもカバーできるでしょ

Linuxでドライバのバージョンを調べる方法を教えてください

もう少し具体的に聞かないと・・・・。

ドライバのバージョンによって、振る舞いを変えたいので、
UbuntuでGPUのドライバのバージョンをプログラム上から調べる方法を教えてください

Ubuntu限定な必要があるかどうかはともかく、ドライバ自体のバージョンを知るAPIは用意されていません。
/usr/lib64のディレクトリでlibcuda.so.* のレギュラーファイルを探すのが確実かと思います。

CUDA実行環境がインストールされてるかどうかを調べてライブラリを遅延ロードできるような仕組みを
標準で用意して欲しいかな

CUDA3.0
ttp://developer.nvidia.com/object/cuda_3_0_downloads.html

Fermi対応版か
肝心のブツが手元にまわってくるかも怪しいのに

deviceQueryがあるじゃん

Fermiは、入手しやすくなるまで松わ。
初物は爆熱で卒倒価格だろうし。

俺は特攻する
470か480かが問題
そもそも手に入るのか、という話もあるが

そうか。取り合えず貼っておかねばなるまい
::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::。:::::::::::::::::::::::::::::::::::::::::::::
:::::::::::::::::::::::::::::::::。::::::...... ...　　 --─-　　:::::::::::::::::::: ..::::: . ..::::::::
:::::::::::::::::...... ....:::::::゜::::::::::..　　 （___ ）（___ ）　::::。:::::::::::::::::　゜.::::::::::::
:. .:::::。:::........　.　.::::::::::::::::: _　i/ ＝　＝ヽi　:::::::::::::。::::::::::: . . . ..::::
:::: :::::::::.....:☆彡::::　　　／/[||　 　 」　　||]　>>683 ::::::::::゜:::::::::: ...:: :::::
　:::::::::::::::::: . . . ..: ::::　/　ﾍ　| |　 ____,ヽ | |　:::::::::::.... .... ..　.::::::::::::::
::::::...゜　.　.:::::::::　　／ヽ ﾉ　 　 ヽ＿＿/　　.......　.　.::::::::::::........ ..::::
:.... .... ..　.　　　　 く　 /　　　 　三三三∠⌒＞:.... .... ..　.:.... .... ..
:.... .... ..:.... .... ..... .... ..　.:.... .... ..　..... .... ..　..... ............. ..　.　........　......
:.... . ∧∧　 　∧∧　　∧∧　 　∧∧　.... .... ..　.:.... .... ..... .... ..　.
... ..:（　　 ）ゝ （　　 ）ゝ（　　 ）ゝ（　　 ）ゝ無茶しやがって… ..........
....　 i⌒　/ 　 i⌒　/ 　i⌒　/ 　 i⌒　/　..　..... ................... ..　.　...
.. 　 三　 | 　 三　 | 　 三　 | 　 三　|　　... ............. ...........　.　.....
...　 ∪ ∪ 　 ∪ ∪ 　 ∪ ∪ 　∪ ∪　............. ............. ..　........　...
　　三三　　三三　　三三　　　三三
　三三　　三三　　三三　　　三三

東工大、気象庁の次世代気象モデルのフルGPU化に成功
http://pc.watch.impress.co.jp/docs/news/20100324_356466.html

120GPUで3.22TFLOPSか。大変なんだろうけど微妙な数字だな。

どう微妙なんだい

なんか膨大な演算能力は殆ど遊んでいるな。
こういう用途なら演算機減らした方が、電力効率的にはマシになるんじゃ。

ha?

え？

ま、近い将来fermiに置き換わるんだろ

>>685
1GPUあたり44.3GFlopsかぁ。維持で対応したって感じだ。

浮動小数点演算速度が3.22TFLOPSなのか？

それでもCPUより80倍の実行性能なんだね

6時間の気象モデルが70分で終了するなら
同じ気象モデルを使えばリアルタイム予測が可能？

偏微分方程式の数値解法をやる大学院生の演習課題みたいなもんですなw

気象問題って、ノード間の影響はどうなの？
TESLAって結局PCI Expressで繋がっているから、
レイテンシが大きそうだな。

とりあえず運用してノウハウ貯めて、
28nmのFermi2で一気にパワーアップってのがいいんじゃないか？
明確なハズレ世代を大量導入するのはちとどうかと。

>>687
そういうのは GPU 以外に期待した方がいいんじゃないの？
電力効率なんて気にしたせいでピーク性能が落ちてゲームユーザが買わなくなったら、GPU の市場自体が崩壊するよ。

>>696
レイテンシを隠蔽するようにcode組んだらしい

>>699
へー、そーなんだ。

>>693
これがよくわからんな。
なんで80倍なんだろう？
G200って倍精度の理論値って80Gflops位じゃなかった？
それに対してCPUが10Gflopsだろ。
8倍の間違いじゃないのか？

http://techon.nikkeibp.co.jp/article/NEWS/20100324/181319/?ST=lsi
> 「最新のマイクロプロセサとの比較ではなく，チューニングの程度の
> 差もあるため，80倍という数字自体は重要ではない。GPUの活用で
> ケタ違いの性能が得られることを確認できた点に意味がある」（青木氏）

だそうで

倍精度もハードウェアで本格対応して500GFLOPSになったんじゃないの
と思ったけどFermi世代からで、G200世代のものは90GFLOPS程度みたいだね。

これで倍精度で44.3GFLOPSなら効率50%近くということになるし
アルゴリズムとかかなり頑張っただろうね。

>>701
シェアードメモリをうまく使ったり、大量のスレッドでノード間のレイテンシを隠蔽したり
帯域がボトルネックになりにくいように出来たからじゃないの。

スカラープロセッサは、流体力学とか多体問題とかの
計算結果を相互に利用しながら並列計算する場合には
帯域がボトルネックになって効率がかなり低くなりがちらしいし、
比較対象のCPUを使ったシステムは効率が10%を下回っていたとかかも。

完全に並列化できて帯域がボトルネックになりにくい演算の場合は
10〜20倍という話だったかな。

>>703
メモリ帯域が問題になるのはよくわかるんだけど、
それも今のNehalemあたりは30GB/sぐらいあるから、
精々teslsaとは5倍くらいだろう。
それだとメモリ転送時間は1/5、計算時間が1/8になるとして、
CPUのメモリ転送時間をA、計算時間をBとすれば、
実行時間は単純にするとC=A+Bになる。
GPUの場合はCg=0.2A+0.125Bになる。
A=10Bとしたとしても5倍程度にしかならない様な気がするなあ。
CPUの場合巨大なキャッシュがあるし、プリフェッチもあるから
その差が縮まるはずなんだけどな。
東工大のクラスのだからCPUのコードがくそだということもなさそうだから、
なんかOpteronのCPU1コアと比較してそうだなあ。

と思っていたら、>>702が書いてくれているね。

比較するCPUが遅すぎたわけか

このcpuの速度は、普通にべたべたfpu演算を書いたときの数字だね。
理研の姫野ベンチ並みのバカコードと対照させても意味がない。

Linpackで最近のcpuをベンチマークすると、理論値の8掛け程度の数字は出る。
TUBAMEのopteronも1コアあたり10GFlops前後。しかも倍精度で。
Linpackに比べて気象エミュは速度が上がらんのは間違いないが、
この青木とやらの記事は全然ダメだ。
nVIDIAが成果としてレファできないレベル。

CPUというのは全然チューニングしてないレファレンスコードだろ。
よく使う手。しかもCPU名すら書いてない。

ベンチマークには２種類ある。嘘か大嘘。

PPTを見ると、44.3GFLOPSは単精度、倍精度では15GFLOPSだね。
比較のCPUは倍精度で0.5GFLOPS以下で、単精度の44.3GFLOPSと比較し、80倍以上といってるんだね。
青木いい加減にしろよｗ

これはひどい。

>>704
この手の計算の場合はCPUのキャッシュやプリフェッチはあまり効果が無いと思うけど
Nehalemだと30GB/s程度メモリ帯域があるのか。

ただCUDAでは単なるメモリ帯域の比較だけでなく、シェアードメモリの利用や
大量のスレッドによるメモリアクセスレイテンシの隠蔽が重要になってくると思う。

>702の記事見たけど、デュアルコア2.4GHzのOpteronの1コアと
TESLA S1070の中の1基の比較なのかな・・・

>>708
倍精度だと思ったら、単精度だったのか・・・
そのpptってどこで公開されているんだろ？

倍精度でも30ばいかい？

なんでみんなプレスリリースくらい読まないの？
不思議

>>708
おいおい・・・まじかよ。
かなり酷いなあ。まだ騙すことは出来るかもしれないが、
これがばれてくると偉いことになるぞ。

そういえばNVIDIAのGPUカンファレンスでCPUの2000倍速くなりましたってのがあったが、
Tesrax4、Opteron 2.4GHz 1Coreのものだった。
かなりつっこまれていたよ。
で、GPUは最適化しましたけど、実はCPUは最適化してませんだった。

>>710
この手の計算はどちらかというとストリーミングに近いから、
キャッシュやプリフェッチは十分効くよ。
特にハードウェアプリフェッチはかなり効果的だよ。
ただ、SSEを使ってしまったりすると、計算時間を隠蔽できなくなって、
今度は転送時間を隠蔽する方法を考えなくてはならなくなる。

Nehalemの中でも2000MHzのメモリに対応したものであれば、
48GB/sにもなるよ。このあたりはオーバークロック気味になるので、
コンシューマレベルでしかないけどね。

GPUを使うのを否定はしないが、いい加減嘘に近い誇張は止めてもらいたいものだ。
数倍でもいいだろう。2倍でも2日かかったものが1日で終わるんだぜ。
あんまり速いと仕事が増えるじゃねぇか！

嘘は言ってないんだよな
比較対象が微妙すぎるだけで
CPUでも効率を出すのが難しい問題はあるから、両方ともきちんとチューニングした上で比較するのが理想ではある
まあXX倍という数字が一人歩きするのはよくないね

>>712
そのプレスリリースってどこにあるの？

>>713
この手の計算って常に帯域を使い切っているわけじゃないの？
それならキャッシュの効果は限定的になりそうだけど、プリフェッチはかなり効果ありそう。

http://journal.mycom.co.jp/news/2010/03/24/055/index.html
> 単一のGPU(Tesla S1070)を用いた際の性能は
> 「単精度ではTSUBAMEに搭載されているOpteron比で最大で100倍を超すレベルを達成、
> 平均でも 44,3GFlopsを達成している。
> 性能が落ちる倍精度でも15GFlops程度を達成しており、
> CPU比では相当高い値を実現した」(青木教授)と説明する。

http://www.gsic.titech.ac.jp/contents/press_release0324.pdf
> スパコンTSUBAMEの単一GPUを使った計算で
> 44.3GFlops（CPUの1コアに対して約80倍）

これってやっぱり、Tesla S1070の4基を単一GPUと言って
Opteronの1コアと比較している同じパターン・・・？

あまりに酷い誇張はやめてもらわないと、有能な開発者がGPUの性能をそれなりに引き出しても
能力が低いとか言われたり、色々と深刻な事態になりかねない・・・

実際、某社でなりかけている。

私の現場では、寧ろ旧世代のCPUを積んだサーバ機をGPUで延命させるってシナリオだから
処理能力は1.5倍でも御の字さw まぁ、実際のところはXeon1core対比で5倍程度は出たから
サーバ機一台で換算してなんとか2倍(5+1*3 vs 4)の数字は出たけど。

>>715
結局実装アルゴリズムによるけど、
プリフェッチはかなり効くよ。

最後の2行はすごく同意だな。
わかっていない人からすると、GPU使えば数10倍になるはずと思ったりするので、
有能な奴が5倍速い書いたコードを書いたとしても相対的に低いと評価されそうだ。

>>716
たぶんこういう使い方が一番いいのかもね。

>>715
CPU比ではそうかもしれんが
AMD比ではかなりしょぼくないか？ｗ

http://nttxstore.jp/_II_EA13119354
NVIDIA Tesla C2070　ETS2070-C6ER　475,545円（税込)　発売開始日 2010/9/30

GTX 470欲しいな。でもどうせ瞬殺なんでしょ？
いいよな秋葉原に住んでるおまえらは

>>716
GPUで延命→鯖予算獲得→新鯖にGPU移植
の流れが理想

GPGPUでの性能評価の指針って明確に定められないのかね
CPU単一コア，レファレンスコードと比較して何倍速くなったと言っても，
そんな条件で計算すること自体なさそうなんで，意味無い気がするんだが

GPUでチューニングするなら，CPUでもチューニングするべきだろう

AMDかintelかでめんどくさいからじゃない

>>722 CPUまでチューニングする必要はないと思うけど、
気象コードがMPIで並列化されてないとも思えないし、1GPU内でも相当の並列化を行ってるんだから、
CPUの1コアのみと比較するのは、ちょっとやりすぎ。
で実際Xeon/X5570と倍精度で比較したら、単一コアで10倍程度、4コアで3倍程度のアドバンテージしかないんじゃない。
ノード単位だと逆に遅くなるとか。120GPUの3.2TFLOPSも単精度だし…、比較自体を都合のいいように、いいとこどりしすぎ。
ベンチマークだけでなく、せっかく実アプリをフルGPU化したのは十分意味があるのに、世間受けを狙いすぎた発表で逆に良識を疑ってしまう。

第三者（？）が書いた記事よりも
NVIDIAの発表のほうがまだ信頼できるな。

詐欺師の言葉のほうが信頼があるとはこれ如何に

詐欺師はお客様に信用されなければいけない。
だから最低限の嘘しかつかない。

大道芸人は嘘をつくたびに金をもらえる。
だからいくらでも嘘をつく。

モックアップＮＶＩＤＩＡとうそつきが多い野心的な学者連中の
たわ言なんて信用できないよな

ま俺は数字しか信じない
ttp://www.anandtech.com/video/showdoc.aspx?i=3783&p=6

数字なんていくらでもいじれるんですよ！
信用のあるデータが欲しいなら自分で実測するしかない

つまりそのデータを公開しても
誰も信じないってわけね

その通り！
よく分かってるじゃん
論文だって他人の実験の結果なんて基本誰も信用しないよ
とりあえず自分で再現実験してみるのは常識

そしてその論文も誰も信用しない

利害を共にしない数十人の人が
方法の妥当性を検証し、追試し、同じ結果が出て初めて
ある程度の信用を得る
それでも「ある程度」なのが学問の厳しい所だ

GPUは1コアだと思いますよ。TESRA内でもGPU同士はPCI越しなんで
遅くなると思います。
3000x3000x50って1GPUにおさまるサイズと思うし

誰も追試できないと思っているから舐めているんだろ。
今後は青木の言うことはハイハイワロスだな。

>>735
それだったら120GPUの意味がないだろ。
あと、Teslaな

いくつかの別な測定を同じ測定かもしれない感じで書くのはテクニックじゃないですか
１GPUで44.3GFLOPSなのに120GPUでなんで3.22TFLOPSなんでしょうか。

Zotac GeForce GTX 480
Amazon.comでPre Order　$499だったから
申し込んでしまったｗ　(米→日転送業者使用)

結局512spじゃなかったねGTX480

Ultraがくる

>>740
最初から、576にしておけばよかったのにね。
32x18で。2ブロックダメでも512で出せるのに。

そんなことしたらますますイールドが悪くなるわけで。

1CUDAコア破損しただけで３２個分が台無しになるFermi
512すべて無事なのはほとんどないんだろうな

PS3と同じ作戦でござる。

　　　　 ＿＿＿＿＿＿＿＿＿＿＿＿
　　　　ヾミ || || || || || || || ,l,,l,,l 川〃彡|
　　　　　V~~''-山┴''''""~ 　　ヾニﾆ彡|　　　　　　　512SPは存在する・・・・・・！
　　　　 /　二ー―''二　　　　　 ヾﾆニ┤　　　　　　 存在するが・・・
　　　 <'-.,　　￣￣ 　 　 _,,,..-‐､　〉ニﾆ|　　　　　　　今回　まだ　全部有効にするとの
　　　/"''-ﾆ,‐l 　　l`__ニ-‐'''""` /ニ二|　　　　　　　指定まではしていない
　　　|　==＝､!　　`＝====､　　l =ｌべ=|
.　　　| `ーﾟ‐'/　　 `ー‐ﾟ―' 　　l.=lへ|~|　　　　　　　そのことを
　　　 |`ー‐/　　　　`ー――　　H<,〉|=|　　　　　　　どうか諸君らも
　　　 |　　/　　　 ､　　 　　 　 　 l|__ノｰ|　　　　　　　思い出していただきたい
.　　　| ／`ー　~　′　　 ＼　　　.|ヾ.ニ|ヽ
　　　 |ｌ　下王l王l王l王lヲ｜　 　|　ヾ_,|　＼　　　　　つまり・・・・
. 　 　 |　　　　≡　　　　　　　　　|　　 `l　　 ＼__　　　我々がその気になれば
　　　　!、　　　　　　　　　　 _,,..-'′ ／l　　 　 |　~'''　 FermiのSP数は
‐''"￣|　`iｰ-..,,,＿,,,,,....-‐'''"　 　 ／　 |　 　 　 |　　　 320SP　384SP　ということも
　-―|　　|＼　　　　　　　　　 ／　 　 |　　　 　 |　　　可能だろう・・・・・・・・・・ということ・・・・！
　 　 |　　 |　 ＼　　　　　　／　 　 　 |　　　　　 |

　 　 　 　 　 ＿＿＿_
　　　　　　　|<三`'ヨ′
　　　　　　_/6|ｰ廿┤
　　　 /l￣　ＫL.三.｣￣h
.　　 /　|　　 レ兮y′/ l
　　〈　 く　　 ∨ l/　,ｲ　|
　　　＼_,.>、　　　 /,L..｣_
.　0ﾆﾆﾆ)而}ﾆﾆﾆﾆニ),ﾘﾘﾆﾆ)
.　　 Ｌ| |＿＿＿__|＿＿__| |
　 　 l | |.＿＿＿＿＿＿_| |　　,:
　,　　l ＼ヽ　l　　|　　 , '／　　;'
　:, ____l＿|_|_;_|＿|＿__|_|＿＿　　　;
　　|＼ﾞ;三三ﾞ';三三三,;ﾞ三三＼　;'
　　|＼＼三三ﾞ三ｼﾞｼﾞ三三,''三;'＼,;'　　;'
　　|、 ＼＼三ﾞ;三三ｼﾞｼﾞ･'三三三;＼　;
　　0ﾄ､＼＼＼;'三三;'三三三;''三三,;'＼
　　　　＼＼＼| 炎炎炎炎炎炎炎炎炎 |
　　　　　 ＼＼| 二I二二I二二I二二I二 |
　　　　　　　＼ＬI二二I二二I二二I二二｣
　　　　　　　　0｣　　　　　　　　　　　　0｣

NVDIAフォーラムでGTX480の倍精度性能はTeslaの1/4という発言が
ありますがどうなんでしょう

>>747
おれ、人柱としてぽちったから待っててくれ。

>>748
俺はお前を待っているぞ
+　　　+
　 ∧＿∧ 　+
　（0ﾟ・∀・）　　　ﾜｸﾜｸﾃｶﾃｶ
　（0ﾟ∪ ∪ +
　と＿_）__）　+

Quadroと同じ戦略か。
妥当っちゃ妥当だな

EECだけでなく倍精度も性能を劣化させたのか
GPGPUの利点がどんどんなくなってるな

>>749
あ、748なんだけど、さすがにTesla買う金は無いのよ。
なんか、GTX480で動かして「明らかに倍精度の性能落としてやがる！！！１１」と
分かるプログラムはどこかにあるかな。

Teslaもないと比べようがなくね？

visual profilerの関数の実行時間を比較くらいしか厳密な計測はできなさそう。

リーロードしてなかった、悪気はないんだｗ

NVIDIA_GPU_Computing_SDK/C/bin/linux/release$ ./matrixMul
Processing time: 0.120000 (ms)
Test PASSED

Press ENTER to exit...
あたりを倍精度化してもらえばいいかと

>>753
かぶって申し訳ないｗｗ
おれ>>739なんだけど、失敗かな。日本で入手可能になるのを
待ったほうが早くて安かったかもしんないよなorz

>>755
おけ、matrixMulでGTX280とGTX480の比較ならやってさしあげられる。
floatのままと、doubleに全部置換した版で。

floatは内部で型変換してるからdoubleの方が早いって聞いた事あるんだけど…

倍精度のFMAをひたすら繰り返す感じのカーネルで計測するのがよいかと。

>>758
どこかのCPUでintをdoubleに変換してたという話かと。

しかし、480SPとか、AMDはリアル12コアだとか、
時代の進歩は速いものだな。

しかし、性能はGTX295から毛の生えた程度
ていうか、もし倍精度の性能がいまいちだったらGTX480/470より
GTX295を買ったほうが安くていいかもね

長崎大のようなことをやられたらNVDIAが東工大から呼出をうけて
「２度目はないからな」とか言われても不思議じゃないよね

>>762
GTX295、5万円くらいだからなぁ。倍精度の性能がGTX280の２倍程度ならば…
GTX295でヨシってことになると。>>757よろしく！

matrixMulはメインメモリの転送が時間に含まれてるしサイズが小さめ
プログラムが書けるなら1000x1000あたりの性能をお願いします

>>760
そうなのか？
普通のCPUの構造上浮動小数点演算は、
64bitでやるから型変換をしてると聞いたんだけど

もしかしてGPUだと話は別なのかな…

>>766
x86アーキテクチャでは、普通は変数型に関わらず浮動小数点演算は80bitの拡張倍精度で行う。
メモリロード/ストアの際に、変数がfloat型ならfloat型に変換されはするが、
別に大した処理でもないので速度にそんな影響は出なかったと思うが。

R3000かなんかで64bitint乗算がなくてとかなかったっけ

>>767
GPU だとどころか、倍精度演算器が載ってない CPU だってあるから。
処理系によって話が別。

device emulationってなくなっちゃうのね

ちょっと整理。

・x86(fpu)
浮動小数点レジスタが80bitなので、floatを突っ込んでも80bitで演算する。
従って、floatとの変換処理が入るのでdoubleの方が速いことがしばしば。

・x86(sse)
MMXレジスタが汎用なので、floatはfloatのまま演算する。
従って、定数や標準関数でdoubleに汚染されないように気をつければfloatの方が多少速くなる。
また、MMXレジスタにfloatの方が2倍詰め込めるのでベクタ化した場合に2倍速くなる可能性もある。
更に、キャッシュ効率もよくなるのでより速くなるかもしれない。

・GPU(cuda)
単精度レジスタと倍精度レジスタが分かれているんだっけ? 倍精度についてよく知らんのでフォローお願い。
演算器の特性上、単精度の方がずっと高速に演算するし、転送量も当然半分にできる。
応用にも拠るけれど、CPUとのI/Fは単精度で内部だけ倍精度にできると最強かも。

64bit版gccだと、-m32オプションをつけない限りもはやx87のコードは吐かないなあ。

>CPUとのI/Fは単精度で内部だけ倍精度にできると最強かも。
これって何回丸めが発生するかによるけど、精度は単精度+αだね。
でもこういうことはよくやるなあ。
データ量を減らしたくて少し精度が欲しい場合は、最終的に欲しいデータはfloatで持っておいて、
計算するときにすべてdouble型のデータにコピーして最後に結果をfloat型に戻す。
この方法をとれば、丸め誤差の発生は一回で済む。
計算の中身が超越関数使ったり、複雑であったりすると結構有効だよ。
その分速度は犠牲になるけど、メモリ帯域が支配的であると余り影響がないな。
G200系だとまだ効果はないだろうけど、Fermiなら結構役に立つんでない？

>・x86(fpu) 
>浮動小数点レジスタが80bitなので、floatを突っ込んでも80bitで演算する。 
>従って、floatとの変換処理が入るのでdoubleの方が速いことがしばしば。 
doubleは64bitなんだから結局まるめが入るのは同じじゃないの？

おいおいあれだけCPU側のコードが糞だといっておいて、CPUのこと知らなさすぎだろｗ

floatが遅くなるのは、MSVCが、floatのときは毎回メモリに書き戻すことで精度を32bitに落としてIEEE互換にするコードを出すから
レジスタ間ならfloat/double/long double関係なく80bit

GPUの種類を取得する関数か何かはありませんか？

>>776
サンプルでついてくるdeviceQueryのソース読んでみれば

>>776
CUDA APIガイドに書いてあるだろ

そうじゃなくて製品名です

製品名って、ASUSかEVGAか判別したいってこと？？

メーカーの判別はどうがんばっても無理だろ

じゃあせめてWindowsかMacかInaxかぐらいわかりませんか？

製品名って言うくらいだから、ELSA GLADIAC 998 GTX Plus V2 512MBみたいなのじゃないのか。

>>782
CUDA以前の問題。てか、そのレベルじゃ絶対無理だろ。

#ifdef _WIN32

まさかOpenCLの話か？
#ifdef _apple
とかやった覚えがある。

#ifdef _appleは豆知識だな

Inax は釣りだろ。782 は偽者じゃないか？
TOTO 向けと別の最適化するのかなｗ

水流をGPUでシミュレーションして最適化するのか

流体シミュレーションはGPGPUのメインテーマだから、まさにうってつけだろう

OpenCL使えよ馬鹿ども

日本の便器メーカーは水量削減に血道を上げているからな
いかに少ない水量で、効率良く、かつきっちり排泄物を流しきるか

便器開発での計算に使用するだけでなく、
便器自体にTeslaを搭載し、
排泄物を画像認識して最適な水流を計算する。

排熱も有効利用できそうだ

>>793
画像認識のコード書くやつは大変だな

確かに、今は節水ということで一般家庭向けの便器は
流れが悪いといくことを感じるな。

現場では流体演算とかして設計しているのだろうか？

TOTO 節水 シミュレーション でググってみた。
http://www.toto.co.jp/saiyo/new/techno/person/person_05_2.htm

流体どころか三相全てシミュレートしてるんだな

GTX470が先に発売されるみたいですね。
購入される方いますか。

どうしてcudaはosと密接なのか？
最新のubuntuをいれたくてもいれられへん
理由を教えてくれろ

>>799
ドライバレベルで提供されているから。

>>797
固体ってのが生々しいな。シミュレーションするためにウンコの物性とか
測定したりしたんだろうな。

壊れ方とかがリアルな模型があるとか聞いたことがある。
でも物体の測定データのファイル名とかはなまなましいだろうなｗ

シミュレーションしているときのSSとかあればいいのに。

たしか法律で、流す水の量は決まってるんだよね。

>>799
たぶん、nvccがgccの進化について行けていないだけだと思う。
バイナリだけなら、最新のUbuntuでも動くよ。

最近のディストリビューションって大抵1年かそこらでサポート終わりじゃん。
1つまえのバージョンにしか入れられないと、半年程度でいれかえなきゃいけないんだよね。
そこらへんを早く何とかしてくれよ。
CentOSとかつかえばいいのかもしれんけど

そこを何とかしてもらいたいね。
俺は古いバージョンのLinuxをVirtualBoxで新しいバージョンのLinux上で動かして、
クロスコンパイルライクなことをしているよ。コンパイルだけなら仮想化環境でも通るからな。

GTX480で倍精度削られたのは本当らしい…

Quadroではどうなるんだろう

今すぐ計算して論文書かなきゃって人以外は↓ここ聞いてからにした方がいいんじゃないかと。
ttp://www.hardocp.com/article/2010/03/26/nvidia_fermi_gtx_470_480_sli_review/7

「当機はまもなく離陸しますw」

>>807
ドイツ語読めないけど、これ？
ttp://www2.hardware-infos.com/news.php?news=3497

Q1.同じGT-240を二枚挿せばCUDAも倍近く早く処理できるようになりますか？
Q2.GDDR5とGDDR3とDDR3でCUDA動画エンコードの処理スピードはどれくらい
違いますか？GDDR3を100として。

>>810
もう少し具体的に知りたいことを書いた方が答えやすい。

例えば動画エンコード用途に限ると「複数枚挿しはどうよ？」とか、「FermiとかATIとかの中でどれが一番速いか」とか。

十分具体的だろ･･･。

ソフトウェアによって、複数挿に対応しているかどうかは変わるので、
使いたいソフトについて調べよう。

メモリの速さは重要だけど、GPGPUでは、メモリにアクセスするときの遅延のほうが問題となっている。
ハードウェアの構造的にもGT200系のほうがメモリアクセスが柔軟なので、GTX260あたりを検討してはどうだろう。

具体的だろと言っておきながらソフトによって変わるとか意味不明

>>811-812
レスありがとうございます。

CUDAの使用目的：動画エンコのみ。 使用するソフト：MediaCoder
使用するかもしれないソフト：TMPGEncKarmaPlus
現状：AVIUTLでロゴ消しとインタレ解除のプロジェクト→
TMPGEnc4で色γクロップ・リサイズしてHUFFYUVで出力→MediaCoderのx264で。
課題：Q6600でVGAでx264エンコが22fpsしか出ない。1080pだと4fps！orz
MediaCoderにCUDAでH.264エンコできる機能があるので使いたい。
GT240がGDDR3で6000円、GDDR5で7500円〜なので、二台組むより二枚入れた
ほうがいいのかも？
GDDR3とGDDR5でＣＵＤＡエンコに殆ど差がないならGDDR3のほうにしたい。
複数枚挿しはどうです？
現状G43/G41/G31なので新しくマザー買わなきゃできないけど…。

>>812
GPGPUやメモリアクセスについて仕組みとか全く知らないのですが、
GTX260はGT240の倍以上の値段だけど、倍の性能あるんですか？
wikiに書いてあるSPとCUDAコア数がどう違うのか分からないので…
もし倍の性能（エンコが倍早くなりそう）ならGT240よりGTX260を選ぶかも
しれません。PXI-EX�Oが1つですむし。

>>810
素直にCorei7にしておけ。
GT240じゃ2枚さそうがCorei7の方がマシだ。

現状CPUを強化したほうがメリット多い
CUDA使うならGTX260以上じゃないとCPUの足を引張る可能性がある

aviutlでの処理時間考えたらCPU強化の方が妥当だな

>>815
http://www.techarp.com/showarticle.aspx?artno=520&pgno=7
でみるとi7はQ6600の1.25倍くらい早い
http://www.katch.ne.jp/~kakonacl/douga/mediacoder/v0.71_cuda.html
でみるとCUDAはQ6600の3.7倍くらい早い
∴CUDAはi7より3倍近く早い…？？？　CUDAは実写向き？

X58+i7買いたいけど、LGA1155待とうかと。

>>818
あ〜、俺もこれ使ってcudaエンコしてるけど、Bフレームが4までだったり2passできなかったり
画質はいまいちですよん。たしかにCPUよりは速いとはおもうけど・・・。

新しいバージョンだと改善されてる可能性はあるけどね。

>>814
GT240の3倍くらいの値段で売られているが、3倍速くなるとはいえないけど、
MediaCoderは複数差対応していなし、GT240を2枚買っても無駄になる。
あたらしくマザー買う予算が削れるならGTX260でもいいかと。

画質もとめるなら、i7ってのは合意。

みなさん色々有難うございます。ググりながら考えてるんですが難しいですね。

>>816>>817
GTX260だと電源交換、i7はM/B&DDR3全部揃えないと…先立つものが…。(ToT)
http://www.katch.ne.jp/~kakonacl/douga/mediacoder/v0.71_cuda.html の人
は　GeForce9600GT(VRAM 512MB)Vista HomePremium(SP2)で実行、CUDAエンコ
でQ6600使用率83％と書いてあるけど、GT240のほうが高速ですよね？
http://pc.watch.impress.co.jp/docs/column/tawada/20091117_329556.html によると
GeForce 9600 GTの1,800MHz/256bit メモリ帯域幅は57.6GB/sec、
GeForce GT 240のGDDR5/3,400MHz/128bit 54.4GB/sec、
GeForce GT 240のDDR3/2,000MHz/128bit 32GB/sec。
DDR2-667（5.3GB/s） DDR2-800（6.4GB/s）デュアルだと倍。
もしかしてDDR2がボトルネックになる？　
GT240 GDDR3とGDDR5でCUDAエンコの速さの違いどうなんでしょう？
そこらへんが一番気になります。

>>819
最新ではBフレームが16まで、Average/Variable/CBR/2pass/3pass できるみたいだよ。
x264ではAverage/Variableは何故かコマ落ちする。

>>820
SP216&DDR3のGTX260とCUDAコア数96のGT240だと、エンコ速度二倍差が出ます？
PHARAOH 500W電源だとGTX260+Q6600はギリギリかなぁ。

(Q6600+GT240)複数台にするほうが経済的かな? G41とDDR2が1組余ってるし…

http://hothardware.com/articleimages/Item1477/san.png
doubleはfloatの1/2*9ってとこか。

Teslaでやって1/2とかだったら、ロックなんだろうな。

これは低すぎるのでOpenCLでDoubleがHWサポート
されてないのではないかと

>>821
GDDR3とGDDR5とでは、メモリ帯域が倍違うので、ストリーム系のアプリでは大きく変わります。
ていうか、このクラスのボードで2枚挿しとか意味がない。

>823−824
むしろ理論値で単精度の5分の1になるはずのRADEONの倍精度が半分程度で済んでる方が気になる

>822
GT240なんてゴミ買うぐらいならいっそ中古のQ9xxxのCPUでも買った方が良くね？

安物買いの銭失いとはこのこと

CUDAに大きな期待を抱かないほうがいい

>>823で、GTX295が285よりもスコア低いのはなぜ？

>>828
単純にGPUを一個しか使っていないからだと思う。
単一GPUなら285の方が速いからね。

>>825
CUDAエンコはストリームと違ってGDDR3もGDDR5殆ど関係無いってさ。
SP数(CUDAコア数)でほとんど性能が決まる。

>>827
最新のMediaCoderでVBRでやってみたらどう？

MLB オバマ始球式

HPC向けGPGPU終わりつつあるな

102 ：Socket774 ：sage ：2010/04/06(火) 23:24:12 ID:n4owrnuu




>>98
> HPC向けはどうなるんだろう？

マキーノの話だとこんなのが。
http://grape.mtk.nao.ac.jp/~makino/journal/journal-2010-04.html#3
> 一枚5万とかで買うのでない限り GPU は価格性能比では CPU に勝てなくなってしまった

以前はGPUの方が同コストのCPU比で10倍↑とか軽く叩き出してたけど、
CPUはマルチコア化が進みまくり値段下がりまくり…
ハイエンドGPGPU買うよりCPUの方がコスパが良くなってしまった。
基本直線番長のGPGPUよりCPUの方が扱いやすいし、プログラミングも先行きわからない
CUDAやらなくても、今までやってきた事そのままで走るし…

確かに、OpenMPとかでそれなりに性能出るならそっちの方が超簡単だもんな。

精度を削って性能を出したマシンで有名になったのにこういう時は倍精度の話だけか

>>832
マルチコア化が進みまくりって、↓か？当面CUDAエンコの半分にもならんだろう
http://akiba-pc.watch.impress.co.jp/hotline/20100403/ni_c6176se.html
http://akiba-pc.watch.impress.co.jp/hotline/20100327/ni_cw3680.html

>>835
CUDAエンコは実用性ゼロと言われているだろ
HPC以外の分野でCUDAを利用している人はいないだろ

>>834
そこ大事だよな。CPUと違って、倍精度・単精度の使い分けで性能違うもんな。
て、マキーノは昔、計算パスの場所によって計算精度が違う計算機を作ってたもんな

>>836
エンコの中の人が面倒くさがっているだけなんじゃないの

>>838
環境を作るのがめんどくさい。
コードを書き換えるのがめんどくさい。
最適化するのがめんどくさい。

>>838
まともなエンコを利用できないという事実が重要

たしかにOpteron12コアX4がFermiと同じ値段ならGPU終わるな。
既存のコードがそのまま動くし。
CUDAでこれ以上の性能が出るアプリは限られている。

fermiの場合ボッタクリなだけのような。
最終的にはfusionみたいな物に落ち着くだろうけど。
GPUのアーキテクチャとしては変に汎用に振るより
コンパクトな割に暗号解読みたいに得意なものが速い
と言う方が良いだろう。

>>838

正直エンコにCUDAを適応出来る処理が少なすぎる

x264の開発者が全員RADE愛好者だったとかいうオチなら面白い

>>839
ペガシスがKarmaPlusに導入したCUDAエンコをTXP4になかなか導入
しないのはそれが原因かｗ

>>841
メニーコア化が進展すればGPUは終わるよね
実際のところ、全然進んでないからNVは助かっているけど

メニーコアのメニーの次元が、GPUとCPUじゃ、全然違うしな

新しいものが出てくるときはチャンスだと思うんだが、このスレではそんな気配かけらもないな。
やっぱり日本人てダメなのかね

だって、HPC向けはコストパフォーマンス悪いし
GeForceは機能削られまくりだし、いまいちなんだよね

ATOM+IONチップセットで
エンコ爆速になったりしないか

>>850
動画データの転送には最低でもPenDは必要。

>>848
バカンスの概念が無い日本人はここぞって時に余力が無い。

>>841
4万のマザーに32,480円の8コア載せるより
5千円のマザーに1.6万円の4コア載せて数万のビデオカード挿す方が数倍早い
んだろうし、どっちもムーアの法則どうり進化すればGPUは当分優勢では？

GPUで縁故するとなんであんなに汚いの?

ソースみたいとわからん。

SpursEngineでエンコして汚いのはハードのせい
CUDAでエンコして汚い場合はソフトが成熟してないから

>>851
適材適所ってことだよなー
y=a*x+b を100万個×10万回　みたいな計算には巨大コアはいらない、
小さいコアがたくさんある方が速い　みたいな。

>>851
なぜ今頃ムーアの法則？
成り立たなくなってかなり経つんだが

え？

>>856

え？

>>856
　　　　　　 　　 _......_
　　 　 __　　 ／.::::::::::｀:.､
　　／, - ｒ, /::::i::::ﾊ:i:::::;::',
　 r-'ヽ./ｲ i::::|ｲ/'　'-ﾍl:::i
　 ｀ ｰ､i　{ l::::l '"´　￣ l:::l　＜またまたｗご冗談を
　　　　l!　l l::::ﾄ､ r_っ ,ィ:::l
　　 　 l　 ﾄヽ::l弋ニ＜l::::l!
　　　　 ゝ- イ`　　 　ｲ^ｲ
　　　　　　　 | 　 　　/-{′

http://wiredvision.jp/news/200709/2007092021.html
http://www.itmedia.co.jp/news/articles/0906/18/news006.html

あと10年ぐらいはどうにかなりそうな気がするけど

>>856はfreelunchは終わった発言と混同しているに1000ペソ

>>860
どんだけ低クロックなんだよ、電算機系分野の１０年ってものすげぇ進歩するんだぞ

クロック？
ムーアの法則って集積密度の話じゃなかったっけ

ムーアの法則なんて、明らかに無理だと分かった時点で
定義の方を変えて無理やり存続させているだけだろ。

IntelはAMDの様子見して出し惜しみ＆殿様商売
「半導体の集積密度は18〜24ヶ月で倍増する」ゆえ
CPUの性能は2年で倍近くになる。1年だと√2倍弱
GPUも同様。

集積密度≠性能ということをきちんと理解しましょう

GPUで無理矢理あれこれするより、
計算専用のユニットを別途開発した方がいいんじゃないの？と思う
GPUより効率よくできる部分もあるだろうし

　 　∩＿＿＿∩　　 　　　　　　｜
　　 | ノ＼　　　　 ヽ　 　　　　　　|
　　/　　●゛　　● |　　　　 　　　|
　 |　∪　　( _●_)　ミ　 　　　　　j
　彡､　　　|∪|　　 |　　　 　　　 Ｊ>>867
/　　　　 ∩ノ ⊃　 ヽ
(　 ＼　／ ＿ノ　|　 |
.＼　“　　／＿＿|　 |
　　＼ ／＿＿＿ ／

x264がCUDA対応してないことと、AVIUTLでCUDA使えないこと、
MediaCoderのcudaH264Enc.exeの画質をx264並みにするにはビットレートを
何割増しにすればいいのか不明なこと、
PowerDirectorは設定がゴミなこと、

が問題かな

http://forums.nvidia.com/index.php?showtopic=165055

Double precision is 1/2 of single precision for Tesla 20-series, whereas double precision
is 1/8th of single precision for GeForce GTX 470/480

倍精度だとteslaの1/4だけと、値段は1/5なんだよね

ノードの数がGTX480は480でteslaは442なんだよね

あれ、HD5870でよくね？

http://pc.watch.impress.co.jp/docs/news/20100405_359261.html

AMDは単精度しかなかったみたいね

IEEE754準拠ではない64bit double floatはR7xxの頃からあった。
俺は使ったことないからわからないけど。
少なくともドキュメントには、Radeon HD4xxx を除外するような文言はなかった。

R8xxはIEEE754準拠の命令が結構揃ってるよ。

なんでラデはFFTのライブラリを出さないのか理由がわかりますか

つくれば？

ですよね

http://ascii.jp/elem/000/000/513/513385/
秋葉に出たみたいです。どうしようかな。

熱的にやめとけ。

売り切れたそうです。GTX480、3枚買った方がいるとか。

>>875
IEEE754準拠の精度になったのがRV770で
RV670のころからdoubleはサポートされているよ。

3枚とかアホとしかww

470なら3台行けるかなー

オークションで売るんじゃないか

PCI-EXｘ16　3つあるマザーでPCIと交互にあるマザーでCUDAエンコに
使うんだろう。

結局、倍精度は削られているのかね？
削られていなければ買いたいのだが。

まだ実測した人はいないみたいですね。スレの人が手に入れるのをまちましょう。

GTX480じゃ倍精度無効になっているだけで
sandraのベンチのやつはEMUですが

無効じゃなくて1/4ですよね
ベンチはOpenGLがそうなってるということで

OpenGLのなにが？

sandraのあれはOpenCLだろ>倍精度無効

sandraはOpenCLだろうがGLだろうがCUDAだろうがCSだろうがvideorenderingだろうが
ハードで使えなきゃエミュでだすよ

http://pc11.2ch.net/test/read.cgi/jisaku/1269857880/740
購入者

>>894
ゲーム目的だったらHD5970を買ったほうがよかったのにね

M/Bとか貧弱なんでゲーマでもないみたいでなんで２枚も買ったのか

480/470の倍精度演算削られたのか…
倍精度演算やる人はぼったくり価格のC2050/2070買えってか

>>897
そのソースはどこ？

>>897
まぁ、このスレ住人が実際にCUDAで試すのを待とうや。

米アマゾンはまだ発売前になってるね

これだけ出回ってるのに倍精度はおろか単精度の演算を流す人もいないとは

出来る人はこんな所に来ない、つまりここは無能の衆が集う所だから

みんな科研費で買うんじゃない？５月まで待たないと．

無駄遣いはやめてもらうようにこのスレのことも仕分け人に伝えとかないといかんね

科研費が無駄に物価を高騰させてるな
さっさと仕分けされろ

無駄な科学者・技術者は農業や林業にまわって効率を上げてやってほしい。

470ポチった。ついでにCUDAの入門書もポチった。
両方到着は14日予定。おまえらよろしく

>>897
倍精度のテストプログラムか何かあれば提示してくれないか？
手元にGeForce系列VGAが一切ないので一切やったことがないんだ

14日、期待してますよ。

人柱になってくれるのに情報薄くて申し訳ないが
このスレの上の方にあるCUDA公式での行列かけ算コードを倍精度化したものを走らせるのが一番簡単かな？
手持ちがあれば提供したいが

未だELSAからボードが来ない……

480のCUDAのベンチマーク結果はどかでみられませんか？
ゲームのベンチだといくらでもみつかるのですが。

　　　　　　　　　　　　　_(こ＾)､_
　　　　　　　　　　　 〃､__ﾉノ､__,ヽ
　　　　　 　 　 　 　 {.っ＞　 ＜っﾄ､
　　　 　 　 　 　 (⌒i　 （千于`ｰ┴'─────┐
　　　　　　　　　（O人　 `ー|　 　 　 　 　 　 　 　 |
　　　　　　　　　　　/⌒ヽ(＾う　見せられ.　 　 　 |
　　　　　　　　　　　`ｧー─イ　　　　ないよ！ 　 |
　　　　　　　　　　　/　　（0::|＿＿＿＿＿_＿_＿_|
　　　　 　 　 　 　 /＼＿＿＿_/
　　　　　　　 　 ／　　　/　　⌒ヽ
　　　　　　___／　 ／￣￣｀)　　ﾉ
　　　　　(__r＿__ノ　　　　 (.＿＿つ