- 1 名前：デフォルトの名無しさん mailto:sage [2007/09/17(月) 14:54:28 ]
- 本家
developer.nvidia.com/object/cuda.html
- 2 名前：デフォルトの名無しさん [2007/09/17(月) 14:55:24 ]
- グプグプー
- 3 名前：デフォルトの名無しさん mailto:sage [2007/09/17(月) 21:03:53 ]
- >>1
GPGPUスレと統合じゃだめなんか？
pc11.2ch.net/test/read.cgi/tech/1188374938/
- 4 名前：デフォルトの名無しさん mailto:sage [2007/09/18(火) 10:06:32 ]
- おお、いつの間にこんなスレが。一応あっちで回答しているけどどうしようか。
- 5 名前：デフォルトの名無しさん mailto:sage [2007/09/19(水) 14:34:55 ]
- 削除依頼出しとけ
- 6 名前：デフォルトの名無しさん [2007/09/21(金) 05:49:47 ]
- WindowsのCUDAとLinuxのCUDAで書式に何か違いはありますか？
- 7 名前：デフォルトの名無しさん mailto:sage [2007/09/21(金) 12:09:20 ]
- 削除依頼出しとけ
- 8 名前：デフォルトの名無しさん [2007/09/21(金) 15:15:23 ]
- 三角関数や指数関数はないの？
- 9 名前：デフォルトの名無しさん [2007/09/21(金) 15:25:30 ]
- 削除依頼出しとけ
- 10 名前：4 mailto:sage [2007/09/21(金) 15:59:09 ]
- >>6
基本的に同一。但し、Win用nvcc(コンパイラ)では64bit版にはできないこととオブジェクトモジュールがMS仕様なことに注意。
更に、nvccはgccで実装されているので、C++のnewやvectorはそのままではリンクできなくなってしまうことにも要注意。
>>8
math.hにある、float系の関数は大抵使えると思っていい。エミュレーションモードなら全部使えた筈。
また、sin(), cos()などの一部の関数は組み込み命令(intrinsic)版が使えるので更に高速。
実際の例は、GPGPUスレ参照で。
pc11.2ch.net/test/read.cgi/tech/1188374938/71-72
- 11 名前：デフォルトの名無しさん mailto:sage [2007/09/21(金) 16:19:38 ]
- CUDAとかって何らかの標準化の動きはあるの？
- 12 名前：デフォルトの名無しさん mailto:sage [2007/09/22(土) 14:43:16 ]
- 笛吹けど踊らず
# まぁ、笛吹いている会社の中でも踊らない連中がいるみたいだし
- 13 名前：デフォルトの名無しさん [2007/09/29(土) 08:08:41 ]
- CUDAは無料ですか？？
- 14 名前：デフォルトの名無しさん mailto:sage [2007/09/29(土) 09:38:17 ]
- ソフトは無料、ハードは有料なんじゃね
- 15 名前：デフォルトの名無しさん mailto:sage [2007/09/29(土) 22:08:38 ]
- CUDAなんてハードはありません。
- 16 名前：デフォルトの名無しさん mailto:sage [2007/09/29(土) 22:10:13 ]
- 処で、テスラc870ってQuadroFXの一番上の奴とどこが違うんだ?
- 17 名前：デフォルトの名無しさん mailto:sage [2007/09/29(土) 23:03:54 ]
- >>15
おマイ馬鹿か、環境の事に決まってるだろ
- 18 名前：デフォルトの名無しさん mailto:sage [2007/09/30(日) 02:33:09 ]
- Compute Unified Device Architecture
~~~~~~~
- 19 名前：デフォルトの名無しさん [2007/09/30(日) 14:58:40 ]
- CompUte Driver Architecture
- 20 名前：4 mailto:sage [2007/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 名前：デフォルトの名無しさん mailto:sage [2007/10/12(金) 18:56:43 ]
- >>20みたいな処理って数値演算よりデータ転送の比重が多すぎて、
並列化させる旨みは非常に少ないと思うんだけどどうなのよ。
本家のフォーラムでいわゆるセマフォを実装して集計処理やらせようとしてる人がいたけど、
そういう問題をCUDAでやる事自体がそもそも間違っとる、みたいに言われてた。
かく言う自分も似たような事しようとしてたけどパフォーマンス出ないんで諦めて、
問題の分割方法から見直したら上手くいったりした。
- 22 名前：デフォルトの名無しさん mailto:sage [2007/10/12(金) 18:57:14 ]
- >>20みたいな処理って数値演算よりデータ転送の比重が多すぎて、
並列化させる旨みは非常に少ないと思うんだけどどうなのよ。
本家のフォーラムでいわゆるセマフォを実装して集計処理やらせようとしてる人がいたけど、
そういう問題をCUDAでやる事自体がそもそも間違っとる、みたいに言われてた。
かく言う自分も似たような事しようとしてたけどパフォーマンス出ないんで諦めて、
問題の分割方法から見直したら上手くいったりした。
- 23 名前：デフォルトの名無しさん mailto:sage [2007/10/12(金) 18:58:08 ]
- 二重orz
- 24 名前：4 mailto:sage [2007/10/12(金) 19:29:53 ]
- >>21
>20の処理は、最初に定数テーブルを仕込んでおいて幾つかパラメータを渡して計算させるのがメイン。
>20に書いているのは最後の集計部なんだけど、結果はarとaiに入るのでそれを転送するだけ。
集計をGPUにやらせられないとなるとtmpを転送しなければならないので、それをなんとかしたかったわけ。
で、未だその後修正掛けていないけど、begin, endのペアをidxの値の分だけを定数テーブルに
仕込んで置けることが判ったからなんとかなりそう。そうすればidxの値の分だけ並列に走らせられることだし。
- 25 名前：デフォルトの名無しさん mailto:sage [2007/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 名前：デフォルトの名無しさん mailto:sage [2007/10/16(火) 23:58:44 ]
- GPUでテクスチャの全てのピクセル値の合計を求めたいってこと？
1x1サイズのミップマップ作って、ピクセル数を掛けるのはどう？
速いかどうかは知らんけど。
- 27 名前：デフォルトの名無しさん mailto:sage [2007/10/17(水) 06:58:08 ]
- opensuse 10.3でCUDAは動く？
- 28 名前：デフォルトの名無しさん mailto:sage [2007/10/19(金) 10:01:25 ]
- Linuxサポートしてるから動くんじゃね？
- 29 名前：デフォルトの名無しさん mailto:sage [2007/10/22(月) 10:32:10 ]
- >>27
ドライバが載らなくても取り敢えずエミュレータは動くでしょうね。
つーか、本家に10.1用と10.2用は置いてあるから試してみては?
ttp://developer.nvidia.com/object/cuda.html
- 30 名前：デフォルトの名無しさん mailto:sage [2007/11/20(火) 15:20:50 ]
- >>27
もしかしたら、標準でインストールされるgccのバージョンが違う所為で動かないかも知れない。
その場合、10.2にインストールされているバージョンのgccをインストールしておけば大丈夫。
# 要はnvccが内部でgccを使っているから、gccのバージョンに依存していると。
# ドライバはOKだった筈。
- 31 名前：デフォルトの名無しさん [2007/11/25(日) 09:41:58 ]
- GeForce 8400GSの安いカードでもCUDAを試すことはできますか？
- 32 名前：デフォルトの名無しさん mailto:sage [2007/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 名前：デフォルトの名無しさん mailto:sage [2007/11/25(日) 18:24:52 ]
- >>31
8800を買うのが無難
- 34 名前：デフォルトの名無しさん mailto:sage [2007/11/25(日) 20:07:13 ]
- 安く済ませるなら、(消費電力の少ない)8600GTもお勧め。
第2世代(と思われる)8800GTは割とパフォーマンスがよさそうだけどね。
- 35 名前：デフォルトの名無しさん mailto:sage [2007/11/25(日) 20:07:50 ]
- >>31
公式にはサポートされていない。
でも、ノートパソコンで動いたという人もいるくらいだから、動く可能性がないわけではないと思う
- 36 名前：デフォルトの名無しさん mailto:sage [2007/11/28(水) 01:39:57 ]
- 8400Mは対応してたはずだけど、普通の8400の方は対応していないんだよね
自作板でサンプルが動いたという報告があったような気みするけど、
対応していない製品を使うのは心配だよね
- 37 名前：デフォルトの名無しさん mailto:sage [2007/11/29(木) 23:46:49 ]
- 8600買ってきたー
試したいのだけど、まず取っ掛かりはどこからいけばいいですか＞＜
- 38 名前：デフォルトの名無しさん mailto:sage [2007/11/30(金) 00:20:36 ]
- >>37
>29にある本家から一式拾って、ドライバとツールキットとSDKを入れて、
サンプル一式をビルドして動かしてみるよろし。
- 39 名前：デフォルトの名無しさん mailto:sage [2007/12/05(水) 18:31:17 ]
- 64bit 版 ubuntu7.10にインストールしたいのですが、どれをダウンロードすればよいのでしょうか？
- 40 名前：デフォルトの名無しさん mailto:sage [2007/12/05(水) 18:48:07 ]
- ttp://developer.nvidia.com/object/cuda.htmlかttp://developer.nvidia.com/object/cuda_1_1_beta.htmlから
x86-64のDriverとx86-64のToolKitとSDKを拾えばいいとして。
ToolKitはどれが使えるかは不明。nvccがgccを使うから、インストールの状況によるので単純じゃない模様。
- 41 名前：デフォルトの名無しさん mailto:sage [2007/12/07(金) 08:32:31 ]
- ubuntu版きてるーーーー
- 42 名前：デフォルトの名無しさん mailto:sage [2007/12/07(金) 12:33:05 ]
- 7.04ってなんだYO
- 43 名前：デフォルトの名無しさん mailto:sage [2007/12/07(金) 13:31:11 ]
- タイムリーなんだか間が悪いんだかw
- 44 名前：デフォルトの名無しさん mailto:sage [2007/12/08(土) 02:28:34 ]
- NECのVALUESTARを使ってるんですが
nVIDIAをインストールしようとすると
NVIDIA setupプログラムは、現在のハードウェアと互換性のあるドライバ
を見つけることができませんでした。
と表示され、インストールができません。
何が原因なんでしょう？
OSはVistaです。
- 45 名前：デフォルトの名無しさん mailto:sage [2007/12/08(土) 08:33:55 ]
- >>44
NECのVALUESTARなんてどうでもいい情報よりもグラボの情報だせよん
- 46 名前：デフォルトの名無しさん mailto:sage [2007/12/08(土) 12:14:51 ]
- なぜココで聞く。
ハードウェアスレに聞け
- 47 名前：デフォルトの名無しさん mailto:sage [2007/12/10(月) 13:04:06 ]
- >>44
そもそもVista対応のCUDA対応ドライバがありません。
- 48 名前：デフォルトの名無しさん mailto:sage [2007/12/13(木) 00:22:06 ]
- >>44
>nVIDIAをインストールしようとすると
って時点で意味がわからんのだが、
とりあえずデバイスエミュレータ(-deviceemu)でどうよ？
- 49 名前：sage [2007/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 mailto:sage [2007/12/17(月) 15:07:07 ]
- 自己解決しました！
- 51 名前：49 mailto:sage [2007/12/17(月) 16:13:39 ]
- 勘違いでした
誰かヘルプ。。。
- 52 名前：デフォルトの名無しさん mailto:sage [2007/12/17(月) 18:50:43 ]
- 何がどう巧くいってないと判断したのかよく分かりませんが。
cudaMallocPitch(), cudaMemcpy2D()じゃなくて、cudaMalloc(), cudaMemcpy()では巧くいったの?
それと、テストしたソースは丸ごとコピペするか、どっかにアプロドして欲しい。
- 53 名前：デフォルトの名無しさん mailto:sage [2007/12/18(火) 11:25:32 ]
- ＞アプロド
なんか可愛いw
- 54 名前：デフォルトの名無しさん mailto:sage [2007/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 名前：デフォルトの名無しさん mailto:sage [2007/12/18(火) 12:14:10 ]
- >>49
＞float* devPtr,an;
このanは只のfloat型。
cudaMemcpy2D()の第1,3パラメータのキャストもおかしい。
もしサンプルにそう書かれていたのなら、そのサンプルがタコ。
- 56 名前：デフォルトの名無しさん mailto:sage [2007/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 mailto:sage [2007/12/19(水) 13:56:13 ]
- 表をそのまま貼ったから見難くなったのは勘弁。
補足しておくと、FX/Plex/NVSはどれもQuadroシリーズ(非OEM)ね。
細かく見てたら、8800GTが追加された他に8800GTSに512MBが追加されている。
しかも、クロックが既存の8800GTSよりも速いだけでなくShaderClockでは最速に。
MemoryInterfaceが512bitだから、8800GTSというより8800GTの系列かしら。
まぁこれで、旧8800GTS(640MB/320MB)は消える方向なんでしょうね。
- 58 名前：デフォルトの名無しさん mailto:sage [2007/12/22(土) 15:09:32 ]
- ほほー
廉価版でCUDA試してみたいだけの奴は8400あたりでもOKということか。
ありがたやありがたや
- 59 名前：デフォルトの名無しさん [2007/12/31(月) 18:53:34 ]
- CUDAを使用できるGPUのうち、
GeForceとQUADROのどちらの方がCUDAに最適なのでしょうか？
- 60 名前：デフォルトの名無しさん mailto:sage [2008/01/03(木) 00:07:49 ]
- >>59
「最適」の基準によります。
入手性、価格、バリエーションではGeForceシリーズに文句なしで軍配が上がりますが、
個別の性能はそれぞれ異なりますので。
因みに、TeslaはQuadroFX5600(-アナログ回路)相当だそうです。
あー、QuadroPlexを買える待遇なら、違う意味で最適かもしれません。
- 61 名前：デフォルトの名無しさん mailto:sage [2008/01/07(月) 12:11:34 ]
- 処で、8800GTXでどこまで速度が出るのか実測してみた。
単純な実数の足し算をループで回しただけなんだが、約43GFlops出た。
一方、ShaderClockは1350MHz、プロセッサ数は128、ということなので計算上は(1クロック1演算なら)172.8GFLOPSとなる。
命令セット上は、積和が1クロックで実行できるから2倍して約350GFLOPSだからカタログスペックはこの値なのだろう。
実際には、ループを構成するには実数の足し算の他に整数加算、整数比較、条件ジャンプの3クロックが必要になる。
従って、約43GFLOPSとなって実測値と見事に一致した。
というわけで、(メモリアクセスが無視できるなら)アセンブリ(ptx)出力から所要時間を割り出すことができる模様。
問題は、メモリアクセスが絡むとどれだけ遅くなるかで、こちらの方は今後調査の予定。
- 62 名前：デフォルトの名無しさん mailto:sage [2008/01/07(月) 21:38:50 ]
- ニーモニックて公開されてるんだっけ？ or 固定長？
- 63 名前：デフォルトの名無しさん mailto:sage [2008/01/07(月) 23:24:57 ]
- 5秒制限とかあるらしいけど
これどうやって回避するんだろう
- 64 名前：デフォルトの名無しさん mailto:sage [2008/01/07(月) 23:36:08 ]
- >>62
ptx出力を睨めっこすれば大体当たりがつく。どっかで公開されてるかも知れんが。
>>63
5秒も計算させたまま帰ってこないような組み方すると、どうせ遅くなるから余り気にならない。
単純なループだけならXeonの方が速いからね。
- 65 名前：デフォルトの名無しさん mailto:sage [2008/01/10(木) 15:13:03 ]
- >>61
おつー
これは目安に使えそうだね
- 66 名前：デフォルトの名無しさん mailto:sage [2008/01/11(金) 11:30:09 ]
- ttp://www.nvidia.com/object/cuda_home.html
CUDAのサイトリニューアルされてるね。
相変わらずfor Vistaはまだだけど。
- 67 名前：デフォルトの名無しさん mailto:sage [2008/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 mailto:sage [2008/01/11(金) 15:23:14 ]
- ずっと旧トップページをブクマしてて今日やっと転送されるようになったので
気づいてなかったｽﾏﾝｺ
1.0になってから遊んでいなかったので久しぶりに遊んでみよう。
- 69 名前：デフォルトの名無しさん mailto:sage [2008/01/11(金) 22:18:42 ]
- >>67
メモリ転送は、PCI-Expressの性能で上限があるはず。
たぶんこれがかなり時間食ってると思う。
4K×4kは画像なんでしょうか。大きいですね。
- 70 名前：デフォルトの名無しさん mailto:sage [2008/01/12(土) 00:56:15 ]
- PCIe 1.1と2.0でどれぐらい差が出るのか気になりますね。
3DMarkのベンチマークやゲームではほとんど差が無いみたいですが
CUDAだとやはりそのあたりがボトルネックになる場合も出てくるでしょうかね。
またメモリ容量に依存する場合もやはり多いのでしょうか？
- 71 名前：67 mailto:sage [2008/01/12(土) 06:35:30 ]
- >>69
FFT解析絡みのことをやっているので、周波数空間像を拡げてから逆FFTなんてことをしばしば。
まぁ、4096x4096は単にテストするのが目的でしたが。
>>70
未だPCIe2を試せる環境がないからなんとも言えませんが、演算の種類によっては影響出るでしょうね。
私のところのプロジェクトでは寧ろ、GPUとCPUを巧く並行動作できるかどうかの方が影響大きいのですが。
- 72 名前：デフォルトの名無しさん mailto:sage [2008/01/18(金) 11:50:14 ]
- むぅ、周波数合成みたいなロジックをCUDAに移植したんだが、余り速度が出ない。
元のソースが三角関数テーブルを使っていたのでメモリアクセスが足を引っ張っているのかと思って
__sincosf()を使ってみたが、今度は何故か実行ごとの所要時間のばらつきが大きい。
速いときはXeon並の速度になるのに遅いときは3倍くらい時間が掛かる。
はてさて、超越関数ユニットの数が少ないのだろうか……
- 73 名前：デフォルトの名無しさん mailto:sage [2008/01/18(金) 15:08:32 ]
- 超越関数は級数近似と思われ
see /cuda/include/math_functions.h
- 74 名前：デフォルトの名無しさん mailto:sage [2008/01/18(金) 17:05:57 ]
- いえいえ、ちゃんとニモニックもあるのですよ。そっちを使うとドキュメントに拠れば32クロックだそうで。
但し、そいつが「To issue one instruction for a warp,」とあるからぶつかるのかなと。
- 75 名前：デフォルトの名無しさん mailto:sage [2008/01/21(月) 00:26:07 ]
- 8800前後の機種の日本での売れ行きがかなり芳しくないようだがこれは何を意味するのだろう．本国でのアメリカでは使用例（論文）が着実に出つつあるから，なんだかそいつらの為だけに開発されてる感じじゃない？？
- 76 名前：デフォルトの名無しさん mailto:sage [2008/01/22(火) 23:13:36 ]
- 8800系って、日本ではNEETゲーマーが買い漁るから
1$=200円くらいのボッタ栗値段が付いてしまって、
貧乏ラボじゃ買えん罠。箱もヲタ臭くてボスの目
が気になる。
- 77 名前：デフォルトの名無しさん mailto:sage [2008/01/22(火) 23:24:09 ]
- つ [玄人志向]
- 78 名前：75 mailto:sage [2008/01/23(水) 00:48:46 ]
- そうか，研究者の間であまり大量に出回っていないと思ったが，そういう事だったのか．＞76
国内で一番CUDA使いこなしているのはどこのチームだろう？
- 79 名前：デフォルトの名無しさん mailto:sage [2008/01/23(水) 01:17:56 ]
- まぁ、CUDAでXeonに勝つ速度を叩き出している漏れが一番だねw
- 80 名前：デフォルトの名無しさん [2008/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 mailto:sage [2008/01/23(水) 16:03:38 ]
- >80を一部訂正。
・但し、冪乗関数はexpf()は1.4427を、exp10f()は3.3219を、それぞれ掛けて(mul.f32)からex2.f32を使うようだ。
- 82 名前：デフォルトの名無しさん mailto:sage [2008/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 名前：デフォルトの名無しさん mailto:sage [2008/01/29(火) 04:47:39 ]
- 3次元格子は扱っていないから一般的な方法は知らん。
で、>82で流れはあっていそうだけど効率は悪いと思うよ。
ブロックごとにdVelをSharedにキャッシュした方がいいかも知れない。
# まぁ、動くものを作る方が先だけど。
- 84 名前：デフォルトの名無しさん mailto:sage [2008/01/29(火) 09:10:24 ]
- >>83
ありがとうございます
吐き出される数値を見て、ぜんぜん計算されてなさそうなので、
心配をしていました。
効率悪いことは覚悟で、まずは動くものにしてみます
- 85 名前：デフォルトの名無しさん mailto:sage [2008/01/29(火) 09:43:00 ]
- >82のコードは断片だと思うから敢えて指摘しなかったけど、一応注意点を列挙。
・(少なくとも安定するまでは)CUT_SAFE_CALL()でAPIを括っておいた方がデバッグしやすい。
# エラーチェックを毎回書くのも面倒でしょ。
・cudaMallocArray()は扱いが難しいので、cudaMalloc()にした方が楽では?
- 86 名前：デフォルトの名無しさん mailto:sage [2008/01/29(火) 12:56:30 ]
- そうそう、忘れてた。
初期の段階ならエミュレーションで動かした方がデバッグプリントできるから楽かも。
- 87 名前：デフォルトの名無しさん mailto:sage [2008/01/29(火) 13:30:21 ]
- >>85,86
どうも、アドバイス感謝です。
なんか、動かすとOSリセットかかる凶悪なプログラムに育ちました
もちろん実行はユーザ権限でやってるのに、そんなことできるんですね
エミュレーションでやってみます
- 88 名前：デフォルトの名無しさん mailto:sage [2008/01/29(火) 13:40:59 ]
- そうなんだよねえ。
昔ビデオカードのドライバ無理矢理作らされてひどい目にあった。
- 89 名前：デフォルトの名無しさん mailto:sage [2008/01/29(火) 14:06:29 ]
- >>87
それはあれだ、CPUでいうところのSegmentation Faultだと思う。
明後日の場所をアクセスすると、そうなりがち。
それと、XWindow環境なら5秒ルールにも注意。
# 仕事なら、移植作業を特価で承りますゼw
- 90 名前：87 mailto:sage [2008/01/29(火) 19:04:21 ]
- コンパイルされたバイナリを実行しないで、GPU使ってるか、エミュレーションしてるか見分ける
方法ってありますか。nvccに-deviceemuつけてコンパイルしてるつもりだけど、どうやら
まだデバイス上で動いているみたいなんです。
まあ、おかげさまでOSは落ちずにXだけ落ちるようになりましたがw
>>89
学業の一部ですので、コツコツ勉強しながら自力でなんとかやらせていただきます
- 91 名前：デフォルトの名無しさん mailto:sage [2008/01/29(火) 21:31:21 ]
- XWindow環境なら5秒ルールに引っかかってしまうのか
知らなかった
- 92 名前：デフォルトの名無しさん mailto:sage [2008/01/30(水) 08:07:06 ]
- CUDAをGentooに乗せた人いますか？
ディストリビューションをのせかえるのは大変なので
他のディストロ用のCUDAをインストールしたら
SDKがまともにコンパイルできないのです…
- 93 名前：デフォルトの名無しさん mailto:sage [2008/01/30(水) 14:06:51 ]
- gccのバージョンが違うんじゃない?
CUDAに使えるgccを別途拾ってこないとダメじゃないかな。
# あと、DirectXとか。
## まぁ、エラー処理のところを作り替えちゃえばいらないけど。
- 94 名前：デフォルトの名無しさん mailto:sage [2008/01/30(水) 16:37:38 ]
- 質問です。
cudaで演算を行うプログラムを作成し、なんとかcore2duoの
80〜100倍の処理速度のものが完成したんですが、
プログラムをdll化してVC++で作成したUIから呼び出したところ
処理速度がcpuの４〜５倍に落ち込んでしまいました。
何が原因かはわかりませんが、タイマーを使って調べたところ
cudaMemcpy関数の処理速度が激しく落ちていました。
これはやはりdll化したことが原因なのでしょうか？
しかしNVIDIAのフォーラムにはdll化は問題なくいけると書いてありました。
dll化もフォーラムのとおり行いました。
cudaMemcpy以外の処理はdll化前と変わらない速度で行われているようです。
いったい何が原因でこのようなことになってしまったのでしょうか？
どなたか解決法、もしくは同じ問題を扱ったことのある方いましたら
アドバイスをよろしくお願いします。
- 95 名前：デフォルトの名無しさん mailto:sage [2008/01/30(水) 18:49:56 ]
- それだけじゃ判らんなぁ。
cudaMemcpy()の転送方向は?
HostToDeviceだったらCPUキャッシュ、DeviceToHostならGPUの処理未終了、って辺りかな。
取り敢えず、デバイス側関数呼び出しから帰ってきてもそれはGPUの処理終了じゃないってことに注意。
- 96 名前：デフォルトの名無しさん mailto:sage [2008/01/30(水) 19:03:48 ]
- >>95さん
足りない説明ですみませんでした。方向はDeviceToHostです。
問題は計測している時間の中にＧＰＵの終わっていない処理にかかってる分も
含まれてしまっているということですね。
だとしたら納得いきます。
カーネル実行直後にメモリーコピーを行っているので。
まずはそれぞれの処理にかかってる時間をちゃんと測れるようにがんばってみます
ありがとうございました！！
- 97 名前：デフォルトの名無しさん [2008/01/30(水) 19:07:29 ]
- CUDAを始めるに当たって、何をインストールすればよいのでしょうか？
とりあえず、Visual C++ 2008はインストールし、グラフィックドライバを最新のものにしました。
CUDA SDK以外に必須のものを教えてください。
- 98 名前：デフォルトの名無しさん mailto:sage [2008/01/30(水) 19:15:15 ]
- 必要なもの。
・英語力
・忍耐
・(cudaの動く)GPU
・(cudaサイトからリンクされている)ドライバ
・ランタイム
・SDK
- 99 名前：デフォルトの名無しさん [2008/01/30(水) 23:54:00 ]
- Linuxでやろうとしてます。どのような
マシンを組んで、どのディストリを入れる
のが良い&安上がりですか？
- 100 名前：デフォルトの名無しさん mailto:sage [2008/01/31(木) 05:28:40 ]
- >>99
安定性を求めるのならメーカー製のものがよいが、自作するつもりなら
少なくともマザーボードは鉄板にすること。
それから、電源が貧弱だとトラブルの元になりやすい。
ビデオカードのお薦めはGeForce 8800 GT。
|
 |
