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にならなくね？
455:デフォルトの名無しさん
10/01/04 16:13:31
257⇒0 258⇒1 … でいいの？
でもそうすると512⇒0がヘンだけどな。
int val;
val = a + 256;
if (val > 512) val = val - 513;
でいいんじゃないの？
というかC言語の疑問をCUDAスレで聞くのはまちがっとる、ここは人口が少ないぞ
456:デフォルトの名無しさん
10/01/04 16:26:48
じゃあCUDAっぽく
const int TABLE[512] = { 256, ... , 512, ... , 0 }
__device__ int f(int a){
return TABLE[a];
}
457:デフォルトの名無しさん
10/01/04 17:09:36
レスありがとうございます
わけわかんないこと書いてすみません
テーブルとif文試して出直してきます
458:デフォルトの名無しさん
10/01/04 17:37:31
>>453
b⇒aとして、
a = (b+256)%768;
でいいのかな？
0⇒256
256⇒512
511⇒767
512⇒0
459:デフォルトの名無しさん
10/01/04 21:45:29
釣りだよな？
460:デフォルトの名無しさん
10/01/05 01:06:37
申し訳ないっす
461:デフォルトの名無しさん
10/01/07 00:21:36
ランダム関数の作成で悩んでます。
・CPUで作成した乱数リストをGPUに転送しない
・1スレッドの中で何度も呼ばれる
・グローバルメモリは使わない。スレッドIDやclock()をseedとして使う
この制約で良い乱数生成アルゴリズムを教えて下さい。
MTGPっていうのを参考にしようとしましたがムズ過ぎて挫折…まずコンパイルが通らず(´・ω・`)
462:デフォルトの名無しさん
10/01/07 01:01:35
おれMTGPコンパイルできたよ。GPU側で作成してそのままGPUで使えるし。
（シンボルコピー）
463:デフォルトの名無しさん
10/01/07 08:15:53
使ったファイルはこれだけ
mtgp32-fast.h
mtgp32-cuda.cu
mtgp32-fast.c
mtgp32-param-fast.c
make_single_random()とかの末尾を適当に書き換える。
464:デフォルトの名無しさん
10/01/07 12:09:10
誰かＯｐｔｉＸやってる人いませんかー？
465:デフォルトの名無しさん
10/01/08 09:15:11
CES 2010: NVIDIA Shows GeForce GF100 Fermi Video Card - Legit Reviews
URLﾘﾝｸ(www.legitreviews.com)
3/2に本当に出るのかしらん
466:デフォルトの名無しさん
10/01/08 09:22:19
気にするな、出たら買う、出るまではCUDAの修行をする。
単純でｲｲじゃないか。
467:デフォルトの名無しさん
10/01/08 09:41:48
>>461
Perlinノイズやってみてよ
URLﾘﾝｸ(mrl.nyu.edu)
Cのソースが有る
468:デフォルトの名無しさん
10/01/08 11:02:05
Fermiもダメそうってどこかのブログで見た
NVIDIAまじやばい
469:デフォルトの名無しさん
10/01/08 11:05:40
ようやくCUDAに慣れてきたところなのに、StreamSDKとかOpenCLとか
DirectComputeとかまた勉強しないといけないの？
辛いのうｗｗプログラマは辛いのうｗｗｗ orz
470:デフォルトの名無しさん
10/01/08 13:36:24
windows7 64bit、Visual C++ 2008 Pro環境に、CUDA toolkit、sdk 2.3をインストールして、
C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK\C\srcにある、
_vc90.slnファイルを開いて、問題なくビルドできます。
exeファイルの出力先は、デフォルトでC:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK\C\bin\win64\ですが、
これ以外の出力先に変更する方法を教えてください。
471:デフォルトの名無しさん
10/01/08 13:42:23
プロジェクトのプロパティで変更できるよ
$(OutDir)\$(ProjectName).exe
これを丸ごとフルパスで指定してしまいなさい。
472:デフォルトの名無しさん
10/01/08 17:14:12
XP,32bitでコンパイルしたプログラムを7,64bitで動かそうとしたら動かなかった･･･
こういうものなのか
473:デフォルトの名無しさん
10/01/08 17:23:41
32/64ビットのcudart.dllが必要
474:デフォルトの名無しさん
10/01/08 19:20:38
>>471
thanks
475:デフォルトの名無しさん
10/01/10 12:33:02
現在CUDA SDKを導入しVisual C++ 2005 Express Editionを使用しCPUでのエミュレーションを行っています。
いくつかのサンプルは動かせたのですがparticlesなど複数のサンプルでコンパイルはできているようですが実行すると以下のようなエラーが出ます
device emulation mode and device execution mode cannot be mixed.
エラーが出ている行は
cutilCheckMsg("cudaMalloc failed");
cutilSafeCall(cudaGLRegisterBufferObject(vbo));
cutilSafeCall( cudaMalloc3DArray(&noiseArray, &channelDesc, size) );
のように「cutil～」になってます。オプションから設定しているのでデバイスエミュレーション自体はうまく行ってるようなのですが
エミュレーションを行う際はソースを書き換える必要があるのでしょうか?
476:デフォルトの名無しさん
10/01/10 12:55:52
メッセージのとおりエミュONOFFまぜてる
またはリビルドしてないんじゃ？
プロジェクト設定じゃなくてビルド構成のほうでエミュレーションにかえられないっけ？
477:デフォルトの名無しさん
10/01/10 13:05:00
メッセージの内容は把握できています。
リビルドも多分できてると思います。
エミュレーションの設定は
プロジェクトのプロパティ→CUDA Build Rule→Emulation Mode はい
にしてます。ほかの方法はよくわからないのですがほかの方法があるのでしょうか?
478:デフォルトの名無しさん
10/01/10 22:13:13
ビルド構成のほうからエミュレーションの設定ができました。問題なくビルド&実行できました。ありがとうございます。
ただFPSが0.1なんでほぼ無意味でした。
479:デフォルトの名無しさん
10/01/10 23:24:17
エミュレーションなんてHello world以上のことをやるもんじゃないな
480:デフォルトの名無しさん
10/01/11 00:20:59
処理する画面領域を１０００分の１ぐらいに限定して、バグを調査するくらいのことはできる。
481:デフォルトの名無しさん
10/01/11 01:26:19
なんだっけ、エミュじゃなくて実機で実デバッグできるやつ
早く来てほしいな
482:デフォルトの名無しさん
10/01/12 14:25:34
CPU版アプリの移植の際に、エミュはデータ検証に使える。
483:デフォルトの名無しさん
10/01/12 20:21:38
-G オプションって使える？WinXPで
「'assembler' は、内部コマンドまたは外部コマンド、
操作可能なプログラムまたはバッチ ファイルとして認識されていません。」
と出てビルドできないけど。
--deviceemuもつけないとだめなの？
484:デフォルトの名無しさん
10/01/13 09:38:31
CUDAのバージョンを書き出しておきたいのですが、
CUDAのバージョンを調べる関数か何かありませんか？
485:デフォルトの名無しさん
10/01/13 09:41:29
>>484
>>436 でどう？
486:デフォルトの名無しさん
10/01/13 09:50:27
こっちを要求している気がする。
cudaRuntimeGetVersion()
cudaDriverGetVersion()
487:デフォルトの名無しさん
10/01/13 09:56:08
WindowsでCUDAを使っています。
ループが一定回数超える度に、coutで経過を表示しています。
Linuxに比べて遅かったので、いろいろ試してみたところ、
DOSウインドウの前で、マウスを動かすと、速くなります。
垂直同期か何かで引っかかっているのかと思い、設定で切ってみましたが
状況が変わりません。
どなたか似たような状況になって解決された方はいないでしょうか？
488:487
10/01/13 11:05:28
追加報告です。
cudaMemcpyでcudaMemcpyDeviceToHostでコピーするところで異常に時間がかかっているようです。
この1行を消すだけで、100秒ほどかかっていたのが2秒にまでなりました。
逆方向のcudaMemcpyHostToDeviceは問題ありません。
489:デフォルトの名無しさん
10/01/13 11:13:17
CUDAプログラムがGPUで計算等をしている間は画面表示の更新が遅くなります。
１回のループにどれくらいの時間がかかるのか、一定回数がどれくらいかによりますが、
経過の表示される間隔が短いとその影響を受ける可能性があります。
あるいは、マウスを動かすことでCPUに負荷をかけないとパワーセーブモードに入ってしまうとか・・・・。
490:デフォルトの名無しさん
10/01/13 11:17:06
>>488
転送に１００秒かかるのは現実的にありえないので、恐らくcudaThreadSynchronize()等でカーネルの実行完了待ちを
していないのではないかと思います。カーネル呼び出しそのものは一瞬で終わりますが、
後続のcudaMemcpy(.....,...,...., cudaMemcpyDeviceToHost)時に暗黙的にそれまでに呼び出したカーネルの実行完了を待っています。
・・・・・ということかどうかは分かりませんが。
491:デフォルトの名無しさん
10/01/13 11:29:27
>>489,490さん
早速の返信ありがとうございます。
cudaMemcpyでcudaMemcpyDeviceToHostを消したプログラムで、
プログラムの一番最後にだけ結果を取り出してみたところ、
正しく計算されていました。
次に、カーネルの実行部分のみをコメントアウトした場合、
やはり非常に時間がかかってしまいました。
やはり、GPUからCPUへのデータ転送に（というよりなにか同期の部分の様な気もしますが）
時間がかかっているようです。
計算用マシンなので省電力の設定はしていません。
492:デフォルトの名無しさん
10/01/13 11:48:15
>>491
> cudaMemcpyでcudaMemcpyDeviceToHostを消したプログラムで、
> プログラムの一番最後にだけ結果を取り出してみたところ、
> 正しく計算されていました。
これを読むと、そもそも何度もDeviceToHostの転送をする必要がない処理という解釈で
よいのでしょうか？
DeviceToHostが遅くなる理由はハードウェア的なものから色々あります。
マザーボードを交換したという人もいました。
SDKに含まれているbandwidthTestの結果はどうなっていますか？
"--memory=pinned" を付けた場合と付けなかった場合をそれぞれ調べてみてください。
493:デフォルトの名無しさん
10/01/13 11:59:12
Device 0: GeForce GTX 285
Quick Mode
Host to Device Bandwidth, 1 Device(s), Paged memory
Transfer Size (Bytes) Bandwidth(MB/s)
33554432 4340.2
Device to Host Bandwidth, 1 Device(s), Paged memory
Transfer Size (Bytes) Bandwidth(MB/s)
33554432 4193.7
Device to Device Bandwidth, 1 Device(s)
Transfer Size (Bytes) Bandwidth(MB/s)
33554432 119962.3
--memory=pinned
Device 0: GeForce GTX 285
Quick Mode
Host to Device Bandwidth, 1 Device(s), Pinned memory, Write-Combined Memory Enabled
Transfer Size (Bytes) Bandwidth(MB/s)
33554432 5758.1
Device to Host Bandwidth, 1 Device(s), Pinned memory, Write-Combined Memory Enabled
Transfer Size (Bytes) Bandwidth(MB/s)
33554432 5471.9
Device to Device Bandwidth, 1 Device(s)
Transfer Size (Bytes) Bandwidth(MB/s)
33554432 119852.3
494:デフォルトの名無しさん
10/01/13 12:00:32
結果はこんな感じです。
特に片方だけ遅いというわけでもないようです。
Linuxで起動した場合、遅くなるという問題は起きませんでした。
DeviceToHostが必要なのは結果を取り出すときだけで、
計算自体には必要ありません。
495:デフォルトの名無しさん
10/01/13 12:03:47
似たような話をnvidiaのforumで見つけました
URLﾘﾝｸ(forums.nvidia.com)
URLﾘﾝｸ(forums.nvidia.com)
返信は結構あるのですが、解決策はよくわかりません。
496:デフォルトの名無しさん
10/01/13 14:46:26
・CUDA内での計算処理、ループ数とループ内のステップ数、スレッド数ブロック数、だいたいどんなもんすか
・HostToDevice、DeviceToHostで転送するメモリのサイズはどのくらいですか
497:デフォルトの名無しさん
10/01/13 15:02:38
とりあえず
タイマー1,2作成
cudaMemCopy(HostToDevice)
タイマー1,2起動
<<<>>>カーネル呼ぶ
cudaThreadSyncronize()
タイマー1停止
cudaMemCopy(DeviceToHost)
タイマー2停止
してほんとにMemCopyなのかカーネルなのか確かめてみる必要は。
498:デフォルトの名無しさん
10/01/13 19:23:55
コピーが細切れでforループでやたら回数呼んで転送してたりしない？
１発で全部転送するのと1/1000ずつ1000回コピーするのとでは、
かかる時間に雲泥の差があるけど。
499:487
10/01/14 04:10:12
PCを再起動したら、上記の問題は出なくなりました。
お騒がせしました。
500:デフォルトの名無しさん
10/01/14 09:16:24
おいｗｗ 終わりかよｗｗ
501:デフォルトの名無しさん
10/01/14 11:44:30
cutil64.dllってどこに置けばいいの？
502:デフォルトの名無しさん
10/01/15 05:05:43
GPU側で配列の合計を求める方法を教えてください
503:デフォルトの名無しさん
10/01/15 07:59:04
サンプルのreduction見れ
504:デフォルトの名無しさん
10/01/15 09:52:56
syncthreadsって異なるブロック間のスレッドも同期されるのでありましょうか？
505:デフォルトの名無しさん
10/01/15 20:54:27
>>504
programing guide よめ
506:デフォルトの名無しさん
10/01/15 23:10:51
Best Practices Guide
もいいこと盛りだくさん書いてあるよ
507:デフォルトの名無しさん
10/01/16 19:28:55
cudaでは多次元配列の使用は推奨されてないの？
508:デフォルトの名無しさん
10/01/17 12:17:08
pthreadとの組み合わせでうまく動作せずに困っています。
引数として渡す構造体に、float * device_Xを用意して
for(int k=0;k<NUM_THREADS;k++){
cudaMalloc((void**)&targ[k].device_X,
sizeof(float)*(hostX.nRow)*(hostX.nCol));
cudaMemcpy((void*)targ[k].device_X,
host_X,
sizeof(float)*(hostX.nRow)*(hostX.nCol),
cudaMemcpyHostToDevice);
pthread_create(&handle[k], NULL,
compute_thread_func,
(void *)&targ[k]);
}
と渡すと、pthread内で
cudaMemcpy(temp,argument->device_X, //argumentは渡されたvoid* argをキャストしたもの
sizeof(float)*(nRow)*(nCol),
cudaMemcpyDeviceToHost);
としてもうまく取り出せません。（中身がぐちゃぐちゃになる）
同様の処理をpthreadで呼び出す関数内で行うとうまく動作します。
スレッドごとにメモリを取るのだから、中でやってもいいのですが気になります。
何か原因になっていそうなことは無いでしょうか。
509:デフォルトの名無しさん
10/01/17 12:59:47
>>508
＃1個のGPUに複数のスレッドからアクセスしたいアプリケーションなのかよく分かりませんが・・・・。
CUDAはスレッド毎にGPUのリソース（コンテキスト）を管理しています。
従ってこの例では
子スレッドのcompute_thread_func内の処理は
親スレッドがcudaMallocした領域にアクセスできません。
510:デフォルトの名無しさん
10/01/17 13:32:14
>>509
ありがとうございます。
複数のスレッドから一台のGPUへのアクセスです。
（後々、スレッドごとに一個のGPUにしたいのですが）
スレッドごとに管理しているのははじめて知りました。
解決しました。
511:デフォルトの名無しさん
10/01/18 20:51:33
「Picture Motion Browser VAIO Edition」を追加。
超解像HDアップコンバートやノイズ除去、手ぶれ補正などの編集や動画作成を、NVIDIA CUDAによるGPU演算を使って行なえる。
512:デフォルトの名無しさん
10/01/18 22:00:36
GT200 Fermi
トランジスタ数 14億 30億
倍精度浮動小数点 30FMA Oｐｓ/clock 256FMA Oｐｓ/clock
単精度浮動小数点 240 512
シェアドメモリ(SM) 16KB 64KB
L1メモリ 無し 48KB or 16KB
L2メモリ 無し 768KB
アドレス幅 32bit 64bit
複数のカーネル実行
無し 16個まで（GigaThredエンジン)
＊L1キャッシュ搭載
GT200では16KBのSMが32個あり、それぞれCUDAコアを8個割り当てられ、L1キャッシュが無し。
Fermiでは64KBのSMが16個。それぞれにCUDAコアが32個割り当てられ、SMから16KBか48KBのどちらかをL1キャッシュとして使用可能。
GT200に対して、3倍のSMと16KBのL1が使用可もしくは同じサイズのSMと48KBのL1が使用できるようになった。これにより、今までCUDA化できなかったプログラムの対応を増やし、さらに高速化もできる。
各CUDAコアに含まれるFPユニットは倍精度浮動少数演算を強化し、GT200に対し8倍の能力。
＊L2メモリの搭載 グローバルメモリに対する高速化。
＊C++をサポート
＊複数のカーネルの動作をサポート
SM内部のパイプラインを強化。SFUが複数に分けられたのでタスクとデータをより効率化。スレッドスケジューラを2個。
＊双方向PCIE通信
GT200ではPCIEバスの送受信をどちらか片方しか一度に実行できず、理論値8GB/s・実測4～5GB/s程度だが
Fermiでは双方向通信が可能になり12GB/sを実現
＊新しいメモリコントローラ FermiよりGDDR5まで対応し、ECCにも対応する。
＊コア内部を各部でモジュール化
設定された価格帯や用途ごとにコアを設計しなおさず、機能をカットオフしたり追加したりできる。
SM単位でCUDAのコアを減らしたり、D3DやOpenGLなどの固定ハードウェアが不要なTeslaなどでオフになったりする可能性もある。
513:,,・´∀｀・,,）っ-○○○
10/01/19 02:32:35
GT200 理想のFermi 現実のFermi
トランジスタ数 14億 30億 30億
倍精度浮動小数点 30FMA Oｐｓ/clock 256FMA Oｐｓ/clock 224FMA Oｐｓ/clock
単精度浮動小数点 240 512 448FMA Oｐｓ/clock
シェアドメモリ(SM) 16KB 64KB 〃
L1メモリ 無し 48KB or 16KB 〃
L2メモリ 無し 768KB 〃
アドレス幅 32bit 64bit 〃
複数のカーネル実行 無し 16個まで（GigaThredエンジン) 14個まで（GigaThredエンジン)
514:デフォルトの名無しさん
10/01/19 04:04:04
OpenCL待ちしてたけど、GPUの対応はCUDAだけっていう開発ばっかな
Fermiで本気出すから
515:デフォルトの名無しさん
10/01/20 02:23:16
GT100ってなんかメモリアクセスのレイテンシがすごくでかくない？
それを隠蔽するために、L2キャッシュが128Kになっているけど、
コヒーレンシとかどうなのかな？
やっぱりGPUはデータ並列なアプリしか向いていないのかね。
516:デフォルトの名無しさん
10/01/20 16:04:40
Fermiのグローバルメモリのレイテンシが遅いってのは何処から来た情報？GDDR5に対応するんだから、帯域は大きくなりそうだけど。
517:デフォルトの名無しさん
10/01/20 17:08:55
>>512
SFUは今のGT200でも複数ある。
ただ、SFUが外部のパイプラインに分かれたからSFUを使っているときに違う種類の演算ができるようになったってことじゃね・。
>>515
最適化しなければいけない
518:デフォルトの名無しさん
10/01/21 03:41:00
>>516
これをみて思った。
URLﾘﾝｸ(pc.watch.impress.co.jp)
でもGT200と一緒なんだね。
URLﾘﾝｸ(pc.watch.impress.co.jp)
いずれにしても、何でGPUがあんなにレイテンシが大きいのかがわかったよ。
だから、coalesced accessが必要なのね。
519:デフォルトの名無しさん
10/01/21 04:03:50
>>518
>いずれにしても、何でGPUがあんなにレイテンシが大きいのかがわかったよ。
俺のために説明してくれ！
520:デフォルトの名無しさん
10/01/21 22:43:04
GPUのメモリのレイテンシが大きいというよりも、CPUの場合と同じぐらいはかかるといった方がいい
521:デフォルトの名無しさん
10/01/22 01:50:01
CPUだとメモリチャンネルが精々2～3チャネル位しか無いが、GT100だと12チャネルもあるから、DRAMのチャネルのスイッチングに
それ以上かかると思っていい。あと、仮にGPC1?がDRAM0チャネル上にあるデータと、
DRAM11チャネル上にあるデータを同時にアクセスする場合、当然レイテンシがかかる。
もし、GPC1がDRAM0、DRAM1だけのデータアクセスですむなら、CPUと変わらない。
最適化が必要というのはそういったところだと思う。
522:デフォルトの名無しさん
10/01/22 05:43:22
MTGPがコンパイルできません。
>>463の言う、4つのファイル使って試したところ100個近くコンパイルエラーが出てしまいます。
あと「末尾を適当に書き換える。」の意味がよく分からんどす…
エラー内容は
1>\略)\mtgp32-fast.h(117) : error C2054: 'inline' の後に '(' が必要です。
1>\略)\mtgp32-fast.h(120) : error C2057: 定数式が必要です。
1>\略)\mtgp32-fast.h(120) : error C2466: サイズが 0 の配列を割り当てまたは宣言しようとしました。
1>\略)\mtgp32-fast.h(120) : error C2085: 'mtgp32_do_recursion' : 仮パラメータ リスト内にありません。
1>\略)\mtgp32-fast.h(121) : error C2061: 構文エラー : 識別子 'inline'
1>\略)\mtgp32-fast.h(122) : error C2054: 'inline' の後に '(' が必要です。
...
WinXP32bit、VC++2005、CUDA SDK2.3
該当箇所のコード見てもどこが悪いのか分からない…助けて…。
523:デフォルトの名無しさん
10/01/22 07:00:08
単にパイプラインステージが何百もあるからレイテンシがでかいのだと思ってました
524:デフォルトの名無しさん
10/01/22 08:15:33
>>522
あ。おれか。
元のソースと自分のソースでdiffしてみた。
・inline を全部削除してみて。
・末尾を適当に書き換える。は、元だとプリントして捨ててしまっているので、
利用したいように書き換えてねと。
おれはmake_uint32_random(d_status, num_data, hostbuff); として
hostbuffを自分で使えるようにしました。
525:デフォルトの名無しさん
10/01/22 20:11:37
>>523
正解
メモリチャンネル云々は殆ど関係ない
CPUはメモリとのクロック差のため
GPUは長大なパイプラインのためレイテンシがデカイ
526:デフォルトの名無しさん
10/01/22 20:30:44
別にパイプラインが深いわけではなくてバッファが大きいだけなんだけど。
527:デフォルトの名無しさん
10/01/22 20:51:46
200-400段は十分に・・
528:デフォルトの名無しさん
10/01/23 00:27:17
>>524
レスサンクス。
inlineを削除したら見事にエラーが亡くなりました(`・ω・´)
これから自分のソースに組み込もうか…と状況に入れます。
重ねて質問申し訳無いですが、hostbuffの名前からしてMTGPの乱数は一度バッファにぶち込んでから使うという事になるのですか？
>>461ようにカーネルの中で特に意識せず使いたいのだけれども…
MTGPがglobalﾒﾓﾘやﾃｸｽﾁｬを使ってるのなら、Cのrand()のように連続で呼び出して使えたら良いなと思ってるんですが無理ですかね？
529:デフォルトの名無しさん
10/01/23 02:00:42
>>528
おれはホスト側で使いたいので、make_uint32_randomの末尾でホスト側のバッファにコピーして使ってる。
デバイス側で使いたいなら、make_xxxxx_randomにd_buffをmain()側でとって渡し、
make_xxxx_random()内ではfreeせずにそのまま。すると次にmain()内でデバイス側関数を呼び出すときに
そのまま渡して使える。 ※スレッドが変わってはいかんので注意
530:デフォルトの名無しさん
10/01/23 05:30:39
CUDAを使って大量のデータのごく簡単な統計値（最小、最大、平均、標準偏差）を計算したいんですが、何かいいサンプルとかありますか？
531:デフォルトの名無しさん
10/01/23 08:51:31
CUBLASにあるんじゃ。
532:デフォルトの名無しさん
10/01/23 11:54:45
>>530
そういう、「最後が一個」な演算は、reductionで探すと有る気がする
並列性を有用に使うのがなかなか面倒らしい
533:デフォルトの名無しさん
10/01/23 12:09:54
OpenGL glutで時系列で変化する数値計算の結果を表示しようとしています。
GPU 2個で計算する場合、コールバック関数内でスレッドを生成することになると思うのですが、
この度（毎フレーム）にホストからグローバルメモリへデータ（時間による変化なし）を
コピーしなくてはいけません。
glutMainLoop(); 呼び出し前に両GPUにコピーした場合は、
コールバック関数内で生成したスレッドにグローバルメモリのポインタを渡しても参照できません。
データのコピーを一度だけですます方法はないでしょうか？ご教示ください。
スレッド生成は URLﾘﾝｸ(tech.ckme.co.jp) を参考にしました。
534:デフォルトの名無しさん
10/01/23 21:06:46
>530
それぐらいだったらたぶんCPUで実行した方が速いよ
535:デフォルトの名無しさん
10/01/24 00:53:03
>>533
GPU0用、GPU1用のスレッドをglutMainLoop()呼び出し前に生成するべきかと。
536:デフォルトの名無しさん
10/01/25 22:22:13
thread数 = N
Mat[N];
id = threadIdx.x;
if(id%2==0){
演算Ａ
Mat[id]にコピー
}else if(id%2!=0){
演算Ｂ
Mat[id]にコピー
}
のようなプログラムを組んでいるのですが、結果をみると最後の2つの要素と同じ計算結果が全体に入ってしまいます。
N=16なら
14の結果が0,2,4…12に
15の結果が1,3,5…13に入ってしまいます。
どこに問題があるのでしょうか
537:デフォルトの名無しさん
10/01/25 22:49:59
>>536
とりあえずC言語から始めようか
538:デフォルトの名無しさん
10/01/26 00:02:22
>>530
俺もこれ知りたい。
N社のSDKを見ても、画像処理のサンプルとかたくさんあっても、単純な総和計算とかないもんだね。やはり向いてないからか・。。。
ご丁寧にCのほうが早いとか教えてくださる人もいるがｗ
マルチパスの画像フィルターとか、
デバイスの中にある中間結果を作って、
その統計値を、次のパスの計算で使ったりするのが常套手段だから
ここでいったんホストのほうにコピーするとボトルネックになってしまう。
デバイスの中のデータの統計値を出すライブラリとか作ってくれると本にありがたいんだが
539:デフォルトの名無しさん
10/01/26 00:30:52
>>530, 538
SDKのreductionサンプルが参考になると思う。 確かpdfのわかりやすいスライドが一緒に入っているはず。
CUDAのアーキテクチャって、総和とかのreduction系演算は
不向きとまでは言わないまでもちょっと頭をひねらなきゃいけない分野だよね。
540:536
10/01/26 01:19:31
演算途中で計算領域が被ってただけでした。失礼しました。
541:デフォルトの名無しさん
10/01/26 01:27:25
>>539
そうだね。
沢山のデータがあるなら、最終はCPUにやらせても良さそうだけど、
転送がボトルネックになるなあ。
それだったらGPU内でやらした方が速いか。
あとI/O系があるとだめだよね。
ダブルバッファやトリプルバッファが使えるようなアプリならいいんだけど。
そうなると画像、映像系ばかりになってしまうなあ。
542:デフォルトの名無しさん
10/01/26 09:14:50
処理に割り当てるmultiprocessorの数を指定とか出来ませんか？
出来ればgridサイズ変えずに
そもそもOSもGPU使う場合割り当てがどうなってるのか分からないんですが
543:デフォルトの名無しさん
10/01/26 10:43:54
>>542
CUDAプログラムが実行中はOSはGPUを使えないので全てのSMを使用しても問題ありません。
544:デフォルトの名無しさん
10/01/26 10:47:22
5秒以上使ったらエラーが出たよ。
545:デフォルトの名無しさん
10/01/26 10:50:22
SMを1個しか使用していなくても5秒以上１つのCUDAカーネルが実行されるとタイムアウト。
546:542
10/01/26 11:15:31
へー、なるほど。
あとgridサイズ小さくする以外に使うSM数を制限できますか？
547:デフォルトの名無しさん
10/01/26 11:41:29
できませんし、その必要もないでしょう。
548:542
10/01/26 11:48:03
>>547
そうですか、ありがとうございます
並列計算の研究の一環でCUDA使ってるんで
並列数の変化でデータが取れたら嬉しいな、という理由です
549:デフォルトの名無しさん
10/01/26 14:26:47
>>548
素直に描画用とは別にGPUを用意しましょう。
総和を取る処理は私も書いているけど、仮想的に二次元にデータを配置して、縦に集計するところまでをGPUでやっている。
残りはCPU側に吐き出してからCPUではSSEを使って集計。
550:デフォルトの名無しさん
10/01/26 14:53:47
>>548
1番最初に実行が始まるCTAの配置は予測可能なので、
その中で使わないSMに割り当てられたCTAはコア内でダミーの演算を長時間繰り返すことでそのSMを占有し続ける。
こうすれば本来の計算は残りのSMでのみ行われるようになる。
通常時間計測できるのは全CTAが完了するまでの時間なので以下のどちらかを。
１）最後のCTAがdevice memoryにマップされたhost pinned memoryに何か書き込む。
２）ダミー演算の繰り返し回数を段々少なくしていき、計測される時間の減少率が変わるところを探す。
なお、全SMを使わないとメモリアクセス効率も落ちるのであまり面白くない。
551:デフォルトの名無しさん
10/01/26 15:20:53
> 仮想的に二次元にデータを配置して、縦に集計するところまで
なるほど。ふつくしい。ｎ次元でCPUでは最後の一次元だけやらせれば、対称になるな
552:542
10/01/26 21:31:38
>>549
別に用意してこの場合メリットってありますか？
あと総和ならReductionで組んだけど今回はGPU単体の話だったんで最後までGPUでやりました
デバイスホスト間の転送時間って馬鹿にならないイメージあるんですが、CPUにやらせたほうが速いんですかね？
まあ最後のちょっとだから誤差の範囲な気がしないでもないｗ
>>550
結構シビアですね、
直接的な方法が無ければgridサイズ縮めてIdx周り弄ろうと思います
↑の方法で弊害無いですよね？；
553:528
10/01/27 17:26:29
少し前にMTGPについて質問した者ですが、どうやら自分の要求とズレた感じでした。
thread 256、block 1でmake_uint32_random()するとバッファに768個のデータが生成されるが、でっていう…。
これはﾒﾙｾﾝﾇﾂｲｽﾀの高周期な乱数列の１部分って事で、本当にあの高周期な乱数を使いたいならその分のメモリが必要だということなのかな。
独自の乱数ジェネレータを作って、その768個の中から１つを取り出して…みたいな事をやり始めるとまた性質が変わってしまうし、本末転倒に。
結局、カーネルの中でシミュレーション目的の使用方法としては微妙だったので絶望。。。
スレッドID毎に使える線形合同法(遷移型)の乱数として使えるようになんとかできないものか…
554:デフォルトの名無しさん
10/01/28 00:49:57
え、おれ100万個単位で作って使えてるよ。
int num_data, にちゃんとでかい数与えてるかな
555:デフォルトの名無しさん
10/01/28 09:22:34
>>554
いや、num_dataの数を変えて生成される個数の事はあまり重要じゃないのよ…
問題はカーネルの中で使おうとした時、バッファに作成された乱数が並んでいる形態では微妙なのです。
例えば、100万個作ったとして256のスレ数で使うなら使用部分を分割する事になりますよね（thID==0は、バッファのindex0～約4000、という感じ）
いや、各スレッドは100万個のうち初めのindexだけseedとして決めて、あとは順次indexを増やして使っていく感じでもいいけど、
両者とも乱数列の周期はバッファのサイズに依存してしまいます。
一方、よくある{x=x*1103515245+12345; return x&2147483647;}このような方法は、アルゴリズムが優秀だとxが4byteなら最大で2^32の周期を持ちますよね。
今の状態のﾒﾙｾﾝﾇﾂｲｽﾀで2^32の周期を得ようとしたら、どんだけ大きいバッファが必要かっていう…
精度の良い乱数という点では利点ありますが、globalﾒﾓﾘを物凄く使う割にはなんだかなぁ…という複雑な気持ち。
MTGPの形態を知らなかった自分が悪いんだけど、要はこれ乱数生成を並列化して高速にしたもので、
実際にパラレルな計算で使う用にはなりえない事が分かりました。
自分の要求としては、GPU Gems3の「CUDAによる効率的な乱数の生成と応用」がﾁﾗｯと見た感じでは合致してたので、今からこっちを参考にしてみます。
長文ｽﾏｿ。>>554にはとても感謝している。こんな結末で申し訳ない。
556:デフォルトの名無しさん
10/01/28 09:38:33
＞乱数列の周期はバッファのサイズに依存してしまいます
横から失礼しますが、
for( ; ; )
{
make_uint32_random(, , d_buff); //デバイスメモリに作らせて残す
my_kernel_func<<<>>>(d_buff, , , , ); //MTGPで作った乱数を消費
}
こんな感じとして、256スレッドが一回に8192個の乱数を使う
→make_uint32_randomは2097152個の乱数を生成する
で良いのでは? make_uint32_random() は複数回呼び出すごとに
前回のmtgp32_params_fast_tの続きから処理するわけで、周期は
2^11213-1 でしょう。÷256しても2^11205とかだ
557:デフォルトの名無しさん
10/01/29 08:52:57
stupid questionですみませんが、VC++ 9.0で
1. .cppと.hのように、.cuファイル内でインクルードしているファイルが更新されたら.cuを再コンパイル対象にしたい。
2. ptxを出力したい。nvcc -ptxでは無理でした。
以上について教えてください。お願いします。
558:デフォルトの名無しさん
10/01/29 19:27:06
>>557
コマンドラインから nvcc -ptx *.cu とやっても駄目？
559:デフォルトの名無しさん
10/01/29 23:45:27
-keepでいいんじゃね
1.についてはSource Dependenciesを個別に設定すれば一応できる
560:デフォルトの名無しさん
10/01/30 15:52:51
>>558
XP Pro 32bitとVista Ultimate 64bitの両環境で、
コマンドラインからnvcc と打つとcl.exeがないと怒られます。
プロジェクトのプロパティを参考にパスとincludeを指定してやっても
エラーは控えていませんがコンパイルできません。
VC使ってる人はどうしてるんでしょう？
561:デフォルトの名無しさん
10/01/30 15:56:22
-ccbin で指定しても駄目ですか？
562:デフォルトの名無しさん
10/01/31 02:17:49
もちろんVSコマンドプロンプトから打ってるよな
563:デフォルトの名無しさん
10/01/31 16:27:27
>>562
ふつうのコマンドプロンプトを使っていました。
どうもお騒がせしました。
564:デフォルトの名無しさん
10/01/31 16:30:07
普通にC++のコードを書けるようになってからじゃないと、学習効率が悪くてどうしようもないぞ。
565:デフォルトの名無しさん
10/01/31 18:38:21
>>559
俺は556じゃないが-keep知らなかった。ありがとう。
566:デフォルトの名無しさん
10/02/01 21:17:18
GPU Challenge 2010
URLﾘﾝｸ(www.hpcc.jp)
自由課題もあるそうな
567:デフォルトの名無しさん
10/02/02 11:27:25
まあ俺は学生だから規定課題でるけどな
しかしCellとかに比べてあんまり最適化するとこないな
568:デフォルトの名無しさん
10/02/02 18:26:32
どのくらい参加するんだろう？
俺もとりあえずエントリーしようかな。
569:デフォルトの名無しさん
10/02/02 22:14:16
自由課題の方で、パターン形成する発展方程式とかの数値計算すると、絵的にも面白そうなの出来そうじゃない？
570:デフォルトの名無しさん
10/02/10 10:54:42
jpegをデコードするライブラリもしくはCUDAのコードはどこかにありませんか？
571:デフォルトの名無しさん
10/02/10 11:54:41
GPU Challengeの課題が増えた
メルセンヌツイスタと言われるとHack the Cellを思い出すな
572:デフォルトの名無しさん
10/02/10 12:30:21
SLI環境で、プログラムを2つ実行した場合、それぞれ別のGPUを
利用する方法を教えてくれ
573:デフォルトの名無しさん
10/02/10 14:52:21
> SLI環境で
えっ、ひとつにしか見えないんじゃないの?? 出来るの？
574:デフォルトの名無しさん
10/02/11 00:45:35
>>572
cudaSetDevice()でそれぞれ0と1を指定する。
575:デフォルトの名無しさん
10/02/11 13:21:15
>>573
東工大青木先生はGeForce4つ並べてた
OpenMPで並列化していたと思う
576:デフォルトの名無しさん
10/02/11 13:39:01
>>575
CUDA版OpenMPてあるの？
577:デフォルトの名無しさん
10/02/11 14:25:56
CUDA版と言えるOpenMPはない
OpenMPのスレッド指定とCUDAのdevice指定を組み合わせただけのこと
578:デフォルトの名無しさん
10/02/11 14:53:48
そういうことか、今度挑戦してみようかな
579:デフォルトの名無しさん
10/02/12 00:56:59
> GeForce4つ並べ
と
> SLI
は違うんじゃね？ ケーブルで繋ぐのがSLI･･･かな
580:デフォルトの名無しさん
10/02/12 10:07:29
OpenMPを使えば複数のGPUを簡単に使えるのですか？
やりかたをおしえてくれろ
581:デフォルトの名無しさん
10/02/12 11:11:35
CUDAで顧客向けプログラムを作成しています。
CUDAプログラムの配布先には本体以外に何を配布すればよろしいのでしょうか？
582:デフォルトの名無しさん
10/02/12 11:46:23
>>581
GeForce人数分
583:デフォルトの名無しさん
10/02/12 11:49:25
TeslaかQuadroを配っておけば良いと思うよ
584:デフォルトの名無しさん
10/02/12 11:56:09
cutil使わなければcudart.dllだけでよろしよ
585:デフォルトの名無しさん
10/02/12 17:57:29
>>579
内部でケーブルでつながってても、
デバイスメモリが共有されるわけじゃないから
CUDA的には関係ない。
586:デフォルトの名無しさん
10/02/13 04:48:44
>>580
SDKにサンプルがある
587:デフォルトの名無しさん
10/02/14 12:30:31
CUDA FORTRANのセミナーが青木先生のとこで開催されるらしいが、おまいら行く？
588:デフォルトの名無しさん
10/02/14 12:42:11
>>587
青木先生か、Cだったら行くんだけどな
589:デフォルトの名無しさん
10/02/15 17:44:43
デバイスエミュレーション時の速度って、実際のCPUとの目安で考えたら
どのくらいスケールして考えればいいですか？
初めてエミュレーションモードを使ってみたんですが、3000倍以上の差が付いて明らかにおかしいと思うんです…
CPU: Core i7 Q720@1.6GHz、 GPU:Tesla C1060
590:デフォルトの名無しさん
10/02/15 17:53:40
Q720って720QMのこと？
ノートPCにTeslaが搭載されているとか、聞いたことがないんだけど
591:589
10/02/15 17:56:06
追記。
grid(2,1,1)、block(256,1,1)でのカーネルで、3000倍になります。
これからgridを増やすと、さらに差が広がっていきます。
カーネルで実装した内容を、CPU版で実装したくないけど速度比較はしたい。
･･･無理でつか？
592:589
10/02/15 18:06:39
>>590
Teslaは別のデスクトップPCので、エミュを動かしたのはノートPCでの方です。
紛らわしくて申し訳ない。
どちらもPCもCUDA使えるんですが、CPU自体はノートの方が性能良かったのでこちらを使いました。
ノートPCのGPU： GeForce GTX 260M
593:デフォルトの名無しさん
10/02/15 18:14:45
そもそも、エミューレーションモードって非CUDA環境でも
CUDAカーネルのデバッグが出来るようにしたものでしょ
あくまでテスト用のもの
594:デフォルトの名無しさん
10/02/15 18:15:25
誤)エミューレーションモード
訂)エミュレーションモード
595:デフォルトの名無しさん
10/02/15 19:10:55
3000倍？そんなもんでしょ
596:デフォルトの名無しさん
10/02/15 19:28:11
効率の良いブロック分けの仕方？がわからず困っています．
実行時にＮ個のデータ系列が与えられて，
それぞれの系列へ施す処理内容自体は同一なのですが，
その処理に用いるパラメタ値が異なります．
例えばＮ＝３の場合，パラメタもp[0]～p[2]の３個があって，
データ系列0の全データ { D[0,0], D[0,1], D[0,2], ..., D[0,m0] } にはp[0]を加算，
データ系列1の全データ { D[1,0], .... , D[1,m1] } にはp[1]を加算，
データ系列2の全データ { D[2,0], .... , D[2,m2] } にはp[2]を加算
という具合です．
全系列のデータ数が同じ(m0=m1=m2)ならば
グリッドの次元の１方向を系列(0～N-1)に割り当てれば良いかと思うのですが，
系列毎にデータ個数がかなり異なる場合はどうすればいいのでしょうか？
データ個数は系列ごとに少なくとも数千個くらいになります．
同じような割り振り方だと何もしないブロックが大量にできてしまいそうです．
597:デフォルトの名無しさん
10/02/15 20:36:27
CにD[0,0]というものはないのでよくわからないけど、
いったん長い配列にまとめて処理して、あとでCPUでばらせばいいのでは。
D[i,j]のjについて、自分はどのpに属するのか覚えさせて。
598:デフォルトの名無しさん
10/02/15 23:36:05
>>592
Nvidia Nexus使えば？
599:デフォルトの名無しさん
10/02/15 23:37:13
追記
Nexus使うとネットワーク経由で、
コード書く用のPCとデバッグするPCを分けられるよ、ってことね。
600:デフォルトの名無しさん
10/02/16 14:48:16
>>598
うーん、デバッグというよりは単にCPUとGPUで速度比較をしたいだけなんです。
うまく並列化して普通は、1～50倍くらいの成果になると思うんですが･･･
目安でいいからエミュレーションモードから大体の速度が分からないのかなと。
601:デフォルトの名無しさん
10/02/16 23:24:57
いまいち意味がわかんないけどCPUコードとGPUコードをデバイスエミュで実行したらCPUコードのが3000倍早いって事？
それだったらそんなもんかと。デバイスエミュは重いし。
違うんだったらごめん。
CPUとGPUで速度比較したいなら普通にCPUとGPUそれぞれ向けのコード書いて実行したらいいんじゃない？？
602:デフォルトの名無しさん
10/02/17 00:50:55
>>601
あ、あれ？ 自分のデバイスエミュの認識自体が間違ってたかな…？
言いたかったのは、実行するハードの方での両者の比較です。
CPUコードと言うのはありません。
カーネルや、その内部で呼ぶ__device__の関数らがGPUコードだよね？それを普通に「GPU」が実行した時の速度と、
デバイスエミュを使ってCPUが実行した場合（内部では逐次計算？）の速度では、普通に「GPU」で実行した方が3000倍速いということです。
>>591の通り、GPUコードが多くなりすぎて、同じ事をするCPUコードを実装するのが面倒なのです。
デバイスエミュはCPUが実行するとの事で淡い期待を抱いてましたが無理そうな感じなんですな…
603:デフォルトの名無しさん
10/02/17 01:32:32
>>602
比較する目的はなんでしょうか？
研究目的であれば面倒であろうがCPU用も実装しなければなりません。
そうでないなら比較なぞしなくてよいのではないかと。
ちなみにGPU：エミュが3000：１程度であればCPU用に普通に実装した方がGPUより速い可能性が十分にあります。
604:デフォルトの名無しさん
10/02/17 15:10:51
面倒でも計算結果の比較しろよｗ
nvidiaのサンプルコードでも結果の比較してるだろ
605:デフォルトの名無しさん
10/02/17 15:41:47
>>602
emulation modeは，名前の通りGPUでの動作を模擬しているだけで，
その計算速度とGPUの計算速度を比較することに意味はない．
emulation modeがある理由は，カーネル内にprintfとかのコードを書いて
debugしたり出来るようにするため．
従って，CPUとGPUの計算速度の比較を行いたいなら，それぞれに最適化した
２つのコードを書いて比較する必要がある．
関係ないけど，
CPUとの比較しているときにCPUのコア一つとの比較が多い気がするけど，
それフェアな比較じゃないよね．せめてOpenMPくらい使おうよと．
まぁ使うとGPUの存在感が薄れるのだけれども…．
606:デフォルトの名無しさん
10/02/17 17:32:51
>>603
研究なんて言えないようなものです。目的としては自己満足になりますね。
ただ、目安程度であれど比較できないとGPUとCUDAを使う意義に関わってきます。
早ければSUGEEEE!!ってなって満足し、遅ければそこできっぱり止めるという選択ができる。
そして3000：1ならまだCPUの方が早そうだというのは参考になりました。
>>605
おっしゃる通りですが、厳密に比較するまでは求めていないんです。
今自分がやってることは無意味なのか？（つまりCPUの方が普通に早い）が分かればいいんです。
grid(2,1,1)でフェアじゃないのは、ｺｰﾃﾞｨﾝｸﾞが糞なので2以上じゃないとGPUで動作しないんです('A`)･･･(ﾒﾓﾘ周りの関係で)
我侭な要求でスマン。
607:デフォルトの名無しさん
10/02/17 18:07:50
逆に考えるんだ
CPU側で動作をきっちり検証したプログラムを、
GPUに移植して、速度を比べる。
GPUに本当に適した問題なら、数十倍出る場合もあるし。
608:デフォルトの名無しさん
10/02/17 18:47:11
>>605
GPUの優位性をうたうような文脈で
比較対象のCPUコードが1スレッドだったら
それを評価する人間は目が余程の節穴でない限り、
CPUのコア数倍した結果と比べるでしょ。
物凄く差がある場合はそれでも十分優位性をうたえるから。
609:デフォルトの名無しさん
10/02/17 21:05:25
>>605
１コアの40倍とあれば4コアの10倍と読み替えればいい訳で。
610:デフォルトの名無しさん
10/02/17 22:05:37
４コアと比較したらどうなるかということではなく、
１コアと比較してる人がCPU版もまともにチューニングしてるとは思えない、ということかと。
まぁ、皆GPU版ばっかりチューニングしてますから。
611:デフォルトの名無しさん
10/02/18 06:16:20
たぶんCPU版はSSEすら使っていないんだろうね。
メモリ帯域がものを言うコードでなくて、CPUがNehalemだったら、
安いGPUなんかじゃ全く優位ないからね。
612:デフォルトの名無しさん
10/02/19 02:31:33
俺はNvidiaちゃんを信じるよ
URLﾘﾝｸ(twitter.com)
613:デフォルトの名無しさん
10/02/19 02:32:33
誤爆った／(^o^)＼
614:デフォルトの名無しさん
10/02/20 04:09:26
うちは理論で「***手法より*%高速化して最速!」とかやってないってのもあるけど
GPUで組んだ手法と既存の手法を比べる場合、既存のほうはベーシックにしろと指導された。
複数CPUだとかSSEを使ってガチガチに最適化した手法と比べちゃうと基準が分からなくなるからだと。
他の高速化との差を知りたければその論文と比較しろということだと思う。
CPU最適化して無いなんて糞というのも分かるけど、こういうところもあるということで。
615:デフォルトの名無しさん
10/02/20 08:30:23
コードの比較もいろいろだよな。
同じアルゴリズムを採用しても、CPUでも書き手によってGPUでも明らかに差が出てくる。
でもGPUを使う場合、多くの場合はCPUよりも速くなりました。というのが目的な訳で、
CPUの方が速いならあえてGPUを使う必要はないからね。
基準が曖昧になるのもわかるけど、そもそも基準が曖昧な気がするなあ。
場合によってはかなり恣意的になることもあるし・・・・。
616:デフォルトの名無しさん
10/02/20 10:07:52
Femiやばいまた延期確定かも
617:デフォルトの名無しさん
10/02/20 11:37:52
一般人が入手できるのは１年後になる可能性もあるらしいね
618:デフォルトの名無しさん
10/02/20 11:44:04
なんでそんな度々延期になるの
619:デフォルトの名無しさん
10/02/20 12:14:45
>>618
ペーパーロンチで実際開発が
行われていないからだよ
620:デフォルトの名無しさん
10/02/20 13:17:36
今回のケースは大きな欠陥があることを知りながら、小手先の改良でなんとかしようとして
「完成品」を大量生産をして、まとにチップが取れなかったのが原因だろ
１％程度とされる歩留まり率で、１チップ当たり5000ドルの原価
これでは商売にならないね
621:デフォルトの名無しさん
10/02/20 13:39:05
3/19に東工大青木先生がCUDA Fortranのセミナやるんだって
622:デフォルトの名無しさん
10/02/20 13:40:57
関係者の宣伝おつ
623:デフォルトの名無しさん
10/02/23 08:21:59
青木先生に集客されたくねえなあ正直
624:デフォルトの名無しさん
10/02/23 22:06:44
nexusをリモートで動かそうとしたが、ブレークポイントでとまらねぇ・・・。
色々試したがどうにも解決しないので教えてください。
状況としては、nexusのユーザーズガイドに沿って設定。
ためしにnexusサンプル動かそうとしたら、
ランタイムAPIプロジェクトはGetDeviceで引数に０が。
で、次の行で落ちる。
DriveAPIは落ちない。んでホスト側の画面右下に青いポップアップ出て、
ターゲットマシンにコンソール画面出てるのでプログラムは正常に動いてるっぽい。
でもカーネル関数内にブレークポイント置いても止まらず。
自分でSDKサンプルのプロジェクトの設定変えて試しても同じ。
マシン環境はこんな感じ。
ホストマシン
Vista 64bit SP2
.Net3.5 SP1
Host nexus1.0(jan 64bit)
GPU FX570
VC++ 2008 SP1
DirectX10 August 09
CUDA SDK2.3 32bit
CUDA ToolKit2.3 32bit