09/10/08 19:32:36
関連サイト
CUDA Zone
URLﾘﾝｸ(www.nvidia.co.jp)
フォーラム
URLﾘﾝｸ(forum.nvidia.co.jp)
CUDAに触れてみる
URLﾘﾝｸ(chihara.naist.jp)
cudaさわってみた
URLﾘﾝｸ(gpgpu.jp)
CUDA のインストール
URLﾘﾝｸ(blog.goo.ne.jp)
NVIDIAの「GeForce 8800 GT(G92)」と次に控える64-bit GPUアーキテクチャ
URLﾘﾝｸ(pc.watch.impress.co.jp)
CUDAを使う
URLﾘﾝｸ(tech.ckme.co.jp)
NVIDIA CUDAを弄ってみた その2
URLﾘﾝｸ(dvd-r.sblo.jp)
CUDAベンチ
URLﾘﾝｸ(wataco.air-nifty.com)
KNOPPIX for CUDA
URLﾘﾝｸ(www.yasuoka.mech.keio.ac.jp)
3:デフォルトの名無しさん
09/10/08 19:33:30
前スレ
【GPGPU】くだすれCUDAスレ【NVIDIA】
ｽﾚﾘﾝｸ(tech板)
忘れてた・・・。
4:デフォルトの名無しさん
09/10/08 23:06:10
pert2という時点で>>1の頭が超低レベルであることがわかっちゃうけどいちおういっておこう乙
5:デフォルトの名無しさん
09/10/09 00:32:17
>>1さんの「超低レベル」というお言葉に甘えて
インストール以前の質問なんですが、
普通にPCを使ってるときのグラフィックボード用のドライバと
CUDAドライバって別物ですよね？
URLﾘﾝｸ(www.a2art.com)
これの「ドライバは特に入れなくてok」ってどゆこと？
6:デフォルトの名無しさん
09/10/09 09:46:30
>>5
リンク先を読まずに当て推量で回答。
リビジョンの違いこそあれ、原則同じもの。
勿論、対応さえしていればどちらを使っても同じ。
逆に、NVIDIAのドライバが全く入っていないのなら入れる必要あり。
7:デフォルトの名無しさん
09/10/09 18:25:44
>>5
取り合えず最新のドライバいれて、CUDAツールキットを入れればおｋ
リンク読んでねえけど
8:5
09/10/09 21:08:40
レスサンクスです
具体的に言うとですね。
GeForce 9500M GT の入ったモニタ一体型VAIOがありまして、
それでCUDAってみようと思ったのですが、
CUDAドライバがうまくインストールできなくて、
とりあえず、ツールキットとSDK入れて、
bandwidthtestしてみたら、
認識できるGPUがないとかなんとか（今手元にないのでスマソ）
言われたんですよね。
やっぱCUDAドライバ入れなきゃいけないのかな。
9:デフォルトの名無しさん
09/10/09 22:29:43
>>8
モバイル版のドライバじゃないとけないんじゃね？
10:5
09/10/09 23:28:23
Mってのがモバイルみたいですね
ノートじゃないから、きづかなかった
ちょっといろいろやってみます
ありがと
11:デフォルトの名無しさん
09/10/14 00:00:52
画像処理なんかのサンプルを見ると
メモリ転送を行っていないのにカーネル呼び出しの引数になり
GPU側で使われている変数があるのですが
どのメモリに格納されているのでしょうか
filter_kernel<<<grid, block>>>(d_result, d_data, width);
こんな感じでwidthのメモリ転送を明示的に行っていないのに使用されています
12:デフォルトの名無しさん
09/10/14 01:08:33
言葉じゃよく判らん。 もちっと例くれ。
あるいは誰でもダウンロードできるサンプルで説明してくれ
13:デフォルトの名無しさん
09/10/14 01:20:42
まぁ，
__device__
__constant__
__shared__
のいずれかの修飾子の変数が予め確保されてるんだろなー･･･
14:デフォルトの名無しさん
09/10/14 20:18:42
>>11
配列でない引数はレジスタに入るんじゃないの
15:デフォルトの名無しさん
09/10/15 00:02:51
>>12
URLﾘﾝｸ(developer.download.nvidia.com)
ここのboxfilerというサンプルに載っています
boxfilter_x、boxfilter_yというカーネルを呼び出しているのですが、width,height,radiusが
メモリ転送されていないにもかかわらず引数として使われています
>>14
レジスタでしたか
希少なのでコンスタントメモリあたりで代用したほうがよさそうですね
16:デフォルトの名無しさん
09/10/15 00:36:03
引数は勝手に転送される
格納場所はLocal Memory
17:デフォルトの名無しさん
09/10/15 00:45:45
>>16
ありがとうございます
助かりました
18:5
09/10/15 15:35:43
相変わらずCUDAれません(;o;
ツールキットとSDKはちゃんとインストールできているようです。
CUDA2.2を入れようとしているのですが、
URLﾘﾝｸ(www.nvidia.co.jp)
ここの
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と言われます
URLﾘﾝｸ(www.nvidia.co.jp)
ここの一番下にVAIOは対象外って書いてあるから、VAIOじゃだめなのかな、、、
19:デフォルトの名無しさん
09/10/15 15:50:05
それかもね。ご愁傷様。
20:5
09/10/15 18:04:53
どんどんダウングレードして、2.0用のCUDAドライバ入れたら入った。
そしたら、今度はなぜか2.2のツールキットとSDKがアンインストールできないヽ(`Д´)ﾉ
エラー6003でぐぐっていろいろやったけど、できない。
アンインスコしないで、2.0インストールしたらまずいですよねぇ
連投スマソ
21:デフォルトの名無しさん
09/10/15 18:26:01
アンインストールしなくても医院で内科医。
22:5
09/10/15 21:28:37
>>21
あなたを信じてインスコしたら、
2.0のインストーラーが自動的に2.2をアンインスコしてくれました。
無事2.0環境整いました。
どうもでした(^^
23:デフォルトの名無しさん
09/10/16 03:15:22
多倍長整数使いたいときは自分で作るしかないん？
24:デフォルトの名無しさん
09/10/16 09:54:23
cudaに多倍長整数ライブラリは付属していないので、そういうことになりますね。
25:デフォルトの名無しさん
09/10/17 01:17:37
for文のネストにどこまで耐えられるのかがわからない＼(^o^)／ﾑﾘﾎﾟ
CUDAのスタックって案外少ない？
ハードウェアによってまちまち？
26:,,・´∀｀・,,）っ-○○○
09/10/17 01:23:49
束胃腸やるのはいいが、そもそもキャリーアップフラグとかあったっけ
27:デフォルトの名無しさん
09/10/18 17:41:26
>>25
スタックはないです
28:デフォルトの名無しさん
09/10/18 21:11:14
>>25
template使った再帰もどきとかやらない限り
心配することはないかと思われる。
29:デフォルトの名無しさん
09/10/18 21:25:31
CUDAカーネルのルーチンはインラインとして扱われるんで、スタックなし。
んでもって、レジスタが使用できないデータタイプだったり、レジスタが枯渇してたりすると、ローカルメモリ(デバイスメモリ)を使用しだすので、パフォーマンスが一気に低下するんで心配してください。
30:デフォルトの名無しさん
09/10/18 21:47:56
25ですが，みなさんありがとうございます．
最近になってサンプルコードをじっくり読み始めましたが，
自分のやろうとしてるくらいのことは素直に出来そうで安心です．
んでもって，もう少しアーキテクチャを知らないといかんな．．．
31:デフォルトの名無しさん
09/10/20 02:09:58
CPU側のプログラミングにインテルのiccを使いたいのだけど、どこかにやり方は載っていませんでしょうか？
OSはWindowsXPです。
32:デフォルトの名無しさん
09/10/20 04:59:45
>>31
ソース分けてオブジェクトリンクすれば？
33:デフォルトの名無しさん
09/10/20 19:12:31
>>31
つか勝手に使えばいいじゃん
34:デフォルトの名無しさん
09/10/22 15:51:00
CUDA2.0入れてVS2005でtemplateビルドしたら
LINK : fatal error LNK1181: 入力ファイル 'cutil32D.lib' を開けません。
って言われます
VISTA64版なので、cutil64D.libはあっても32はないのですが、、
どこを修正すればいいのかもわかりません。
ちなみにVSも使ったことない初心者です
どうしたらよいでしょう？
35:デフォルトの名無しさん
09/10/22 19:38:02
>>34
上のツールバーの真ん中の方にWin32となっているところをx64に変更。
そもそも64bit用のツールをインストールしているかどうか・・・・。
36:34
09/10/22 23:00:43
>>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:デフォルトの名無しさん
09/10/23 02:28:17
スレ違い、と言うか人生の道を踏み外したな…。
38:34
09/10/23 02:37:36
>>37
せめて正しいスレに導いておくんなまし
とりあえずコンパイルできないことには段ボーラーです。
踏み外す以前に舗装道路なんて歩いたことございません。
39:デフォルトの名無しさん
09/10/23 21:18:44
メモ、とりあえず報告しておきます。
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ビルドを行う。
まずはURLﾘﾝｸ(www.sharkpp.net)
で64bitビルド出来るようにする。
このまま進めてくと、
nvcc fatal : Visual Studio configuration file '(null)' could not be found
みたいなエラーが出てくると思う。
これを解決する方法が→URLﾘﾝｸ(forums.nvidia.com)
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のコマンドラインとかちゃんと設定しておけばいける！はず・・・
40:デフォルトの名無しさん
09/10/23 21:41:37
NVIDIAが目指す究極のプロセッサへと進む「Fermi」アーキテクチャ
URLﾘﾝｸ(pc.watch.impress.co.jp)
41:デフォルトの名無しさん
09/10/23 21:46:13
>>39 の意味はわからないけど、64ビットで使うときは、コンパイルとリンクのオプションに
-ccbin "C:\Program Files (x86)\Microsoft Visual Studio 9.0\VC\bin"
をつければたいていうまくいく。
42:デフォルトの名無しさん
09/10/27 10:37:15
>>36
その仕事、うちの会社にくれ。あんたがやるより速いものが早くできるぞ。
43:デフォルトの名無しさん
09/10/29 13:36:59
> VSも使ったことない初心者です
でいきなりCUDAか… なんというかどういう会社だと…
「まず32bitで試させてくださいお願いします」
で後から64bitのことを考えたほうがいいんじゃね？
44:デフォルトの名無しさん
09/10/29 14:05:12
>>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で遊ぼうかなと
45:デフォルトの名無しさん
09/10/31 10:11:38
>>44
追記... Win7は64bitです。
ToolKitとSDKは2.3です。
46:デフォルトの名無しさん
09/10/31 19:19:26
ボリュームテクスチャをルックアップテーブルとして使いたいのだが
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);
で実行時間に１０倍以上の差がついてしまうんだがなぜだろう。
メモリ上の配置でかなりパフォーマンスが落ちるということはありますか？
47:デフォルトの名無しさん
09/10/31 20:16:28
>>46
テクスチャキャッシュにヒットするかどうかでパフォーマンスは大きく変わります。
48:デフォルトの名無しさん
09/11/01 06:52:25
win7x64にCUDA_VS_Wizard_W64.2.0.1入れてもテンプレに表示されるようになる
けどエラーでプロジェクトが生成できない・・・
Vistax64の時はこれで一発だったんだが・・・
7でも32bitのSDK入れてCUDA_VS_Wizard_W32使ったらすんなり入ったが
49:デフォルトの名無しさん
09/11/01 07:30:31
>>48
使えると便利だから使えるようになってほしいな。
フォーラム見に行ったら10/31までmergeでdownだよと書いてあったorz
復活したらエラーメッセージをポストしてください
50:デフォルトの名無しさん
09/11/02 07:33:45
誰かtesla使ってる人いる？
tesla一枚挿しのマシンに、リモートからwinデフォのtelnet使って動かそうとすると、
tesla自体認識しないんだが・・これどうやって使うの？
51: ◆0uxK91AxII
09/11/02 07:57:43
認識してすらいない、っと。
52:デフォルトの名無しさん
09/11/02 09:31:50
>>50
TeslaがささっているマシンがWindowsなのかどうか読み取れませんが、
リモートからWindowsマシンのTesla使うのであればVNCを使うことが多いですね。
リモートからの利用中心であればLinuxにしてしまった方が・・・。
53:デフォルトの名無しさん
09/11/02 12:21:30
>>47
詳しく！ お願いします 助かります
54:デフォルトの名無しさん
09/11/02 12:23:50
逆に、デバイスから高速かつランダムな位置に参照したいときはどのメモリ使えばいい？
55:デフォルトの名無しさん
09/11/02 12:28:34
>>54
サイズも示さないで高速かつランダムって、馬鹿なの? 間抜けなの?
レジスタに乗るなら1クロックでアクセスできる。
共有メモリに乗るなら4クロックでアクセスできる。
グローバルメモリに乗るならランダムと言う前提からcoalscedであることは期待できないから最長数百クロック掛かる。
56:デフォルトの名無しさん
09/11/02 12:31:27
>>50
んじゃ私はLinuxサーバ機であるとして。
一部のメーカのサーバ機(BIOS依存かマザボ依存かは不明)ではTeslaをGPUとして認識する。
その場合は(オンボのGPUが切り離されて涙目になることはあっても)なんとかなる。
処が、一部のサーバ機ではGPUとして認識できないので自前でデバイスを叩かないと認識しない。
57:デフォルトの名無しさん
09/11/02 17:33:33
確かにそうだな サイズは20メガバイトくらいだ
58:デフォルトの名無しさん
09/11/04 18:14:36
8800GTXがあるんだが四倍率早く処理するにはなにがある？
59:デフォルトの名無しさん
09/11/04 21:40:42
ちょっと日本語でＯＫ
60:デフォルトの名無しさん
09/11/04 23:27:30
8800GTXを4本さす
61:デフォルトの名無しさん
09/11/04 23:38:39
クロックを４倍に
62:デフォルトの名無しさん
09/11/05 09:05:00
>>52
レスありがとう。
winでGTX280とTeslaさした状態だと、両方のデバイス認識してて、GTX280外してTesla単体にすると、
Teslaを認識しなくなるんだけど、原因わかる人いる？ドライバではないと思うけど・・
ちなみに、winデフォのリモートデスクトップで確認しました。
63:デフォルトの名無しさん
09/11/05 09:10:33
DeviceQueryでどうなるのかはっきりして欲しい
64:デフォルトの名無しさん
09/11/05 10:22:46
>>62
Tesla単体のときって画面は何に表示されるのでしょうか？
TeslaはGPUを搭載していますが、ビデオカードではないことに注意してください。
Windowsの制約で複数ベンダーのビデオドライバを同時にロードできないので
オンボードにATIやIntel製のビデオチップとかがあると駄目です。
どうしてもその構成にしたいのであればLinuxに。
というかWindowsのバージョンによるけどまず無理です。
65:デフォルトの名無しさん
09/11/05 13:05:40
馬鹿みたいに高いのに、すごい仕様だな
66:デフォルトの名無しさん
09/11/05 15:11:24
>64
Win7からはWDDM1.1対応ドライバであれば普通に混載できる
67:デフォルトの名無しさん
09/11/05 18:11:37
なんかメモリアロケートのアルゴリズムいじったら五倍以上速くなった、、
あやうく三枚発注するとこだったよ。よかった
68:デフォルトの名無しさん
09/11/06 00:56:55
CPUの100倍速い、になるまで満足しちゃいかん
69:デフォルトの名無しさん
09/11/06 02:21:26
将来性は銅なのよ?
70:デフォルトの名無しさん
09/11/06 02:27:56
GeForce 8600 GT使って、3日目の初心者です。
素朴な疑問なのですが、GeForce 8600 GTには512MBのメモリが実装されているのですが、
cudaMalloc を行って、alloc領域が512MBを超えた場合には、どの様になるのでしょうか？
ご存知の方がいらっしゃれば、事象や回避策を教えてください。m(_ _)m
71:デフォルトの名無しさん
09/11/06 03:25:19
>>70
素朴な通り取れない。
エラーが返ってくるはず。
72:デフォルトの名無しさん
09/11/06 03:27:14
>>67
これはCUDAの話？
73:70
09/11/06 08:35:41
>>71
ありがとうございます。
まだ、某サイトにあった手順でHellow…が出せた段階です。
もしや実装メモリ量を超えると、HDDとの間でガリガリswapでもするのかなぁ?っと思った次第です。
74:73
09/11/06 08:52:35
うっ…
w 綴り違うし orz...
75:デフォルトの名無しさん
09/11/06 10:32:36
>>73
大丈夫、swapなんてするはずもない。そもそも512MiB全部使えるわけでさえない。
その世代で512MiB搭載のボードなら、実際に使えるのは高々511.69MiBだけだ。
しかも、そのボードを実際に描画にも使っているならそこから更に画面解像度に依存した分使える量が減る。
逆に、目一杯cudaMalloc()で確保した状態で画面解像度を上げたりすると面白いことになる。
76:デフォルトの名無しさん
09/11/06 23:30:47
>>73
つかQueryつかってデバイスメモリの上限チェックしないと
77:デフォルトの名無しさん
09/11/07 00:27:47
>>63,64
レスありがとう。どうやら、winではディスプレイ出力できるグラボを最低ひとつ積んでないと、
cuda使えないみたいです。つまりTesla単体では動かない。。deviceQueryしても、
there is no device supporting cudaしか返ってきませんでした。倍精度の数値計算したくて
Teslaシングルで動かしたいときは、nvidiaの公式が推奨してるようにLinuxでやるのがベスト
みたいです。いろいろとありがとうございました～
78:デフォルトの名無しさん
09/11/07 03:56:44
snow leopardじゃつかえないの？
79:デフォルトの名無しさん
09/11/07 06:48:09
>>74
皮肉のつもりじゃなかったのか。
80:34
09/11/09 21:08:31
いまだにtemplateのコンパイルもできません。
32bitコンパイルしようとするとlibがないといわれ、
64bitコンパイルしようとするとスキップされる。
人に聞いたら、ソースが32bit用なんじゃないかって言われた。
ちゃんと64bit用SDKインスコしたつもりなんだけど、
64bit用のソースでなかったりするのでしょうか？
81:デフォルトの名無しさん
09/11/09 21:49:20
>>80
リンカが言うようにlibのファイルがないんでしょ？
パスが通ってるか、ファイルがあるかどうかくらい自分でなんとかしなよ。
82:デフォルトの名無しさん
09/11/09 22:40:46
>>81
レスサンクスです。
でも、>>34に書いたとおり、32bit用のlibはもともとないんです。
まったく初心者なので、勘違いしてたらすんません。
83:デフォルトの名無しさん
09/11/09 22:45:41
>>82
それなら、リンカに指定するライブラリをcutil64D.libに変えたらいいじゃん。
84:デフォルトの名無しさん
09/11/09 22:55:47
>>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が入っていないんじゃないかな？詳しく確認はしていないけれどね。
85:デフォルトの名無しさん
09/11/09 23:03:07
>>83,84
ありがとっす～
試してみます！
86:73
09/11/09 23:08:04
>>75､76
どうもです。お返事遅れました。
先週はXPでやってましたが、今日からはFedoraに入れ直し実施してます。
結果ですが、画面が点滅したり、真っ黒になったり、となりました。
復旧には、sshで入ってプロセス消してもだめで、結局rebootで…。
ちゃんとsizeみて使う様にします。
m(_ _)m
87:74
09/11/09 23:11:42
済みません。連投です。
>>79
はぃ、
天然ってより、単に英語は苦手です。
88:34
09/11/09 23:12:49
exe消してもだめでした、、、
つか、なんかソリューションエクスプローラーのところで
template.cuを右クリック→プロパティにした時点で、
「操作を完了できませんでした。エラーを特定できません。」
ってなる(´･ω･`)だめぽ
リンカへの指定はちょっと勉強してみます。
89:デフォルトの名無しさん
09/11/11 17:55:34
どうにもうまくいかないので質問なのですが
short intでGPUにmemcpyしたデータをFFTしたいのですが、
short intのデータをfloatにキャストするには
どうしたらいいのでしょうか？
int→floatやfloat→intはプログラミングガイドに載っていた通りできたのですが、
short intだとうまくいかないのは何故だろう…？
90:デフォルトの名無しさん
09/11/11 18:45:15
>>89
何を試したのかよく分かりませんが、short2型をfloat2型に変換する感じがよさそうな。
91:デフォルトの名無しさん
09/11/11 21:49:44
プログラミングガイド読んでたらﾊｧﾊｧ(´д`*)してきた
なにこの気持ち（*´д`*）ﾊｧﾊｧ
92:デフォルトの名無しさん
09/11/12 17:01:33
>>90
おかげさまで出来ました。ほんとにありがとう
ちなみに試したのは
__int_as_float(int)のような形式のキャストと
__int2float_rn(int)のような形式のものです。
どちらもshortをキャストしたり、shortにキャストすることが出来ませんでした。
93:デフォルトの名無しさん
09/11/12 17:28:04
GPUに型変換するのは遅いと思うのだが。
94:デフォルトの名無しさん
09/11/12 20:44:37
その通りだな、そこで相当なロスが出ると思う
95:デフォルトの名無しさん
09/11/12 23:30:59
voidでコピーして、floatのポインタで計算すれば良いんじゃねーの?
何か勘違いしてる?
96:デフォルトの名無しさん
09/11/12 23:47:32
>>91
興奮してるなぁｗ
97:デフォルトの名無しさん
09/11/13 00:31:33
int -> floatのキャストなら1命令だからむしろ速いんじゃないのか
アライメントは考慮したほうがいいけど
98:デフォルトの名無しさん
09/11/13 10:27:25
いやだから、gpuでshort -> intはダメだろ。
intで渡しておいて、int -> floatはイントリンシックで。
99:デフォルトの名無しさん
09/11/13 10:54:01
別にshortで渡しても問題ないかと。
GPU内でshort -> intの変換をしてから int -> floatの変換が起こりますが、
そんなに頻繁に変換をするわけでなければメモリ転送量の削減効果の方が大きいでしょう。
＃32-bitアクセスにするためにshort2型をお薦め。
100:デフォルトの名無しさん
09/11/13 11:55:38
GPGPUのハードのいいベンダーだか機種だかオススメある？
予算は潤沢にあるとして
101:デフォルトの名無しさん
09/11/13 12:11:15
NVIDIAのTesla。つーか、Teslaのラック筐体マジお勧め。
漏れなくNVIDIAのサポートがついてくるから。
102:デフォルトの名無しさん
09/11/13 12:36:58
>>101
CUDAしたいからtesla c1060を搭載した機種を買うのは
当然なんだけど、teslaを組み込んだ
ハード全体としてのオススメを知りたいのでした。
ようするにフラクティカだとかELSAだとか爆速だとか、
どこがいいのよっって話です。
teslaのラック筐体ってなに？
タワー型じゃなくてユニット型がいいってこと？
それともteslaの名を冠したサーバーがあるの？
すまんあんま詳しくないんでわからん
103:デフォルトの名無しさん
09/11/13 12:53:43
HPC用にTeslaC1060相当を4枚入れた、TeslaS1070っていうラック筐体があるですよ。
当然、PC筐体は別途必要。
PC本体なら、QuadroPlex2200S4ってーのがTeslaS1070にマザボをつけたような仕様だったかと。
ラック筐体じゃなければ、QuadroPlex2200D2がタワー型でGPU2枚挿しのPCになる。
NVIDIAの営業曰く、「GeForceなんてアキバ的発想はやめましょう」ということなので。
# GeForceでいいならELSA辺りがリファレンスボードをそのまま使うから安定しているけどね。
## つーか、QuadroPlex使うような予算があるならソフト開発受注したいぞw
104:デフォルトの名無しさん
09/11/13 13:05:58
>>103
ありがと～
調べてみる
105:デフォルトの名無しさん
09/11/13 14:48:43
URLﾘﾝｸ(www.supermicro.com)
これ、自分はアメリカで買った。
国内の代理店でも売ってるとこあるよ。
106:デフォルトの名無しさん
09/11/13 15:01:35
>105のPCすげぇ。TeslaC1060が4台も載っている。なのにVGAはMatrox。
107:デフォルトの名無しさん
09/11/13 15:20:48
おまえらどんな仕事してんだよ
108:デフォルトの名無しさん
09/11/13 15:25:42
自宅の治安を守る仕事
109:デフォルトの名無しさん
09/11/13 15:31:22
>>106
だってサーバーだから。
すんごくうるさいよこのマシン。標準でリモート管理機能が付いているから、
リモートから電源のON/OFFやキーボード、マウス、VGA、IDEポートの
ネットワークリダイレクトができる。
だから管理用IPアドレスだけ設定してサーバールームに入れっぱなし。
110:デフォルトの名無しさん
09/11/13 21:37:09
>>105
国内のページ見つけたけど
URLﾘﾝｸ(www.able.across.or.jp)
スペックが書いてあるようで書いてないような、、、
この値段は一体なんの値段なんだろう
111:デフォルトの名無しさん
09/11/14 00:42:44
GeForceとQuadroってどう違うんですか？？
112:デフォルトの名無しさん
09/11/14 00:49:20
シールが違う。
113:デフォルトの名無しさん
09/11/14 01:02:41
>>93-99
色々参考になります。ありがとう
とりあえずそれぞれのやり方で試して処理時間比較してみます。
114:デフォルトの名無しさん
09/11/14 01:16:16
>>110
書いてあるじゃん。
7046GT-TRF-TC4はサーバーのベアボーン。
箱、電源、マザーボード、ドライブエンクロージャ、そしてTesla C1060 * 4だけ。
あとはCPU、メモリ、ハードディスクを買ってきて刺してやれば動くってこと。
115:111
09/11/14 01:20:18
>>112
それだけですかｗ
ありがとうございました m(_ _)m
116:デフォルトの名無しさん
09/11/14 02:51:33
>>109
自宅にサーバールームか。アニメみたいww
117:デフォルトの名無しさん
09/11/14 03:30:00
やはりワレワレはコストパフォーマンスの良いものをと考えて…
i7-920、12GB
ASUS P6T7 WS nForce200が二個乗り
とりあえず GTX275あたりを二枚位かな
1000W電源
これでざっと20万コース？
118:デフォルトの名無しさん
09/11/14 04:37:05
randってつかえないのですか？
119:デフォルトの名無しさん
09/11/14 04:42:01
>>114
そっか単純にCPUもメモリもついてないのか
安いわけだ
120:デフォルトの名無しさん
09/11/14 05:12:47
サーバー用のメモリも安くなってきたから、サーバー用でいいかも。
Kingstonの4GB×3本で$350とかだ。
121:デフォルトの名無しさん
09/11/14 05:32:24
CUDA_SAFE_CALLって必ず必要なのですか？
122:デフォルトの名無しさん
09/11/14 05:36:18
> nForce200が二個乗り
おれの場合、O(n^2)だと、使うメモリ帯域<<計算量で、
x16で有る必要は無いなと思ったよ。
123:デフォルトの名無しさん
09/11/14 07:48:22
カーネルを実行する際に、引数で、変数を渡すことができますが、多数の数値を配列で渡したい場合、
配列をデバイスにコピーしてから、配列へのポインタを渡すしかないのでしょうか？
オーバーヘッドの少ない、数値の渡し方を差がしているのですが・・・
124:デフォルトの名無しさん
09/11/14 08:32:24
引数で配列のポインタを渡したって、どっちみちホストメモリ→デバイスメモリの
コピーは要るし？
125:デフォルトの名無しさん
09/11/14 15:41:12
すみません、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"
126:続き
09/11/14 15:43:23
以上を実行後、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が無いように思えます。
これはどのようにすればいいのでしょうか。
127:デフォルトの名無しさん
09/11/15 15:27:35
en_curr_regionがないってエラー出るんで、調べてたらbreakやcontinueのところでエラー出てるっぽいんだけど、CUDAってbreakやcontinueとの相性悪いのだろうか？
128:デフォルトの名無しさん
09/11/15 16:54:01
分岐は罪
129:デフォルトの名無しさん
09/11/15 20:02:42
質問です
CUDAを使ったプログラムをCUDA toolkitをインストールせずに実行する方法ってありますか？
130:デフォルトの名無しさん
09/11/15 20:56:22
試した事はないけど、driverと実行ファイルがあれば出来るだろうな
131:デフォルトの名無しさん
09/11/15 21:09:53
よ～し、パパ、CUDAを覚えて会社で活躍しちゃうぞ～！
132:デフォルトの名無しさん
09/11/16 01:31:49
VCランタイムがあるとして、cudart.dllと、cutil32.dll だけ必要(Win32)。
133:デフォルトの名無しさん
09/11/16 01:44:32
>>132
それって再配布できるの？
134:デフォルトの名無しさん
09/11/16 03:58:07
cudatoolkitのEULA、cudart.dll はredistributableと書いてあるな
CUTILは、cutil.cppの頭に「何の保証もしねぇよ」って書いてある。CUDAとは別で単なるサポート用なのかな。
ていうかcutil32.dllは無くても動くように書けると思う。ソースみたらくだらないぞこれ
nVidiaのドライバ入れればCUDAも動く、ようにはまだなってないのかな。
135:デフォルトの名無しさん
09/11/17 00:52:04
>>134
cudart.dll はCUDAの上位互換を保つために、配布するんだそうだ。
ちなみに、cutil32 or 64 .lib は使わないほうが。。。
ともかく、
cudart's version <= cuda driver's version
の条件が必須
136:,,・´∀｀・,,）っ-○○○
09/11/17 01:46:28
cutilはクソすぐる
137:デフォルトの名無しさん
09/11/17 01:55:12
いつの間にか団子が２ちゃんのスポンサーになってる・・・
【２ちゃんねる10週年】神秘の「乳団子」の秘密とは？
ｽﾚﾘﾝｸ(tech板)
138:デフォルトの名無しさん
09/11/17 05:06:18
Tesla C2050 $2,499 発表になったじゃない。倍精度500-630GFLOPS。
倍精度なら10倍速い、ってどかすか買うお金持ちな人居るんだろうなぁ。
139:デフォルトの名無しさん
09/11/17 08:14:48
今まで散々売られてもいないLarrabeeを盾にして現存のプロセッサと
それを擁護する人たちを馬鹿にしていた癖して
いざ完成が近づいてくると
ディスクリート版には期待していないｷﾘｯ
とか馬鹿にしてるよな。
140:デフォルトの名無しさん
09/11/17 09:20:00
おまえは誰と戦って(ry
141:デフォルトの名無しさん
09/11/17 20:19:11
書き込むスレ間違えたが分かっている人はいると思うから問題ない。
142:デフォルトの名無しさん
09/11/17 20:52:26
>>138
倍精度性能だけ見るとHD5870と同じ性能だというのは本当なのか？
単精度の1/2の性能でいいって言ってるけど
HD5870の単精度性能が数倍早いだけだという噂が
143:デフォルトの名無しさん
09/11/18 01:33:01
>>142
AMDはなにかを犠牲にしていると思うんだけど何を犠牲にしているの？
144:,,・´∀｀・,,）っ-○○○
09/11/18 02:05:20
>>143
たとえばプロセッサエレメントはFP演算とデータの読み書きは排他実行だったりとか
それでなくとも各プロセッサエレメントにぶら下がってる5つの演算ユニットが平均2～3程度しか稼動してないとか
145:デフォルトの名無しさん
09/11/18 06:05:13
ららびーは倍精度で１TFLOPSとか行くのかな。
Core i7で60GFLOPSくらいだっけ。
146:デフォルトの名無しさん
09/11/18 08:27:53
>>144
データ読み書きは別スレッドの分が並列に動くだろ。
メモリアクセス性能低いからALU命令の比率が余程高くないと
読み書き速度がネックになるけど。
147:デフォルトの名無しさん
09/11/18 10:21:58
初歩的な疑問なのですが
URLﾘﾝｸ(tech.ckme.co.jp)
のコードのように
__constant__ float g_idata[10000];
とか書いた時，このconstantメモリ領域は
いつ確保されるのでしょうか？
宣言だけ書かれてあっても
host,device共にその領域を使うコードを実行しない限り，
領域確保は行われないと思っていいのでしょうか？
148:デフォルトの名無しさん
09/11/18 12:05:28
>>147
GPUにはそもそも領域確保と言う概念がありません。
定数メモリ領域も共有メモリ領域もただそこにあるだけなので、宣言はあくまでもコード上のものでしかありません。
149:デフォルトの名無しさん
09/11/18 15:16:17
一応タイムシェアリングできてるみたいだけどGPUのメモリ領域の退避はしないの？
150:デフォルトの名無しさん
09/11/18 18:43:48
CUDAでCPU、GPUを並列に動作させられますか？
CPU GPU 並列 CUDA あたりで検索しても出てきません。
原理的にできそうな気もするのですが、
GPUで操作させている間はドライバを働かせているので無理なのでしょうか？
151:デフォルトの名無しさん
09/11/18 18:52:14
スレッドを使えばできますよ。GPUのプログラムは殆んどCPUに負荷をかけません。
152:デフォルトの名無しさん
09/11/18 22:50:37
>>150
GPUを起動したスレッドを待機状態のまま放置しておけば、他のスレッドでCPU資源を遣り繰りできます。
但し、SPの個数を超えるようなグリッドを食わせると途中で転送処理が入るので要注意。
153:デフォルトの名無しさん
09/11/19 06:44:52
CentOSで動きますか？
154:デフォルトの名無しさん
09/11/19 14:51:15
FermiでC言語が使えるってどーいうこと？
155:デフォルトの名無しさん
09/11/19 17:13:26
>>153
動きます。
156:デフォルトの名無しさん
09/11/19 19:37:25
fermi用のCコンパイラを用意しましたってことじゃねーの？
157:デフォルトの名無しさん
09/11/20 19:04:36
スレッド内部でレジスタがどのように使われているかわかりません
変数の数＝レジスタの数という認識をしていたのですが、そうではないようです
どなたかご教授ください
158:デフォルトの名無しさん
09/11/21 10:30:39
すいません、テクスチャメモリで疑問があります。
よろしければどなたか教えて頂けませんか？
コードは
URLﾘﾝｸ(gpu.fixstars.com)
のテクスチャメモリを参照するていう項目のソースをとりあえず勉強がてらコンパイルしてみました。
配列に値を3ずつ入れていって、バインドして、それにtex1Dでアクセスするというだけのものです
13行目にインデクスに3.14fを用いている事から線形補間なしの場合、小数切捨てで3になり
配列[3]の値を読んで、9が返ってくるのは理解できたんですが、線形補間を有りにすると
9と12の間を補間してるはずなのに7.921875なんていう値が帰ってきています。不思議に思ってv2.3のPrograming guideを読んだところ、P.139に記述があり、どうも与えた値が-0.5されているようです。
ので、はじめから+0.5シフトしてやれば正しい値になりそうですが、そもそも何故-0.5されているのかがわかりません。
検索してもぜんぜん出てこないので当たり前な事なのかもしれませんがドツボにハマってしまってわかりません。よろしければどなたか教えて頂けませんか？長文失礼しました
159:デフォルトの名無しさん
09/11/21 10:32:29
158です。
上記質問で訂正があります。programming guideのP.139ではなくP.137でした。すいません。
160:デフォルトの名無しさん
09/11/21 11:49:04
>>158
テクスチャだからとしか言い様が無いような。
配列の[0]の値は0 ～ 1が守備範囲で中心は0.5
...
配列の[i]の値はi ～ i+1が守備範囲で中心はi+0.5
...
配列の[N-1]の値はN-1 ～ Nが守備範囲で中心はN-0.5
N個の値を0～Nの範囲に均等にマッピングするためにこうなっている。
161:デフォルトの名無しさん
09/11/21 12:40:17
>>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
あれ？ちょっと違う値だね
162:158
09/11/22 15:50:57
>>159,160
ありがとうございます。
なんとなく理由が理解できた気がします。
が、こちらでも計算しましたが
buffer[3.141592653589]=7.924777960767
となって件のサイトの値7.921875とは合わないですね・・・
buffer[9.80665]=27.91995となって27.925781ではないですね・・・
低精度の線形補間というのは有効数字3桁って事なんですかね？
163:デフォルトの名無しさん
09/11/23 05:07:25
リニアフィルタつきのテクスチャでは、
座標(0,0)てのは左上端のドットが持つ四角い領域の左上端を意味するからね。
ドットの真ん中は座標(0.5, 0.5)になる。
真ん中にそのドットの本来の色（値）がくる。
精度はシラネ。リファレンスマニュアルに何か書いてあると思うけど。
164:デフォルトの名無しさん
09/11/23 14:39:37
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
165:デフォルトの名無しさん
09/11/23 18:16:06
>164
コンパイル時に -arch=sm_13 はつけてる？
166:デフォルトの名無しさん
09/11/23 19:09:46
>>165
それだっ！動作しました。
cublasは倍精度で動いたんで、いろいろ調べたんですが。
ありがとうございます。
167:デフォルトの名無しさん
09/11/24 05:01:59
>>164
CUDAって3．2までいっているの？
つい最近2．3が出たような気がする。
ひょっとして2．3の間違い？
168:デフォルトの名無しさん
09/11/24 12:40:41
おそらく。
今は3.0のβだね。
169:デフォルトの名無しさん
09/11/25 03:38:40
newでメモリを確保するのは反則ですか？
170:デフォルトの名無しさん
09/11/25 06:36:25
>>169
どうやってnewで確保された領域をGPUに転送するの？
171:デフォルトの名無しさん
09/11/25 10:28:10
>>169
CPU側のメモリでしたら反則ではありません。但し、VCで使う場合は*.cuでnewしてもmsvcrtにリンクできません。
172:デフォルトの名無しさん
09/11/25 19:49:02
>>169
パフォーマンス求めるなら、論外です。
173:デフォルトの名無しさん
09/11/25 22:00:21
newでメモリ確保するのがだめならどうやってCPU側のメモリ確保するの?
全部静的確保?
それともcudaHostAllocを使えって話?
174:デフォルトの名無しさん
09/11/26 06:26:18
メモリ転送ってＳＳＥで高速化されたりするんかな？
だとしたらnewより専用で用意されたものを使ったほうがいいかもね
175:デフォルトの名無しさん
09/11/26 06:58:15
>>174
メモリ転送ってCPUメモリ間？それともHOSTｰGPU間？
前者だったら高速化はされるけど、後者はDMAでPCIEにダイレクトに転送されるだろうから、
SSEは関係ないんじゃない？
176:デフォルトの名無しさん
09/11/26 17:12:50
これは・・・
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
177:デフォルトの名無しさん
09/11/26 19:45:48
対称行列になるものをGPUに送信したいのだが、うまい方法はない？
一般の場合は、GPU上では対称ではないとみたいのだが
送る際対称になる場合が結構あって、転送時間無駄だなあと。
178:,,・´∀｀・,,）っ-○○○
09/11/27 00:42:44
>>175
これじゃないの？
URLﾘﾝｸ(gpu.fixstars.com)
179:,,・´∀｀・,,）っ-○○○
09/11/27 00:47:50
このコードはいろいろ酷いからそのまま使えると思っちゃ駄目よ。
srcのアドレスはこの場合64バイトでアラインメントされてないといけない。
destのほうも最低限16バイトアライン
180:デフォルトの名無しさん
09/11/27 01:21:10
>>176
なかなか良い説明だなｗ
181:デフォルトの名無しさん
09/11/27 01:24:20
>>176
あと、競争するコースに分岐があったら、
その分岐に用のある人が２，３人だったとしても全員一応付き合わされた後
本流に戻るというか、そんな感じだな。
182:デフォルトの名無しさん
09/11/27 14:02:37
>>180,181
あれ、みんな読んでないのかな。
最近ようやくCUDA本がでたわけだけど、
まんまこの文章書いてあるんだよね。
青木先生乙
183:デフォルトの名無しさん
09/11/27 14:41:15
>>182
絵もついてるしなｗ
184:デフォルトの名無しさん
09/11/27 19:44:26
地球シミュレータを蹴散らし一躍脚光を浴びたというのに
中身スカスカなスレだなｗ
185:デフォルトの名無しさん
09/11/28 02:15:17
CUDAスパコンってそのときだけのものなんじゃない？
研究機関のスパコンは定期的に更新されるもので、更新したら前の代の
ソフトウェア資産はさっぱり使えなくなりましたじゃとても困るだろう。
186:,,・´∀｀・,,）っ-○○○
09/11/28 03:50:50
まあG80～GT200世代のコードも一応次のFermiでは動くし
NVIDIAが父さんしない限りはずっと続くんじゃない？
187:デフォルトの名無しさん
09/11/28 12:16:54
>>182
もうある程度把握しちゃってるからいまさら入門書買ってもとも
思ったけど、どうなんだろ。
188:デフォルトの名無しさん
09/11/28 12:19:29
>>185
IBMがHPC向けにはCELLやめてOpenCLに舵を切った的なことが書いてある。
ま、TheInquirerの記事は眉唾ではあるけど。
URLﾘﾝｸ(www.theinquirer.net)
あ、ここCUDAスレだけどまあ似たようなもんということで。
189:,,・´∀｀・,,）っ-○○○
09/11/28 14:03:35
×CellやめてOpenCL
○Cellベッタリのコード書くのやめてOpenCL
段階的には切り捨てることも考えられるが
いきなりOpenCLで他のデバイスとかいっても、資産が無いじゃん
190:デフォルトの名無しさん
09/11/28 20:47:07
はじめてのCUDAプログラミング
買った人いる？
どうだった？
191:デフォルトの名無しさん
09/11/28 22:01:41
図書館に頼んだら陳列は来月からと言われてしまった
192:デフォルトの名無しさん
09/11/28 23:23:30
本なんか読まなくても分かるだろ
193:デフォルトの名無しさん
09/11/29 00:55:57
日販は使えない会社だ。
194:デフォルトの名無しさん
09/11/29 01:10:09
>>190
思ったより安かったのでぽちってきた。
ASCII.techの特集も買ったけどアクセスの最適化あたりで苦戦中なレベルなので、
どうだった？とか言われても答えられないかもしれない。
195:デフォルトの名無しさん
09/11/29 12:11:38
>>184
これ、とんでもない誤報
浮動小数点演算を理解していないバカコミの馬鹿記事
科学技術立国 日本の恥を世界に晒したもの
「ふざけたやつがペンもつな、馬鹿野郎」だ。
196:デフォルトの名無しさん
09/11/29 14:01:50
浮動小数点数
197:デフォルトの名無しさん
09/11/29 15:13:49
GTX295と電源買った～
さて、何に使おうw
198:デフォルトの名無しさん
09/11/29 15:40:35
っBOINC
199:デフォルトの名無しさん
09/11/29 15:46:59
>>176
おにいさん、もっと！
200:197
09/11/29 18:00:24
smokeParticles.exe の動作が 8600GTS より遅く感じるんだがなぜだー
201:デフォルトの名無しさん
09/11/29 18:14:08
>>200
２個同時使用に対応してないんだろ
202:197
09/11/29 18:19:17
>>201
それでも8600GTSよりは早くなるはずじゃない？
あまり詳しくないんで間違ってたらすまん。
203:デフォルトの名無しさん
09/11/29 18:26:30
>>202
言われて見ればそうだな、約２倍に性能アップしてるはずだし
いつからか分からんけど古いＣＵＤＡと最近のＣＵＤＡのサンプルプロジェクトが入れ替わってるからな
パーティクル関係の数字が増えてるけど同じプログラムでやってみた？
204:197
09/11/29 18:37:06
>>203
同じプログラムでやってる。
ベンチマークでもやってみるかな。
205:デフォルトの名無しさん
09/11/29 19:47:47
文字列処理をさせてはみたものの
遅すぎて使い物にならねーぞこんちくしょー
206:デフォルトの名無しさん
09/11/29 19:49:56
>>205
単精度な数値計算に変換すれば良いんでないかい
207:デフォルトの名無しさん
09/11/29 20:06:22
>>206
UTF-8をどうやって数値計算にすればいいぉか？
208:デフォルトの名無しさん
09/11/29 20:14:33
UTFの全領域を使う分けじゃなければ、必要な部分だけを数値にマップするとか。
209:デフォルトの名無しさん
09/11/29 20:57:44
文字列処理って言ってもいろいろあるだろ
210:197
09/11/29 22:02:50
ベンチマークやったら電源が落ちた・・・
780Wじゃ足りないのかな
211:デフォルトの名無しさん
09/11/29 22:28:55
>>210
よっぽどの詐欺電源でも買ってない限りは、さすがに足りないってことは無いと思うけど
初期不良じゃないか？
212:デフォルトの名無しさん
09/11/29 22:35:29
熱落ちじゃないの？
213:デフォルトの名無しさん
09/11/29 22:39:18
Geforceは80 PLUS シルバー以上の
電源じゃないとまともに高負荷に耐えられないぞ
214:197
09/11/30 06:48:15
80 Plusって書いてないわ・・・
CORAZON ってやつです。
215:デフォルトの名無しさん
09/11/30 08:07:19
>>214
URLﾘﾝｸ(www.keian.co.jp)
12Vが25AだなHDDとかいっぱい着いてたら微妙だな
最小構成なら問題ない範囲だと思うけど
てかこれ安いな、これ買えばよかったｗ
216:197
09/11/30 17:40:28
25Aが2つ書いてあるのはどういう意味？
217:デフォルトの名無しさん
09/11/30 20:42:55
定員25人のエレベーター2基と、定員50人のエレベーター1基は違うというのは分かるな？
218:デフォルトの名無しさん
09/11/30 20:46:53
ＨＤＤとかPCIE用の電源が２本電源の中にあるってことだ
ちゃんと分割して接続すれば12x25=300WあるからGTX295のTDP300Wは支えられる
HDDとかつけるともう足りない
ちなみにうちの700W表記の電源は一本36Aある
219:197
09/11/30 21:20:06
IDE電源にハードディスクは接続しているけれども、
PCIEの電源は2本ともGTX295に挿してるぜ。
そういう意味じゃない？
ごめん詳しくないんだ。
大人しくこのあたり買っておくかな。
URLﾘﾝｸ(www.scythe.co.jp)
220:デフォルトの名無しさん
09/11/30 21:25:34
>>219
検索すりゃわかるが今の電源でGTX295を動かしてる人が居るから動くって
ＨＤＤ１個にしてみて周辺機器もはずしてＯＣしてるならデフォルトにして
配線を入れ替えたりしてみてダメなら初期不良だろう
221:デフォルトの名無しさん
09/11/30 21:30:13
つまり、安いもんじゃないし保障が切れたらもともこもないから
電源が原因だとしても保障があるうちに買った店に持っていって確認してもらったほうがいい
お店の人が電源がだめだと言うなら電源を買えばいいし
お店にあるちゃんとした電源でもダメだったら初期不良で交換してもらえるからね
222:197
09/11/30 21:50:02
そうしてみる。
電源も一緒に買ったものだし、持って行ってみるわ。
めっちゃ勉強になった。ありがとう。
223:197
09/12/01 06:40:06
追記
HDを1つにしたら、落ちるまでの時間が長くなりました
224:デフォルトの名無しさん
09/12/01 07:27:11
排熱もやばいんだろw
きちんとしたケースとFAN買えw
225:デフォルトの名無しさん
09/12/01 17:12:58
多次元配列の領域確保、コピーってcudaMallocとcudaMemcpyでできる？
226:デフォルトの名無しさん
09/12/01 18:48:14
できる。
てか普通のCみたいにコピーできる。
ちなみにCudaだと多次元配列だと面倒だから1次元配列として扱うことがおおい。
cudaMemcpyAsyncってのもある
227:デフォルトの名無しさん
09/12/01 20:22:02
面倒っていうか１次元しか扱えないし
228:デフォルトの名無しさん
09/12/01 21:14:19
1次元だけだったか
自分の技量が足りなくてできないのかと思ってた。
229:デフォルトの名無しさん
09/12/02 19:12:31
今日から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
230:デフォルトの名無しさん
09/12/02 19:56:18
>>229
CUDA_Func<<<1, 512, 0>>>(dest, src, nPos);を2ｋ回も呼んでるのがまず悪いんじゃね？
あとは詳しい人に任せた。
俺も勉強中。共有メモリのバンクコンフリクトがわけわからねえ。
231:デフォルトの名無しさん
09/12/02 21:07:48
>>229
そりゃカーネルをキックするコストはかなりでかいから、ループで何回も呼んだらCPUにまけるだろ。。。。
>>230
shared memoryのbank conflictは
ものすごーーーーく大雑把にいうと、thread_id順でshared memoryのアドレスにアクセスすると各バンクのチャネルがぶつからなくて、パラレルに出来るよってお話
232:229
09/12/02 21:33:41
レスありがとうございます。
リニアに1次元配列を処理するような事は意味がないと言う事でしょうか？
例えば、ある程度の長さの、サンプリング単位のPCMの演算や、ピクセル単位の画像の演算とか…
233:229
09/12/02 21:40:38
連投スマソ
なんかレスを書いてて、やっとピンと来たんですが、
例えば各スレッドでさらにループで回して、CUDA_Funcを減らせば良い的な話だったりします？
>>229 のコードで言うと
CUDA_Func内で1KB分ループさせて、各スレッドへは1KBのオフセットを渡す。
その分、CUDA_Funcの起動回数を減らす。
違う…？
234:初心者
09/12/02 21:55:19
>>299
通常CPUなら、forで何回もやるような処理を
CUDAのカーネルを一発たたくことによって処理させるっていうのが基本的な考え方じゃないの？
あと、メモリは一度になるべく大きくとって転送したほうが効率がいいらしいよん
235:デフォルトの名無しさん
09/12/02 22:17:38
>>233
そのコードを見た感じ、1024*1024*512スレッドつかってることになってるけど、
何をしてるの？
236:デフォルトの名無しさん
09/12/03 06:31:18
いや>>229のやり方は正しいよ
画面描画とＣＵＤＡは同期処理だから大きい単位でやると画面がタイムアウト起こす
これ以上の最適化はＣＵＤＡでは不可能
これで遅いというならそれがそのカードの性能限界だと考えるしかないな
237:デフォルトの名無しさん
09/12/03 06:39:52
ちなみに>>229のCPUとカードの具体的な名前と周波数と
PCIEの速度とx16 gen2とかね
遅いって実際にどれくらい遅かったのは知りたいね
上位のクアッドCPUと8400GSなんかじゃ勝負にならないのは当たり前だから
238:デフォルトの名無しさん
09/12/03 06:59:13
>>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に書き出し
するともっと速くなる
239:デフォルトの名無しさん
09/12/03 07:11:45
>>238
早けりゃいいってもんじゃないぞ
そんなもん低クラスのカードで動かしたら一発で画面真っ暗だわ
240:デフォルトの名無しさん
09/12/03 07:16:32
あとsharedメモリはそんな使い方するもんじゃないだろう
毎回コピーしてたらそのコストの方がデカイわな
241:デフォルトの名無しさん
09/12/03 07:22:16
sharedメモリってあれだ
１スレッドでやる計算が複雑な時に頻繁に変数の値を更新するだろ
そういう時にグローバルメモリよりもアクセスが早い一時領域として利用するもんだ
こういう計算自体が単純なケースでは効果はない
242:デフォルトの名無しさん
09/12/03 07:30:33
>>239
別人だけど8400GSくらいだとそうなの？経験上何msを超えるとハングする？
ググルとOSにより2秒や5秒でタイムアウトとあるがギリギリまでやるのはまずそうな気はする。
243:デフォルトの名無しさん
09/12/03 07:38:09
あ、失礼 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;
}
}
実際書いて動かさないと良く分からないすな。グレーのビットマップ二つ用意してやってみる形かな。
244:デフォルトの名無しさん
09/12/03 08:17:42
>>242
タイムアウトが何秒とか議論することですらない
マウスすら動かない状態が２、３秒も続くようなアプリはアプリ失格だろ
245:デフォルトの名無しさん
09/12/03 09:05:21
CPUより遅いくらいの8400GSで動かそうとしたなら2,3秒のフリーズ程度なら止む無し。
クラッシュしてデータを失わせるかもしれないリスクを犯すよりは
起動時に簡単なベンチ走らせて遅いGPUはハネちゃうのもありかな。
8400GSを考慮したせいでミドル以上のグラボの足を引っ張るとか馬鹿すぎる。
246:デフォルトの名無しさん
09/12/03 09:14:35
プログレスバー表示したら遅くなるから表示しないで画面を固まらせるなんて
そんなものは個人で使うだけにするんだなｗ
247:デフォルトの名無しさん
09/12/03 10:20:49
ちょっと待て、みんな一台のGPUであれこれやろうとしているのか?
それじゃ出るスピードも出ないぞ。
248:229
09/12/03 10:22:20
皆様おはようございます。そして、レス感謝です。
朝イチで打ち合わせがあるので、結果だけ取り急ぎ報告します。
前のコード
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が動きませんでした。ちょっとヤバイかも…
詳細は追ってフォローさせて下さい。
249:デフォルトの名無しさん
09/12/03 10:25:19
cudaMemcpyは同期を取ってから転送するから、結果の利用のタイミングぎりぎりまで実行を遅らせられれば
見掛け上の処理時間を短縮できるよ。
250:デフォルトの名無しさん
09/12/03 11:27:12
>>248
RDP経由でCUDAが動かないのは仕様です。
251:デフォルトの名無しさん
09/12/03 12:36:54
俺もやってみたらこうなりましたが。
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();
}
}
252:デフォルトの名無しさん
09/12/03 12:40:39
//ホストでこうして、
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でやってみるといいかな？
253:デフォルトの名無しさん
09/12/03 12:46:40
4096x4096でやってみたらだいぶ差が出てきましたよ。
GPU threads : 2048
Processing time GPU: 33.380219 (ms)
Processing time CPU: 260.214355 (ms)
254:デフォルトの名無しさん
09/12/03 13:12:34
accessAddressの計算おかしくありません？
255:デフォルトの名無しさん
09/12/03 13:34:08
コードちゃんと見て無いけど、floatならGPUでパラで動かした方が速いに決まってる
athlon系の浮動小数点演算のコストは、整数演算のざっと平均70倍
intel系は知らないにゃぁ
256:デフォルトの名無しさん
09/12/03 13:49:01
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側コードが酷いのが多すぎる。
257:デフォルトの名無しさん
09/12/03 13:57:09
ぶたぎりですが、
geforce gtx285 と tesla c1060 って
どのくらい違います？
c2050がどうせでるから、c1060と285が
対してかわらないなら、285がいいかと思うのですけど。
ところで
URLﾘﾝｸ(www.nvidia.co.jp)
いつのまにかnvidiaのページに
「Tesla GPU コンピューティングプロセッサは発売されています:
- Tesla C2050/C2070
- Tesla C1060」
こんなことが書いてあるけど、まだ出てないですよね。
258:デフォルトの名無しさん
09/12/03 14:25:16
>>257
そんなこと言うと、NVIDIAの営業に「そんなアキバ的発想はダメですよ」って言われちゃいますよw
259:デフォルトの名無しさん
09/12/03 14:27:55
CPU側は、やっぱりCore i7で8スレッド並列とか動かしてあげないとだめじゃね？
そんでGTX285と勝負するみたいな。
260:デフォルトの名無しさん
09/12/03 14:29:52
率直に言って、なんで値段が5倍のTeslaが売れるのか良く分からないのです＜アキバ的発想ですが
Fermiアーキテクチャも、GeForceが先に出るんじゃないですか？ 来年2月くらいでしょうかね。
261:デフォルトの名無しさん
09/12/03 14:33:23
>>260
そりゃぁ、GeForceが叩き台だモノ。NVIDIA曰く、GeForceよりも信頼性が10倍高いから5倍の価格差はペイできるって発想でしょ。
262:デフォルトの名無しさん
09/12/03 14:37:11
>>259
HPC用によく最適化されたコードだとHT切ったほうが速い
>>260
サポートが手厚い
263:デフォルトの名無しさん
09/12/03 14:49:44
Engadgetに、米軍が旧型PS3を2000台買ったとか書いてあったよな
$200ドルで256MB、100GFlops、て安いのかな??
10倍して$2000ドル、2.5GB、1TFLOPS、て。どうなんだろ。GPUのほうが安そうな。びみょう？
264:デフォルトの名無しさん
09/12/03 16:25:21
GTX285とか、240SPに、512とか1024とか2^nで放り込んでもうまく処理してくれるのですか？
半端な部分の処理がどうなるのか心配していると脳みそかゆくなります。
265:デフォルトの名無しさん
09/12/03 16:37:48
>>264
cudaライブラリが「適当に」割り振ります。グラフィック兼用だと、そもそも全部使う保証さえありません。
266:デフォルトの名無しさん
09/12/03 19:13:14
信頼性とか保証とかサポートとかの違いか
267:デフォルトの名無しさん
09/12/03 19:36:45
いや、メモリ搭載量全然違うし。
268:デフォルトの名無しさん
09/12/03 20:35:09
そうだね、GeForceの方が安いメモリを使っているしね。
269:デフォルトの名無しさん
09/12/03 22:47:05
はじめてのCUDAプログラミング買ったが
テクスチャには触れてないも同然だな。
270:デフォルトの名無しさん
09/12/03 22:52:47
>>263
ゲーム機としても使うつもりだから
271:デフォルトの名無しさん
09/12/04 01:08:38
>>251
実コード動かしたのは乙だけど、なんでデータがfloatになったの？
整数演算だったら、結局ちゃんとコードが書かれていればボトルネックは転送速度であって、
DDR2-CPUとDDR2-(PCI-Ex16)-GPUの比較になりそうな気がする。
272:デフォルトの名無しさん
09/12/04 01:42:05
CUDAの処理速度を計りたいならまずプロファイラにかけろ。話はそれからだ
273:デフォルトの名無しさん
09/12/04 01:47:26
>>263
コードの資産の問題なんでしょ。
Cell用に書いていたので、GPUにポーティングする手間を考えたら、
PS3を大量に買った方が安いのでしょう。
もちろんアプリによるけど、PS3は実効性能で100Gflopsはでるけど、
GPUは実効性能で1TFLOPSも出ないでしょ。
自分の経験では、Cｅｌｌに対して、G200は大体2倍～3弱倍という感じだった。
Cellはカリカリにチューンしたけどね。
274:デフォルトの名無しさん
09/12/04 02:53:46
アキバ的発想ではやっぱり、GTX295を二枚というのが現時点ではさいきょうCUDA環境ですか
275:デフォルトの名無しさん
09/12/04 05:25:55
>>273
GPU内部だけで完結させる処理なら１Tいくかもしれないが
PCの場合はデータ転送がボトルネックになって１Gｆlopsとかになるが
276:デフォルトの名無しさん
09/12/04 06:04:38
>>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...
ほう
277:デフォルトの名無しさん
09/12/04 08:54:02
結局さぁ、「ゲーム」だったらCPUで全部演算しようなんて考えるやついないんだし、
別にGPUでやるほどのもんじゃねぇ、ってだけなんじゃないの＞229の例
278:デフォルトの名無しさん
09/12/04 11:40:15
>>274
アキバ的にはGTX295を４枚が最強だけど、
そこまでするならtesla買った方が幸せだわな。
279:デフォルトの名無しさん
09/12/04 12:09:47
Larrabeeがなんかよさそう
280:デフォルトの名無しさん
09/12/04 12:33:49
>>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遅い…のはこれはこれでよいのかな
281:デフォルトの名無しさん
09/12/04 14:40:50
>>278
電源がそれだけで1200Wくらい要るなぁ。
1000W電源二台って感じか。うちのブレーカーは間違いなく飛ぶだろう
282:デフォルトの名無しさん
09/12/04 14:50:01
電子レンジとかドライヤーつけなければいけるだろう
たぶん
283:デフォルトの名無しさん
09/12/04 15:05:13
夏は地獄を見るだろう
284:デフォルトの名無しさん
09/12/04 15:55:44
CUDAで倍精度演算をしたいと思っているのですが、
・GeForce GTX 285
・GeForce GTX 295
なら、どっちがいいですか？
また、もっといいものがあれば教えてください。
Teslaは高いらしいのでそれ以外で、お願いします。
285:デフォルトの名無しさん
09/12/04 16:02:55
>>284
Fermiが出るまで待つ
でなければGTX295でしょうねぇ。(メモリ・帯域幅は小さくても倍精度演算リソースが二倍)
286:デフォルトの名無しさん
09/12/04 17:01:27
>>285
ありがとうございます。
Fermiがいつ出るかわからない状況なので、GTX295の方向で考えます。
いつ出るか発表されてませんよね？
287:デフォルトの名無しさん
09/12/04 19:46:16
cudaMallocHostってmallocのラッパーではない？
ライブラリつかってて、なぜかcudaMallocHostを指定していて
mallocで大丈夫だろうと思ったら動作せず
cudaMallocHostなら動作したので。
どう違うの？
288:デフォルトの名無しさん
09/12/04 19:58:37
>>287
つ page lock
289:デフォルトの名無しさん
09/12/04 20:54:29
>>285
１個づつ付いてるのを２倍って呼ぶのは詐欺くさいですな
290:デフォルトの名無しさん
09/12/04 21:02:18
>>287
メモリコピーのハードウェアアクセラレートって決まった区切りで始まる決まったサイズのメモリ単位でしか扱えなくて
そういうのを考慮してメモリを確保するんでしょう
291:,,・´∀｀・,,）っ-○○○
09/12/04 23:12:59
>>280
byteこそSSE使わないと。
最大16倍程度は速くなる
292:,,・´∀｀・,,）っ-○○○
09/12/04 23:25:37
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);
}
}
293:,,・´∀｀・,,）っ-○○○
09/12/04 23:54:33
>>279
最終的にはメモリ帯域ネックなんで同等のメモリ帯域を持つGTXと変わらんと思うよ。
ただ2つの配列の平均取るだけならほぼ帯域ベンチ。
PCIeかGPUのL/S帯域で律速。
キャッシュメモリが256KB/coreほどあって16コアあれば、1024x1024ならキャッシュにほぼ収まるので
それを有効に使って処理を連結していくなら、そこではじめてLarrabeeの旨味が出てくる。
個人的にはN-Bodyとかが良い勝負しそう。
294:デフォルトの名無しさん
09/12/05 01:12:39
>>293
だんごさんてきにはGrape-DRが
お勧めということですね
295:デフォルトの名無しさん
09/12/05 01:17:38
多体問題だけやりたいんならいいんじゃね？
296:デフォルトの名無しさん
09/12/05 02:09:51
おいおまいら、こんなのがあるぞ。
URLﾘﾝｸ(www.acceleware.com)
リッチな会社にいたら旅行がてらいってみるといいぞ。
これ見るとＦｉｘｓｔａｒｓがいかに良心的かわかるぞ。
297:,,・´∀｀・,,）っ-○○○
09/12/05 12:49:04
>>295
Larrabeeは机上計算だと1コアあたり16並列で15サイクル（＋α）くらいで回るんで、かなり理想的なデバイスなんだけどね。
（ただしvrcppsの精度補完ありならもう少しかかる）
298:デフォルトの名無しさん
09/12/05 13:54:20
GTX260 で決まりすよ。18000円なんだから4枚位買えばいい。
299:デフォルトの名無しさん
09/12/05 14:08:17
ららびーはもう終わったみたいな話が出てますけど
URLﾘﾝｸ(venturebeat.com)
300:,,・´∀｀・,,）っ-○○○
09/12/05 15:09:19
読んで字のごとくコンシューマ向けディスクリートはね。
何のためにスパコン向けカンファレンスで発表したと思う？
まあなんにせよ「パソコン」に刺さるカードは無くなったな。
PTX2.0の資料請求しとかなきゃ。
301:284
09/12/05 18:47:17
>>289
っていうことは、285の方がよさそうでしょうか？
メモリの帯域が大きく、こちらの方が使いやすそうな気がします。
デュアルGPUのコーディングも大変そうなので……
302:,,・´∀｀・,,）っ-○○○
09/12/05 18:58:27
>>301
待てるならFermiを待った方がいいよ
303:デフォルトの名無しさん
09/12/05 19:32:04
倍精度性能はGT200世代はおまけで１個演算機がついてるだけだから期待するだけ無駄だよ
Fermi世代から各コアに演算機付けるって言ってるけど、
一般向けには倍精度削ってGF100の名前で出すとか言ってるからどうなるか分からん
CUDAに拘らないならRADEON HD5xxxって選択もある
GT200世代より倍精度演算能力は圧倒的にHD5xxxの方が高いし
ただRADEONはドライバが糞だしCUDAも動かないしいろいろ中途半端
Larrabeeがどうなるかってところか
現状実用的なものは無いから、実験的な目的以外では買わない方がいいし
実験的な目的なら安いやつでいいじゃないかという話になる
304:,,・´∀｀・,,）っ-○○○
09/12/05 20:29:53
Larrabeeはコンシューマ向けディスクリートGPUとしてはキャンセル。
HPC向けには続投しそうだから100万くらい出せば手に入るようになるかもよ？
305:デフォルトの名無しさん
09/12/05 23:36:41
RADEONはFFTを出せないところを見ると行列積が精一杯のようだよ。
GTX280は512bitのバンド幅がどうもよろしくないのでGTX260を奨めます。
306:デフォルトの名無しさん
09/12/05 23:41:53
もしRADEONで遊ぶなら、現時点ではDXCSでSM5.0のスレッド制御を
使ってどこまでできるかだろうなぁ。
DirectX 11 SDKにFFTのサンプルコードなんかもあるから、持ってる人は色々
ベンチマーク取ってみて欲しいな。
307:デフォルトの名無しさん
09/12/05 23:43:57
>>304
Tesla by Fermiの値段にぴったり張り付いて売るのがIntel式な気が
しないでもない。そのためにもnVidiaには頑張ってもらわないと！
308:デフォルトの名無しさん
09/12/05 23:44:40
295だろ、普通に考えて。
スレッド数を多く使える分だけ、高速化が容易
メモリのバンド幅とかよりも重要だと思うが？
309:デフォルトの名無しさん
09/12/05 23:51:04
迷っている段階なら、とりあえずGTX260を買って試すのがオススメかな
310:,,・´∀｀・,,）っ-○○○
09/12/06 00:24:51
>>307
流石にIntelシンパの俺でも30万は出せない。
Fermi出た時点で一番コストパフォーマンスいいの選ぶわ。
311:デフォルトの名無しさん
09/12/06 01:23:34
試行錯誤するのにcompute capability 1.3に対応した
ファンレスカードとかあれば良いんだけどな。
260や295を付けっぱなしってのは、なんか精神衛生に良くない。
312:デフォルトの名無しさん
09/12/06 01:25:46
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で発生していませんでした。
意図的にバンクコンフリクトを発生させると更に１０％程遅くなるのでバンクコンフリクトは原因ではなさそうです。
共有メモリは速いはずなのにグローバルメモリよりなぜ遅くなるのか悩んでいます。初歩的なミスだろうとは思うのですが。
313:デフォルトの名無しさん
09/12/06 07:23:36
>>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上)に行ってないか心配だ。
314:デフォルトの名無しさん
09/12/06 07:53:46
vram2[i] がレジスタのってたりしないかな。
親から vram3 として渡してみるとか、参照を i+1 にするとか。
315:312
09/12/06 10:18:05
>>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__ の構文がなぜあるのか不思議に思っていたのですがこういう理由だったのかと。
しかし大きな共有メモリで遅くなるのは厄介な仕様ですね。
316:313
09/12/06 16:21:39
>>315
おつ。
勘が鈍ってなくて安心したｗ
共有メモリは多数のコアで共有される有限リソースだから、そこも一定以下に抑えないといけないって話ね。
あと本当にsharedを使う必要があるのか、一時変数（レジスタ）で済ませられないか？
sharedを使うべきなのは、同時に動くコア同士のやりとりがある場合と、
どう処理を分配してもcoarescedアクセスにできなくて、でもsharedを使えばできるという場合ぐらいかと思う。
あ、そういや、逆にレジスタが多くてOccupancyが上がらない場合の退避用とかにも使えるだろうか？
317:312
09/12/06 17:53:52
訂正 *((uint1 *)&px) = shared[i] を px.x = shared[i].x;の部分は勘違いでした。
１バイトしか書き換えてないので等価ではありません。無視してくださいです…
>>315
実はどういう理屈で遅くなるのかよくわかってないのですがお蔭様で対処の方法がわかってひと安心です。
最終的にはいろいろな画像処理に使う予定なのでキャッシュ的な使い方をsharedでさせてみるテストでした。
比較のためテクスチャ版とshared版作ろうとしてはまってしまいました。
318:デフォルトの名無しさん
09/12/07 12:02:53
自宅でGPU4枚とかって人はなにに使うの？
319:デフォルトの名無しさん
09/12/07 12:36:53
エンコ
320:デフォルトの名無しさん
09/12/07 13:11:16
ベンチ見てほくそえむ
321:デフォルトの名無しさん
09/12/07 16:15:39
>>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出力を読めるかどうかは割りと重要かも知れず。
# つーか、メモリアクセスの個数を数えるくらいのことは普通にやってる。
322:デフォルトの名無しさん
09/12/08 01:19:36
すみません、かなり初心者です。
行ごとに要素数の違う配列をデバイス側に渡したいんですけど、
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;
}
というようにはできないんでしょうか？
323:デフォルトの名無しさん
09/12/08 06:39:08
>>322
>CUDA_SAFE_CALL(cudaMemcpy(mat_d,mat,size,cudaMemcpyHostToDevice));
これがいらないんじゃない？
ホスト側のポインタ列をデバイスに渡しても使いようがない。メモリ空間が違う。
324:デフォルトの名無しさん
09/12/08 09:12:09
開発環境の入っていないマシンで動かすには、どのファイルを持って行けばよいのでしょうか？
325:デフォルトの名無しさん
09/12/08 10:25:10
>>291
個人的な興味なんですけど、CUDA使ってるcodecってなんで速いんでしょうね？
BYTEアクセスじゃ無いとしても、どの部分をGPUでやらせれば爆速になるのか不思議ﾁｬﾝﾃﾞｽ
326:デフォルトの名無しさん
09/12/08 14:53:55
SDKのドキュメントの通りに，VS2005で
サンプルプロジェクト template を元に使っているのですが
このプロジェクトに新しい.cuファイルを追加した場合に
そのcuファイルへの更新がビルド時に察知されません
毎回リビルドすれば一応反映できるのですが
通常のビルドで反映させるにはどうすればよいのでしょうか？
327:322
09/12/08 15:52:18
>>323
すみません。
抜けてましたが、
>CUDA_SAFE_CALL(cudaMemcpy(mat_d,mat,size,cudaMemcpyHostToDevice));
は、
ホスト側であらかじめ計算を行い、
その結果を初期値としてポインタ列、matに与えた上で、
その値をデバイス側のポインタ列、mat_dにコピーし、
その値を使ってデバイス側で計算をする。
というつもりで書きました。
このようにホスト側で計算した値をデバイス側に渡す時は
どのように記述するのがよいのでしょうか？
328:デフォルトの名無しさん
09/12/08 16:15:43
{1個、2個、4個、8個、16個、…} みたいなデータを渡したいのかな??
固めてまとめて渡しちゃったほうがオーバーヘッドが少ないと思います。
トータル何列あるよ、は別にパラメータで渡す。
(実際にCUDAのルーチン書く前に、コピー/戻しの時間を色々計ってみるといいです)
cudaMallocしたデータにはホストからは触れないので、
ホストでmallocしたデータ(mat)に計算結果格納
→同じサイズでcudaMalloc(mat_d)
→cudaMemcpyHostToDeviceで渡す
なのでそれでいいです
329:323
09/12/08 17:35:58
>>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の言うように
固めたほうがベターではある。
330:323
09/12/08 17:48:04
>>329
ちなみに、ポインタサイズが32bitか64bitかはホスト側に依存してることを確認した。
URLﾘﾝｸ(forums.nvidia.com)
331:デフォルトの名無しさん
09/12/08 17:59:16
>>325
あいまいな質問だけど答えてみる。
CPUよりGPUのほうが計算能力の総量はでかいから、本来GPUが速いほうが当然なんだけど、
以下のようなものは遅くなる場合もある。
・計算負荷よりもメモリ転送がボトルネックになるもの
・処理を細かく分解して並列化することが難しいもの
codecなどは、メモリ転送がボトルネックになりやすい傾向はあるものの、
だいたい画面の領域ごとに分解できる処理が多いから、多少とも速くはできるでしょう。
あと速くならない処理はCPU側でやればいい。
332:デフォルトの名無しさん
09/12/08 20:37:12
>>328>>329
参考になりました。ありがとうございます！
まだ不慣れなので色々試行錯誤しながらやってみます！
333:デフォルトの名無しさん
09/12/08 21:48:05
>>325
CUDAだけ使って早いと思ってる？ｗ
334:,,・´∀｀・,,）っ-○○○
09/12/08 23:12:01
動き検出なんかは当然CPUにやらせたほうがいいぞ
SpursEngineなりCellなりがあるならそれでもいいが
335:デフォルトの名無しさん
09/12/09 00:11:13
>>333-334
流石にGPUだけとは思ってないけど、atomでh264なんかエンコすると日が暮れる勢いだけど、ionプラットフォームだと実用範囲になる。
エンコ中でも完全にCPUを使い切ってないところを見ると、やっぱGPUをかなり使ってるんだなと勝手な想像
336:デフォルトの名無しさん
09/12/09 06:34:40
IONはメインメモリをグローバルメモリとして使う、つまりマッピングしておけば
CPUからもGPUからもフルスピードでアクセスできるメモリ空間を作ることができる。
ION2は買おうかな・・・。
337:デフォルトの名無しさん
09/12/09 08:29:47
>>336
実際はそんなに甘くはなく、GPUからフルスピードでアクセスするためにはWriteCombinedモードに
する必要があって、それを付けるとCPUからのアクセスが死ぬほど遅くなります。
通常は素直に転送（実際にはメモリコピーですが）したほうがマシです。
低価格で4GB近いデバイスメモリが使えるというメリットはありますが。
338:デフォルトの名無しさん
09/12/09 09:47:42
そのためのMOVNTDQA/MOVTDQだろう。
まあ前者はSSE4.1でないと使えないがな。Atomは仕様不可。あしからず。
339:デフォルトの名無しさん
09/12/09 10:11:54
それらを使っても十分に遅いと思うのですが。
340: ◆0uxK91AxII
09/12/09 10:22:13
WCにしつつ、CPU->GPUはmovnt*でcopyすれば良さそうだね。
CUDAとか知らないけど:b
341:デフォルトの名無しさん
09/12/09 10:26:15
>>339
だからCPU専用のワーク領域と分けて読み書きを最小限にする。
342:デフォルトの名無しさん
09/12/09 19:25:43
>>337
そか、CPUのキャッシュが使えなかったね。
死ぬほどまで遅くなるとは思ってなかった。
343:デフォルトの名無しさん
09/12/10 10:12:49
GTX295 Vista x64 CUDA2.3 の環境で、CUDA VS WIZARDで作成したプロジェクトを使っています。
質問1. ブロックあたりのスレッド数を512にすると計算結果が正しくなりません。256なら正しいです。なぜでしょう？
shared memory はグローバル変数として宣言して、各スレッドで共用（リードオンリー）しています。スレッド内で__shared__の宣言はありません。
質問2. 完成したので、SDKのsimpleMultiGPUを参考にGPU 2個使おうとしています。
とりあえずマルチスレッド用の空の関数を書いたところ、コンパイルはできるのですが、リンクできません。
cutStartThreadなどが外部で定義されてないと言われます。
simpleMultiGPUはビルドできるので、プロジェクトのプロパティを見比べてみましたが特に違いは見あたりません。
何が悪いんでしょうか？
どなたかお助けください。
344:デフォルトの名無しさん
09/12/10 10:45:27
>>343
・sharedにはCPUから書けませんが、各スレッドがリードオンリーならどこで書き込んでいるのでしょう。
・cutで始まる名前の関数はcutilライブラリ内にあります。cutilライブラリもビルドし、あなたのプロジェクトからリンクできるようにしないといけません。
345:デフォルトの名無しさん
09/12/10 12:13:48
>>344
上については、言葉足らずでした。カーネルの最初に
if( threadIdx.x == 0 ) { /* デバイスメモリからコピー */ }
__syncthreads();
でコピーしています。
将来的にはコピーも並列化する予定です。
下についてはありがとうございます。既存のlib(dll?)だけでは駄目だってことですね。
346:デフォルトの名無しさん
09/12/10 12:40:10
threadIdx.xはスレッドの識別子で、実行順番とは無関係では？
347:デフォルトの名無しさん
09/12/10 13:15:43
共有メモリのつくりに関する知識が思いっきり欠如していると思われ。
threadIdxが0の人だけが書こうとしたら、他の人はみんな何もできなくて時間が無駄。
おまけに、スレッド数が一定数を超えると(GTX295の場合は256か)、別の実行単位になってしまうから(表に見えない)
共有メモリ間転送が行なわれてしまってその点でも時間の無駄。
尤も、>345の様にthreadIdx0でしかコピーしない場合は転送は発生しない代わりに結果がご覧の通りなわけで。
どうせthreadIdx0がコピーしている間他はなにもできないんだから、一斉に同じものを書いてしまえばいいんでない?
348:デフォルトの名無しさん
09/12/10 13:33:51
>>347
> threadIdxが0の人だけが書こうとしたら、他の人はみんな何もできなくて時間が無駄。
開発途中なので、とりあえずこうしているだけです。
将来的にはコアレスにします。
> スレッド数が一定数を超えると(GTX295の場合は256か)、別の実行単位になってしまうから(表に見えない)
512で上手くいかないのはこのあたりが原因のようですね。
349:デフォルトの名無しさん
09/12/10 13:49:31
>>347
たびたびすみません。
「別の実行単位」の意味がよく分からないのですが、
要は1つのブロックが2つのSMに割り当てられるって意味ですか？
350:デフォルトの名無しさん
09/12/10 14:03:14
>>345
Sharedは文字通り、共有できるメモリなので、
各スレッドが１ワードずつ協力して作成→どこでも読み書き可能、
ですよ。0番スレッドだけに任せなくてもいい
351:デフォルトの名無しさん
09/12/10 17:12:55
自分は>>343じゃないけど256スレッドでOKで512スレッドで計算結果がおかしくなる理由がわからない。
>おまけに、スレッド数が一定数を超えると(GTX295の場合は256か)、別の実行単位になってしまうから(表に見えない)
>共有メモリ間転送が行なわれてしまってその点でも時間の無駄。
というのは時間の無駄を指摘してるだけで結果がおかしくなる理由ではないという認識なのですがあってますか？
この時の共有メモリ間転送というのは別の実行単位に転送しているということ？
>>343
256スレッドでやればOKなのでもう興味はないかもしれないが、シェアードメモリの代わりにグローバルメモリで
共用して計算した場合は512スレッドの時も結果は正しくなるのですか？
352:デフォルトの名無しさん
09/12/10 19:27:33
>質問1. ブロックあたりのスレッド数を512にすると計算結果が正しくなりません。256なら正しいです。なぜでしょう？
カーネルコールの後でエラーキャッチのコード入れてる？エラー出てない？
353:デフォルトの名無しさん
09/12/10 21:05:48
質問の傾向をみると、
CUDAって、普通のCのように自由度高く書けるけど、
実際は、サンプルソースからは読み取れない
あれは駄目、こうしたら駄目、的な制約が多すぎるのかな？
DXAPI+HLSLのような、自由度の少ない環境の方が
むしろ、良質なソースが書けるのかも。
354:デフォルトの名無しさん
09/12/10 21:54:43
参考資料
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)