【GPGPU】NVIDIA CUDA ..
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
上のほうにもあるけど、総和の計算って実際どーすればいいんだ？
スレッドブロックごとに小計を求めて、各ブロックの結果をまとめるとかいう
方法になりそうなのは予想できるが、実際どんな風に書けばいいのかわからん。
具体的にはベクトルの長さが出したいんだが、各要素の２乗の和がうまく出せない。
ヒントでもいいので分かる方教えてください。
次ページ最新レス表示スレッドの検索類似スレ一覧話題のニュースおまかせリスト▼オプションを表示暇つぶし2ch
4325日前に更新/252 KB
担当:undef