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

【GPGPU】くだすれCUD ..

231:デフォルトの名無しさん
08/08/25 09:16:36
だからお前はマルチプラットフォームの仕事をしたことあるのかと

232:デフォルトの名無しさん
08/08/25 11:23:15
サターンとプレステで同じゲームを同時に作ったぞ
同時に作っただけだけど

233:デフォルトの名無しさん
08/08/25 12:41:48
>>231
Firefoxはマルチプラットフォームじゃないのかね。

234:デフォルトの名無しさん
08/08/25 12:43:22
マルチプラットフォームの要件はオープンソースではかなり多い。
オープンソース系の企業ならそういう仕事していてもおかしくないね。

235:デフォルトの名無しさん
08/08/25 12:48:10
日本語読めないの？
聞かれてるのは「お前は仕事でマルチプラットフォームをやったことがあるのか？」だぞ

236:デフォルトの名無しさん
08/08/25 12:51:51
>>235
話のすりかえ。
俺自身は学生時代にFirefoxの開発に携わったが、
今の話題は「マルチプラットフォームの仕事があるかないか」。
お前が間違っていることに気づけ。

237:デフォルトの名無しさん
08/08/25 12:53:16
そもそも、ここの住人が開発に携わったことがあるかどうかなんて誰も興味を持っていないし、
信憑性もない。
もっと有益は議論をすべきだと思うけどね。

238:デフォルトの名無しさん
08/08/25 12:54:24
スレ違い　お前らネットイナゴは佃煮にされて死んでしまえ

239:デフォルトの名無しさん
08/08/25 12:58:38
派遣先企業で派遣先企業のプラットフォームに合わせたものを実装するのがおまえら派遣プログラマの仕事だろ
派遣先正社員からマルチプラットフォーム対応なんて指示されてないだろ
くだらないこと考えて余計な工数使うなよ

240:デフォルトの名無しさん
08/08/25 13:13:32
>派遣先企業で派遣先企業のプラットフォームに合わせたものを実装するのがおまえら派遣プログラマの仕事だろ
派遣先がパッケージベンダってこともあるだろう
なにムキになってるの？

241:デフォルトの名無しさん
08/08/25 13:23:54
>>239
何故派遣にこだわる？

242:デフォルトの名無しさん
08/08/25 13:49:24
ID出すとしょうもない素人以下の奴がいろんなスレに粘着してた

243:デフォルトの名無しさん
08/08/25 14:17:22
>>242
この板はID出ないじゃん

244:デフォルトの名無しさん
08/08/25 14:18:21
なんだみんな派遣だったのか

245:デフォルトの名無しさん
08/08/25 15:59:39
派遣に異常なコンプレックスを持っている人がいるようですが、
池沼の方なので基本放置でおねがいします。

246:デフォルトの名無しさん
08/08/25 16:06:10
夏の間だけ2chもアカウント認証式にすればいいと思う。

247:デフォルトの名無しさん
08/08/25 21:59:49
正社員にコンプレックスでしょ？
派遣のどこにコンプレックスをもてと

248:デフォルトの名無しさん
08/08/25 23:45:38
ひどい脱線ぶりだな。

249:デフォルトの名無しさん
08/08/25 23:49:36
派遣野郎はスレ違い

250:デフォルトの名無しさん
08/08/26 00:13:24
>>232
それありそう。PS3 と Xbox360 とか。どう似せるかに力が注がれる・・・

251:デフォルトの名無しさん
08/08/26 00:59:55
ここはなんのスレなんだ

252:デフォルトの名無しさん
08/08/26 04:09:43
くだすれ

253:デフォルトの名無しさん
08/08/26 06:23:56
CUDAはマルチプラットフォーム。

254:デフォルトの名無しさん
08/08/26 07:35:34
アホかい。

255:デフォルトの名無しさん
08/09/25 14:29:37
━━━━━━━━━━━━━━━━━━━
■ 今週の名言
───────────────────
　　　　　　　　　「HPCは“ハイ・プロダクティビティ・コンピューティング”」
　　　　　　日本SGIのHPC・サービス事業本部で本部長を務める田坂隆明執行役員）

256:デフォルトの名無しさん
08/09/28 22:50:51
CUDAってSIMD並列といいうことは、過去のスパコン用のコード
移植したら激速ってことでOK?

257:デフォルトの名無しさん
08/09/29 11:43:54
>>256
SIMDじゃないよ。仮にそうだとしても、そんなに単純な話じゃない。

258:デフォルトの名無しさん
08/09/30 05:16:33
SIMDだよｗ

259:デフォルトの名無しさん
08/09/30 16:35:15
SIMDじゃないなら、何なのか問いたい。

260:デフォルトの名無しさん
08/09/30 17:38:27
SIMDだったりMIMDだったり

261:デフォルトの名無しさん
08/09/30 21:12:21
実行はSIMDだけど
コードはスカラだから面倒臭いよ

262:デフォルトの名無しさん
08/10/01 10:37:52
NVIDIAはSIMTと言っているね。SIMD＋αくらい。
オンボメモリはSX8i並だからやっぱ速いんじゃねーの。

263:デフォルトの名無しさん
08/10/01 13:58:23
理屈はいいからなんか作れよ
＞＞過去のスパコン用のコード移植したら激速ってことでOK?
OKの一言で終わることをなにぐだぐだやってんだか

264:デフォルトの名無しさん
08/10/01 18:46:54
要するに、「移植できるもんならしてみろ」ってことだろｗ

265:デフォルトの名無しさん
08/10/01 23:50:07
だね。単純に移植できるとは思わない方がいい。

266:デフォルトの名無しさん
08/10/02 00:27:50
たしかにブロック切り分けの所もう少しスマートにならないのかと思うのだが
あんなもんデバイスの種類で最適化した組み合わせを自動算出すりゃいいだろ

267:デフォルトの名無しさん
08/10/02 00:29:52
>>266
CUDA2でストリームプロセッサ数なんかも取得できるようになったから、
是非とも自動算出関数を作ってみてくれ。

268:デフォルトの名無しさん
08/10/02 01:08:48
いや作るのは簡単なんだけどCUDAってクラスも使えないしラップ出来ないからだめじゃん

269:デフォルトの名無しさん
08/10/02 01:11:17
あいや出来るのか
使うのなら作るぞ

270:デフォルトの名無しさん
08/10/06 20:37:38
CUDAに興味があっていろいろ調べてるんだけど
URLﾘﾝｸ(slashdot.jp)
こんなにメモリアクセスがきついのか…
へたれの俺はCPUで頑張ろうって気になってきた
このあたりの条件が和らぐのはどれくらい後だろう？

271:デフォルトの名無しさん
08/10/06 21:10:43
それでもビデオカードは一世代速いメモリを積んでいるから、レイテンシを誤魔化せれば、CPUより多少はいいんじゃないかな。
CPUもGPUも、演算能力よりメモリのスピードがネックなんだよねえ。

272:デフォルトの名無しさん
08/10/06 22:18:57
>>270
そうそう、グローバルメモリはコヒーレントな読み出しなら4クロックなのに
ランダムアクセスすると100倍のクロックを喰ってしまう。
スラッシュドットの記事の書き方はちょっと言葉が足りてなくて、
実際にはいかにメモリをコヒーレントな読み書きで済ませるかが鍵ということになる。

例えば、CUFFTの2Dを1000回回すのと1Dバッチと自前の転置を2回ずつ動かすのとでは、
64x64辺りだと余り変わらないのに128x128になると途端に2Dの方が速くなる。
プロファイラで調べると、CUFFTの内部で用意されている転置関数はコヒーレントな読み書きしかしていない。
一方、自前の転置はコヒーレントな読み出しとインコヒーレントな書き出しになっている。
つまり、1000回呼び出すコストをインコヒーレントな書き出しコストが上回ってしまうということだ。

273:デフォルトの名無しさん
08/10/07 01:53:43
CPUより遅くなるということはない
電気代の割りに効果が少ないというのはあるｗ

274:デフォルトの名無しさん
08/10/11 06:33:45
TMPGEncで軽い処理はCPUより遅い。
VRAMとのメモリ転送あるし。

275:デフォルトの名無しさん
08/10/15 22:02:44
メインメモリと共有するタイプのGPUで処理した方が速い事も状況によってはあるのかな

276:デフォルトの名無しさん
08/10/15 22:26:03
CUDAでは有り得ないからスレ違い。GPGPUスレにでもどうぞ。

277:デフォルトの名無しさん
08/10/26 18:43:49
すいません、超低級な質問です。

dim3 threads(128, 1, 1);
dim3 grid(128, 1, 1);

hoge_kernel <<< grid, threads >>> (d_ptr, 128);

とかで関数を呼び出したんですが、ホスト側のスレッド数が128個生成
されるみたいです。これって、GPU内にスレッドが出来るんだと思って
たんですが、違うんでしょうか？

ちなみに、-deviceemu はつけてません。コンソールに以下のように
出るので、GPUにはいってると思います。

Using device 0: GeForce 8600 GT

278:デフォルトの名無しさん
08/10/26 20:58:06
>>277
＞ホスト側のスレッド数が128個生成されるみたいです。
そう判断した根拠は?

それとは別になるが、128x128が妥当かどうかは検証が必要だと思う。

279:デフォルトの名無しさん
08/10/27 17:11:09

質問です。

__device__ void hoge_kernel(void)
{
}

extern "C" __global__ void hoge(void)
{
dim3 threads(16, 1, 1);
dim3 grid(16, 1, 1);

hoge_kernel <<< grid, threads >>> ();
}

上をコンパイルすると、
「error: call can not be configured」
となるんですが、これはどういうことなんでしょうか。

280:デフォルトの名無しさん
08/10/27 17:47:00
__global__な関数から__device__の関数を<<<>>>で呼ぶことはできない。

281:デフォルトの名無しさん
08/11/01 09:40:52
ちょっと質問させてくらさい。。

８ストリーミングプロセッサ（SE)単位で構成されている
ストリーミングマルチプロセッサ（SM)でのことなのですが、
グローバル（ローカル）メモリ上へのデータのロード、
ストア命令は100サイクル以上かかってしまいますよね。

その間、8つのストリーミングプロセッサはおお休みしている
のでしょうか？

それとも、他のWarpのインストラクションが割り当てられたり
するのでしょうか。そうであるといいのですが、
そうするとSharedメモリがあとからのWARPにのっとられてしまって
まずそうですよね???

レジスタ的には32bit×2048本あるようなので(280の場合)
OKそうですが、Sharedメモリは16KBytesしかないし。

実際LD、STに160サイクルぐらいくったら、
160命令ぐらい無駄にしちゃって効率をかなり落としてる気がするので、
なんかやです。

24WARPぐらい、バッファリングしていて、割り当てられるWARPを
out of oderで実行するという記事も、どっかで見たような
気がするのですがSharedメモリがじゃまして無理があるような・・・

なんか良い方法があるのかな???

っていうか、なにか根本的な部分で俺の勘違い？

282:デフォルトの名無しさん
08/11/01 10:12:52
・グローバルメモリアクセスは、最大400(?)クロック掛かるが、最短では4クロックで済む。
# そのためには、coalescedにアクセスできるように工夫する必要がある。
・各ストリーミングプロセッサは、独立して動作する。Sharedメモリも同様。

例えば、行列の転置のような処理の場合、普通に書くとcoalescedに読んでincoherentに書かざるを得ない。
# 或いはその逆か。
そこで、CUFFT内で行なっている転置処理では、(プロファイルで見る限り)一旦共有メモリにおいて同期を取ることで、
読み書き共にcoalescedアクセスを維持しているようだ。

283:デフォルトの名無しさん
08/11/01 18:40:52
>>282
早速のレス、さんくす。

なるほど、最短４クロックですか。coalescedにしてもレイテンシー
だからだめかなって思ったけど、よくよく考えると、DRAMへの直接
アクセスに数百クロックってのはおかしいことに気づいた。

各ストリーミングプロセッサは、独立だけど１インストラクション単位で
同期してるんじゃないんかな（んなことぁない？）

・・・っておもったけどブランチがあるから、TPXレベルから見ると独立か???

最初見たときSharedメモリも他のSPのregも０レイテンシで使える
ようだったんで、演算と独立にグローバルメモリとのロード、ストアができん
１６K程度のSharedメモリ、あまり意味ないじゃんとかおもったけど、
そんなことなさそだね。

なかなか面白そうなチップだ。
分岐粒度が32って実際、汎用でどうなんだろ。
ベクトルつかったことないんでわからんが、
適当にセルオートマタでも乗っけて遊んでみまつ。

284:デフォルトの名無しさん
08/11/12 22:42:34
CUDA-Zなんて便利なものがありました。
forum.nvidia.co.jp

285:デフォルトの名無しさん
08/11/12 22:55:32
>>284
kwsk

286:デフォルトの名無しさん
08/11/12 23:19:24
URLﾘﾝｸ(cuda-z.sourceforge.net)

287:デフォルトの名無しさん
08/11/14 19:39:30
あげてすまん。

CUDAのためにGTX280つきのPCかったのだが
ほかにもう一枚グラボ買わないといけんのかな？

あとPCIバスしか残ってないのだが・・・
LINUXで使う予定。

288:デフォルトの名無しさん
08/11/14 22:57:10
>>287
別に一枚でも大丈夫。二枚挿しても(開発には)却って面倒だったり。
まぁ、二枚あれば表示に足引っ張られる心配なくなるから安定はするけどね。

289:デフォルトの名無しさん
08/11/15 00:41:03
>>288
そうなのか、安心した。

ありが㌧

290:デフォルトの名無しさん
08/11/15 17:39:36
質問です。

float にてゴリゴリ計算して、結果を返すプログラムを書いてみたのですが、
普通のアルゴリズム、CUDA+GPU、CUDA+CPU(deviceemu) の2つで比較して
みたところ、思ったより差が大きいのですが、こんなもんですか？

もしくは、機種依存があったりするんでしょうか？

291:290
08/11/15 17:40:23
>>290

すいません、間違えました。
「2つで比較して」->「3つで比較して」

292:デフォルトの名無しさん
08/11/15 19:43:01
何の差?
演算精度なら、そりゃぁあるさ、floatなんだもの。

293:デフォルトの名無しさん
08/11/15 20:51:43
ホストは、デフォルトではfloatをfloatで計算してくれないからな。

SSEを使うなりで、ホストがちゃんとfloatで計算すれば、結果は一致するんジャマイカ？

294:デフォルトの名無しさん
08/11/15 20:57:47
>>293
一致する保証はないよ。CUDAのドキュメントにもあるけれど、超越関数はGPU内部の組み込み版を使うと若干誤差が残る。
いずれにしてもfloatの想定の範囲内だから、実用上は問題にならないけどね。

295:,,・´∀｀・,,）っ-●◎○
08/11/15 22:26:10
x87 SSE CellB.E. CUDA の浮動小数サポートの対比表みたいなのがCUDAのマニュアルにあったな。
確かに完全じゃない。
糞まじめに準拠してるのはx87くらいだ

296:290
08/11/16 02:56:28
>>292-295
なるほど。出てきた結果が、CPU上の double で計算した結果と、
GPUの float で計算した結果が、最大1%程度違ったから、正直
驚きました。普段doubleしか使って無くて、誤差なんかほとんど
気にしなくてよかったので。

誤差が蓄積してくようなタイプのアルゴリズムではないと思っていた
だけに、少し驚きました。

297:デフォルトの名無しさん
08/11/20 18:53:51
Cで使っていた自作ライブラリは、
nvccでコンパイルし直さなきゃダメなの？

298:,,・´∀｀・,,）っ-○◎-
08/11/20 18:58:40
食べてみたけどおいしかったよ。
一つどうぞ、あーんして

299:,,・´∀｀・,,）っ-●◎○
08/11/20 19:44:47
飾りだ、食うなボケ。

300:,,・´∀｀・,,）っ----
08/11/20 21:36:00
僕のもうないです
>>299から貰ってね

301:デフォルトの名無しさん
08/11/20 23:10:32
>>297
ホスト(CPU)上で実行させるものならVC++と一緒にリンクできる。
（cppIntegrationサンプル参照）

GPU上で実行させたいならそれなりにいじくらないと無理。

302:デフォルトの名無しさん
08/12/01 15:16:29
-deviceemuでは動くプログラムが、GPUを使うと誤動作します。
__device__関数に引数が一部しか渡らなくなります(float4のz成分だけしか渡せない)
ループ回数を極端に減らすと改善されるのですが、これはregisterメモリがパンクしてしまったということでしょうか？

303:デフォルトの名無しさん
08/12/01 17:15:56
ソースを見ないとなんとも言えませんが、ループ回数とregisterは依存関係にはありません。
nvcc -ptxでptxファイルを出力して眺めてみては如何でしょう。

304:デフォルトの名無しさん
08/12/01 23:02:30
ptxファイルですか。。。
正直どこをどう見たらいいのかわからないので敬遠していましたが、
遂に避けられないとこまで来てしまいましたかねえ。

見るべきポイントとかありましたら、よろしければ教えてください。

305:,,・´∀｀・,,）っ-●◎○
08/12/01 23:09:56
ptxもネイティブコードじゃなしに中間コードの断片だからな。
レジスタが何本割り当てられるかとかそういった情報すら持ってないから困る。

その点Larrabeeは16本+16本+αとかレジスタ本数が決まってて
コード生成時点で割り当てが決まってしまう。
結果的に静的な最適化がしやすい。

良し悪し。

306:デフォルトの名無しさん
08/12/02 13:02:50
ﾘｮｳﾃﾆ●ヽ( ・∀・)ﾉ● CUDA！

307:デフォルトの名無しさん
08/12/03 17:58:11
.oとかにしてgccでリンクすることってできますか？

308:デフォルトの名無しさん
08/12/04 00:09:51
>>307
何を? どこで? なんで質問もろくにできないの?
ちなみに、エスパー募集ならお門違い。

309:デフォルトの名無しさん
08/12/04 00:52:30
>>307

普通にできますよ。
-l/usr/local/cuda/lib -lcudart
等のオプションは当然自分で追加する必要があります。

310:デフォルトの名無しさん
08/12/04 06:13:34
Windowsで、とかいう話だったりしてね。

311:デフォルトの名無しさん
08/12/04 10:05:39
.objじゃなくて.oって書いてあるんだからそれはないんじゃない？

312:デフォルトの名無しさん
08/12/04 22:31:14
それだったらわざわざ聞くかねぇ。まぁいいか。

313:デフォルトの名無しさん
08/12/12 18:47:36

【GPGPU】CUDA/ATI　STREAM　速度・画質検証スレ
ｽﾚﾘﾝｸ(jisaku板)
　

314:デフォルトの名無しさん
08/12/14 19:30:29
これからCUDAを始めようといろいろサイトを巡回しているのですが、
グリッドサイズとブロックサイズはどのようにして決めたらいいのでしょうか?
同期をとるかとらないかで変わると思いますが、とりあえずとらない場合でお願いします。

315:デフォルトの名無しさん
08/12/14 20:41:30
>>314
BlockSize(スレッド数)は32の倍数で余り大きくない辺り。GridSize(ブロック数)は充分大きく。

316:デフォルトの名無しさん
08/12/14 21:31:49
>>315
ありがとうございます。
もう一つ質問させていただきたいのですが、たとえば、512x512のようにブロックサイズ、グリッドサイズ
ともに最大値未満の場合、<16, 32, 1>のように複数の次元に分けるのと、<512, 1, 1>のように1つの次元
にまとめるのとどちらがいいのでしょうか?
多分通常はブロックサイズは各次元の最大値<512, 512, 64>をオーバーすることが多いと思うので、こちらは
分けることになると思いますが、グリッドサイズの最大値は<65536, 65536, 1>もあるので、
次元を分けなくてもいいときは分けない方がいいのでしょうか?

317:デフォルトの名無しさん
08/12/14 22:01:51
threadサイズが実際の分散処理の単位になる
つまりthread数だけ並列演算が行われる
例えばthread(1)にしてblock(10)とかだと順次処理と変わらない
ただしthreadは256までしか出来ないのでblockという単位が容易されてる
blockは単にプログラムがやりやすいように用意されただけのもので
実際に分散処理には影響しない

318:デフォルトの名無しさん
08/12/14 22:15:42
>>316
言いたいことがよく判らんが、演算量が512x512ならブロックを8x8にしてグリッドを64x64でいいんでない?
或いはxyの区別が重要でないならブロックを64、グリッドを4096とか。
私は後者の方針でptxファイルの出力とプロファイルの結果を睨みながら決定することが多いけど。
# 横着するときは前者だなぁw

319:デフォルトの名無しさん
08/12/14 22:17:30
>>316
単に内部ループが
for(x=0;x<?;x++){}
_global_
か
for(x=0;x<?;x++){
for(y=0;y<?;y++){
_global_
}
}
になるかの違い

320:デフォルトの名無しさん
08/12/14 22:24:34
元の位置を特定する為にglobal関数でインデックスを計算することになるから
もしxやyのどちらかで収まるなら１つだけ使うのが良い
global関数内部の計算コストが一番ネックになる部分だからね
インデックスの計算部分が一番邪魔で削れるだけ削った方が良い

321:デフォルトの名無しさん
08/12/14 22:35:25
>>317,319
ありがとうございます。ようやく喉のつっかえが取れました。

>>318
試しに昔書いた画像処理のコードを移植しようとしていたのですが、入力によって
いろいろサイズが変わるので、512x512で固定するのはちょっとやりづらいかなと…
ただ、ここまでの話からするとchannel(RGB)*height*widthの三次元配列に画像をつっこんで、
grid = <height, channel, 1>
block = <32, width / 32, 1>
で処理を回せそうですね。画像の幅をを32の倍数に合わせる必要はありそうですが。

>>320
ありがとうございます。実際に使うときは上のような感じにブロックとグリッドの数を決めようかなと
思いますが、どうでしょうか?

322:デフォルトの名無しさん
08/12/14 22:59:47
heightは超えてもいいの？
範囲は内部でインデックスが範囲かどうか判定するよりは
端数を含めたメモリを確保して必要無い部分も計算させてしまった方が速いよ

323:デフォルトの名無しさん
08/12/14 23:22:38
>>322
そこまで大きな画像を入れるかどうかは分かりませんが、確かに高さが65536を
超える場合は分割してやる必要がありますね。

>端数を含めたメモリを確保して必要無い部分も計算させてしまった方が速いよ
ということは、配列にコピーするときに幅を合わせてコピーするのではなく、後ろに空白部分を
まとめた方がいいということでしょうか?

#これって多次元配列を引数に渡すのってどうするんだろう…

324:デフォルトの名無しさん
08/12/15 00:01:31
いや、最終的には全て一次元の配列でそれが適切な間隔でアクセスされるように並んでいるのが一番いいことになる。
だから、画像なんかの場合だと座標値に意味がないなら詰め込んで構わない。
実際には、コンボリューションフィルターなどで隣の画素を参照しないといけないことが多いだろうけど、
その場合はどうせ共有メモリを使ってアクセスを減らすなどの工夫が必要だからどの途32の倍数に拘る必要はない気がする。

325:デフォルトの名無しさん
08/12/15 00:50:40
論文用のならともかく普通の画像処理で画像サイズが固定されてることはまずないし
サイズも共有メモリじゃとても足りない
共有メモリを使うならCPUサイドであらかじめ画像を分割して渡すという方法しかないけど効率悪いよ
毎回隣接した画素情報も含めないといけないから

326:デフォルトの名無しさん
08/12/15 01:06:12
>>324
カーネルサイズにもよりますが、sharedメモリはちょっとサイズ的にきつそうですね…
constantメモリならブロックサイズを調整すれば何とかなるかも?

>>325
単なる趣味グラムなので、ぶっちゃけ全部globalメモリからとってきてもいいような気がしたのですが、
逆にそれだとCPUで処理するよりも遅くなりそうですね

327:デフォルトの名無しさん
08/12/16 01:54:49
遅くなりそうというかどんだけがんばっても最終的にはglobalメモリを使うのが一番早いという結論になるよ
まあそれがGPGPUの一番のネックなんだけどね
CPUサイドのメモリを共用するようにいずれはなるんだろうけど今の限界はそこ

328:デフォルトの名無しさん
08/12/16 06:11:12
一番(手っ取り)早いのは同意するが、一番速いとは限らない。
それこそ、思考放棄じゃないかな?

329:デフォルトの名無しさん
08/12/19 00:31:55
詰まってます、助けて・・・
私はVS2005で開発しようとカスタムウィザードを入れ、ツールキットとSDKをインスコし
パスも通したのにLNK1181エラーでるんです・・・存在しないオブジェクトファイルを指定するのはなぜでしょう？
ライブラリパスがみつからないってどういうことなんでしょうか？

330:デフォルトの名無しさん
08/12/20 01:37:50
"Program" "Files"と認識してるとか

331:デフォルトの名無しさん
08/12/20 02:38:31
Windowsってなんで重要なディレクトリ名にスペースなんか入れるんだろう。
MSって頭悪いの？

332:デフォルトの名無しさん
08/12/20 04:16:06
低脳を排除するためかもねｗ

333:デフォルトの名無しさん
08/12/20 05:06:25
プログラマお断り。ずっと消費者で居てくだちい。

334:デフォルトの名無しさん
08/12/20 06:27:59
デスクトップやマイドキュメントのスペースや半角カナフォルダはなくなったが、Program Filesはそのままだな。

335:デフォルトの名無しさん
08/12/20 09:43:22
>>331
その昔１６ビットＯＳが主流だったころＭ＄は８文字ファイル名しか扱えなかった
そしてUnix系ＯＳは既にロングファイルネームに対応していた
そこで登場したのが32bit Window 95だった
長いファイル名が使えるんだぞ、凄いんだぞと新発明でもしたかのように宣伝していた
そして自慢げにフォルダ名を無駄に長くした

336:,,・´∀｀・,,）っ-●◎○
08/12/20 12:23:53
VistaはUNIXerには優しくなったな。
上位エディション限定だけどSUAも使えるし

337:デフォルトの名無しさん
08/12/20 20:39:06
CUDAをさ、VS2005で使えてる人っているの？
つか、なんで英語のマニュアルしかないんだよ。わかんねーよ

338:,,・´∀｀・,,）っ-○◎●
08/12/21 00:44:02
むしろ2005以外でどうやって使うんだ？
2.1βで2008対応したけど

339:デフォルトの名無しさん
08/12/21 00:58:06
>>338
おお、それはいいこと聞いた
やっと2005アンインストールできる

340:デフォルトの名無しさん
08/12/21 01:37:02
>>338
どこか日本語で環境設定を教えてくれるサイト知ってたら教えてくれ・・・

341:デフォルトの名無しさん
08/12/21 01:55:48
環境設定なんか何も要らんがな。スクリプト走らせるだけ。
bashのバージョンによってはウォーニング出るけど、ちゃんと動く。
スクリプト猿にも成れないとは、どんだけゆとり脳なんだよ?スポンジなの?

342:,,・´∀｀・,,）っ-●◎○
08/12/21 10:35:26
どのみちあの程度の英語が読めない人に優れたプログラミング書くのは無理だと思うんだ。
俺の私的な業務の手伝いしてくれるなら日本法人さんに代わりにチュートリアルとかの要求
してあげてもいいんだけど、職権乱用か。

>>340 ここで解らなかったらもう諦めたら？
URLﾘﾝｸ(chihara.naist.jp)

343:デフォルトの名無しさん
08/12/21 11:55:12
どうでもいいけど、↑のURLは「disp_content」を削らないとエラーになった。

344:デフォルトの名無しさん
08/12/21 12:58:13
>>340
サンプルプロジェクトを開いて設定を見て勉強するのがいいよ

345:デフォルトの名無しさん
08/12/24 12:51:08
英語読めないのにプログラミング言語してる男の人って・・・

346:デフォルトの名無しさん
08/12/24 18:35:50
英語に優越感もってるやつって、さもしいな

347:デフォルトの名無しさん
08/12/24 19:03:01
>>345
可愛いもんだよねｖ

348:デフォルトの名無しさん
08/12/24 23:58:39
英語なんて誰でも１ヶ月もあればある程度は読めるようになるよ
出来ないと思い込んでやろうとしないだけで

349:デフォルトの名無しさん
08/12/26 12:46:32
vs2005の環境で拡張子cuの場合、定義参照はできるのですが、
インテリセンスが機能しません。freeとかvcのランタイムも同様に
機能しないのですが、やり方分かる方いますか?
ツール->オプション->テキストエディタ->ファイル拡張子にcu c++として
登録するとできるのは上記だけです

350:デフォルトの名無しさん
08/12/26 22:27:18
いっそ、nvccの拡張機能以外は全部cppでやったら?

351:デフォルトの名無しさん
08/12/27 05:03:22
それがスマートだね

352:デフォルトの名無しさん
08/12/27 17:53:02
質問です。
VC++2005で開発する際に、コマンドプロンプトにデバイス名も含め一切文字を表示させたくないのですが、どのようにすればよいでしょうか？

353:デフォルトの名無しさん
08/12/27 18:38:27
>>352
長いパス名が邪魔なだけなら set PROMPT=$G とか？

354:デフォルトの名無しさん
08/12/27 18:47:25
>>352
CUTILを使うと標準出力があるって話なら、リダイレクトするかCUTILを使わなければいいと思うのだが。

355:デフォルトの名無しさん
08/12/28 01:33:08
プロンプト出したくないって話なら_tmainじゃね

356:352
08/12/28 02:53:51
>>353-355
回答ありがとうございます。

>>353
質問があいまいですみません。
パス名が邪魔、ってわけではないです。

>>354
おそらくこの指摘にあてはまるのだと思います。
＞リダイレクト
リダイレクトしても、コマンドプロンプトに表示が出てしまいます。
＞CUTIL
cutil.hを使わずに、cutil.hで提供されている関数を使わない方法、もしくは、代替可能なヘッダーファイルはありますか？

>>355
_tmainを使ってしまうとWindowsでしか使えないソースコードになってしまうので、別の方法があれば教えてください。

357:352
08/12/28 02:55:28
自己レスですｗ

誤：cutil.hで提供されている関数を使わない方法
正：cutil.hで提供されている関数を使う方法

の間違いです。

358:デフォルトの名無しさん
08/12/28 03:10:48
system("cls");
とか

359:デフォルトの名無しさん
08/12/28 10:49:28
>>356
> _tmainを使ってしまうとWindowsでしか使えないソースコードになってしまうので、別の方法があれば教えてください。
Windows かどうかで #ifdef するくらいは許されるんじゃないの？

360:デフォルトの名無しさん
08/12/28 14:29:28
プロジェクトの設定でWindowsアプリにすればいいだけでしょ

361:デフォルトの名無しさん
08/12/28 14:32:32
>>_tmainを使ってしまうとWindowsでしか使えないソースコードになってしまうので、別の方法があれば教えてください。
mainだろうがWindows依存部分のコードが発生するのは当たり前でしょ
CUDAの部分とC++の部分とWindows依存部分とファイルを分離するのが鉄則だよ
複数の環境で使う場合は#ifdefプラグマを使って分離したWindows依存部分の
#includeを切り替えるようにするだけ

362:デフォルトの名無しさん
08/12/28 17:29:43
コンパイラのオプションで指定できるでしょ
環境ごとにmakefile書くだけじゃね

363:デフォルトの名無しさん
08/12/28 17:48:49
>>361
#ifdefはいつからpragmaになりましたか?

364:デフォルトの名無しさん
08/12/28 17:53:47
>>363
#ifdefディレクティブと言いたかったんじゃないの?
で、使い慣れない言葉なんでつい間違えたと。
句読点もちゃんと使えないようだし、日本語に慣れていない三国人なんでしょ。

365:デフォルトの名無しさん
08/12/28 22:36:27
cudaのMDコードどこかに追ってませんか？

366:デフォルトの名無しさん
08/12/29 12:11:23
たった1行でここまで意味不明なのも凄いな。

367:デフォルトの名無しさん
08/12/29 12:24:23
×追ってませんか？
○オッドアイいませんか？

368:デフォルトの名無しさん
08/12/29 13:39:38
>>365氏はcudaで書かれたMD(分子動力学）コードどこかに落ちてませんか
と聞いてるのかと。当方も知らないので答えられませんが。
Fortranで書いたものを全部はアレなのでGPU上で実行したいサブルーチンだけ
Ｃに変えてCUDAで動かしたいのですが、そんな例とかは落ちてないですかね。
mainルーチンその他関係ないところまで全部Ｃに移植するのが嫌ってだけな
んですが。あ、当方はIntelFortran使用。
当方まだCUDA触りたて、試しにSTREAM BMTのtriadだけ手で適当に書いて
GeForce9600GTで40GB/s弱（効率7割弱）のメモリバンド幅。あ、じゃもうでき
るじゃんとか勝手に思い、Fortranのコードに挑もうとしてあえなく止まってますorz

369:デフォルトの名無しさん
08/12/29 14:08:06
つURLﾘﾝｸ(www.easize.jp)

370:368
08/12/29 14:20:59
>>369
ありがとうございます。m(_ _)m
多次元配列はどのみちGPU上では一次元化するつもりでしたがなるほど。
参考にさせていただきます。

371:デフォルトの名無しさん
08/12/29 14:45:46
>>369
そのページ、ポインタの扱いがメタメタだな
わかってるなら別にいいが

372:デフォルトの名無しさん
08/12/29 14:49:43
うん、ぶっちゃけ「"cuda fortran"でぐぐって一番上に来たリンク貼り付けただけ」なんだ、すまない

373:デフォルトの名無しさん
08/12/29 14:50:01
何か普段使うものをCUDAに移植しようと思いつつ適当な物が見当たらない

374:デフォルトの名無しさん
08/12/30 02:15:09
ぶっちゃけ、俺もそうなんだよね。おまけで付いて来たサンプル書き換えて
動かしてFLOPSベンチしてるだけ。
CUDA動かす環境をコスパ良く構築するには?と考えて色々やってみたが、
構築した環境で動かすモノって、結局サンプルの改造ばっか。まんまと
nVidiaの販促戦略に乗せられたぜw

375:デフォルトの名無しさん
08/12/30 08:50:29
zip圧縮とかjpg圧縮とかを移植したらライセンスの関係はどうなるの？

376:デフォルトの名無しさん
08/12/30 12:52:04
圧縮アルゴリズムに関しては問題ないんじゃね？

377:デフォルトの名無しさん
09/01/19 15:08:44
CUDAプログラミングモデルの概要
URLﾘﾝｸ(http.download.nvidia.com)
CUDAプログラミングの基本（パート I - ソフトウェアスタックとメモリ管理)
URLﾘﾝｸ(http.download.nvidia.com)
CUDAプログラミングの基本（パート II - カーネル)
URLﾘﾝｸ(http.download.nvidia.com)

378:デフォルトの名無しさん
09/01/21 03:21:23
デバイス側に確保したメモリにホスト側から
cudaMemcpyのように一括でコピーするのではなく、
一部分だけコピーする、又は書き換える良い方法がありましたら教えてください。

379:デフォルトの名無しさん
09/01/21 03:28:19
・cudaMemcpy()で一部分だけコピーする。
・cudaMemset()で（ｒｙ
・そういうカーネル関数を用意して呼び出す。

380:デフォルトの名無しさん
09/01/21 12:27:16
一部分だけってのはたいてい順次処理になるからCPUにやらせたほうが有利だよ

381:デフォルトの名無しさん
09/01/22 01:33:54
CUDAのプログラミングでZIPアーカイブのパスワード解析とか、早くなりませんかねぇ。
エンコード/デコードに使えるんだから、どうなんでしょう？

382:デフォルトの名無しさん
09/01/22 06:38:19
CUDAのfortranサポート予定ってGPUカーネルの部分をfortranライクに
書けるようになるって理解でいいのかな？2.0からって見かけたけど、それっぽい記述が全く無い…

383:デフォルトの名無しさん
09/01/22 08:47:26
>>379-380
ありがとうございました。
参考にさせていただきます。

384:デフォルトの名無しさん
09/01/22 19:18:17
>>381
総当りか辞書型か知らんが演算能力よりI/Oの問題だろ普通。

385:デフォルトの名無しさん
09/01/22 20:12:54
>>381
早くなるよ

386:デフォルトの名無しさん
09/01/23 00:20:22
>>384
ファイル数が多い場合はI/Oも問題になるかもしれませんが、
ある程度のサイズであればWindowsでもメモリに読み込むので、
それほど問題にならないのでは？
毎回物理的に読みに行くわけではないし。

387:デフォルトの名無しさん
09/01/26 09:28:38
ファイルI/OじゃなくメモリI/Oなのでは？

388:デフォルトの名無しさん
09/01/28 11:49:27
>>387
お前は何を言っているんだ

389:デフォルトの名無しさん
09/01/28 12:47:12
>>388
pciバスつかうんだろ？
立派なioとおもうが

つかクーダってほんとマイナーなんだね。

390:デフォルトの名無しさん
09/01/28 13:18:14
そりゃNVIDIA限定だからだろ？グラボを選ぶのなら汎用的なアプリは作れない。
CUDAよりもっといろんなグラボに対応しているOpenCLへのつなぎに過ぎないよ。

391:デフォルトの名無しさん
09/01/29 00:08:14
ZIPのパスワード解析にZIPファイル全体へのアクセスが必要だと勘違いしてる馬鹿ｗ

392:384
09/01/29 08:00:22
>>381-391
ああ、俺がぼけてた。なんか知らんけど自分でもなんであんなレスつけたんだろ orz
わけわからない流れにしちまって、すまん。

正しくはこうだな。↓
パス解析では総当りにしても辞書にしてもforなどのループ速度が求められるだけで、
CUDAによる計算は意味があるのかい？と思った。
MD5とかデータから導き出す数値ならば本当の意味での解析だから意味ありそうだけど。

393:デフォルトの名無しさん
09/01/29 23:02:34
こんな記事もあるしな

GPUを利用した無線LANのパスワードクラッキングが主流に
URLﾘﾝｸ(itpro.nikkeibp.co.jp)

394:デフォルトの名無しさん
09/01/30 03:18:20
>>393の言ってるのはまさに本当の意味での解析ってやつだしな。
WPA/WPA2-PSKからの復元は総当り/辞書比較と違ってGPGPUにする意味がある。

395:デフォルトの名無しさん
09/02/01 15:41:49
Catalystだけど、トリップ検索は随分速くなるみたい
URLﾘﾝｸ(sourceforge.jp)

396:デフォルトの名無しさん
09/02/02 21:51:40
ホスト側でconstantメモリを動的に確保できないのでしょうか？
方法がありましたら教えてください。

397:デフォルトの名無しさん
09/02/02 22:44:30
>>396
目的が良く分からないのでなんとも言えませんが、
どうせ64KBしかないのだから静的に64KB確保しておくのではだめなのでしょうか。

398:デフォルトの名無しさん
09/02/02 22:54:45
>>396
あるけど、なんでprograming guide読まないの？？

399:デフォルトの名無しさん
09/02/02 23:59:28
>>397
静的に確保しておいてもよかったんですが，綺麗にかけるならとw
>>398
一度目を通したつもりでしたが，素通りしていたようです。
もう一度読み直してきますm(_ _)m

400:デフォルトの名無しさん
09/02/06 00:08:13
誰かZIPの暗号解析ツール作って＞＜
それか解析のループ部分だけうｐしてくれたら、カーネルは俺が書きます。

401:デフォルトの名無しさん
09/02/06 06:54:02
総当りするだけならどっかのオープンソース実装の
パスワードチェックと復号ルーチンとってくればいいだろ

402:デフォルトの名無しさん
09/02/06 10:20:56
ただでさえ消費電力がデカイのにパスワード解析に何週間も動かしてられっかｗ

403:デフォルトの名無しさん
09/02/07 16:57:19
超解像のような考え方は昔からあるようだな

フリーソフトImageD2
URLﾘﾝｸ(www.tiu.ac.jp)

これは時間軸方向にも参照するものだ。
NVIDIAのMOTIONDSPと同じ考え方だな。
西川善司の大画面☆マニア第113回超解像
URLﾘﾝｸ(av.watch.impress.co.jp)

この方法は数フレームを参照することで低解像な映像のブレから情報量をアップさせるようだが、
この方法だとシーンチェンジやめまぐるしく動く映像では逆効果でめちゃくちゃな映像になるのではないか？
（これは上記のフリーソフトの別ページでも注意点として載っていた）
でも東芝などの日本の各社がやってる1フレームだけで行う超解像はそもそも無理がある。だから不自然な画質になったり、情報量が逆に消失したりする。
Lanczosなどでそのままアプコンしたほうがずっと情報量あるし自然な画質だ。比較してみれば一発で分かる。
URLﾘﾝｸ(plusd.itmedia.co.jp)
ここの元画像を720×480にし、それをAVIUTLなどでLanczosでフルHDにしたもののほうがずっと綺麗。
超解像は単純な処理だから柵とか崩れてるし、文字も駄目だし、元からあった情報を処理によって消しちゃう副作用のほうが強い。
超解像、超解像と目新しくいって盛り上げようとしたいのは分かるが、こんなのはまやかしだよ。

URLﾘﾝｸ(www1.axfc.net)
比較用画像もうｐしておいた

一方、数フレームでやる方法は計算が大変だが、シーンチェンジや盛大な動きの問題さえクリアすればかなり使えそうではある。
MOTIONDSPや↓はそのあたりちゃんとクリアしているのだろうか？
URLﾘﾝｸ(www.flashbackj.com)

404:デフォルトの名無しさん
09/02/07 17:01:19
10年ぐらい前にカーネギメロン大学で勉強してた頃に超解像のプログラム作れっていう課題を出されたことがある

405:デフォルトの名無しさん
09/02/07 17:34:44
ちょー解像は判ったからマルチすんなや。

406:デフォルトの名無しさん
09/02/07 20:13:09
>>403
スレ違いじゃない？
GPGPU 向きな処理なのは確かだし、そもそも Cell なり SpursEngine は GPU じゃないし。
画像処理スレとかあったと思うよ？

407:デフォルトの名無しさん
09/02/07 23:45:41
アメリカのメロン大学は一応名門だけど日本校はどうなんだ？

408:デフォルトの名無しさん
09/02/08 01:05:45
>>404
なにその高レベルな課題
俺の大学時代と雲泥の差

409:デフォルトの名無しさん
09/02/08 02:58:06
超解像の原理が未だに理解出来ません
複数の映像フレームの同じポイントの色の出現頻度に一番高い色を適用するってことですか？

410:デフォルトの名無しさん
09/02/08 13:45:42
は？今の超解像って、時間軸の補完までやってるのか？

411:,,・´∀｀・,,）っ-○◎●
09/02/08 13:52:05
MPEGの原理は主にフレーム間差分をとってJPEG圧縮なんで、それを改めてチェックしたところで・・・

412:,,・´∀｀・,,）っ-○◎●
09/02/08 21:06:32
>>407
神戸のハーバーランドの近くにあるあれか？
兵庫県の財政圧迫してるらしいよ。
学生少なくて撤退の噂もあります

同じビルの同じ階に神戸電子専門学校の学校法人が経営する大学院がある。
何を勘違いしたのか兵庫県立大学も情報科学の大学院を神戸市内に置いてる。

神戸という都市に何を求めてるのか、理解不能である。

たしかに古くからの工業都市で組み込みソフト屋の数がそれなりにいるのはわかる。
最大の勘違いは、日本のITドカタは職場を離れて大学院に通えるほど裕福ではないことだ。

413:デフォルトの名無しさん
09/02/08 21:49:03
俺はてっきり、ラーメン大学とか、洗車大学と同じノリで、
神戸電子専門学校が作ったヒト寄せパンダだと思ってた。
ってか別法人なのか?
学生少なくて、って言う前に、派遣のワーキングプア、
デジタル土方に成る為にわざわざ苦労して大学通う馬鹿
が神戸のどこに居るんだよっ?

俺が居るよ…orz

414:,,・´∀｀・,,）っ-○◎●
09/02/08 22:51:22
兵庫県もネギメロンなんて誘致するくらいなら神戸大に寄附講座でも作ったほうがよかったんじゃね？
現状兵庫県内で優秀な学生がいるんだし。

おっと、「学生」の枠で括ると近くの中高一貫校のほうがよっぽど・・・

415:デフォルトの名無しさん
09/02/09 09:41:15
情報系の大学院は多すぎるよな
旧帝でも定員割れする世の中なのに

このうえ西和彦が秋葉に大学院作るんだから馬鹿としか

416:デフォルトの名無しさん
09/02/09 12:50:30
>>411
誰がMPEGの話をしてるの?
誰も君の知識を披露してくれなんて頼んでないよ?

圧縮で壊れるったって圧縮率によるし、位相が生き残ってりゃ理論的には解像度は上げられる。
圧縮率によっては壊れた分が回復するだけかも知れないけどな。

417:デフォルトの名無しさん
09/02/09 14:21:54
誰がMPEGの話をしてるの?
誰も君の知識を披露してくれなんて頼んでないよ?

418:デフォルトの名無しさん
09/02/09 15:40:30
それを言うなら、なんでいきなり超解像の話になったんだろ？

419:デフォルトの名無しさん
09/02/09 15:43:20
メロン日本校のカリキュラムを見たけど、ウンコすぎて話にならねぇｗ
こりゃ学生もこねーわ

420:,,・´∀｀・,,）っ-○◎●
09/02/09 18:02:06
>>416-417
だれも連投してくれなんて頼んでないよ？

421:,,・´∀｀・,,）っ-○◎●
09/02/09 18:06:41
GPGPU向けの画像・動画関連のソフトってどれも速さばかり求めて品質は二の次だな
ウンコなエンコーダだとRGB/YUV変換で腐る
一番腐るのは量子化だろうけどな

422:デフォルトの名無しさん
09/02/09 18:25:28
Geforceは何気にGTXシリーズで初めて64bitに対応してるけど
演算装置は１個しかないｗ
並列処理が得意な分野でfloatだけで済むようなものはほぼ無いだろｗ

423:デフォルトの名無しさん
09/02/09 19:05:29
>>420
連投じゃなくて別人
2ch歴長いんだからそれくらい分かれ

424:デフォルトの名無しさん
09/02/09 20:42:10
甘いな。floatでも速ければ使い物になる用途は結構あるもんだよ。
例えば、近似計算なんかはfloatで近似させてからdoubleで更に近づけることもできるしね。

425:デフォルトの名無しさん
09/02/09 21:56:34
そういや西和彦もこっちの出身だな。確か甲陽から和田大
だから、落ちこぼれか。そもそも博士余りで高学歴ワープアPD
が社会問題に成ってる今日この頃に、大学院新設とかもうね
アボガドバナナ…。
鴨葱メロンと言えば、金出教授もこっちの出身だったな。
沢山カメラ使って超解像みたいな論文も有ったような無かったような。

426:デフォルトの名無しさん
09/02/13 13:06:41
>>425
お前、自分の考えまとめるの下手だな。

427:デフォルトの名無しさん
09/02/13 13:31:44
まあ、CUDAは英文のドキュメントが読めてある程度知能がないと
使えないからな。

スレ違いの話で盛り上がるのも分かるな。バカばっかりだしw

428:デフォルトの名無しさん
09/02/13 17:06:53
CUDA版のTripcode Explorerができたみたいですね。最適化に期待します。
URLﾘﾝｸ(tripper.kousaku.in)
URLﾘﾝｸ(download.kousaku.in)

429:,,・´∀｀・,,）っ-●◎○
09/02/13 18:13:51
ようこそ、バーボンハウスへ。
このテキーラはサービスだから、まず飲んで落ち着いて欲しい。

うん、「絶対に動かない」んだ。済まない。
仏の顔もって言うしね、謝って許してもらおうとも思っていない。

でも、このネタプログラムを見たとき、君は、きっと言葉では言い表せない
「ときめき」みたいなものを感じてくれたと思う。
殺伐とした世の中で、そういう気持ちを忘れないで欲しい、そう思って
５分ででっちあげたんだ。

じゃあ、注文を聞こうか。

--------------------
1M超のバイナリファイルに何が詰まってるか疑問な人は、テキストエディタで開いてみればいいよ＞＜

430:デフォルトの名無しさん
09/02/13 23:27:26
CUDAはコアレスと分岐の扱いを把握すれば、やりたいことは大体クリア出来ると思われ。

431:,,・´∀｀・,,）っ-○◎●
09/02/14 00:05:15
結局プレディケートなんだよね。
SMあたりの命令デコーダは1基だからSP毎に別々のフローを実行することができない。

Larrabee（Ct）なら分岐は容易に表現できる。

432:デフォルトの名無しさん
09/02/14 01:58:51
>>431
プレディケードって？？

433:,,・´∀｀・,,）っ-●◎○
09/02/14 11:48:56
えーとね、たとえば
if (cond) {
　　funcA();
} else {
　　funcB();
}
なんてコードがあるとしよう。
普通のCPU向けのCだと、 cond の条件にしたがって、funcA()を実行するブロック、
あるいはfuncB()を実行するブロックにジャンプする。
すなわち命令ポインタを操作してコードを飛ぶ。

CUDAにおいてはシェーダマルチプロセッサ一つにたいし、命令デコーダは1つしかない。
にもかかわらず、condは要素ごとに変わるわけで、条件分岐先はSPごとにバラバラになる可能性がある。
んで、そこで使うのがプレディケートなわけだけど、簡単にいえば、ifとelseの両方を通るようにする。
funcAとfuncBをインライン展開して、条件ビットで選択的に実行するコードに展開する。
んで、各要素に対して、実行するか実行しない（あるいは実行結果を反映しない）かを選択的に行うわけだ。
並列度の高いプロセッサではよく使われる方法だ。

んで、こっからはこのアプローチの弱点。
問題はif-elseを何重にも組み合わせたり、switch文を多用する場合、総当たりにかかる計算時間量が
並列化によるパフォーマンスメリットを相殺し、逆に遅くなるケースもある。
並列処理を諦めて素直に要素ごとに逐次処理をさせてくれたほうがかえって効率がいいかもしれない。

しかしCUDAってそのへんの融通がきかないんだよね。基本的に【並列処理しか記述できない】から。
正確には逐次処理は専用のプリミティブなんかを使って限定的に逐次処理はやれるけど記述面では
かえって面倒になる。
GPUで不得意な処理はCPUでやれってアプローチだからそのへんの融通を利かせる気は無いらしい。

434:デフォルトの名無しさん
09/02/14 11:57:31
CUDAで困るのはその点のほかに、並列数を途中で変えられないこともあるよね。
一度ホストに処理を戻すと遅くなりかねないし、共有メモリが失われてしまうし。
私の関わっているプロジェクトでは演算処理が中心なので、ある程度融通が利いてくれないとね。

次ページ