07/09/17 14:54:28
本家
URLﾘﾝｸ(developer.nvidia.com)
2:デフォルトの名無しさん
07/09/17 14:55:24
グプグプー
3:デフォルトの名無しさん
07/09/17 21:03:53
>>1
GPGPUスレと統合じゃだめなんか？
ｽﾚﾘﾝｸ(tech板)
4:デフォルトの名無しさん
07/09/18 10:06:32
おお、いつの間にこんなスレが。一応あっちで回答しているけどどうしようか。
5:デフォルトの名無しさん
07/09/19 14:34:55
削除依頼出しとけ
6:デフォルトの名無しさん
07/09/21 05:49:47
WindowsのCUDAとLinuxのCUDAで書式に何か違いはありますか？
7:デフォルトの名無しさん
07/09/21 12:09:20
削除依頼出しとけ
8:デフォルトの名無しさん
07/09/21 15:15:23
三角関数や指数関数はないの？
9:デフォルトの名無しさん
07/09/21 15:25:30
削除依頼出しとけ
10:4
07/09/21 15:59:09
>>6
基本的に同一。但し、Win用nvcc(コンパイラ)では64bit版にはできないこととオブジェクトモジュールがMS仕様なことに注意。
更に、nvccはgccで実装されているので、C++のnewやvectorはそのままではリンクできなくなってしまうことにも要注意。
>>8
math.hにある、float系の関数は大抵使えると思っていい。エミュレーションモードなら全部使えた筈。
また、sin(), cos()などの一部の関数は組み込み命令(intrinsic)版が使えるので更に高速。
実際の例は、GPGPUスレ参照で。
ｽﾚﾘﾝｸ(tech板:71-72番)
11:デフォルトの名無しさん
07/09/21 16:19:38
CUDAとかって何らかの標準化の動きはあるの？
12:デフォルトの名無しさん
07/09/22 14:43:16
笛吹けど踊らず
# まぁ、笛吹いている会社の中でも踊らない連中がいるみたいだし
13:デフォルトの名無しさん
07/09/29 08:08:41
CUDAは無料ですか？？
14:デフォルトの名無しさん
07/09/29 09:38:17
ソフトは無料、ハードは有料なんじゃね
15:デフォルトの名無しさん
07/09/29 22:08:38
CUDAなんてハードはありません。
16:デフォルトの名無しさん
07/09/29 22:10:13
処で、テスラc870ってQuadroFXの一番上の奴とどこが違うんだ?
17:デフォルトの名無しさん
07/09/29 23:03:54
>>15
おマイ馬鹿か、環境の事に決まってるだろ
18:デフォルトの名無しさん
07/09/30 02:33:09
Compute Unified Device Architecture
~~~~~~~
19:デフォルトの名無しさん
07/09/30 14:58:40
CompUte Driver Architecture
20:4
07/10/01 12:05:40
ちょっとメモ。
device側メモリは共有メモリであろうとグローバルメモリであろうと、
並列動作で一箇所に書き込もうとすると(当然の如く)破綻する。
従って、ヒストグラム作成のようなロジックはそのままでは実装できない。
生データをソートしてから集計するロジックを書いてみたが、さてこいつをどう並列化するか。
--
__global__ void devSum(int begin, int end, int idx, float2 * tmp)
{
int ic = begin;
float2 sum = make_float2(tmp[ic].x, tmp[ic].y);
for (++ic; ic < end; ++ic) {
sum.x += tmp[ic].x;
sum.y += tmp[ic].y;
}
ar[idx] = sum.x;
ai[idx] = sum.y;
}
--
単純に考えると開始位置を検索しなければならないが、それ自体にコストが掛かりそうだ。
21:デフォルトの名無しさん
07/10/12 18:56:43
>>20みたいな処理って数値演算よりデータ転送の比重が多すぎて、
並列化させる旨みは非常に少ないと思うんだけどどうなのよ。
本家のフォーラムでいわゆるセマフォを実装して集計処理やらせようとしてる人がいたけど、
そういう問題をCUDAでやる事自体がそもそも間違っとる、みたいに言われてた。
かく言う自分も似たような事しようとしてたけどパフォーマンス出ないんで諦めて、
問題の分割方法から見直したら上手くいったりした。
22:デフォルトの名無しさん
07/10/12 18:57:14
>>20みたいな処理って数値演算よりデータ転送の比重が多すぎて、
並列化させる旨みは非常に少ないと思うんだけどどうなのよ。
本家のフォーラムでいわゆるセマフォを実装して集計処理やらせようとしてる人がいたけど、
そういう問題をCUDAでやる事自体がそもそも間違っとる、みたいに言われてた。
かく言う自分も似たような事しようとしてたけどパフォーマンス出ないんで諦めて、
問題の分割方法から見直したら上手くいったりした。
23:デフォルトの名無しさん
07/10/12 18:58:08
二重orz
24:4
07/10/12 19:29:53
>>21
>20の処理は、最初に定数テーブルを仕込んでおいて幾つかパラメータを渡して計算させるのがメイン。
>20に書いているのは最後の集計部なんだけど、結果はarとaiに入るのでそれを転送するだけ。
集計をGPUにやらせられないとなるとtmpを転送しなければならないので、それをなんとかしたかったわけ。
で、未だその後修正掛けていないけど、begin, endのペアをidxの値の分だけを定数テーブルに
仕込んで置けることが判ったからなんとかなりそう。そうすればidxの値の分だけ並列に走らせられることだし。
25:デフォルトの名無しさん
07/10/15 12:09:23
>>24
状況はわかったけど最後の方で言ってる意味が良くわからん…俺の頭がヘタレなのか。
今のとこ、並列で集計処理するのに一番早そうかなって思ってるのはデータを半分ずつ加算してく方法。
例えば処理したいデータがnコなら、n/2コのスレッドで2コずつデータ加算。後はリカーシブにやる。
(めんどいからとりあえずnは偶数としてね)
これならlog_2 nのオーダーで集計出来んじゃね？ってかんじ。
実装したことないがスレッド数が変わるのが厄介かな…今は
if(tx==0) {
for(i=0;i<BLOCKSIZE;i++)
sum+=data[i];
}
みたいに横着してる。他で十分高速化してるんでとりあえず後回し中。
26:デフォルトの名無しさん
07/10/16 23:58:44
GPUでテクスチャの全てのピクセル値の合計を求めたいってこと？
1x1サイズのミップマップ作って、ピクセル数を掛けるのはどう？
速いかどうかは知らんけど。
27:デフォルトの名無しさん
07/10/17 06:58:08
opensuse 10.3でCUDAは動く？
28:デフォルトの名無しさん
07/10/19 10:01:25
Linuxサポートしてるから動くんじゃね？
29:デフォルトの名無しさん
07/10/22 10:32:10
>>27
ドライバが載らなくても取り敢えずエミュレータは動くでしょうね。
つーか、本家に10.1用と10.2用は置いてあるから試してみては?
URLﾘﾝｸ(developer.nvidia.com)
30:デフォルトの名無しさん
07/11/20 15:20:50
>>27
もしかしたら、標準でインストールされるgccのバージョンが違う所為で動かないかも知れない。
その場合、10.2にインストールされているバージョンのgccをインストールしておけば大丈夫。
# 要はnvccが内部でgccを使っているから、gccのバージョンに依存していると。
# ドライバはOKだった筈。
31:デフォルトの名無しさん
07/11/25 09:41:58
GeForce 8400GSの安いカードでもCUDAを試すことはできますか？
32:デフォルトの名無しさん
07/11/25 12:05:37
>>31
ドキュメントには書かれていないので微妙。
時期的には、CUDA1.0に間に合うタイミングで出たと思うけど……
# 1.0のドキュメントに書かれているのはこれら。
--
GeForce 8800 Ultra
GeForce 8800 GTX
GeForce 8800 GTS
GeForce 8600 GTS
GeForce 8600 GT
GeForce 8500 GT
Quadro FX 5600
Quadro FX 4600
Tesla C870
Tesla D870
Tesla S870
--
# 1.1betaのリリースノートには8800GTに対応とある。
33:デフォルトの名無しさん
07/11/25 18:24:52
>>31
8800を買うのが無難
34:デフォルトの名無しさん
07/11/25 20:07:13
安く済ませるなら、(消費電力の少ない)8600GTもお勧め。
第2世代(と思われる)8800GTは割とパフォーマンスがよさそうだけどね。
35:デフォルトの名無しさん
07/11/25 20:07:50
>>31
公式にはサポートされていない。
でも、ノートパソコンで動いたという人もいるくらいだから、動く可能性がないわけではないと思う
36:デフォルトの名無しさん
07/11/28 01:39:57
8400Mは対応してたはずだけど、普通の8400の方は対応していないんだよね
自作板でサンプルが動いたという報告があったような気みするけど、
対応していない製品を使うのは心配だよね
37:デフォルトの名無しさん
07/11/29 23:46:49
8600買ってきたー
試したいのだけど、まず取っ掛かりはどこからいけばいいですか＞＜
38:デフォルトの名無しさん
07/11/30 00:20:36
>>37
>29にある本家から一式拾って、ドライバとツールキットとSDKを入れて、
サンプル一式をビルドして動かしてみるよろし。
39:デフォルトの名無しさん
07/12/05 18:31:17
64bit 版 ubuntu7.10にインストールしたいのですが、どれをダウンロードすればよいのでしょうか？
40:デフォルトの名無しさん
07/12/05 18:48:07
URLﾘﾝｸ(developer.nvidia.com)かURLﾘﾝｸ(developer.nvidia.com)から
x86-64のDriverとx86-64のToolKitとSDKを拾えばいいとして。
ToolKitはどれが使えるかは不明。nvccがgccを使うから、インストールの状況によるので単純じゃない模様。
41:デフォルトの名無しさん
07/12/07 08:32:31
ubuntu版きてるーーーー
42:デフォルトの名無しさん
07/12/07 12:33:05
7.04ってなんだYO
43:デフォルトの名無しさん
07/12/07 13:31:11
タイムリーなんだか間が悪いんだかw
44:デフォルトの名無しさん
07/12/08 02:28:34
NECのVALUESTARを使ってるんですが
nVIDIAをインストールしようとすると
NVIDIA setupプログラムは、現在のハードウェアと互換性のあるドライバ
を見つけることができませんでした。
と表示され、インストールができません。
何が原因なんでしょう？
OSはVistaです。
45:デフォルトの名無しさん
07/12/08 08:33:55
>>44
NECのVALUESTARなんてどうでもいい情報よりもグラボの情報だせよん
46:デフォルトの名無しさん
07/12/08 12:14:51
なぜココで聞く。
ハードウェアスレに聞け
47:デフォルトの名無しさん
07/12/10 13:04:06
>>44
そもそもVista対応のCUDA対応ドライバがありません。
48:デフォルトの名無しさん
07/12/13 00:22:06
>>44
>nVIDIAをインストールしようとすると
って時点で意味がわからんのだが、
とりあえずデバイスエミュレータ(-deviceemu)でどうよ？
49:sage
07/12/15 19:52:50
2次元配列を使用したいのですが
以下のサンプルコードを参考にしてみましたがうまく使えません
よろしければご教示お願いします
/* 2D array test */
#include<stdio.h>
#include<cutil.h>
#define width 4
#define height 4
// device code
__global__ void myKernel(float* devPtr, int pitch)
{
for (int r = 0; r < height; ++r) {
float* row = (float*)((char*)devPtr + r * pitch);
for (int c = 0; c < width; ++c) {
float element = row[c];
}
}
devPtr[1][1]=0.0;
}
int main(){
float* devPtr,an;
size_t pitch;
cudaMallocPitch((void**)&devPtr, &pitch,
width * sizeof(float), height);
myKernel<<<100, 512>>>(devPtr, pitch);
CUDA_SAFE_CALL(cudaMemcpy2D((void**)an,pitch,(void**)&devPtr,pitch,4*sizeof(float),4,cudaMemcpyDeviceToHost));
50:49
07/12/17 15:07:07
自己解決しました！
51:49
07/12/17 16:13:39
勘違いでした
誰かヘルプ。。。
52:デフォルトの名無しさん
07/12/17 18:50:43
何がどう巧くいってないと判断したのかよく分かりませんが。
cudaMallocPitch(), cudaMemcpy2D()じゃなくて、cudaMalloc(), cudaMemcpy()では巧くいったの?
それと、テストしたソースは丸ごとコピペするか、どっかにアプロドして欲しい。
53:デフォルトの名無しさん
07/12/18 11:25:32
＞アプロド
なんか可愛いw
54:デフォルトの名無しさん
07/12/18 12:07:39
>>49
＞devPtr[1][1]=0.0;
float * devPtrなんだから、エラーになる。
＞CUDA_SAFE_CALL(cudaMemcpy2D((void**)an,pitch,(void**)&devPtr,pitch,4*sizeof(float),4,cudaMemcpyDeviceToHost));
cudaMemcpy2D()で転送する先のanが初期化されていないから絶対にまともにコピーできない。
サンプルをよっぽど誤読しているのか、Cの経験が全く足りていないのか、いずれにしても阿呆過ぎる。
55:デフォルトの名無しさん
07/12/18 12:14:10
>>49
＞float* devPtr,an;
このanは只のfloat型。
cudaMemcpy2D()の第1,3パラメータのキャストもおかしい。
もしサンプルにそう書かれていたのなら、そのサンプルがタコ。
56:デフォルトの名無しさん
07/12/19 13:38:10
処で、サイトが更新されて、CUDA ZONEって目立つようになったね。
1.1は未だbetaなのか正式リリースなのか、中味は変わっていないみたいだけど。
ついでに、使えるチップの一覧も更新されているので転記。
--
GeForce Tesla Quadro
8800 Ultra C870 FX 5600
8800 GTX D870 FX 4600
8800 GTS S870 FX 1700
8800 GT FX 570
8600 GTS FX 370
8600 GT NVS 290
8500 GT FX 1600M
8400 GS FX 570M
8700M GT FX 360M
8600M GT Quadro Plex 1000 Model IV
8600M GS Quadro Plex 1000 Model S4
8400M GT NVS 320M
8400M GS NVS 140M
8400M G NVS 135M
NVS 130M
57:56
07/12/19 13:56:13
表をそのまま貼ったから見難くなったのは勘弁。
補足しておくと、FX/Plex/NVSはどれもQuadroシリーズ(非OEM)ね。
細かく見てたら、8800GTが追加された他に8800GTSに512MBが追加されている。
しかも、クロックが既存の8800GTSよりも速いだけでなくShaderClockでは最速に。
MemoryInterfaceが512bitだから、8800GTSというより8800GTの系列かしら。
まぁこれで、旧8800GTS(640MB/320MB)は消える方向なんでしょうね。
58:デフォルトの名無しさん
07/12/22 15:09:32
ほほー
廉価版でCUDA試してみたいだけの奴は8400あたりでもOKということか。
ありがたやありがたや
59:デフォルトの名無しさん
07/12/31 18:53:34
CUDAを使用できるGPUのうち、
GeForceとQUADROのどちらの方がCUDAに最適なのでしょうか？
60:デフォルトの名無しさん
08/01/03 00:07:49
>>59
「最適」の基準によります。
入手性、価格、バリエーションではGeForceシリーズに文句なしで軍配が上がりますが、
個別の性能はそれぞれ異なりますので。
因みに、TeslaはQuadroFX5600(-アナログ回路)相当だそうです。
あー、QuadroPlexを買える待遇なら、違う意味で最適かもしれません。
61:デフォルトの名無しさん
08/01/07 12:11:34
処で、8800GTXでどこまで速度が出るのか実測してみた。
単純な実数の足し算をループで回しただけなんだが、約43GFlops出た。
一方、ShaderClockは1350MHz、プロセッサ数は128、ということなので計算上は(1クロック1演算なら)172.8GFLOPSとなる。
命令セット上は、積和が1クロックで実行できるから2倍して約350GFLOPSだからカタログスペックはこの値なのだろう。
実際には、ループを構成するには実数の足し算の他に整数加算、整数比較、条件ジャンプの3クロックが必要になる。
従って、約43GFLOPSとなって実測値と見事に一致した。
というわけで、(メモリアクセスが無視できるなら)アセンブリ(ptx)出力から所要時間を割り出すことができる模様。
問題は、メモリアクセスが絡むとどれだけ遅くなるかで、こちらの方は今後調査の予定。
62:デフォルトの名無しさん
08/01/07 21:38:50
ニーモニックて公開されてるんだっけ？ or 固定長？
63:デフォルトの名無しさん
08/01/07 23:24:57
5秒制限とかあるらしいけど
これどうやって回避するんだろう
64:デフォルトの名無しさん
08/01/07 23:36:08
>>62
ptx出力を睨めっこすれば大体当たりがつく。どっかで公開されてるかも知れんが。
>>63
5秒も計算させたまま帰ってこないような組み方すると、どうせ遅くなるから余り気にならない。
単純なループだけならXeonの方が速いからね。
65:デフォルトの名無しさん
08/01/10 15:13:03
>>61
おつー
これは目安に使えそうだね
66:デフォルトの名無しさん
08/01/11 11:30:09
URLﾘﾝｸ(www.nvidia.com)
CUDAのサイトリニューアルされてるね。
相変わらずfor Vistaはまだだけど。
67:デフォルトの名無しさん
08/01/11 14:40:01
>>66
一応>56で指摘済み。
処で、cufftを使ってみた。
SDKのサンプル(convolutionFFT2D)では100MPix/sの処理速度があるような実行結果が得られたけれど、
実際に試してみたら4096x4096の大きな画像を使っても10MPix/sしか出ない。
どうも、オーバヘッドが大きくて速度が出にくいみたいだ。
まぁ、サンプルと違ってテストしたコードは「プラン作成」「メモリ確保」「メモリ転送」「FFT実行」
「メモリ転送」｢メモリ破棄」「プラン破棄」を全部実行しているからだとは思うけど。
と言うことで、来週辺りは実際のプログラムに組み込んだ形でテストしてみる羽目になりそうだ。
そうそう、8800GTXをCUDA1.0で動かした場合と8800GTをCUDA1.1で動かした場合で
convolutionFFT2Dの所要時間が殆ど変わらなかった。GPUの性能差を埋める程度にはCUDA1.1で改善されたのかな?
68:66
08/01/11 15:23:14
ずっと旧トップページをブクマしてて今日やっと転送されるようになったので
気づいてなかったｽﾏﾝｺ
1.0になってから遊んでいなかったので久しぶりに遊んでみよう。
69:デフォルトの名無しさん
08/01/11 22:18:42
>>67
メモリ転送は、PCI-Expressの性能で上限があるはず。
たぶんこれがかなり時間食ってると思う。
4K×4kは画像なんでしょうか。大きいですね。
70:デフォルトの名無しさん
08/01/12 00:56:15
PCIe 1.1と2.0でどれぐらい差が出るのか気になりますね。
3DMarkのベンチマークやゲームではほとんど差が無いみたいですが
CUDAだとやはりそのあたりがボトルネックになる場合も出てくるでしょうかね。
またメモリ容量に依存する場合もやはり多いのでしょうか？
71:67
08/01/12 06:35:30
>>69
FFT解析絡みのことをやっているので、周波数空間像を拡げてから逆FFTなんてことをしばしば。
まぁ、4096x4096は単にテストするのが目的でしたが。
>>70
未だPCIe2を試せる環境がないからなんとも言えませんが、演算の種類によっては影響出るでしょうね。
私のところのプロジェクトでは寧ろ、GPUとCPUを巧く並行動作できるかどうかの方が影響大きいのですが。
72:デフォルトの名無しさん
08/01/18 11:50:14
むぅ、周波数合成みたいなロジックをCUDAに移植したんだが、余り速度が出ない。
元のソースが三角関数テーブルを使っていたのでメモリアクセスが足を引っ張っているのかと思って
__sincosf()を使ってみたが、今度は何故か実行ごとの所要時間のばらつきが大きい。
速いときはXeon並の速度になるのに遅いときは3倍くらい時間が掛かる。
はてさて、超越関数ユニットの数が少ないのだろうか……
73:デフォルトの名無しさん
08/01/18 15:08:32
超越関数は級数近似と思われ
see /cuda/include/math_functions.h
74:デフォルトの名無しさん
08/01/18 17:05:57
いえいえ、ちゃんとニモニックもあるのですよ。そっちを使うとドキュメントに拠れば32クロックだそうで。
但し、そいつが「To issue one instruction for a warp,」とあるからぶつかるのかなと。
75:デフォルトの名無しさん
08/01/21 00:26:07
8800前後の機種の日本での売れ行きがかなり芳しくないようだがこれは何を意味するのだろう．本国でのアメリカでは使用例（論文）が着実に出つつあるから，なんだかそいつらの為だけに開発されてる感じじゃない？？
76:デフォルトの名無しさん
08/01/22 23:13:36
8800系って、日本ではNEETゲーマーが買い漁るから
1$=200円くらいのボッタ栗値段が付いてしまって、
貧乏ラボじゃ買えん罠。箱もヲタ臭くてボスの目
が気になる。
77:デフォルトの名無しさん
08/01/22 23:24:09
つ [玄人志向]
78:75
08/01/23 00:48:46
そうか，研究者の間であまり大量に出回っていないと思ったが，そういう事だったのか．＞76
国内で一番CUDA使いこなしているのはどこのチームだろう？
79:デフォルトの名無しさん
08/01/23 01:17:56
まぁ、CUDAでXeonに勝つ速度を叩き出している漏れが一番だねw
80:デフォルトの名無しさん
08/01/23 12:23:21
なるほど、確かに>79が一番阿呆かも知れない。
大してネタがないのだけれど、保守代わりにメモage。
--
※nvccでコンパイル時に、制限に注意しつつ-use_fast_mathを指定する方がいい。
・三角関数や対数などの超越関数は、>74が指摘したようなニモニック(sin.f32など)を使ったコードを出力するようになる。
・また、巨大な数(数字忘れた)で割る割り算の精度を確保する為の回避ロジック(分母分子とも0.25を掛ける)も省略されるようになる。
# つまり、割り算(dif.f32)には制限があると言うこと。
・但し、冪乗関数はeの冪も10の冪も2の冪のニモニック(ep2.f32)を使い、それに定数を掛けるようだ。
※__sincosf()は最適化を阻害するので、__sin(), __cos()を使って実装した方が良さそう。
# __sincosf()はどの道sin.f32とcos.f32を別々に使うので、別々にしても遅くなることはない。
--
あとで整理してどっかに上げとくかな。
81:80
08/01/23 16:03:38
>80を一部訂正。
・但し、冪乗関数はexpf()は1.4427を、exp10f()は3.3219を、それぞれ掛けて(mul.f32)からex2.f32を使うようだ。
82:デフォルトの名無しさん
08/01/28 23:07:10
すみません。質問です。
3次元の格子上の各点で計算をするプログラムを書きたいのですが、
各点のインデックスを作るとき、たとえば以下のように書くのは間違いですか？
CUDAで3次元格子上のシミュレーションをやるときに使う一般的な方法が
ありましたら、教えていただければさいわいです。
#define IMAX 64
#define JMAX 64
#define KMAX 64
#define IDX(i,j,k) ((k)*IMAX*JMAX+(j)*IMAX+(i))
float *dVel;
int main() {
cudaMallocArray(dVel,0,sizeof(float)*IMAX*JMAX*KMAX);
初期化＆デバイスにデータ転送の関数();
dim3 grid(IMAX-2,JMAX-2,1);
dim3 thread(KMAX-2,1,1);
while(1) {
問題のfunction<<<grid,thread>>>(dVel);
if(収束) break;
}
}
__global__
void 問題のfunction(float *dVel) {
int idx = blockIdx.x+1; int idy = blockIdx.y+1;int idz = threadIdx.x+1;
float cx = (dVel[IDX(i+1,j,k)]+dVel[IDX(i-1,j,k)])*0.5;
float cy = (dVel[IDX(i,j+1,k)]+dVel[IDX(i,j-1,k)])*0.5;
float cz = (dVel[IDX(i,j,k+1)]+dVel[IDX(i,j,k-1)])*0.5;//こんな感じの計算がやりたいです
この後、四則演算が続く。
}
83:デフォルトの名無しさん
08/01/29 04:47:39
3次元格子は扱っていないから一般的な方法は知らん。
で、>82で流れはあっていそうだけど効率は悪いと思うよ。
ブロックごとにdVelをSharedにキャッシュした方がいいかも知れない。
# まぁ、動くものを作る方が先だけど。
84:デフォルトの名無しさん
08/01/29 09:10:24
>>83
ありがとうございます
吐き出される数値を見て、ぜんぜん計算されてなさそうなので、
心配をしていました。
効率悪いことは覚悟で、まずは動くものにしてみます
85:デフォルトの名無しさん
08/01/29 09:43:00
>82のコードは断片だと思うから敢えて指摘しなかったけど、一応注意点を列挙。
・(少なくとも安定するまでは)CUT_SAFE_CALL()でAPIを括っておいた方がデバッグしやすい。
# エラーチェックを毎回書くのも面倒でしょ。
・cudaMallocArray()は扱いが難しいので、cudaMalloc()にした方が楽では?
86:デフォルトの名無しさん
08/01/29 12:56:30
そうそう、忘れてた。
初期の段階ならエミュレーションで動かした方がデバッグプリントできるから楽かも。
87:デフォルトの名無しさん
08/01/29 13:30:21
>>85,86
どうも、アドバイス感謝です。
なんか、動かすとOSリセットかかる凶悪なプログラムに育ちました
もちろん実行はユーザ権限でやってるのに、そんなことできるんですね
エミュレーションでやってみます
88:デフォルトの名無しさん
08/01/29 13:40:59
そうなんだよねえ。
昔ビデオカードのドライバ無理矢理作らされてひどい目にあった。
89:デフォルトの名無しさん
08/01/29 14:06:29
>>87
それはあれだ、CPUでいうところのSegmentation Faultだと思う。
明後日の場所をアクセスすると、そうなりがち。
それと、XWindow環境なら5秒ルールにも注意。
# 仕事なら、移植作業を特価で承りますゼw
90:87
08/01/29 19:04:21
コンパイルされたバイナリを実行しないで、GPU使ってるか、エミュレーションしてるか見分ける
方法ってありますか。nvccに-deviceemuつけてコンパイルしてるつもりだけど、どうやら
まだデバイス上で動いているみたいなんです。
まあ、おかげさまでOSは落ちずにXだけ落ちるようになりましたがw
>>89
学業の一部ですので、コツコツ勉強しながら自力でなんとかやらせていただきます
91:デフォルトの名無しさん
08/01/29 21:31:21
XWindow環境なら5秒ルールに引っかかってしまうのか
知らなかった
92:デフォルトの名無しさん
08/01/30 08:07:06
CUDAをGentooに乗せた人いますか？
ディストリビューションをのせかえるのは大変なので
他のディストロ用のCUDAをインストールしたら
SDKがまともにコンパイルできないのです…
93:デフォルトの名無しさん
08/01/30 14:06:51
gccのバージョンが違うんじゃない?
CUDAに使えるgccを別途拾ってこないとダメじゃないかな。
# あと、DirectXとか。
## まぁ、エラー処理のところを作り替えちゃえばいらないけど。
94:デフォルトの名無しさん
08/01/30 16:37:38
質問です。
cudaで演算を行うプログラムを作成し、なんとかcore2duoの
80～100倍の処理速度のものが完成したんですが、
プログラムをdll化してVC++で作成したUIから呼び出したところ
処理速度がcpuの４～５倍に落ち込んでしまいました。
何が原因かはわかりませんが、タイマーを使って調べたところ
cudaMemcpy関数の処理速度が激しく落ちていました。
これはやはりdll化したことが原因なのでしょうか？
しかしNVIDIAのフォーラムにはdll化は問題なくいけると書いてありました。
dll化もフォーラムのとおり行いました。
cudaMemcpy以外の処理はdll化前と変わらない速度で行われているようです。
いったい何が原因でこのようなことになってしまったのでしょうか？
どなたか解決法、もしくは同じ問題を扱ったことのある方いましたら
アドバイスをよろしくお願いします。
95:デフォルトの名無しさん
08/01/30 18:49:56
それだけじゃ判らんなぁ。
cudaMemcpy()の転送方向は?
HostToDeviceだったらCPUキャッシュ、DeviceToHostならGPUの処理未終了、って辺りかな。
取り敢えず、デバイス側関数呼び出しから帰ってきてもそれはGPUの処理終了じゃないってことに注意。
96:デフォルトの名無しさん
08/01/30 19:03:48
>>95さん
足りない説明ですみませんでした。方向はDeviceToHostです。
問題は計測している時間の中にＧＰＵの終わっていない処理にかかってる分も
含まれてしまっているということですね。
だとしたら納得いきます。
カーネル実行直後にメモリーコピーを行っているので。
まずはそれぞれの処理にかかってる時間をちゃんと測れるようにがんばってみます
ありがとうございました！！
97:デフォルトの名無しさん
08/01/30 19:07:29
CUDAを始めるに当たって、何をインストールすればよいのでしょうか？
とりあえず、Visual C++ 2008はインストールし、グラフィックドライバを最新のものにしました。
CUDA SDK以外に必須のものを教えてください。
98:デフォルトの名無しさん
08/01/30 19:15:15
必要なもの。
・英語力
・忍耐
・(cudaの動く)GPU
・(cudaサイトからリンクされている)ドライバ
・ランタイム
・SDK
99:デフォルトの名無しさん
08/01/30 23:54:00
Linuxでやろうとしてます。どのような
マシンを組んで、どのディストリを入れる
のが良い&安上がりですか？
100:デフォルトの名無しさん
08/01/31 05:28:40
>>99
安定性を求めるのならメーカー製のものがよいが、自作するつもりなら
少なくともマザーボードは鉄板にすること。
それから、電源が貧弱だとトラブルの元になりやすい。
ビデオカードのお薦めはGeForce 8800 GT。
101:デフォルトの名無しさん
08/01/31 10:41:05
安西先生、8800は安くないっす…。
102:デフォルトの名無しさん
08/01/31 12:56:52
目的が判らんから何とも言えんが、安さ優先なら8600GTでいいんじゃね?
余裕があるなら8600GTSで。8800GTは消費電力も発生熱量も大きいから、
筐体の空冷性能が悪いと巧くないかも知らんことだし。
103:デフォルトの名無しさん
08/02/01 19:37:13
Visual Studio 2008ではコンパイルできないのでしょうか？
104:デフォルトの名無しさん
08/02/01 20:45:43
現在
C:\Program Files\NVIDIA Corporation\NVIDIA CUDA SDK\projects
以下にあるtemplateを書き換えて使っています。
他のディレクトリに移すと、コンパイルできなくなってしまいます。
他のディレクトリに移す場合、どこを書き換えたらよいのでしょうか？
105:デフォルトの名無しさん
08/02/02 09:05:14
>>103
できない。
>>104
Visual Studioの新規作成で作成すればよい。
106:デフォルトの名無しさん
08/02/05 09:19:22
プロファイラがリリースされたようですね
107:デフォルトの名無しさん
08/02/05 17:45:28
SLI構成でCUDA使ってる人はいますか？
SLIにするには、ボード増設すればいいだけ的なことが書いてあったけど、
一枚から二枚に移行するにあたって、気をつけることってなにかあったら
教えてください。
108:デフォルトの名無しさん
08/02/05 19:12:00
>>106
どこ? ぱっと見で見つからない……
>>107
ボード増設だけではSLIにならない気もす。ボード完直結のケーブルが要るような……
逆に、CUDA使う分には別段困らないと思うけど。
109:デフォルトの名無しさん
08/02/05 20:42:22
>>107
SLIは同一型番のカードでないとうまくいかないことが多いみたいですね。
最低でもメモリ容量やクロックは合わせないとだめなんじゃないでしょうかね。
それからチップセットによってはx16の2枚のフルレーンではなくx8の2枚となるのも要注意でしょうか。
>>108
SLIと非SLIのマルチGPUでパフォーマンスに差が出るのかも気になりますね。
110:デフォルトの名無しさん
08/02/06 21:05:28
ホスト側のメモリの確保はmallocを使わないといけないのでしょうか？
サンプルコードを見るとmallocが使われていますが、newじゃ駄目なのでしょうか？
111:デフォルトの名無しさん
08/02/06 21:57:45
newは使えないって言うのをどこかで見たなぁ
112:デフォルトの名無しさん
08/02/06 23:18:31
Linuxなら大丈夫。Windowsの場合、nvccはgccベースなのにVCのリンカを使うからおかしなことになる。
113:デフォルトの名無しさん
08/02/07 06:27:16
iostreamをインクルードすると、
"C:\Program Files\Microsoft Visual Studio 8\VC\INCLUDE\xlocinfo", line 77: error
:
support for exception handling is disabled
throw runtime_error("bad locale name");
^
と行ったエラーが大量に出てコンパイルできません。何が問題なのでしょうか？
OSはWindowsXPです。
114:デフォルトの名無しさん
08/02/07 06:43:58
あんたの頭。
115:デフォルトの名無しさん
08/02/07 07:49:52
>>113
ロケールが悪いと書かれてるが？
116:デフォルトの名無しさん
08/02/07 07:55:34
ロケールが悪いんじゃなくて、
ロケールが悪いというエラーを返すために例外を投げてるのが悪いんでしょ
117:デフォルトの名無しさん
08/02/07 13:53:04
floatを4つ使うよりfloat4を使った方が高速なの？
float3とかfloat4の利点がいまいちわからないのだけど
118:デフォルトの名無しさん
08/02/07 14:14:17
globalメモリからのアクセスなら、floatはld.global.f32が使われる。
float2ならld.global.v2.f32が使われる。float3, float4も恐らく同等。
尤も、struct {float x, y}だとしてもfloat2扱いしてくれるから使いたくなければ使わなくてもいい。
119:デフォルトの名無しさん
08/02/07 15:19:37
GPGPU#2スレから来ました。
extern __shared__ int shared[];
でsharedメモリを使えますが、複数の配列を使いたい場合はどのようにすればよいのでしょうか？
120:デフォルトの名無しさん
08/02/07 17:53:32
>>119
<<<>>>の呼び出しで共有メモリを確保する方法を使う場合、
<<<>>>で共有メモリのサイズを指定して何らかの__shared__ポインタで場所を指定するだけ。
つまり、複数の配列を使いたければ自分で監理するしかありません。
例えば、
extern __shared__ char sharedTop[];
__shared__ int * sharedInt1 = (int *) (sharedTop + 0);
__shared__ int * sharedInt2 = (int *) (sharedTop + elemOfInt1 * sizeof(* sharedInt1));
というように。
勿論、
func<<<blocks, threads, elemOfInt1 * sizeof(* sharedInt1) + elemOfInt2 * sizeof(* sharedInt2))>>>();
のように呼ぶ必要があります。
121:デフォルトの名無しさん
08/02/07 18:21:09
<<<>>>に入れるブロックの数とかスレッドの数ってどうやって決めればいいのでしょうか。
いまは、適当に縦横の格子点数で決めてるんですが。
ちなみに、使っているGPUは以下のものです。
Device 0: "Quadro FX 4600"
Major revision number: 1
Minor revision number: 0
Total amount of global memory: 804585472 bytes
Total amount of constant memory: 65536 bytes
Total amount of shared memory per block: 16384 bytes
Total number of registers available per block: 8192
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: 1188000 kilohertz
122:デフォルトの名無しさん
08/02/07 18:26:18
ちなみにsharedが速いということはほとんどないから、安心してglobalを使え。
むしろ、コピーする分遅くなる。
まあ、CPUの様に高速なL1キャッシュでも積むようになれば別なのだろうけど。
123:デフォルトの名無しさん
08/02/07 18:27:00
>>121
プログラムによる
124:デフォルトの名無しさん
08/02/07 19:57:12
>>121
ドキュメントを読む限り、スレッド数はワープ単位にするべきかと。
それと、できる限りCPU側でループさせない方がいいと言う点からいけばブロック数は
1000-10000位は欲しいということになるらしい。
まぁ、実測してみればいいんじゃね?
125:デフォルトの名無しさん
08/02/08 19:52:20
GPUとCPUに分散させて計算させることはできませんか？
126:デフォルトの名無しさん
08/02/08 20:03:31
できます。CUDAで言えば、GPU側を起動した後CPUはGPUを待たずに処理できます。
127:デフォルトの名無しさん
08/02/09 03:18:24
CUBLASとCUFFTをソースコードからコンパイル出来た方がいたら、ぜひやり方を教えてください
ソースはココからダウンロードできます
URLﾘﾝｸ(forums.nvidia.com)
128:デフォルトの名無しさん
08/02/09 05:42:34
サンプルがなかったっけ？
129:デフォルトの名無しさん
08/02/09 08:06:18
>>128
何のサンプル?
CUBLASやCUFFTのサンプルならその他のCUDAのサンプルごとSDKに同梱されていると思うけど。
>>127
ソースからコンパイルする目的は?
130:デフォルトの名無しさん
08/02/11 21:01:30
>>122
再帰的につかう変数をsharedにもってくると、断然速いだろう？
131:デフォルトの名無しさん
08/02/12 03:11:08
CUDAで再帰って使えるんですか？
132:デフォルトの名無しさん
08/02/12 08:06:11
CUDAで再帰って使えないと思うような根拠があるんですか?
133:デフォルトの名無しさん
08/02/12 13:37:22
ところで、CUDAって何て読むの?
キューダ? シーユーディーエー?
134:デフォルトの名無しさん
08/02/12 13:42:45
>>133
お好きにどうぞ。
一般的な英語の発音ルールに乗っ取るなら「カダ」になるのかも知れませんが、
barracudaからの連想で「クーダ」と発音する方がそれっぽいかも知れません。
135:デフォルトの名無しさん
08/02/12 13:43:07
s/乗っ取る/則る/
136:デフォルトの名無しさん
08/02/12 20:31:08
クーダの方が語感的にいいな
137:デフォルトの名無しさん
08/02/12 22:39:15
Kirk氏の講演を聴いたことがあるが
クーダと言っているように聞こえた
138:デフォルトの名無しさん
08/02/13 00:30:47
管
139:デフォルトの名無しさん
08/02/13 14:20:22
Windows版の開発もLinux版のようにコマンドラインだけでやれる方法はありませんか？
140:デフォルトの名無しさん
08/02/13 15:56:54
VisualStudioを使わなければOK!
cl.exeやnmake.exeの使い方は自分で調べてくれ。
141:デフォルトの名無しさん
08/02/14 15:21:59
祭りだ祭りだ
Japan CUDA カンファレンス 3/6（木）＠本郷
URLﾘﾝｸ(www.loopinc.jp)
142:デフォルトの名無しさん
08/02/15 19:10:36
ブロックやグリッド数を2つ以上必要な場合ってどのような場合なのでしょうか？
143:デフォルトの名無しさん
08/02/15 20:09:39
GPUのプロセッサ数はどこでわかりますか？
144:デフォルトの名無しさん
08/02/15 23:24:30
>>142
GPUの性能を活かすためには、並列度を上げる必要があります。
逆に、ブロックやグリッドが1だとプロセッサを使い切れなくてパフォーマンスが出ません。
>>143
GPUの仕様か、CUDAのドキュメントに書いてあります。
アプリケーションから動的に取得することはできないようです。
145:142
08/02/16 19:59:42
ブロック数が1つだと、並列には実行されることはないのでしょうか？
プログラミングガイドを見ると、ブロックの中に複数のスレッドがあるのですが、
このスレッドというのは、CPUのスレッドのように並列には実行されないのでしょうか？
146:デフォルトの名無しさん
08/02/16 20:07:15
だから用いるThreadの個数を指定するの
例のソースコードなんかを読んだ方が良いよ
147:デフォルトの名無しさん
08/02/16 20:16:33
そう言えば次の廃エンドになるらしい9800GTXではCUDAもやっと2.0とかになるのか?
148:デフォルトの名無しさん
08/02/17 10:16:54
CUDAは範囲外のメモリを読み込もうとしてもエラーを返さないのですか？
149:デフォルトの名無しさん
08/02/17 10:44:07
>>148
APIはエラーを返すけど、デバイス側の関数は必ずしもエラーを返さないね。
下手すると落ちるけど。
150:デフォルトの名無しさん
08/02/17 14:34:08
blockIdxとthreadIdxは0から始まるのですか？
151:デフォルトの名無しさん
08/02/17 14:48:10
はい。例えばfunc<<<m, n>>>()と呼び出した場合は、
それぞれ0からm-1, 0からn-1の値をとります。
二次元型、三次元型の場合も同様です。
152:142
08/02/17 15:08:14
ブロック数を1つから2つに増やして、スレッド数を100から50に減らして実行してみたのですが、
速くなるどころか逆に2倍くらい遅くなってしまいました。そういうもんなのでしょうか？
dim3 grid( 1, 1, 1);
dim3 threads(100, 1, 1);
↓
dim3 grid( 2, 1, 1);
dim3 threads(50, 1, 1);
153:デフォルトの名無しさん
08/02/17 15:44:38
プログラムにもよると思うけどありえる話だろうね
154:デフォルトの名無しさん
08/02/17 15:53:10
何スレッドが同時に動くかはGPUによっても違うので一概には言えませんが、
少なくとも8の倍数にはなるのでスレッド数50や100は効率が落ちます。
例えばスレッド数を64にしても50のときと所要時間は変わらないかも知れません。
又、デバイス側の関数呼び出しは完了していても演算自体は終わっていないかも知れません。
cudaThreadSynchronize()してから時間を計った方がいいかも知れません。
155:デフォルトの名無しさん
08/02/19 17:46:37
研究室のPCにLinux入れてCUDA動かしてますが、
ユーザ権限実行でもOSが死ぬプログラムを余裕で書けてしまいますね
これはCUDA使う以上、仕方ないことなんですよねきっと
バグのあるデバイスドライバ使ってるのと同じことと、あきらめるしか
ないのでしょうか
156:デフォルトの名無しさん
08/02/19 18:03:53
vista対応と自動回復が働いてくれることを渇望しましょう
157:デフォルトの名無しさん
08/02/19 19:31:36
そして危険なプログラムが走らないような検証系に進むのであった。
といってもDirectX動かす権限があればOS死ぬのは容易いような。
158:デフォルトの名無しさん
08/02/19 19:36:50
vista対応マダー？
159:デフォルトの名無しさん
08/02/20 05:07:39
vistaは要らない子。
*BSD対応マダーー？？
160:デフォルトの名無しさん
08/02/20 21:32:15
いやむしろTRON対応
161:デフォルトの名無しさん
08/02/21 13:54:22
*BSD対応は確かに欲しいな・・・。
っていうか、オプソにするくらいの気持ちじゃないと
この規格は普及せんような・・・。
将来的にnVIDIAのGPUが携帯電話などに使われて行った時に
なかなか面倒じゃないかなぁ。
162:デフォルトの名無しさん
08/02/24 00:15:33
NVIDIAはしばらく内部仕様公開しないだろうね。
ptxは読めてもその先の動作がわからないから辛いわ。
AMDには圧倒的に差をつけていそうだし。
GeForce9600出ましたな。
8800はStreamProcessorが128個×1.35GHzだったけど，今度は64個×1.625GHzみたい。
並列度が低い場合でも早くなるっていうこと？
163:デフォルトの名無しさん
08/02/24 00:31:43
だめっぽい
URLﾘﾝｸ(www.4gamer.net)
164:デフォルトの名無しさん
08/02/24 00:44:42
>>157-158
ユーザースペースで動作させた場合の性能が気になりますね。
>>159-161
*BSDやgentoo等への対応も欲しいですね。
欲をいえばVista以外でのHybrid SLIのHybrid Power対応も w
自動切換えとかGeForce Boostとかは無理でも、コマンド一発でオンボとカードを切り替えられるようになるといいのですが。
165:デフォルトの名無しさん
08/02/24 09:31:22
doubleにはいったいいつ対応するのか
166:デフォルトの名無しさん
08/02/24 09:38:24
ｽﾚﾘﾝｸ(tech板:137番)
167:デフォルトの名無しさん
08/02/24 14:10:04
URLﾘﾝｸ(www.nvidia.com)
9600GT，CUDAも実行できなければ遅い。doubleが使える世代への場繋ぎでしょう。
168:デフォルトの名無しさん
08/02/24 14:17:49
>>167
NVIDIAは伝統的に新しいGPUでCUDAが使えるようになるまでに少なくともドライバを更新してくるから、
今のところ未だ、9600GTでCUDAが使えるようにならないとは言えない。
169:デフォルトの名無しさん
08/02/24 14:28:20
普通に使えるようになるでしょ。つかヘッダとか弄ったら使えたりして。
で、VC2008+cuda1.1の食べ合わせって駄目なのかよ。1.0はいけるらしいけど orz
URLﾘﾝｸ(forums.nvidia.com)
170:デフォルトの名無しさん
08/02/26 07:10:16
Vistaのbeta版が３月の終わりごろに出るらしい
URLﾘﾝｸ(forums.nvidia.com)
171:デフォルトの名無しさん
08/02/26 11:05:55
>>170
㌧
172:デフォルトの名無しさん
08/03/02 23:09:30
ずっと上の方で誰か質問しててまだ回答無いけど、
ゲフォ8400GSでCUDA動きますか？
WinXPでは動いたんだけど、Fedora Linuxでは
認識しなくてエミュレーションに成ってしまった。
何かと金の掛かるWinではCUDAやりたくないです。
173:デフォルトの名無しさん
08/03/02 23:27:05
>>172
一覧にはあるから使えるんじゃないかというレスがあった希ガス。
問題は、ドライバをインストールできるかどうかでそれは端末とBIOS次第かな。
ToolKitはFedora7用のがあるそうだから、逆にそれ以前のはダメかもしれない。
あーでも、エミュレーションでコンパイルできているなら大丈夫か。
コンパイルオプションを指定すれば、ちゃんとGPUコードのモジュールもできるんじゃないかな。
デバイス認識させるのは癖があるから要注意で。
Xが動く状況じゃないと、巧く認識できなくて苦労した記憶がある。
巧く認識できていれば/dev/nvidia0と/dev/nvidiactlができている筈なんで、
それがないならroot権限で正しくコンパイルできたGPUコードのサンプルを動かせば
ドライバの.koを使ってデバイスのエントリができる筈。
Linuxのデバイス周りは詳しくないんで、頓珍漢なことを書いているようなら失礼。
174:デフォルトの名無しさん
08/03/03 10:30:40
しかし、VisualStuido2005EEと組み合わせれば無料でできるのに「金の掛かる」とはまた意味不明な……
WindowsならGPUのプロファイリングツールも使えるし、それはそれで悪くない選択だと思うんだけどね。
# といいつつ、未だ自宅で環境整備してないけどw
175:デフォルトの名無しさん
08/03/03 10:32:18
8600GT WIN2Kの俺に謝れ！！！
XPが無い…
176:デフォルトの名無しさん
08/03/03 10:36:43
>>175
Win2kで動かないのであれば、WinXP買うのも馬鹿馬鹿しいね。ゴメン、配慮が足りなかった。
177:デフォルトの名無しさん
08/03/03 18:37:47
XPが嫌いってだけのもあるんだけどなー
ムダに重たくて
178:デフォルトの名無しさん
08/03/03 18:50:25
>>177
テーマをクラシックにするか、システムプロパティでパフォーマンスを優先にすればWin2Kに較べて重くないと思いますよ。
179:デフォルトの名無しさん
08/03/03 18:58:12
構ってちゃんの相手するなよ・・・
180:デフォルトの名無しさん
08/03/04 02:14:10
そもそもWindowsが糞高いし、ソース付いて来ないし。
不自由過ぎてハックする気が萎える。
181:デフォルトの名無しさん
08/03/04 17:32:33
こんちわす。
.NET環境からCUDAとデータの受け渡しをする例、なんてのは
無いもんでしょうかね。既存のVB.NET、C#.NETの統計解析を
高速化致したく。
相関係数＝Σ(xj-xk)^2÷√Σ(xj-xjの平均)^2*(xk-xkの平均)
なんてのが多発し、j,kは1～1000、それぞれ1万点ずつ、x1～x1000の全ての
組み合わせで計算、という具合です。でも「総和を取る」とかはGPUは
苦手でしょうかね・・・。
182:デフォルトの名無しさん
08/03/04 17:44:36
総和は苦手なので、適当に分割して小計を求めてから最後に総計を求めることになると思う。
それはさておき、先ずはその相関係数の式自体を変形してみてはどうだろう。
分子も分母のそれぞれも平均値を使わない形にできるはずだが。
# ついでに言えば分母の最後は2乗が抜けている希ガス。
それをやると、平均を使う2パス処理ではなくΣxj, Σxj^2, Σxk, Σxk^2,...などを求める1パスですむ筈。
あ、肝腎のVB/C#からの呼び出し方はわかんね。CUDA側は"*.obj"になるからリンクさえできればなんとかなるのかな?
まぁ、計算量が充分大きいならCUDA側を"*.exe"にして起動するようにしても大差ない気もするけど。
183:デフォルトの名無しさん
08/03/05 12:43:43
どうもどすえ。自分なりにお勉強してみました。
.NET←→CUDAは、UnmanagedMemoryStream/共有メモリ/セマフォ、とかで
なんとかなりそうです。典型的な処理は以下のようにすればよいのかなと。
・データは、(128×8)センサー{測定データ64個単位×256｝のように
CPU側で整形する。平均0、分散1に正規化し、欠損・異常値の前処理も
しておく。これはNオーダーの処理なので十数秒ですむ。
・GPU側には、C(j,k)＝(A(j,k)－B(j,k))^2 (j,kは0-7)
つまり8×8画素の画像の、差の二乗？を繰り返し計算させる。
D(j,k)＝D(j,k)＋C(j,k)で集計する。
センサー(0～1023番)対センサー(0～1023番)なんで、N^2オーダー。
今はここに数時間掛かってます。
・終わったらホスト側でD(j,k)を取り出して、画素を全部足す。
数時間が数分になってくれると嬉しいのですがｗ
「GPU Gem」第三版を注文したので、もうちょっとお勉強してみるす。
184:デフォルトの名無しさん
08/03/05 13:27:36
だから、相関係数を求める式を変形しろって。
CUDA云々以前だぞ。
185:デフォルトの名無しさん
08/03/05 14:02:20
え、相関係数を求める式にはもう「平均」は無いすよ。
ごめんどこ怒られてるのかわかんね・・・
186:デフォルトの名無しさん
08/03/05 14:29:04
「相関係数」でぐぐれかす
187:デフォルトの名無しさん
08/03/05 14:33:41
お前ら両方とももういい加減にしろ
188:デフォルトの名無しさん
08/03/05 16:49:01
統計の基本:
相関係数は普通ピアソンの積率相関係数を指す。
これは、r = Sxy / √(Sxx * Syy)で求められる。
このとき、Sxyはx,yの共分散、Sxxはxの分散、Syyはyの分散。
分散はSxx = Σ(xi - x￣)^2。Syyについても同様。
# ここではxの平均を便宜上x￣で現わす。
この式は、Sxx = Σxx - Σx*Σx/nに変形できる。
このとき、Σxxはxiの平方和、Σxはxiの総和。
また、共分散はSxy = Σ(xi - x￣)*(yi - y￣)。
この式は、Sxy = Σxy - Σx*Σy/nに変形できる。
このとき、Σxyはxiとyiの積の総和。
つまり、相関係数はr = (Σxy - Σx*Σy/n)/√((Σxx - Σx*Σx/n)*(Σyy - Σy*Σy/n))。
従って、線形時間で求められる。また、1回のスキャンで充分なので、各要素を保存する必要がない。
189:デフォルトの名無しさん
08/03/06 01:21:05
上のほうにもあるけど、総和の計算って実際どーすればいいんだ？
スレッドブロックごとに小計を求めて、各ブロックの結果をまとめるとかいう
方法になりそうなのは予想できるが、実際どんな風に書けばいいのかわからん。
具体的にはベクトルの長さが出したいんだが、各要素の２乗の和がうまく出せない。
ヒントでもいいので分かる方教えてください。
190:デフォルトの名無しさん
08/03/06 02:29:35
1つのブロックを総和の計算に割り当てる…とか
そんなことしか思い浮かばん
191:上の方の人
08/03/06 07:58:00
後で暇があったら実際のコードの集計部分だけ晒してみるか。
192:デフォルトの名無しさん
08/03/06 19:32:29
CUDAを使ったjpegの展開ライブラリはありませんか？
193:デフォルトの名無しさん
08/03/06 19:55:22
IDCTはそこそこ効率よさそうだけど前段のビットストリーム分解が不向きっぽいぞ。
194:デフォルトの名無しさん
08/03/06 22:20:39
>>191
そうしてくれると助かります。
よろしくお願いします。
195:デフォルトの名無しさん
08/03/06 23:43:14
デイビッド・カーク氏来日，CUDAカンファレンス2008開催
URLﾘﾝｸ(www.4gamer.net)
196:デフォルトの名無しさん
08/03/07 04:38:20
やっとできたらしい日本語版のCUDAZONEとプログラミングガイドの日本語版。
URLﾘﾝｸ(www.nvidia.co.jp)
しかし、機械翻訳そのままなのかな。余りに訳が酷くて原文読まないと意味が判らんとは……
197:デフォルトの名無しさん
08/03/07 15:34:38
CUT_EXIT(argc, argv)はなんでargc, argvを与えなくちゃいけないのでしょうか？
198:デフォルトの名無しさん
08/03/07 15:54:16
>>197
引き数にnopromptが指定されていない場合は
"Press ENTER to exit..."のメッセージを出して入力待ちにするためです。
詳しくは、cutil.h参照で。
199:デフォルトの名無しさん
08/03/07 18:23:26
floatであった変数をいくつかまとめてfloat4を使うようにしました。
そうしたところプログラムの実行速度が10%ほど落ちたのですが、float4よりfloatの方が
実行効率がよいのでしょうか？
200:デフォルトの名無しさん
08/03/07 23:32:28
よく知らないが、float4ってSIMD用の型じゃないの？
201:デフォルトの名無しさん
08/03/08 00:20:55
単にまとめただけじゃ速くならないと思う。一番いいのは-ptxして出力眺めることなんだけどね。
202:デフォルトの名無しさん
08/03/08 09:24:11
なんで、nvidiaのフォーラムには韓国語はあるのに日本語はないの？
203:デフォルトの名無しさん
08/03/08 11:42:50
流石にそれはnvidiaに聞いてくれ。需要が少ないと言うか、要望が少ないのだろ。
半島人と違って日本人は奥床しいからw
204:デフォルトの名無しさん
08/03/08 11:57:24
アメリカの地下鉄とかの案内板や自販機にハングルはあるけど
日本語はないとかいう話を思い出しました w
東大キャンパスであった、CUDAカンファレンス2008の動画とかは公開されないのでしょうかね？
205:デフォルトの名無しさん
08/03/08 21:17:16
ずいぶん前の話だが、20-26あたりで議論されてる問題って結局どんな風に
並列処理するのが最適なんだ？
総和の計算の話が出てたから思い出したのだが。
前に試行錯誤したが、高速化できなかったり、値がおかしくなったりして
あきらめてしまっていた。
どうぞ皆様知恵をわけてください。
この前のカンファレンスでも言われていたがチューニングにかける時間を
惜しんではいけないのだね。
206:デフォルトの名無しさん
08/03/08 21:27:59
総和みたいな計算は向いてないとも言ってたけどね
207:デフォルトの名無しさん
08/03/08 21:59:03
>>205
結局>20は、>24で書いたようにインデックスで並列にした。
要はこんな感じ。
--
大きな配列tmpがあるとき、下図の+で区切られた範囲ごとに集計したい。
tmp+-----+-----+--------+-----+---+...
つまり、1番目のセクションは最初の(先頭の)+から次の+まで、2番目のセクションはその次の+まで……
ここで、こんな構造体を考える。
struct sections {short begin, end;}
こいつの配列を作って次のように値をセットする。
{{0, 6}, {6, 12}, {12, 21}, {21, 27}, {27, 31}, ...}
これをデバイス側に転送しておいて、一つのデータスレッドが一つのセクションを担当する形で集計した。
この方法の問題点：
・セクションのサイズが不均衡なので、ワープ内でも不均衡だと無駄なからループが発生してしまう。
# 1ワープ内では同じインストラクションが走ってしまうため。つまり、使用効率が落ちる状態。
・セクションサイズがワープ内では極力等しくなるようにソートすると、今度はアクセス場所がランダムになる。
# 先程の配列が、例えば{{0, 6}, {6, 12}, {21, 27}, {301, 307}, ...}のようになってしまう。
いずれにしても、綺麗に並べることができない。
但し、今回は前段のtmp配列への演算結果の格納が所要時間の大半を占めたために適当にチューニングして放置。
尚、予告した集計のロジックは、総和ではなく↑とも違う小計算出だったので条件が違うと言うことでパス。
総計は興味があるので、その内テストできたら改めて晒そうと思う。
208:デフォルトの名無しさん
08/03/09 13:44:37
いきなりですが質問です
xp64でVisualStudio2005と8800GTS(G92)がありCUDAドライバー、
ツールキット、SDK、と入れたんですが、結局ビルド方法とか
どのPDFに書いてあんだろ、というのがよくワカンネですよ。
VS2005のプロジェクト設定をx64に変えて、cutil32D.libを
拾ってきたZIPから解凍して入れたら、mandelbrotとか一応
動いたけど。他のmandelbrot描画ソフトより劇速なんで、
動いてはいるみたいです。
209:デフォルトの名無しさん
08/03/09 16:10:18
うーん、IDEは使ってないから判らないなぁw
pdfじゃなくて、ReleaseNoteか何かにビルド方法かいてなかったっけ?
後で見てみるわ。
210:205
08/03/09 23:15:54
>>206
確かに言っていた。まぁそれは言われなくとも自明なことなんだが。
ついでに、CUDAがあるからGPUに向いていない処理も容易に書けるようになってしまった
とも言っていたな。
だが、膨大なデータをホストに書き戻して、総和を計算し、その結果をまたデバイスに
送って使うということは出来れば避けたかった。
GPUが苦手な処理でもデバイス側でやらせるというのがいいのか、転送コストを覚悟の上で
ホストに戻して計算したほうがいいのか実際確かめるために試行錯誤してみている。
>>207
説明ありがとう。
おかげで24でかかれていたことが理解できたよ。
その考え方は思いつかなかったな。
俺ももうちょい考えてみていい感じにできたら晒そうかと思う。
211:デフォルトの名無しさん
08/03/10 02:56:03
ベクトルの総和を出すscanlarge.pdfを見てみると、
elementsが65536ならCPU－GPUでほぼ一緒、
100万elementsでGPUがCPUの5倍速い、ようですね。
なるほどGPU向きではないな・・・
212:デフォルトの名無しさん
08/03/10 03:54:26
HostがCore2Duo辺りなら未だいいんだけど、Woodcrest辺りだと厳しいんだよねw
超越関数を大量に使う演算だと、GPUが随分有利になるんだけど。
# SSEには超越関数がない→ベクタ化のためには関数テーブルが必要→メモリ消費→キャッシュミス
213:デフォルトの名無しさん
08/03/10 11:09:00
cudaでfloat4*float4ってできないの？
各要素ごとに4行に分けて書かないとだめ？
214:デフォルトの名無しさん
08/03/10 18:33:07
>>213
float4 * float4 はないみたいだ。
ちなみに、
--
__device__ float4 a, b, c;
__device__ static void func()
{
c = make_float4(a.x * b.x, a.y * b.y, a.z * b.z, a.w * b.w);
}
--
こんなコードを書いてみたけど、
--
.global .align 16 .b8 a[16];
.global .align 16 .b8 b[16];
.global .align 16 .b8 c[16];
:
:
ld.global.v4.f32 {$f1,$f4,$f7,$f10}, [a+0]; //
ld.global.v4.f32 {$f2,$f5,$f8,$f11}, [b+0]; //
.loc 4 12 0
mul.f32 $f3, $f1, $f2; //
mul.f32 $f6, $f4, $f5; //
mul.f32 $f9, $f7, $f8; //
mul.f32 $f12, $f10, $f11; //
st.global.v4.f32 [c+0], {$f3,$f6,$f9,$f12}; //
--
こうなった。
215:デフォルトの名無しさん
08/03/11 19:17:03
FreeBSDのLinuxエミュでCUDAを使っているんだけど、
Linuxと比較してどれくらい速度が落ちているのかな?
ほぼネィティブと同程度ならこのまま使っていこうかと
思っているんだけど。
216:デフォルトの名無しさん
08/03/12 05:34:03
Windows上でコンパイルしているのですが、nmakeは使えないのでしょうか？
cygwinのmakeコマンドであればコンパイルできるようなのですが、nmakeを使うと
NMAKE : fatal error U1077: 'C:\CUDA\bin\nvcc.EXE' : リターン コード '0xc0000005'
Stop.
となってコンパイルできません。このエラーは何を意味しているのでしょうか？
217:デフォルトの名無しさん
08/03/12 05:43:10
float4ってsimdで計算してくれないの？
218:デフォルトの名無しさん
08/03/12 07:57:11
>>217
そもそもsimdなんてありませんが何か。つーか、>214で回答ついてるっしょ。
>>216
nvccが終了ステータスに5を積んでいるようですが、
ちゃんと動いた結果ならnmake側で無視することができると思います。
そもそもちゃんと動いてますか?
>>215
4種類ほどボード別に速度を測ったデータがあるからそれでも載せてみましょうか。
219:216
08/03/12 08:30:07
>>217
早速のご回答ありがとうございます。
Makefileの内容ですが、非常に初歩的な
test.exe: test.cu test_kernel.cu
nvcc test.cu
です。直接、コマンドプロンプトから打ち込んだ場合、問題なくコンパイルできます。
220:デフォルトの名無しさん
08/03/12 08:49:24
>>219
レス番違いますぜ。
取り敢えず、nvccがエラーステータスを返しているかどうかと、
nmakeでそれを無視できないか調べてみてはいかがでしょう。
221:デフォルトの名無しさん
08/03/12 12:49:19
そういうわけで、各ボードの比較表を作ったので貼ってみる。
--前半
"name","Tesla C870","GeForce 8800 GTX","GeForce 8800 GTS","GeForce 8800 GT","GeForce 8600 GTS","GeForce 8600 GT","使用コマンド"
"totalGlobalMem","0x5ffc0000","0x2ff50000","0x27f50000","0x1ffb0000","0xffb0000","0xffb0000","deviceQuery他"
"[MiB]",1535.75,767.31,639.31,511.69,255.69,255.69,
"sharedMemPerBlock","0x4000","0x4000","0x4000","0x4000","0x4000","0x4000",
"[KiB]",16,16,16,16,16,16,
"regsPerBlock",8192,8192,8192,8192,8192,8192,
"warpSize",32,32,32,32,32,32,
"memPitch","0x40000","0x40000","0x40000","0x40000","0x40000","0x40000",
"[KiB]",256,256,256,256,256,256,
"maxThreadsPerBlock",512,512,512,512,512,512,
"maxThreadsDim",512,512,512,512,512,512,
,512,512,512,512,512,512,
,64,64,64,64,64,64,
"maxGridSize",65535,65535,65535,65535,65535,65535,
,65535,65535,65535,65535,65535,65535,
,1,1,1,1,1,1,
"totalConstMem","0x10000","0x10000","0x10000","0x10000","0x10000","0x10000",
"[KiB]",64,64,64,64,64,64,
"major:minor",1:00,1:00,1:00,1:01,1:01,1:01,
"clockRate","1350000[kHz]","1350000[kHz]","1188000[kHz]","1512000[kHz]","1458000[kHz]","1188000[kHz]",
"textureAlignment","0x100","0x100","0x100","0x100","1x100","0x100",
222:デフォルトの名無しさん
08/03/12 12:49:57
--後半
"with Xeon",,,,,,,"使用コマンド"
"convolutionFFt2d","145.05[MPix/s]","107.99[MPix/s]","99.57[MPix/s]","107.00[MPix/s]",,,"convolutionFFT2D"
"convolutionRowGPU","1100.00[MPix/s]","1071.34[MPix/s]","717.34[MPix/s]","1457.11[MPix/s]",,,"convolutionTexture"
"cudaMemcpyToArray","1209.69[MPix/s]","1141.85[MPix/s]","1049.30[MPix/s]","1483.92[MPix/s]",,,
"simpleTexture","1180.83[MPix/s]","1125.08[MPix/s]","799.22[MPix/s]","910.22[MPix/s]",,,"simpleTexture"
"clockTest",24766,24930,25346,22396,,,"clock"
"with Core2Duo",,,,,,,"使用コマンド"
"convolutionFFt2d",,,,"106.64[MPix/s]","41.21[MPix/s]","32.93[MPix/s]","convolutionFFT2D"
"convolutionRowGPU",,,,"1433.39[MPix/s]","467.43[MPix/s]","375.54[MPix/s]","convolutionTexture"
"cudaMemcpyToArray",,,,"1475.25[MPix/s]","893.47[MPix/s]","596.14[MPix/s]",
"simpleTexture",,,,"873.01[MPix/s]","370.00[MPix/s]","285.30[MPix/s]","simpleTexture"
"clockTest",,,,24066,65006,71736,"clock"
223:デフォルトの名無しさん
08/03/12 13:58:27
これってプログラミングガイドの最後にあるやつ？
224:デフォルトの名無しさん
08/03/12 14:29:37
CUT_EXITを実行せずにプログラムを終了した場合、どのような悪影響がありますか？
225:デフォルトの名無しさん
08/03/12 15:01:43
>>221
SharedMemoryって16KBしか割り当てられないのですか？？
226:デフォルトの名無しさん
08/03/12 15:10:28
幾ら何でも同じスレの中くらい検索してくれたっていいじゃないか(TT
>>223
いいえ、サンプルプログラムの出力です。
>>224
なーーんにも。どうしても気になるならcutil.hを眺めてください。
>>225
この表では、1block辺りのSharedMemoryのサイズを示しています。
227:デフォルトの名無しさん
08/03/12 18:15:17
Cのmallocの場合、NULLかどうかでメモリが確保できたか判別できますが、
cudaMallocでメモリが確保されたかどうかを調べる方法を教えてください
228:デフォルトの名無しさん
08/03/12 18:31:03
CUT_CHECK_ERROR
229:デフォルトの名無しさん
08/03/12 18:31:16
cudaMalloc()
cudaError_t cudaMalloc(void** devPtr, size_t count);
allocates count bytes of linear memory on the device and returns in *devPtr a pointer to the allocated memory.
The allocated memory is suitably aligned for any kind of variable.
The memory is not cleared. cudaMalloc() returns cudaErrorMemoryAllocation in case of failure.
230:228
08/03/12 18:42:05
ああ、あとデバッグモードじゃないとだめだよ
231:デフォルトの名無しさん
08/03/13 03:10:17
warpサイズって簡単に言うと何？
232:デフォルトの名無しさん
08/03/13 06:48:58
同一インストラクションが同時に走る、演算ユニットの集まり?
233:デフォルトの名無しさん
08/03/13 13:50:16
CUDA初心者です。
CUDA難しいよ
難しいよCUDA
・<<<Dg, Db, sizeof(float)*num_threads*4>>> て書いたら動かないんすかね。
int sharemem_size = sizeof(float)*num_threads*4;
<<<Dg, Db, sharemem_size >>> て書いたら動いた。
・畳み込み演算させたら、CPUでやるより1.7倍しか速くないの(´・ω・`)
・動いてる間、システム全部固まってる。最初焦った。
CUDAの道険しいﾅﾘ
234:デフォルトの名無しさん
08/03/13 15:37:14
初歩的な質問かとは思いますが、各ストリームプロセッサにそれぞれ1つずつスレッドが割り当てられるのですか？
235:デフォルトの名無しさん
08/03/13 18:22:14
__syncthreads()ではブロック内でしか同期できませんが、全ブロックを同期したいときは
カーネルを抜けていったんCPU側に戻すしかないのでしょうか？
236:デフォルトの名無しさん
08/03/14 02:24:24
cudaMallocHostとCのmallocの違いを教えてください
237:デフォルトの名無しさん
08/03/14 05:10:46
>>235
確かカーネルは非同期で実行されたはずだから、
カーネルを連続で呼んだら同期されない・・・と思う。
<<<カーネルの呼び出し>>>
cudaThreadSynchronize();
<<<カーネルの呼び出し>>>
ならいけるんじゃないの？自信ないけど。
238:デフォルトの名無しさん
08/03/14 11:06:04
>>236
cudaMallocHostでメモリ確保すると、ページングなしスワッピングなし
物理メモリに連続して常駐、のメモリ領域が確保されるはずなのだ！
と思ってcudaMallocHost使ってます。mallocとmlockの合せ技？
239:デフォルトの名無しさん
08/03/14 16:02:07
Windows, VC2005なんですが、PTXを出すためのオプションはどこに
書けばいいんどすか( ・ω・`)
苦労してPhenomの6倍速いまで持ってきたんですが、さすがに
この先はアセンブラ見てみないと何をどうすればいいか分からんどす。
モノは信号処理用の畳み込み(4096)どす。
240:デフォルトの名無しさん
08/03/14 16:41:32
>>239
プロパティでカスタムビルドであることをチェックできない?
オプションは言うまでもなく -ptx
241:デフォルトの名無しさん
08/03/14 16:56:28
グラフィックカードを2枚挿した状態で、メインメモリからGPU側に転送する際に
どのようにしてGPUを指定すればよいのでしょうか？
カーネルから見られるのは、1枚のGPUだけという認識でよろしいのでしょうか？
242:デフォルトの名無しさん
08/03/14 23:48:11
>>241
cudaGetDeviceCount()からcudaChooseDevice()あたりでは。プログラミング
ガイドのD.1。二枚刺しﾃﾗｳﾗﾔﾏｼｽ
>>240
プロジェクトのプロパティのカスタムビルドステップ、等にもコマンドライン
指定が何もなかとです。次はVSディレクトリに潜って.cuのRuleファイルを
探してみるです。
243:デフォルトの名無しさん
08/03/15 03:32:18
RuntimeAPIとDriverAPIの違いを教えてください。
244:デフォルトの名無しさん
08/03/15 12:17:00
CUDA素晴らしすぎです。ようやくグリッドとブロックとスレッドと
globalとsharedとconstantを理解して、今「Xeonより100倍速いぜ！」
な速度を叩き出しています。もっとチューニングがんばってみよう。
>>243
RuntimeよりDriverの方が低レベルとして書いてありますけど、
似たようなの一杯有るし、違い良く分からないですよね。
245:デフォルトの名無しさん
08/03/15 13:20:14
恐らくは、Driverの方はC++で使うことを想定してない。
実際問題として、nvccはC++としてコンパイルするのでRuntimeAPIの方が使いやすいと思う。
246:デフォルトの名無しさん
08/03/15 17:14:57
>>244
segdmm のC800では120GFlops程度と言われていますが
4coreのXeonでも75GFlopsです。
どのくらい違いますか教えてください。
247:244
08/03/16 00:22:24
>>246
やってみたのが4096単位データ×128個に対して4096のフィルタで
コンボリューション掛ける、なんですが、Xeon2GHzの1スレッド
だと250M操作/秒位しか出来ないですよ。キャッシュミス連発？
GeForce8800ですと、28G操作/秒で動きます。
248:244
08/03/16 04:08:51
あ、そんで、Xeonでやるとこうだってのはもれじゃ無くて既存の
なので、弄れないし、作りはよく分からんのです。とりあえず
「このルーチン速くするライブラリ作れんか？CUDAとか言う奴で」
と言われてやってみたら100倍速いので怒りが有頂天になった所です。
でも冷静に考えると、CPU側もカリカリに書けばあと10倍位速くなる
可能性はありますわな。たぶんSSEとか使ってないし。
でもそこから10倍にするには、つまりもう何台かパソコン追加しないと
ならんわけなので、やっぱりCUDA可愛いよ可愛いよCUDAｗ
249:デフォルトの名無しさん
08/03/16 08:32:33
私のところじゃ、Xeonの3-10倍だなぁ。単体テストで。
# 他の演算と組み合わせると、1.5-2倍が限度（TT
250:デフォルトの名無しさん
08/03/16 09:07:29
「Xeon」てみなさんが言ってるのは、DPマシン8コアのこと？それとも1コアのこと？
1コアに負けるんじゃ確かにちょっと悲しいよね。
251:デフォルトの名無しさん
08/03/16 09:14:14
1スッドレ、って書いてるじゃん
252:デフォルトの名無しさん
08/03/16 09:34:49
つまりこういう暑苦しい感じが最強だと。
喧嘩してる場合じゃねぇ！Quadコア×デュアルCPU＋8800GTXの2枚刺し、
の友情で最強を目指そうじゃないかお前ら！夕日に向かって海まで走れ！
敵は巨大行列のLU分解だ！手ごわいぞO(N^3)！
253:デフォルトの名無しさん
08/03/16 12:57:56
グリッド数、ブロック数、スレッド数の決め方がよくわかりません。
例えば独立な300個のスレッドがあった場合、
一つのブロックに300のスレッドを割り当てるのがよいのでしょうか？
それともスレッド数は1つで300のブロックを作成するのがよいのでしょうか？
ワープ数などはどのように考慮すればよいのでしょうか？
254:デフォルトの名無しさん
08/03/16 13:54:06
>>253
あくまでも漏れの場合ですが、
ブロック数＝16。これは使っているカードの「Multi Processor」に合せて。
違うグリッドに属するデータのやり取りは出来ない事に注意。
スレッド数＝512。目安で。
で、カーネル関数でループで処理を回す場合、最外では
for (i = 0; i < maxcalccount; i=i+スレッド数×ブロック数)
のようにします。これで8192単位での並列化になります。
ループ内部では、グローバルメモリからたとえば１ワード読むなら
sharedメモリ[スレッドID]
＝グローバルメモリ[スレッド数×ブロック数＋ブロックID×スレッド数＋スレッドID]
となります。
floatを使う場合、スレッドが512なら、各スレッドで最大8ワードもてますが、
sharedメモリは使い切らないほうが良いようです。
255:デフォルトの名無しさん
08/03/16 14:57:39
ブロック数は、どうせ多く割り当ててもCUDA内部で直列に並べるだけだから
非同期で少しでもCPUと並列にしたい場合を除けば大目に割り当ててOK。
スレッド数についても多い分はどうせ別のワープに割り当てられるから多めでOK。
但し、同期を取る場合には多過ぎるとダメ。
手元のデバイス関数の場合、ブロック数*スレッド数は少なくとも1024か2048以上必要(8800GTXの実測で)。
これらを踏まえると、スレッド数が32ならブロック数は64以上、スレッド数が64ならブロック数は32以上くらいか。
ブロック数の上限は、実測しながら適当に調整するとして、大体1024を超えるといくつでも変わらないと思う。
# これも、具体的なテスト用のサンプル用意したいところだね。
256:デフォルトの名無しさん
08/03/16 16:42:50
>>188みたいな話が載っている本で使っている本があれば教えてほしいなあ。
入門書の理屈を素直になぞっているような人間には、普通のcpuとgpuとの差を出せなさそう。
257:デフォルトの名無しさん
08/03/16 17:01:32
>>256
>188は統計学を知らんと自分で導けないし、統計学学ぶ香具師がプログラマになるとも限らんしね。
Webでも知識があることを前提としているか、律儀に平均値との差を求める方式しか載ってなかったり。
私の場合は、電卓の統計機能が個別のデータと平均値の差を使わずにどうやって標準偏差を得るのか
不思議に思って調べて以来の知識かな。
ってことで、統計学を知らない私も>188みたいな話の解説があるなら見たい希ガス。
258:デフォルトの名無しさん
08/03/16 17:05:10
CUDAというかGeforceってSIMDなの？
SSEは確かに1つの命令で4つとかの足し算が行われるけど、
CUDAの場合、どれがそれに当てはまるの？
259:253
08/03/16 17:17:24
>>254>>255
早速のご回答ありがとうございます。
基本的にはスレッド数は目一杯であとはブロックすうを調整すると言うことでいいみたいですね。
ところで、処理する数がスレッド数×ブロック数で割り切れない場合、最後のループはどうするのがいいのでしょうか？
やはりifでやるしかないのでしょうか？
260:255
08/03/16 17:38:35
>>259
んにゃ、私はブロック数もスレッド数も実測で決定している。
で、割り切れない場合はループ数で制御するのがいいみたい。
# やっぱり説明するためにはさんぷるが必要かw
>>258
GPUの場合は、SIMDじゃなくてMIMDということになるのかな。
CUDAで考えると、Warp単位で同じインストラクションを走らせて、
複数のデータスレッドを処理するイメージ。
だから、条件分岐をCUDAで書いても、Warp単位で同じインストラクションなので
データスレッドごとに条件が違うと無駄が生じてしまう罠。
詳細は手ごろなサンプルがないので割愛w
261:253
08/03/16 18:08:52
>>260
やっぱり基本は実測なのですね。
ところでループ数で制御するというのはどういうことなのでしょうか？
処理する数/(スレッド数×ブロック数)=ループ数
だと思うのですが・・・
262:255
08/03/16 18:27:15
例えばこんな感じ。
__global__ void func(float const * a1, float * a2, unsigned num)
{ // 実際には一時変数を使った方がいい希ガス
for (unsigned idx = blockIdx.x * blockDim.x + threadIdx.x; idx < num; idx += gridDim.x * blockDim.x) {
a2[idx] = a1[idx] * a1[idx];
}
num=ブロック数*スレッド数ならどのスレッドでもループは一回だけ回る。
num<ブロック数*スレッド数なら、暇なスレッドが発生する(から効率は宜しくない)。
num/(ブロック数*スレッド数)が11と12の中間なら一部のスレッドは12回回って残りは11回回る。
このループ回数が1や2じゃなければ、暇なスレッドの割合が相対的に少なくなる寸法。
# これが理由で、ブロック数は無闇に増やせばいいというわけでもないということになる。
あー説明が下手な私（：；
# サンプル作ったら誰か買わない?w
こういう回し方をすることによって、近くのスレッドが近くのメモリをアクセスする状態のままループが進行する。
つまり、ProgrammingGuideで言うところの"coalesced"。これをプログラミングガイドでは「結合した」と訳しているけど……
263:253
08/03/16 20:27:25
>>262
ありがとうございます。
なるほど非常にうまい手ですね。
プログラミングガイドはなんか日本語訳がめちゃくちゃで非常に読みにくいですが、
がんばって解読したいと思います。
264:デフォルトの名無しさん
08/03/16 20:55:07
shared memoryを使ったら、global memoryを直接使うより遅くなってしまいました。
どうも、カーネルを実行する際のshared memoryの確保に時間がかかっているようなのですが、
カーネルを実行するたびに毎回確保するのではなく、常時確保するような方法はないでしょうか？
265:255
08/03/16 21:20:30
>>264
その現象は確認していないので不明です。
が、保持する方法は多分ないでしょう。
実は一つ問題があって、shared memoryは16バイト同時アクセスができないので
その点でglobal memoryより遅いのは確認済みです。
# ptx出力を見れば見当がつきますが。
266:デフォルトの名無しさん
08/03/17 00:53:35
>>261
もれは、CUDAに投入する前に、CPU側で後詰めゼロとかやって、
「CUDA内では一切条件判断は不要」を原則にしています。
>>264
global、sharedへのアクセスがcoalescedじゃないと、めがっさ遅くなります。
それと、sharedに確保した後で、ブロック内の各スレッドでデータを
shareしない場合には、そもそもsharedを使う理由無いです。レジスタ
に順番にグローバルから読んで処理すればいいです。
267:デフォルトの名無しさん
08/03/17 06:10:01
ブロック間の同期をとる方法を教えてください
268:デフォルトの名無しさん
08/03/17 09:17:39
「ブロック」は、そもそも同時にブロック001とブロック002が動いているか
どうかも分からないんで、同期出来ないんでは？
CPU側で、カーネル関数その１呼び出して、cudaThreadSyncronizeして、
カーネル関数その２呼び出す。ではないかと。
269:255
08/03/17 09:22:38
一応、__syncthreads()という手はありますよ。
但し、全ブロックが同時に動いているわけではない点には注意。
つまり、上の方に書いた分割方法で書いたような大きなブロック数の時には
恐らく巧くいかないかかなり待たされることになるかも知れず。
私の手元のロジックでは、ブロック数が64の時には巧いこと同期が取れています。
# C870, 8800GTX, GTS, GTでしか確認してないけど。
270:デフォルトの名無しさん
08/03/17 14:41:06
>>269
それって各ブロック内のスレッドの同期じゃないの？
271:255
08/03/17 14:55:21
あー、そうそう。失敬、その通りです。
>269は全て、「ブロック」を「スレッド」に読み替えてください。
# どっかでブロック数に切り替わってしまった……_/￣|○
と言うことで、やっぱりブロック間の同期は取れないのでした。
272:デフォルトの名無しさん
08/03/17 15:29:11
>>268
カーネル2はカーネル1が終了してから実行されるんじゃないの？
273:デフォルトの名無しさん
08/03/17 15:56:55
>>272
そこが良く分からんのです。
サンプルではcudaThreadSyncronizeでカーネル呼び出しを挟んでるんで、
一応もれも呼び出すことにしておるのですが。
274:デフォルトの名無しさん
08/03/17 16:04:25
>>273
今、挟んだのと挟んでないのを比べてみたけど、全く同じ結果だった。
ifがあるので、各ブロックはばらばらのタイミングで終わるはずなのだが。
275:255
08/03/17 16:10:18
前にも書いたかもしれないけれど、カーネル終了を明示的に知る手段はないので注意。
つまり、<<<>>>による呼び出しから復帰しても、GPUの処理が終了したわけではないので。
なので、次のような呼び出しをした場合、k1()とk2()は並列に動く可能性があるわけ。
--
k1<<<1, 8>>>();
k2<<<1, 8>>>();
--
終了を保証するには、cudaThreadSynchronize()を呼ぶかcudaMemcopy*()を呼ぶかすればOK。
276:デフォルトの名無しさん
08/03/17 19:44:31
カーネルのループの段数を増やしたら、スレッド数が512だとカーネルの実行ができなくなった。
256ならできる。257はだめ。なんでだろう・・・
277:デフォルトの名無しさん
08/03/17 19:51:18
1BlockのThread数の制限512ってのと関係があるのでは？
278:デフォルトの名無しさん
08/03/17 19:59:35
ちょっと舌足らずだったかもしれませんが、カーネル内のループの段数を増やしたところ、そのカーネルが実行できなくなってしまいました。
そこで、スレッド数を512から256に減らしてみたところうまくいきました。
スタックか何かがあふれてしまったのではないかと疑っているのですが・・・調べる方法がわからず行き詰まってしまいました。
279:デフォルトの名無しさん
08/03/17 20:10:15
実行できないってことは何かエラーが出るの？
280:デフォルトの名無しさん
08/03/17 23:59:13
使っているレジスタの総数が8192を超えたとか。
unrollしすぎだとか。
281:デフォルトの名無しさん
08/03/18 05:17:35
カーネルを2つ同時に実行できますか？
2枚のカードでそれぞれ別のタスクを割り当てて、並列で実行したいのですが。
282:デフォルトの名無しさん
08/03/18 07:49:07
ブロック数が少なければ、勝手に同時実行するな。
283:デフォルトの名無しさん
08/03/18 08:14:14
Maximum memory pitchって何？
284:デフォルトの名無しさん
08/03/18 08:33:44
VCのnmakeって使えないの？
cudaをコンパイルできるサンプルMakefileがあったら教えてください。
285:デフォルトの名無しさん
08/03/18 11:56:28
>>281
サンプルのMultiGPU、8800GTS＋8400で試したら動いたよ。
単にこれ、WINAPIのCreateThreadして中でカーネル関数呼んでるだけですな。
でWaitForMultipleObjectsでWindowsスレッドの終了待っている。
>>283
2次元配列でGPUのメモリを使うときの幅じゃないすか。
cudaMemcpy2Dとか使うときに気にするんだと思われ。
ごめん二次元配列使ってないので良く分からない。
8800GTSも8400も262144bytesです。
286:デフォルトの名無しさん
08/03/18 14:42:22
演算性能がボトルネックになっているのか、バンド幅がボトルネックになっているのか知る良い方法はないものか
287:デフォルトの名無しさん
08/03/18 15:15:30
>>286
演算抜きでcudaMemcpyのHostToDevice、DeviceToHostだけやって
時間を測って見る、そして演算入りの時間と比べる、でどうだろう。
288:デフォルトの名無しさん
08/03/18 16:42:29
__constant__ floatとconst floatの違いって何？
__constant__の方が速いの？実測してみたけど、全然違いが出ないのだけど。
ちなみに何故か定数を直接、式の中に書いたら唯一5％ほど遅くなった。
289:デフォルトの名無しさん
08/03/18 18:11:38
プログラミングガイドの5章が全然わかんねー
結局のところcoalescedって何よ
Figure 5-2の右側は何でダメなの？？？？
教えて偉い人
290:デフォルトの名無しさん
08/03/18 23:54:19
>>288
__constant__は全部で64Kバイトある定数格納用メモリ、ホストから
memcpyすることが出来る。const floatは8192個のレジスタを使って
処理される。なので、普通のコード内定数はconst、ホストからどさ
っと書き込みたい固定データは__constant__でいいと思い。
>>289
5-2の右側は、132=33×4、だから？ 各スレッドは、
16×sizeof()×何か＋スレッドID、にアライメントされた
アドレスにアクセスしないとcoalescedしてくれないと。
291:289
08/03/19 04:35:23
>>290
coalescedを辞書で調べても「合体した」とかしか出てこないので、いまいち意味がよくわからないのですが、
結局0番スレッドは
a[0]
a[16]
a[32]
:
１番スレッドは
a[1]
a[17]
a[33]
:
以外のアドレスにアクセスしないようにすればよいと言うことでいいのでしょうか？
292:デフォルトの名無しさん
08/03/19 04:51:32
要は、配列のアクセスはarray[m * blockDim.x + threadIdx.x]のような形で
blockDim.xは16の倍数になればよろしって話なんじゃないかと。
293:デフォルトの名無しさん
08/03/19 11:08:25
cuMemGetInfoフリーメモリと合計メモリを取得したのだけど、４MBにしかならない。
使い方は
unsigned int f,m;
cuMemGetInfo(&f,&m)
でいいんだよね？
294:デフォルトの名無しさん
08/03/19 11:35:28
>>291
つまり「coalesced」は・・・もれの理解では・・・
・sharedメモリには入出力が32bit×8ポートずつあるかも。
でも開始アドレスは一個かも。
・各スレッドからのアクセス要求アドレスが「0,4,8,12,16,20,24,28」の場合、
GPU側で「ではsharedメモリの出力を綺麗にストリーミングプロセッサに
一個ずつ割り当てればいいのであるな、開始アドレスは0だな」と判断し
てくれるのかも。
>>293
もれの場合は、debugでは動くけどReleaseでは動かないのね。
debugでは FreeMem 465529088 TotalMem 536543232
と取れた。8800GTS/512MB。理由分かったら教えてください。
295:デフォルトの名無しさん
08/03/19 19:26:24
>>290
__constant__って書き換えられるの？
296:デフォルトの名無しさん
08/03/19 22:06:01
>>295
__constant__は定数メモリと考えればいい。
つまり、GPUからは書き換えられないのでCPUから書き換えると。
# constな変数は初期化はできるけど代入できないのと似てるかもしれない。
>>294>>293
>294のdebugでの数値は多分正しい値が出ているみたいだね。
>293の方は、もしかしたらエラーが返っていないかな?
私のところでも、整数にして201だったかのエラーが返ってきていたのだけれど、
その状態では値をセットしてくれないようだった。
つまり、もしかしたら>293の4MBとというのはfとmが未初期化で偶々入った不定値かも。
# 気になるなら、0で初期化してみればいい。
>>294>>291
coalescedなアクセスパターンにすることは、sharedだけでなくglobalでも重要なので
慣れておくことをお勧め。まぁ、例えば>262に書いたようにすればいいだけだけど。
297:デフォルトの名無しさん
08/03/19 23:54:01
>>295
cudaMemcpyToSymbol()でCPUからセット出来るよ！
速いし64KBあるんで結構便利です。
298:デフォルトの名無しさん
08/03/20 00:01:09
問題は、ベクタアクセスできる__global__と違って1データずつのアクセスになるから
逆に遅くなるケースもあることだな。
299:デフォルトの名無しさん
08/03/20 12:55:05
ややこしさはCellとあんまり変わらなそうだな。
300:デフォルトの名無しさん
08/03/20 14:25:54
>>299
とにかくメモリのコピーがヤヤコシイんだわ。ホスト側、GPUの
グローバルメモリ、GPUのチップ内メモリで転送しまくらないと
いかんで。やり方間違えると全然性能出ないし。
Cellはその辺どうなの？
301:デフォルトの名無しさん
08/03/20 14:28:32
256kしかないLSでやりくりするのが大変って聞いたな
302:デフォルトの名無しさん
08/03/20 15:10:07
両方試した私に言わせて貰えば、どっちもどっち。
確かにCBEは256KiBの壁がねぇ。GPUも64KiBの壁やcoalescedの沼があるけど。
超越関数を使える点ではCUDAが有利。ホストの性能でもPPEじゃ結構泣けるし。
303:デフォルトの名無しさん
08/03/20 16:02:10
PPEはひどいよな。ホステスにC2D使ってるけど、ifがいっぱいあるような場合はC2Dの方が速いしね。
304:デフォルトの名無しさん
08/03/20 16:14:13
>>298
それ実験してみたんだけど、変わらないみたい。
・global→shared(行列多数)と__constant__に置いた定行列で行列積
・global→shared(行列多数)とsharedの一部に置いた定行列で行列積
で、後者が1%遅いくらいだった。リードオンリーなだけで、コアとの距離や
所要クロック数はconstantもshared・レジスタも同じなのかもと？
305:デフォルトの名無しさん
08/03/20 16:28:58
CUDAで逆行列を求める方法を教えてください
306:デフォルトの名無しさん
08/03/20 16:49:55
page-lockedってどういう状態を表すのでしょうか？
307:298
08/03/20 17:42:11
>>304
-ptxでニモニックを出力してみれば違いが判るかと。
例えば、ld.global.v2.f32はあるけどld.const.v2.f32はないからld.const.f32が2回になってしまう。
coalescedなアクセスができるglobalは4クロックでアクセスできるからsharedやconstと変わらないわけで。
尤も、一旦そのパターンから外れるとglobalは数百クロックだそうだから途端に劇遅になるけど。
308:298
08/03/20 17:48:04
>>305
先ずは、逆行列を求めるプログラムをCで書いてみてください。
それの、ループの部分を分割する方向でGPU関数を作るのが第一歩になります。
つーかさぁ、一行質問する人達ってなんなの?
情報交換する気がなくて、単にくれくれの精神なんだったら勘弁して欲しいんだけど。
# なんか回答者が数人とそれより少し多いレス要員だけで持ってる気がするよ……
>>304
あーそうそう、書き忘れたけどレジスタアクセスは1クロックじゃなかったカナ。
確実にメモリよりも速いみたい。
309:デフォルトの名無しさん
08/03/20 19:43:17
nvccにはO2とかの最適化オプションは無いのでしょうか？
310:デフォルトの名無しさん
08/03/20 20:26:00
-Wallが使えないみたいね
たぶん無さそうな予感
311:298
08/03/20 21:19:57
一部のオプションはgccと共通ですよ。
例えば、-O3(恐らく-O2なども)や-pgは使えます。
312:デフォルトの名無しさん
08/03/20 23:15:10
>>299
どっちもどっちで、用途や慣れの問題だと思う。302がいうように、SFUの有無は
大きいが、それも用途しだい。
自分的には比較的情報がオープンなCellのほうがいろいろいじりやすいと思うが、
最近人気なくて心配。たぶんそのうちLarrabeeが全部かっさらってくんじゃないかな。
313:デフォルトの名無しさん
08/03/22 18:23:49
GeForce 9800 GX2で現バージョンのCUDAは動きますか？
314:デフォルトの名無しさん
08/03/22 18:28:28
ストリーム番号の同じカーネルやメモリコピーは順番に、違うものは並列に実行される？？の認識でいい？？？
315:デフォルトの名無しさん
08/03/22 18:43:27
テクスチャメモリの容量はどのようにすれば調べられるのでしょうか？
どこかに資料があるのでしょうか？それとも何かAPIがあるのでしょうか？
316:デフォルトの名無しさん
08/03/22 23:25:42
いつも適当に答えてるおじさんです。まことにもうしわけない。
>>313
もれ、8800GTS(G92)で動いてるから、×２で二個見えるんじゃないかと
予想しますけど、一応どっちも対応外だから。漢は度胸でひとつ人柱に。
>>314
ストリームが二つあっても、実行順は関係ないのではと予想。ストリーム
１の半分実行した後ストリーム２の半分実行、もありえるかと。
cudaStreamSynchronizeでストリームの実行完了を待たないといけないです。
>>315
テクスチャメモリはグローバルメモリのキャッシュですが、実サイズが
不明ですね。cudaGetDevicePropertiesでは「alignment=256Bytes」と
読めます。これ以下と以上で実行時間を調べてみると、キャッシュミス
してるかしてないかわかるのではないでしょうかと。
317:デフォルトの名無しさん
08/03/23 13:19:47
openSUSE 10.3でCUDAは動きますか？
318:デフォルトの名無しさん
08/03/23 13:33:44
動くらしい
319:デフォルトの名無しさん
08/03/23 17:02:07
次のバージョンっていつ頃出るの？
新しいカードも出ているわけだしそろそろ、SDKの方もアップデートしてほしいものなのだが
320:デフォルトの名無しさん
08/03/23 23:52:41
Vista64と9600,9800対応版が欲しいすね。
早く出してくれないとAMD/ATIに浮気したくなりますよ。
321:デフォルトの名無しさん
08/03/24 10:55:41
サンプル作ると言いつつなかなか暇が取れない私が来ましたよ。
>>313
ドライバが対応しているかどうかが問題。ドライバが認識さえすれば、動くとは思いますが。
>>314
ストリームのご利用は計画的に。ちなみに、CUDA1.0世代の(要はG8xかな)GPUだと
処理の隠蔽ができないらしく、全く速くならないという噂もあります。
>>317
>27-30
荒らしじゃないんだったら、スレくらい検索してほしいと思う今日子の頃。
# いや、荒らしなら検索如何に関わらずご遠慮願いますが。
>>319
取り敢えず、待つしか。下手すると夏の新製品までお預けの可能性も。
>>320
え～、待っていれば倍精度も来るのに。
322:デフォルトの名無しさん
08/03/26 23:04:49
英wikiでは、9600GTでもcudaできることになっている
悪い子と漢は居ね゛ーがー
323:デフォルトの名無しさん
08/03/26 23:38:52
>>322
もれ9600GT買おうと思ってる。
ARCTICのS1Rev2が付く、という報告待ち・・・。ファンレスで使いたくて。
324:デフォルトの名無しさん
08/03/31 21:17:49
URLﾘﾝｸ(www.nvidia.com)
9800GX2 CUDA OKだとでてるよん
325:デフォルトの名無しさん
08/04/01 00:04:20
おー、更新されている。情報THX。
ドライバは変わっていないから、チップ依存なんですかね。
326:デフォルトの名無しさん
08/04/04 22:15:16
ML115祭りに参加してOpetron使用をポチりました。
で、nbody.exeを実行すると・・・・ 160～170GFLOPSという値が。
うぅ・・速過ぎる。初めて見るそのパワーに感動ものですた。
327:デフォルトの名無しさん
08/04/04 22:15:54
忘れてた GPUは8800GTですぅ。
328:デフォルトの名無しさん
08/04/04 22:24:35
そろそろまじめに実用的実装に入らないのかな？
圧縮解凍ソフトに組み込むとか
暗号化ソフトに組み込むとか
一番簡単なのは画像処理ソフトだろうけど
329:デフォルトの名無しさん
08/04/05 02:08:09
大真面目に実用的な実装で悪戦苦闘していますが何か。
330:デフォルトの名無しさん
08/04/05 09:56:17
画面の解像度によって使えるメモリ量は変わりますか？
331:デフォルトの名無しさん
08/04/05 12:08:56
CUDA使うアプリをインストーラで配るにはどうしたらいいもんだろうか？
というところが。Vista対応してないし。
332:デフォルトの名無しさん
08/04/05 13:45:15
CALならドライバ要らずなんだけどなぁ。
なんで別途ドライバが必要なんだろ。
つかVista版ドライバは3月の終わりごろじゃなかったのかよ！
333:デフォルトの名無しさん
08/04/05 15:56:48
XP用のビデオドライバをインストールできるはずだから、その状態でXP用のCUDAドライバ入れられれば動くんじゃないかな
と無責任なことを言ってみる
334:デフォルトの名無しさん
08/04/05 16:12:49
>>332
たぶん時差
335:デフォルトの名無しさん
08/04/05 19:34:29
URLﾘﾝｸ(japanese.engadget.com)
cudaなんかやってる場合じゃないよな
336:デフォルトの名無しさん
08/04/05 19:52:25
戦犯はマイクロソフトだろ･･･
337:デフォルトの名無しさん
08/04/05 20:08:59
黒歴史確実なOSだからスルーでいいよ
338:デフォルトの名無しさん
08/04/05 20:15:04
え、俺パフォーマンス100倍UPにちょー期待してんだけど
339:デフォルトの名無しさん
08/04/05 22:09:08
普通にopenGL使ったオープンソースのキット使ったほうがいいだろ
340:デフォルトの名無しさん
08/04/05 23:52:39
今月リリース予定のCUDA 2.0ベータ版でVistaサポートだってさ
341:デフォルトの名無しさん
08/04/06 00:55:39
CUDAを有効に使うのなら、gccとの相性から考えてもLinuxで使うだろ。JK
342:デフォルトの名無しさん
08/04/06 01:32:22
仮想メモリが使ってみたいんだもの。
343:デフォルトの名無しさん
08/04/06 06:31:34
ローカル変数はどこに格納されるのですか？
shared memoryでしょうか？
344:デフォルトの名無しさん
08/04/06 08:35:00
スタックフレームはないようなので、全てレジスタに格納されると思って宜しいかと。
ローカルに大きな配列取ろうとしたらどうなるのかはしらんなぁ。
ptx出力して読んでみたら?
345:デフォルトの名無しさん
08/04/06 09:32:54
>>341
商売にならんじゃないか。建前はどうでもいいんですよ。
売れるかどうかだけです。
346:デフォルトの名無しさん
08/04/06 09:49:44
「商売にできない」の間違いですね。
347:デフォルトの名無しさん
08/04/06 11:51:20
エンコソフトでCUDA使ったやつ売れれば儲かりそうね。
動作検証が非常に大変だろうけど。
Vista対応不可って所で今の所どうにもならないですなぁ。
漏れがいま一番熱望してるのは、XilinxがCUDAに対応して、
FPGAのコンパイル・シミュレーションが劇速になること。
ああ夢のようだ・・・
348:デフォルトの名無しさん
08/04/06 16:11:58
GPUで計算させているのだけど、その間のCPU使用率って100%なんだけど、これって正常なの？
349:デフォルトの名無しさん
08/04/06 16:28:42
正常です。
ブロック数がPE数より多い場合は待たされることがありますし、
メモリ転送や同期を取るときには当然待たされます。
実は待っている間もしっかりCPU時間を消費するのです。
従って、GPUとCPUを巧く連携させて高パフォーマンスを狙うには
いかに待たずに済ませるかが鍵になるので、Streamの使用は必須になってくるかと。
350:デフォルトの名無しさん
08/04/06 16:37:26
くそ速いglReadPixelsとして使えるかと思ったが
CUDAで扱うデータしかWrite Readできないのね
ぐすん
351:デフォルトの名無しさん
08/04/06 16:53:37
CUDAにもGLサポート関数があるから、もしかしたら連携できるんじゃないの?
サンプルの、simpleGL辺りになんかない?
352:デフォルトの名無しさん
08/04/06 19:06:19
PBO使え
353:デフォルトの名無しさん
08/04/06 19:22:37
GT200と99GX2はどっちが高性能でしか？
354:デフォルトの名無しさん
08/04/07 10:28:11
このスレの住人なら知っていますね、あの糞開発ツールのことを
・自分のプログラムのバグなのかコンパイラのバグなのかわからない
・他の仕事に応用できない糞開発ツールの独自世界を必死に学習している
・テキストエディタで書いたほうが効率的なのに糞UIツールを懸命に使っている
・糞開発ツールを批判すると「性格が悪いから糞ツールを批判するんだ」と言われる
糞だけど、政治的な理由で無理やり使わされているんですよね。
もう、あんな厨の作った糞ツールを我慢して使うのはやめましょう。
・糞開発ツールを部下に押し付ける上司の命令は無視しましょう。
上司は糞開発ツールが使われる実績を作ることであの会社のごきげんをとっているのです。
・糞開発ツールを使わせる上司の下では働けません、と上司の上司に直訴しましょう。
・あの糞開発ツール提供会社には「おたくの糞開発ツールは話にならない」と突き放しましょう。
バグレポートなどしてはいけません。改善要求などもってのほかです。
あの会社はあなたたちのことをテスター/モルモットとしか思っていません。
・あの会議で「糞開発ツールを使ったら生産性がxx％アップしました」
なんて話が出たら力強く机を叩き、会議室を出ましょう。
あの人たちは糞開発ツールをマンセーすることで立場を確保しているのです。
糞な開発ツールを糞だと言える、そんな当たり前の環境をみんなの力で取り戻しましょう。
355:デフォルトの名無しさん
08/04/07 14:17:07
GX2はSLIと同じように2つのgridを使う必要があるのでしょうか？それとも、1つのgridで両方のGPUを使えるのでしょうか？
356:デフォルトの名無しさん
08/04/07 14:26:08
CUDAにSLIは関係ありません。
二枚挿しのことをSLIと呼ぶのは間違いです。
何故なら、SLI用コネクタを挿さなくてもCUDAではGPUを二つ使うことができるからです。
で、肝腎な9800GX2ですが予想ではGPUが二つ見える筈なので、一つの時のままだとダメな気がします。
357:デフォルトの名無しさん
08/04/07 15:24:49
2つのGPUを使うには、CPU側もマルチスレッドにする必要があるのですか？
358:デフォルトの名無しさん
08/04/07 16:17:21
Streamを使えばシングルスレッドでも何とかなるんじゃないかと思うので、是非とも試してみてください。
# 因みに、StreamはCUDA1.1からの機能なのでサンプルのmultiGPUでは使っていません。
359:デフォルトの名無しさん
08/04/08 18:48:35
ビデオカード2枚刺しの場合、2枚のビデオカード間のデータのやりとりはPCIe経由なのですか？
360:デフォルトの名無しさん
08/04/08 22:08:50
>>359
CUDAにはチップ間の転送なんてなかったと思いますが。