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)
355:デフォルトの名無しさん
09/12/10 22:17:49
>>353
CUDAをただの言語だと思っていてGPUのアーキテクチャを理解していないからかと。
後半については同意。限られた記述で実装可能なアプリしか書かないからというのがより正確かも。
356:343
09/12/11 00:35:49
>>351
スレッド数512、デバイスメモリ使用でやってみましたが、やはり正しくありませんでした。
Emuではちゃんと動くので、原因はよく分かりません。
コンスタントも使って多のでそっちが原因かも知れません。
分からん。もういい。
>>343 の質問2については、multithreading.cppをプロジェクトに追加してないだけでした。
俺死ね。
357:デフォルトの名無しさん
09/12/11 01:23:14
>353
テクスチャはまじで鬼門。
凝った使い方をしようとすると、コンパイル通っても動かない・動いても結果間違ってるがざらにある。
358:デフォルトの名無しさん
09/12/11 01:42:00
>>357
あるある。
deviceemuともまた動作違うこと多いしね。
NEXUS出ればデバッグもうちょっと楽になりそう。
359:デフォルトの名無しさん
09/12/11 09:51:22
>>343
歩留まり向上のため、重い処理をさせると計算が狂うことがあるらしい
どこかのPDFで見た
360:デフォルトの名無しさん
09/12/11 10:57:49
>>359
それちょっと表現が微妙に誤解を招くような・・・。
出荷されている状態で、計算がまともじゃないコアがたまにまぎれこんでいる。
グラフィック用途と違って、メーカーの想定より負荷が偏った計算をさせることがあるから、
やや耐性の低いコアの耐性がちょっと足りない部分がエラーを起こす。
高負荷プログラムなんかを使ってそういうコアというか製品をはじいてやれば、あとはほぼ安定稼働する、という話だと思う。
361:デフォルトの名無しさん
09/12/11 11:17:51
はじめてのCUDAがいつまでたっても届かない
一体どこに発注したんだよ上の人
362:デフォルトの名無しさん
09/12/11 12:42:26
>>360
CUDAのスパコンで1/6が不良品だったと言ってたが
363:デフォルトの名無しさん
09/12/11 12:44:30
えっ Teslaとかでもそうなのか？？
364:デフォルトの名無しさん
09/12/11 13:49:03
Teslaは選別品
365:デフォルトの名無しさん
09/12/11 14:52:34
低価格帯で一番安定してるGPUってなに？
366:デフォルトの名無しさん
09/12/11 15:53:13
URLﾘﾝｸ(slidesha.re) のP.26
長崎大の人に言って選別プログラムを貰うことだなｗ
367:デフォルトの名無しさん
09/12/11 16:12:17
CUDAは保障外だからCUDAやりたい奴は動くのが当たるまで買えってことなんだろう
368:デフォルトの名無しさん
09/12/11 16:17:50
仮に2/3超の確率で完動しないとしても、自前で問題無く選別できるなら
Tesla買うよりコンシューマGPU買うほうがまだ安い位だから、たくさん
買う人は自前で選別できるように頑張るのが正解だよなぁ。
369:デフォルトの名無しさん
09/12/11 16:40:06
多少クロック落とせば動くんじゃないのかな。CPUのオーバークロックと一緒で。
370:デフォルトの名無しさん
09/12/11 20:54:29
>>361
自分は11/27楽天でぽちって 12/01に到着。
ぼちったあとすぐ納期の表示が延びたから焦ったけど滑り込みセーフだったらしい。
書いてあることはわかり易い。しかし意外なことが触れられていなかったりする。
371:デフォルトの名無しさん
09/12/11 22:49:23
>>360
その状態が続く限り、コンシューマ用パッケージソフトじゃGPGPUを使えないね。
「動かない」サポートの爆発でたいへんなことになる。
372:デフォルトの名無しさん
09/12/11 23:34:51
てか>>356の作ってるソフトが選別ソフト代わりになるんじゃね？
373:デフォルトの名無しさん
09/12/12 13:58:47
おれらまだまだ、「主メモリ側が主役」って固定観念なくね？
いかんな。どかどかデバイスメモリにロード、大量のブロック×スレッドを駆使、
CPUはたまにお手伝いをさせていただく、位に思わないとな。俺まだまだ
374:デフォルトの名無しさん
09/12/12 15:52:58
メモリ転送と計算を非同期で出来るのかね？
375:デフォルトの名無しさん
09/12/12 15:58:25
>>374
マルチスレッドで、CPU側は「GPU様が計算し終わるまで寝とけや」くらいの扱いにするわけよ
376:デフォルトの名無しさん
09/12/12 16:43:09
>>374
できる。
基本GPUとCPUでは非同期。
ストリームを設定すれば、GPUで演算中にデータ転送させることも可能。
ちゃんと同期させるための関数もある。
377:デフォルトの名無しさん
09/12/13 15:51:31
> NEXUS
いつ出るんだっけ
VisualStudio2008でいいんだっけ
378:デフォルトの名無しさん
09/12/15 06:00:58
カーネルで計算した結果をCPU側に返すにはどうしたらいいの？
return aのように簡単にはできないのですか？
379:デフォルトの名無しさん
09/12/15 06:22:49
>>378
ちょｗ
カーネルfunctionの引数で float * outdata
とかやってカーネルからこれに書き込む outdata[i] = result;
ホスト側でcudaMemCopyのdeviceToHostで、ホスト側でcudamallocしたfloat * に戻す
みたいになる。ややこしいぞ？
380:デフォルトの名無しさん
09/12/15 19:31:04
thrust便利だなしかし
381:デフォルトの名無しさん
09/12/16 02:52:17
thrust、自分で書いたカーネル関数はどう使うのっすか
382:デフォルトの名無しさん
09/12/16 15:26:45
__shared__ int sh[num*2]; /* numはスレッド数 */
という共有メモリの配列をソートし、最大or最小のデータのみグローバルメモリに書き込みたいのですが
これをカーネルの中でやるいい方法を教えてください。
383:デフォルトの名無しさん
09/12/16 17:36:07
ソートは不要で、reductionで半分を繰り返せばいいのでは
384:デフォルトの名無しさん
09/12/16 23:18:35
>>383
言われてみればそうですね。
解決しました。thx
385:デフォルトの名無しさん
09/12/17 03:27:35
どこかにfortranでの使用例って
ありませんか
386:デフォルトの名無しさん
09/12/18 16:42:19
そういえば見かけたことねえな
387:デフォルトの名無しさん
09/12/18 17:17:27
CUDAはそもそも実行速度がシビアで
高級言語向けではないと思うけどな
388:デフォルトの名無しさん
09/12/20 02:21:29
共有メモリは16個のバンクに分割されて、半ワープのスレッドが同じバンクにアクセスしなければ競合は起こらない。
みたいな事の意味がイマイチ分かりません。
例えばブロックの共有メモリが8KBなら0.5KBづつに分けられて、半ワープのスレッドが0.5KBの中ならどこにアクセスしても良いって意味なの？
それとも、先頭から4ﾊﾞｲﾄとか8ﾊﾞｲﾄとかの固定の領域に連続して分けられてるの？
でもそれだと分けられる領域が4ﾊﾞｲﾄなのか8ﾊﾞｲﾄなのか分からないじゃまいかヽ(`Д´)ﾉ
カーネルの引数にvoid*とデータサイズを渡して任意のデータ長で処理させたいのですが、
こういう事をしようとすると元からバンク競合は避けられないのですか？教えて、エロイ人(´・ω・`)
389:デフォルトの名無しさん
09/12/20 02:33:00
今あるMPIのプログラムをGPUに対応させようと考えているんだけど、
CUDAでMPIを使うとき、データの転送は一端ホスト側のメモリに転送してから
他のノードへ転送することになるの？
だとするとレイテンシがすごく大きそうだね。
それとも専用のライブラリなんかあるのかな？
GPU側のグローバルメモリがホスト側のメモリにマップされていれば、
GPUを意識しないで転送ができそうなんだが。
390:デフォルトの名無しさん
09/12/20 08:02:33
>>388
おれの解釈だと、共有メモリには16個のアドレス入力と16個のデータ入出力が有る
と思ってるんだけど 違うかも
391:デフォルトの名無しさん
09/12/20 09:55:19
>>388
連続する64バイトの領域（先頭アドレスが４の倍数）について、4バイト(32bit)のバンク１６個に分かれています。
例えばバンク０に属するアドレスは０、６４、１２８、１９２、２５６、・・・・・から始まる４バイトの領域。
任意のデータ長が４バイトより大きいのか小さいのか分かりませんが、テンプレート等を使う場合には
型に合わせて場合分けをする必要があるでしょう。
>>390
「１個のアドレス入力と１個のデータ入出力ができるバンク」が１６個あるというのが適切な表現かと。
392:デフォルトの名無しさん
09/12/20 10:00:29
>>389
device memory on node 0 -> host memory on node 0 -> host memory on node 1 -> device memory on node 1
となるのでレイテンシは大きくなります。今のところCUDAではデバイスメモリをメモリマップする手段はありません。
ただし十分大きなデータを転送する場合にはパイプラン化すれば問題なくなると思います。
むしろpinnedメモリとMPIライブラリの干渉が一番の問題・・・・。
393:デフォルトの名無しさん
09/12/20 10:47:33
>>392
レスサンクス
やはりレイテンシは大きいのだね。
目的のアプリは数10kbyteから数100kbyteの転送を頻繁に行うから、
せっかくGPUで速くできても、転送ボトルネックで詰まりそう。
転送するサイズも小さいので、page lockさせない方がよいのかも。
394:デフォルトの名無しさん
09/12/20 14:34:19
GPUの購入で悩んでいるのだが、
Tesla C1060, GTX295,GTX285のうち結局どれが
速いんですか？ちなみに流体に使います。
GTX295ってGPU2基積んでるけど並列プログラミング
しないと2基機能しないとか？
素人質問で申し訳ない。
Teslaの保証とサポートも魅力的だが。
395:デフォルトの名無しさん
09/12/20 14:44:32
>>394
メモリ量でC1060になっちゃう なんてことないかな？
295はちゃんとマルチスレッドでプログラムしないと二基使えないです。
396:デフォルトの名無しさん
09/12/20 15:45:47
>>390 >>391
なるほど…よく分かりました。ﾚｽ感謝。
処理させたい任意のデータ長は5ﾊﾞｲﾄだったり、11ﾊﾞｲﾄだったり半端な数もきます。
任意のデータ長に対してコピーやビット演算を行なうんですが、データ長が4ﾊﾞｲﾄより大きいと
もうバンク競合は避けられない感じなんですね。
プログラムも汎用にさせたかったけどここは我慢するしかないのか…(´・ω・`)
397:デフォルトの名無しさん
09/12/20 16:09:09
VCでプロジェクト作るのが面倒なんだけど、
なんかウィザード的なものはないのかな
398:,,・´∀｀・,,）っ-○○○
09/12/20 16:29:23
URLﾘﾝｸ(forums.nvidia.com)
HTML+JavaScriptだから好きに書き換えられるだろ
399:デフォルトの名無しさん
09/12/20 16:30:40
おいらも流体で使う予定だが
GTX260を二つ買ってCUDAのお勉強しつつFermi待ち
C1060とか中途半端に高いからねえ
400:デフォルトの名無しさん
09/12/20 16:35:47
Nvidia寄りのとある企業のエンジニア曰く、
スペックはまったく同じでGTX285のチップの選別品だといっていた。
クロックがGTX285のほうが高いし、GTX285のほうが若干早いかも。
でもこのスレでもあるように、なんか計算結果がずれる可能性があるし、
メモリの多いTeslaが使えるなら使いたいよね。
401:デフォルトの名無しさん
09/12/20 16:53:55
DodでさえIBM CellじゃなくてPS3を購入してるんだからGTX285で十分でしょう。
402:デフォルトの名無しさん
09/12/20 17:01:10
>>398
まだ完成してないんじゃん
403:デフォルトの名無しさん
09/12/20 18:24:06
>>396
ASCII本の絵がわかりやすいけど、
共有メモリーは[16][17]にすれば問題ないと思う。
共有メモリーは無駄飯食っている用心棒みたいなもんで、
冗長にしてでも使えるときに使わない手はない。
404:394
09/12/20 22:43:43
皆さん、レス感謝です。
なるほど、GTX285とC1060差はさほどないんですね。
メモリと耐久性とサポート,計算結果
に関してはTeslaが有利というわけですな。
とりあえずGTX295のマルチスレッドはちょっと遠慮しようかなと思います。
もうちょい1人で悩んでみます。
405:デフォルトの名無しさん
09/12/21 00:30:00
個人的には285の2GBメモリ版がお薦めなのです。
406:デフォルトの名無しさん
09/12/21 06:47:28
【SIGGRAPH Asia 2009レポート】
東工大、スクウェアエニックスがCUDA実装事例を紹介
URLﾘﾝｸ(pc.watch.impress.co.jp)
407:デフォルトの名無しさん
09/12/21 07:04:05
【SIGGRAPH Asia 2009レポート】
NVIDIAがGPUによるグラフィックスワークスタイルの変革をアピール
URLﾘﾝｸ(pc.watch.impress.co.jp)
408:デフォルトの名無しさん
09/12/21 07:14:00
>>405
どっか、285の4GB版とか安く出してくれないもんかね。
409:デフォルトの名無しさん
09/12/21 07:15:58
...|￣￣ | ＜ Fermiはまだかね？
／:::| ___| ∧∧ ∧∧
/::::＿|＿＿_|＿ （ ｡_｡）. （ ｡_｡）
||:::::::（ ・∀・） ／<▽> ／<▽>
||::／ <ヽ∞/>＼ |::::::;;;;::/ |::::::;;;;::/
||::| <ヽ/>.- | |:と),__｣ |:と),__｣
＿..||::| o o ...|_ξ|:::::::::| .|::::::::|
＼ ＼＿_(久)__／＿＼::::::| |:::::::|
.||.i＼ ､__ﾉﾌ ＼| |:::::::|
.||ヽ .i＼ _ __ ＿＿＿＿ __ _.＼ |::::::|
.|| ﾞヽ i ﾊ i ﾊ i ﾊ i ﾊ | し'_つ
.|| ﾞ|i～^~^～^~^～^~^～
410:デフォルトの名無しさん
09/12/21 11:01:34
めるせんぬついすた、GPU版の使い方＠Windowsがやっとわかったわー
411:デフォルトの名無しさん
09/12/21 15:33:23
|┃三 /::::::::ハ､＼､::::::::＼＼::::::::::::',
|┃ ｉ:::::::ｲ ｀> ー─--ミ::::::::::::|
|┃ {::::::::| ::＼:::／:::: ＼:::ﾘ-}
ｶﾞﾗｯ. |┃ ',::r､:| <●> <●> !> ｲ
|┃ ﾉ// |:､`{ ｀> .:: ､ __ﾉ
|┃三 |::∧ﾍ ／､__r)＼ |:::::|
|┃ |::::::`~', 〈 ,＿ｨｪｧ 〉 l::::::》 ＜フェニミストはまだかね 辛抱たまらん
|┃ |:::::::::::::'､ `=＝'´ ,,ｲ::ﾉﾉ从
|┃三 ノ从､:::::::::`ｉ､,, ... ..,,／ |::::://:从
412:デフォルトの名無しさん
09/12/21 16:19:40
おれの物欲も辛抱たまらん、Core i7＋GTX260Mノートが欲しくて。
でもFermiのモバイル版が乗ったノートが出るのを松
413:デフォルトの名無しさん
09/12/21 21:07:30
故あってCUDAを使うことになったのですが自宅にはRadeonを積んだものしかありません。
コンパイルオプションでCPUでのエミュレーションができると何かで読んだのですが
これを利用すればRadeonの環境でも一応の動作はするのでしょうか？
414:デフォルトの名無しさん
09/12/21 21:24:33
CPUでのエミュレーションなので
radeon関係なくCPU上で動くよ
415:デフォルトの名無しさん
09/12/21 21:44:22
即レスどうもです。学校には設備があるので問題ないのですが自宅でもやりたかったので
あともう一つ質問よろしいでしょうか。VineLinuxでの動作は可能でしょうか?
気になったのですがどうもサポートしてないようなので
416:デフォルトの名無しさん
09/12/22 05:22:26
>>412
悪魔：「単精度なら大して変わらないんだから、買っちゃえよ！Fermi出たらまた買えよ！」
417:デフォルトの名無しさん
09/12/22 06:31:17
>>415
Linuxでサポート期待しちゃいかんだろう
あとエミュレーションはテクスチャ周りなど微妙に？おかしい模様なんで、あんまり期待しないほうがいい。
5000円くらいで適当なカード買うのが一番無難かもよ。
418:デフォルトの名無しさん
09/12/23 09:13:22
なんか、GeForceの在庫がどんどん無くなっていってない？
まもなくFermi出るな！これは！
419:デフォルトの名無しさん
09/12/23 11:06:25
C2050/2070よりも先にFermi搭載のGeForceが出ると発表しているしね
NVIDIA、“Fermi”採用第1弾GPU「Tesla 20」シリーズ発表 - ITmedia +D PC USER
URLﾘﾝｸ(plusd.itmedia.co.jp)
まあ、実際に2010Q1に出るかは怪しいわけだが
C2050が2499ドルってのは自宅用で用意するには結構きつい値段なので、
Fermi搭載のGeForceがどの程度の値段なのかが今から気になっている
420:デフォルトの名無しさん
09/12/23 11:07:45
>>415
とりあえず玄人思考のGTX260でも買ってみよう
でかいからPCケースに入るか確認してからね
421:デフォルトの名無しさん
09/12/26 03:28:46
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個くらいになるんだけど…
こんなんだと何万、何千のスレッドとか無理じゃね？
みんなカーネル以外の関数は作らずにカーネルの中に処理を直書きしてるのですか？
422:デフォルトの名無しさん
09/12/26 04:01:38
>>421
レジスタとスタックとかバイトとレジスタ個数とかごっちゃになってないか色々と
423:デフォルトの名無しさん
09/12/26 11:11:36
4Gamer.net ― Fermi時代の製品展開が分かってきたNVIDIAのGPUロードマップ。DX11世代が出揃うのは2010年Q2以降に（Fermi（開発コードネーム））
URLﾘﾝｸ(www.4gamer.net)
Fermi搭載のGeForceは3月くらいかなあ…？
この記事自体が2ヶ月前だからなんともだけどね
424:デフォルトの名無しさん
09/12/26 17:54:34
あんどうさん更新
URLﾘﾝｸ(www.geocities.jp)
425:デフォルトの名無しさん
09/12/26 18:03:48
>>421
>Device 0: "GeForce 8600 GTS"
>Total number of registers available per block: 8192
直訳すると、「ブロックあたり使用可能なレジスター数：8192」
ブロックあたりのスレッド数は数千・数万もいらない。
128～256程度でだいたいパフォーマンスは出る。（それを複数コアで何セットも同時に動かす）
あとカーネル１個でなんでもかんでもするもんでもない。（直列的な）処理ごとに分割すればいい。
426:421
09/12/27 23:15:38
>>422 >>425
うん…見直すと自分とても恥ずかしい事言いましたね(´・ω・`)
当初の疑問は理解して解決しました。
で、また別の話のGPGPUに関して質問です。
今DirectX使ってゲームを作っていて処理の重い部分(当たり判定とか)をGPUにやらせようと思っているんだけど、
ゲームって元々グラフィック描画でGPU使うので、それに加えてCUDAやプログラマブルシェーダーでGPGPUしたら相当パフォーマンス落ちますか？
目的はCPU使用率を抑えたいんだけど、ゲームのような1フレーム毎にGPGPUをするのは実用で使えうるものなんでしょうか。
427:デフォルトの名無しさん
09/12/27 23:28:32
当たり前の話だけど、GPUをフルに使うようなレンダリングをしていたら遅くなる
フレーム枚にGPGPUの処理を入れるのはそんなに問題ない
ただし処理の重い部分ってのが上手く並列化できるアルゴリズムであることや、そこそこの量あることが求められる
分岐ばっかだったり、処理よりメインメモリ-VRAM間の転送のほうが時間かかると逆に遅くなる
ようは、本当にGPGPUを使ったほうがいい処理なのかを吟味する必要がある
428:デフォルトの名無しさん
09/12/27 23:54:43
質問です。
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の物を入れました。
改善策がありましたら、ぜひお願いします。
429:デフォルトの名無しさん
09/12/28 00:17:23
>>428
Expressは64bitでコンパイルできないかも？
URLﾘﾝｸ(www.microsoft.com)
430:デフォルトの名無しさん
09/12/28 00:25:04
>>429
みたいです。やはりproを買うしかないのでしょうか…
レスありがとうございました。
431:デフォルトの名無しさん
09/12/28 00:48:50
>>428 >>430
答えるのに十分な情報量を持った質問を久しぶりに見たような気がしたｗ
proが良いけどねぇ。
今は最適化コンパイラもついてるし、standardで良い気もする。
（standard使ったことはないので参考まで）
432:デフォルトの名無しさん
09/12/28 02:01:56
PlatformSDKにx64コンパイラあったような
433:デフォルトの名無しさん
09/12/28 20:43:04
落ち着いて、なぜこういうエラーがでるのか？
と、考えるしかないよ。
434:デフォルトの名無しさん
09/12/29 11:29:33
CUDAが使えるGPUが載ってるか否かの判断をプログラム上で行うには
どうするといいのでしょう？
435:デフォルトの名無しさん
09/12/29 11:57:15
OpenCL入門 - マルチコアCPU・GPUのための並列プログラミング
が出版されるそうな
アスキードットテクノロジーズの記事書いた人たちらしいが
GPGPUでもっともメジャーなのはCUDAだと思うけど
OpenCLがこの本の売り文句通りスタンダードになるんだろうか…？
よくわからないけど、無関心でいるわけにも行かないのでとりあえずポチってみる
感想あったらアマゾンのレビューに書く
>>434
プログラム上でその判断をやる必要というのがよくわからない
コマンドラインでやるのでは駄目なのか
436:デフォルトの名無しさん
09/12/29 11:57:22
>>434
それはやはり
cudaGetDeviceCountして、CUDAデバイス個数調べ
cudaGetDevicePropertyをDeviceCountまわして、
.major、.minorでバージョンチェック
.multiProcessorCountでコアの個数調べ
して使うしかないんじゃ。
437:434
09/12/29 12:21:50
>>436
どうも
>>435
CUDA依存部分を共有ライブラリの形でプログラム本体から切り出しておいて、
実行時に動的にリンク出来たらとか考えてます。
…CUDA版バイナリと、CUDA無し版バイナリを用意して、インストール時に
選ばせればいいのか。
438:デフォルトの名無しさん
09/12/29 13:05:23
>>435
スタンダードはどう考えてもDirectComputeだろ。
いまのWindowsのシェアとDirectXの普及率から考えてCUDAがスタンダードになるにはポテンシャルが違いすぎる。
439:438
09/12/29 13:06:36
CUDAはスタンダードになれる可能性はあるけど、OpenCLはたぶん廃れる。
OpenCLって所詮はCUDAのラッパでしょ？
ラッパライブラリがスタンダードになった事例ってあんまり知らないんだけどさぁ。
440:デフォルトの名無しさん
09/12/29 13:16:26
Appleはラッパ好きだよな。
ObjCもラッパではないにしろマクロ駆使してるだけで、骨はCに変わりはないでしょ？
Appleは声は大きいけど技術が無いから、OpenCLもあんまり期待してないよ。
所詮ラッパライブラリだから、コンパイルしたらどのGPUでも使えるってわけではないし。
OpenCLのDLLロードしたらどのGPUでも同じバイナリでおｋみたいになったら使いやすいけどね。
441:デフォルトの名無しさん
09/12/29 13:19:23
てか俺が作っちゃう？ｗｗｗ
442:デフォルトの名無しさん
09/12/29 13:27:41
>>437
436でCUDA使えるかどうか確認して、
プログラムとしてはCPU用ルーチンとGPU用ルーチンを両方持っておく
でいいんじゃね？(やはり、動作確認用にCPU/GPU双方で試すのは必要と思うし)
443:デフォルトの名無しさん
09/12/29 13:30:38
>>438
おれ、
URLﾘﾝｸ(openvidia.sourceforge.net)
見て「うざっ！！」と思って挫折した。 ダメな俺
444:デフォルトの名無しさん
09/12/31 03:50:21
pyCUDAってどうかな
DriverAPIを使ってモジュールを実行時に作成、ロードするみたいですね。
445:デフォルトの名無しさん
10/01/03 16:34:54
3次元のデータをデバイス側に送って計算したいんですけど、
>>328の
＞固めてまとめて渡しちゃったほうがオーバーヘッドが少ないと思います。
＞トータル何列あるよ、は別にパラメータで渡す。
というのはどこにどのように記述すればいいのでしょう？
446:デフォルトの名無しさん
10/01/03 16:57:35
>>445
もちろん、少ないならデバイス関数の引数として渡す。パラメータ複数個まとめた配列で渡しても良いし。
計算の対象、計算の結果、に加えて、計算のパラメータ、も同様にホスト側で確保→cudaMemCopyするわけ。
447:デフォルトの名無しさん
10/01/03 17:38:23
>>446
ありがとうございます。
物分かりが悪くて申し訳ないのですが、
cudamemcpyは確保した領域にデータをコピーする使い方しか思いつかないので、
＞計算の対象、計算の結果、に加えて、計算のパラメータ、も同様にホスト側で確保→cudaMemCopyするわけ。
を簡単に例で説明していただけないでしょうか？
448:デフォルトの名無しさん
10/01/04 00:23:45
えっ なぜ分からないんだ？ たとえばこんな感じだよ
__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();
449:デフォルトの名無しさん
10/01/04 11:06:25
>>448
ありがとうございます。
なんだか変な方向に考えすぎていました。
450:デフォルトの名無しさん
10/01/04 12:52:14
CUDAってGPUカーネルの中で関数読んで値をreturnさせることできたっけ？
あと__global__にしないとホストから関数呼べなくね？
451:448
10/01/04 13:31:34
calculateが__global__で、func_A()、func_B()が__device__ですたorz
452:デフォルトの名無しさん
10/01/04 14:00:48
Fermi搭載GeForceの発売がまた不透明になったとか
453:デフォルトの名無しさん
10/01/04 16:05:01
⇒を何らかの演算として 0⇒256 256⇒512 512⇒0 といった感じにしたいのですが
a = (a+256)%513;
とする以外思いつきません。けど剰余は遅いので違う計算にしたいんです。
いっそのことテーブル作ったりif文にしたほうが早いのでしょうか？
馬鹿すぎてわかりません。誰か助けて
454:デフォルトの名無しさん
10/01/04 16:10:45
>>453
Cudaだと四則演算はほとんど処理時間には影響しないし、その式をつかうといいんじゃない？
っていおうとおもったけど、a=512のとき0にならなくね？