【GPGPU】NVIDIA CUDA質問スレッド

本家
http://developer.nvidia.com/object/cuda.html

グプグプー

>>1
GPGPUスレと統合じゃだめなんか？

http://pc11.2ch.net/test/read.cgi/tech/1188374938/

おお、いつの間にこんなスレが。一応あっちで回答しているけどどうしようか。

削除依頼出しとけ

WindowsのCUDAとLinuxのCUDAで書式に何か違いはありますか？

削除依頼出しとけ

三角関数や指数関数はないの？

削除依頼出しとけ

>>6
基本的に同一。但し、Win用nvcc(コンパイラ)では64bit版にはできないこととオブジェクトモジュールがMS仕様なことに注意。
更に、nvccはgccで実装されているので、C++のnewやvectorはそのままではリンクできなくなってしまうことにも要注意。

>>8
math.hにある、float系の関数は大抵使えると思っていい。エミュレーションモードなら全部使えた筈。
また、sin(), cos()などの一部の関数は組み込み命令(intrinsic)版が使えるので更に高速。
実際の例は、GPGPUスレ参照で。
http://pc11.2ch.net/test/read.cgi/tech/1188374938/71-72

CUDAとかって何らかの標準化の動きはあるの？

笛吹けど踊らず
# まぁ、笛吹いている会社の中でも踊らない連中がいるみたいだし

CUDAは無料ですか？？

ソフトは無料、ハードは有料なんじゃね

CUDAなんてハードはありません。

処で、テスラc870ってQuadroFXの一番上の奴とどこが違うんだ?

>>15
おマイ馬鹿か、環境の事に決まってるだろ

Compute Unified Device Architecture
        ~~~~~~~

CompUte Driver Architecture

ちょっとメモ。
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;
}
--
単純に考えると開始位置を検索しなければならないが、それ自体にコストが掛かりそうだ。

>>20みたいな処理って数値演算よりデータ転送の比重が多すぎて、
並列化させる旨みは非常に少ないと思うんだけどどうなのよ。
本家のフォーラムでいわゆるセマフォを実装して集計処理やらせようとしてる人がいたけど、
そういう問題をCUDAでやる事自体がそもそも間違っとる、みたいに言われてた。
かく言う自分も似たような事しようとしてたけどパフォーマンス出ないんで諦めて、
問題の分割方法から見直したら上手くいったりした。

>>20みたいな処理って数値演算よりデータ転送の比重が多すぎて、
並列化させる旨みは非常に少ないと思うんだけどどうなのよ。
本家のフォーラムでいわゆるセマフォを実装して集計処理やらせようとしてる人がいたけど、
そういう問題をCUDAでやる事自体がそもそも間違っとる、みたいに言われてた。
かく言う自分も似たような事しようとしてたけどパフォーマンス出ないんで諦めて、
問題の分割方法から見直したら上手くいったりした。

二重orz

>>21
>20の処理は、最初に定数テーブルを仕込んでおいて幾つかパラメータを渡して計算させるのがメイン。
>20に書いているのは最後の集計部なんだけど、結果はarとaiに入るのでそれを転送するだけ。
集計をGPUにやらせられないとなるとtmpを転送しなければならないので、それをなんとかしたかったわけ。
で、未だその後修正掛けていないけど、begin, endのペアをidxの値の分だけを定数テーブルに
仕込んで置けることが判ったからなんとかなりそう。そうすればidxの値の分だけ並列に走らせられることだし。

>>24
状況はわかったけど最後の方で言ってる意味が良くわからん…俺の頭がヘタレなのか。

今のとこ、並列で集計処理するのに一番早そうかなって思ってるのはデータを半分ずつ加算してく方法。
例えば処理したいデータがnコなら、n/2コのスレッドで2コずつデータ加算。後はリカーシブにやる。
(めんどいからとりあえずnは偶数としてね)
これならlog_2 nのオーダーで集計出来んじゃね？ってかんじ。

実装したことないがスレッド数が変わるのが厄介かな…今は
if(tx==0) {
　for(i=0;i<BLOCKSIZE;i++)
　　sum+=data[i];
}
みたいに横着してる。他で十分高速化してるんでとりあえず後回し中。

GPUでテクスチャの全てのピクセル値の合計を求めたいってこと？
1x1サイズのミップマップ作って、ピクセル数を掛けるのはどう？
速いかどうかは知らんけど。

opensuse 10.3でCUDAは動く？

Linuxサポートしてるから動くんじゃね？

>>27
ドライバが載らなくても取り敢えずエミュレータは動くでしょうね。
つーか、本家に10.1用と10.2用は置いてあるから試してみては?
ttp://developer.nvidia.com/object/cuda.html

>>27
もしかしたら、標準でインストールされるgccのバージョンが違う所為で動かないかも知れない。
その場合、10.2にインストールされているバージョンのgccをインストールしておけば大丈夫。
# 要はnvccが内部でgccを使っているから、gccのバージョンに依存していると。
# ドライバはOKだった筈。

GeForce 8400GSの安いカードでもCUDAを試すことはできますか？

>>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に対応とある。

>>31
8800を買うのが無難

安く済ませるなら、(消費電力の少ない)8600GTもお勧め。
第2世代(と思われる)8800GTは割とパフォーマンスがよさそうだけどね。

>>31
公式にはサポートされていない。
でも、ノートパソコンで動いたという人もいるくらいだから、動く可能性がないわけではないと思う

8400Mは対応してたはずだけど、普通の8400の方は対応していないんだよね
自作板でサンプルが動いたという報告があったような気みするけど、
対応していない製品を使うのは心配だよね

8600買ってきたー
試したいのだけど、まず取っ掛かりはどこからいけばいいですか＞＜

>>37
>29にある本家から一式拾って、ドライバとツールキットとSDKを入れて、
サンプル一式をビルドして動かしてみるよろし。

64bit 版　ubuntu7.10にインストールしたいのですが、どれをダウンロードすればよいのでしょうか？

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を使うから、インストールの状況によるので単純じゃない模様。

ubuntu版きてるーーーー

7.04ってなんだYO

タイムリーなんだか間が悪いんだかw

NECのVALUESTARを使ってるんですが
nVIDIAをインストールしようとすると

NVIDIA setupプログラムは、現在のハードウェアと互換性のあるドライバ
を見つけることができませんでした。

と表示され、インストールができません。
何が原因なんでしょう？
OSはVistaです。

>>44
NECのVALUESTARなんてどうでもいい情報よりもグラボの情報だせよん

なぜココで聞く。
ハードウェアスレに聞け

>>44
そもそもVista対応のCUDA対応ドライバがありません。

>>44
>nVIDIAをインストールしようとすると
って時点で意味がわからんのだが、
とりあえずデバイスエミュレータ(-deviceemu)でどうよ？

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));

自己解決しました！

勘違いでした
誰かヘルプ。。。

何がどう巧くいってないと判断したのかよく分かりませんが。
cudaMallocPitch(), cudaMemcpy2D()じゃなくて、cudaMalloc(), cudaMemcpy()では巧くいったの?
それと、テストしたソースは丸ごとコピペするか、どっかにアプロドして欲しい。

＞アプロド
なんか可愛いw

>>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の経験が全く足りていないのか、いずれにしても阿呆過ぎる。

>>49
＞float* devPtr,an;
このanは只のfloat型。

cudaMemcpy2D()の第1,3パラメータのキャストもおかしい。
もしサンプルにそう書かれていたのなら、そのサンプルがタコ。

処で、サイトが更新されて、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

表をそのまま貼ったから見難くなったのは勘弁。
補足しておくと、FX/Plex/NVSはどれもQuadroシリーズ(非OEM)ね。
細かく見てたら、8800GTが追加された他に8800GTSに512MBが追加されている。
しかも、クロックが既存の8800GTSよりも速いだけでなくShaderClockでは最速に。
MemoryInterfaceが512bitだから、8800GTSというより8800GTの系列かしら。
まぁこれで、旧8800GTS(640MB/320MB)は消える方向なんでしょうね。

ほほー
廉価版でCUDA試してみたいだけの奴は8400あたりでもOKということか。

ありがたやありがたや

CUDAを使用できるGPUのうち、
GeForceとQUADROのどちらの方がCUDAに最適なのでしょうか？

>>59
「最適」の基準によります。
入手性、価格、バリエーションではGeForceシリーズに文句なしで軍配が上がりますが、
個別の性能はそれぞれ異なりますので。
因みに、TeslaはQuadroFX5600(-アナログ回路)相当だそうです。
あー、QuadroPlexを買える待遇なら、違う意味で最適かもしれません。

処で、8800GTXでどこまで速度が出るのか実測してみた。
単純な実数の足し算をループで回しただけなんだが、約43GFlops出た。

一方、ShaderClockは1350MHz、プロセッサ数は128、ということなので計算上は(1クロック1演算なら)172.8GFLOPSとなる。
命令セット上は、積和が1クロックで実行できるから2倍して約350GFLOPSだからカタログスペックはこの値なのだろう。
実際には、ループを構成するには実数の足し算の他に整数加算、整数比較、条件ジャンプの3クロックが必要になる。
従って、約43GFLOPSとなって実測値と見事に一致した。

というわけで、(メモリアクセスが無視できるなら)アセンブリ(ptx)出力から所要時間を割り出すことができる模様。
問題は、メモリアクセスが絡むとどれだけ遅くなるかで、こちらの方は今後調査の予定。

ニーモニックて公開されてるんだっけ？ or 固定長？

5秒制限とかあるらしいけど
これどうやって回避するんだろう

>>62
ptx出力を睨めっこすれば大体当たりがつく。どっかで公開されてるかも知れんが。

>>63
5秒も計算させたまま帰ってこないような組み方すると、どうせ遅くなるから余り気にならない。
単純なループだけならXeonの方が速いからね。

>>61
おつー

これは目安に使えそうだね

ttp://www.nvidia.com/object/cuda_home.html

CUDAのサイトリニューアルされてるね。
相変わらずfor Vistaはまだだけど。

>>66
一応>56で指摘済み。

処で、cufftを使ってみた。
SDKのサンプル(convolutionFFT2D)では100MPix/sの処理速度があるような実行結果が得られたけれど、
実際に試してみたら4096x4096の大きな画像を使っても10MPix/sしか出ない。
どうも、オーバヘッドが大きくて速度が出にくいみたいだ。
まぁ、サンプルと違ってテストしたコードは「プラン作成」「メモリ確保」「メモリ転送」「FFT実行」
「メモリ転送」｢メモリ破棄」「プラン破棄」を全部実行しているからだとは思うけど。
と言うことで、来週辺りは実際のプログラムに組み込んだ形でテストしてみる羽目になりそうだ。

そうそう、8800GTXをCUDA1.0で動かした場合と8800GTをCUDA1.1で動かした場合で
convolutionFFT2Dの所要時間が殆ど変わらなかった。GPUの性能差を埋める程度にはCUDA1.1で改善されたのかな?

ずっと旧トップページをブクマしてて今日やっと転送されるようになったので
気づいてなかったｽﾏﾝｺ

1.0になってから遊んでいなかったので久しぶりに遊んでみよう。

>>67
メモリ転送は、PCI-Expressの性能で上限があるはず。
たぶんこれがかなり時間食ってると思う。
4K×4kは画像なんでしょうか。大きいですね。

PCIe 1.1と2.0でどれぐらい差が出るのか気になりますね。

3DMarkのベンチマークやゲームではほとんど差が無いみたいですが
CUDAだとやはりそのあたりがボトルネックになる場合も出てくるでしょうかね。

またメモリ容量に依存する場合もやはり多いのでしょうか？

>>69
FFT解析絡みのことをやっているので、周波数空間像を拡げてから逆FFTなんてことをしばしば。
まぁ、4096x4096は単にテストするのが目的でしたが。

>>70
未だPCIe2を試せる環境がないからなんとも言えませんが、演算の種類によっては影響出るでしょうね。
私のところのプロジェクトでは寧ろ、GPUとCPUを巧く並行動作できるかどうかの方が影響大きいのですが。

むぅ、周波数合成みたいなロジックをCUDAに移植したんだが、余り速度が出ない。
元のソースが三角関数テーブルを使っていたのでメモリアクセスが足を引っ張っているのかと思って
__sincosf()を使ってみたが、今度は何故か実行ごとの所要時間のばらつきが大きい。
速いときはXeon並の速度になるのに遅いときは3倍くらい時間が掛かる。
はてさて、超越関数ユニットの数が少ないのだろうか……

超越関数は級数近似と思われ
see /cuda/include/math_functions.h

いえいえ、ちゃんとニモニックもあるのですよ。そっちを使うとドキュメントに拠れば32クロックだそうで。
但し、そいつが「To issue one instruction for a warp,」とあるからぶつかるのかなと。

8800前後の機種の日本での売れ行きがかなり芳しくないようだがこれは何を意味するのだろう．本国でのアメリカでは使用例（論文）が着実に出つつあるから，なんだかそいつらの為だけに開発されてる感じじゃない？？

8800系って、日本ではNEETゲーマーが買い漁るから
1$=200円くらいのボッタ栗値段が付いてしまって、
貧乏ラボじゃ買えん罠。箱もヲタ臭くてボスの目
が気になる。

つ [玄人志向]

そうか，研究者の間であまり大量に出回っていないと思ったが，そういう事だったのか．＞76
国内で一番CUDA使いこなしているのはどこのチームだろう？

まぁ、CUDAでXeonに勝つ速度を叩き出している漏れが一番だねw

なるほど、確かに>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を別々に使うので、別々にしても遅くなることはない。
--
あとで整理してどっかに上げとくかな。

>80を一部訂正。
・但し、冪乗関数はexpf()は1.4427を、exp10f()は3.3219を、それぞれ掛けて(mul.f32)からex2.f32を使うようだ。

すみません。質問です。
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;//こんな感じの計算がやりたいです
この後、四則演算が続く。
}

3次元格子は扱っていないから一般的な方法は知らん。
で、>82で流れはあっていそうだけど効率は悪いと思うよ。
ブロックごとにdVelをSharedにキャッシュした方がいいかも知れない。
# まぁ、動くものを作る方が先だけど。

>>83
ありがとうございます
吐き出される数値を見て、ぜんぜん計算されてなさそうなので、
心配をしていました。

効率悪いことは覚悟で、まずは動くものにしてみます

>82のコードは断片だと思うから敢えて指摘しなかったけど、一応注意点を列挙。
・(少なくとも安定するまでは)CUT_SAFE_CALL()でAPIを括っておいた方がデバッグしやすい。
# エラーチェックを毎回書くのも面倒でしょ。
・cudaMallocArray()は扱いが難しいので、cudaMalloc()にした方が楽では?

そうそう、忘れてた。
初期の段階ならエミュレーションで動かした方がデバッグプリントできるから楽かも。

>>85,86
どうも、アドバイス感謝です。
なんか、動かすとOSリセットかかる凶悪なプログラムに育ちました
もちろん実行はユーザ権限でやってるのに、そんなことできるんですね

エミュレーションでやってみます

そうなんだよねえ。
昔ビデオカードのドライバ無理矢理作らされてひどい目にあった。

>>87
それはあれだ、CPUでいうところのSegmentation Faultだと思う。
明後日の場所をアクセスすると、そうなりがち。
それと、XWindow環境なら5秒ルールにも注意。
# 仕事なら、移植作業を特価で承りますゼw

コンパイルされたバイナリを実行しないで、GPU使ってるか、エミュレーションしてるか見分ける
方法ってありますか。nvccに-deviceemuつけてコンパイルしてるつもりだけど、どうやら
まだデバイス上で動いているみたいなんです。

まあ、おかげさまでOSは落ちずにXだけ落ちるようになりましたがw

>>89
学業の一部ですので、コツコツ勉強しながら自力でなんとかやらせていただきます

XWindow環境なら5秒ルールに引っかかってしまうのか
知らなかった

CUDAをGentooに乗せた人いますか？
ディストリビューションをのせかえるのは大変なので
他のディストロ用のCUDAをインストールしたら
SDKがまともにコンパイルできないのです…

gccのバージョンが違うんじゃない?
CUDAに使えるgccを別途拾ってこないとダメじゃないかな。
# あと、DirectXとか。
## まぁ、エラー処理のところを作り替えちゃえばいらないけど。

質問です。
cudaで演算を行うプログラムを作成し、なんとかcore2duoの
80〜100倍の処理速度のものが完成したんですが、
プログラムをdll化してVC++で作成したUIから呼び出したところ
処理速度がcpuの４〜５倍に落ち込んでしまいました。

何が原因かはわかりませんが、タイマーを使って調べたところ
cudaMemcpy関数の処理速度が激しく落ちていました。
これはやはりdll化したことが原因なのでしょうか？
しかしNVIDIAのフォーラムにはdll化は問題なくいけると書いてありました。
dll化もフォーラムのとおり行いました。
cudaMemcpy以外の処理はdll化前と変わらない速度で行われているようです。
いったい何が原因でこのようなことになってしまったのでしょうか？

どなたか解決法、もしくは同じ問題を扱ったことのある方いましたら
アドバイスをよろしくお願いします。

それだけじゃ判らんなぁ。
cudaMemcpy()の転送方向は?
HostToDeviceだったらCPUキャッシュ、DeviceToHostならGPUの処理未終了、って辺りかな。
取り敢えず、デバイス側関数呼び出しから帰ってきてもそれはGPUの処理終了じゃないってことに注意。

>>95さん
足りない説明ですみませんでした。方向はDeviceToHostです。
問題は計測している時間の中にＧＰＵの終わっていない処理にかかってる分も
含まれてしまっているということですね。
だとしたら納得いきます。
カーネル実行直後にメモリーコピーを行っているので。

まずはそれぞれの処理にかかってる時間をちゃんと測れるようにがんばってみます
ありがとうございました！！

CUDAを始めるに当たって、何をインストールすればよいのでしょうか？
とりあえず、Visual C++ 2008はインストールし、グラフィックドライバを最新のものにしました。
CUDA SDK以外に必須のものを教えてください。

必要なもの。
・英語力
・忍耐
・(cudaの動く)GPU
・(cudaサイトからリンクされている)ドライバ
・ランタイム
・SDK

Linuxでやろうとしてます。どのような
マシンを組んで、どのディストリを入れる
のが良い&安上がりですか？

>>99
安定性を求めるのならメーカー製のものがよいが、自作するつもりなら
少なくともマザーボードは鉄板にすること。
それから、電源が貧弱だとトラブルの元になりやすい。
ビデオカードのお薦めはGeForce 8800 GT。

安西先生、8800は安くないっす…。

目的が判らんから何とも言えんが、安さ優先なら8600GTでいいんじゃね?
余裕があるなら8600GTSで。8800GTは消費電力も発生熱量も大きいから、
筐体の空冷性能が悪いと巧くないかも知らんことだし。

Visual Studio 2008ではコンパイルできないのでしょうか？

現在
C:\Program Files\NVIDIA Corporation\NVIDIA CUDA SDK\projects
以下にあるtemplateを書き換えて使っています。
他のディレクトリに移すと、コンパイルできなくなってしまいます。
他のディレクトリに移す場合、どこを書き換えたらよいのでしょうか？

>>103
できない。
>>104
Visual Studioの新規作成で作成すればよい。

プロファイラがリリースされたようですね

SLI構成でCUDA使ってる人はいますか？
SLIにするには、ボード増設すればいいだけ的なことが書いてあったけど、
一枚から二枚に移行するにあたって、気をつけることってなにかあったら
教えてください。

>>106
どこ? ぱっと見で見つからない……

>>107
ボード増設だけではSLIにならない気もす。ボード完直結のケーブルが要るような……
逆に、CUDA使う分には別段困らないと思うけど。

>>107
SLIは同一型番のカードでないとうまくいかないことが多いみたいですね。
最低でもメモリ容量やクロックは合わせないとだめなんじゃないでしょうかね。

それからチップセットによってはx16の2枚のフルレーンではなくx8の2枚となるのも要注意でしょうか。

>>108
SLIと非SLIのマルチGPUでパフォーマンスに差が出るのかも気になりますね。

ホスト側のメモリの確保はmallocを使わないといけないのでしょうか？
サンプルコードを見るとmallocが使われていますが、newじゃ駄目なのでしょうか？

newは使えないって言うのをどこかで見たなぁ

Linuxなら大丈夫。Windowsの場合、nvccはgccベースなのにVCのリンカを使うからおかしなことになる。

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です。

あんたの頭。

>>113
ロケールが悪いと書かれてるが？

ロケールが悪いんじゃなくて、
ロケールが悪いというエラーを返すために例外を投げてるのが悪いんでしょ

floatを4つ使うよりfloat4を使った方が高速なの？
float3とかfloat4の利点がいまいちわからないのだけど

globalメモリからのアクセスなら、floatはld.global.f32が使われる。
float2ならld.global.v2.f32が使われる。float3, float4も恐らく同等。
尤も、struct {float x, y}だとしてもfloat2扱いしてくれるから使いたくなければ使わなくてもいい。

GPGPU#2スレから来ました。

extern __shared__ int shared[]; 
でsharedメモリを使えますが、複数の配列を使いたい場合はどのようにすればよいのでしょうか？ 

>>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))>>>();
のように呼ぶ必要があります。

<<<>>>に入れるブロックの数とかスレッドの数ってどうやって決めればいいのでしょうか。
いまは、適当に縦横の格子点数で決めてるんですが。
ちなみに、使っている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

ちなみにsharedが速いということはほとんどないから、安心してglobalを使え。
むしろ、コピーする分遅くなる。
まあ、CPUの様に高速なL1キャッシュでも積むようになれば別なのだろうけど。

>>121
プログラムによる

>>121
ドキュメントを読む限り、スレッド数はワープ単位にするべきかと。
それと、できる限りCPU側でループさせない方がいいと言う点からいけばブロック数は
1000-10000位は欲しいということになるらしい。
まぁ、実測してみればいいんじゃね?

GPUとCPUに分散させて計算させることはできませんか？

できます。CUDAで言えば、GPU側を起動した後CPUはGPUを待たずに処理できます。

CUBLASとCUFFTをソースコードからコンパイル出来た方がいたら、ぜひやり方を教えてください

ソースはココからダウンロードできます
http://forums.nvidia.com/index.php?showtopic=59101

サンプルがなかったっけ？

>>128
何のサンプル?
CUBLASやCUFFTのサンプルならその他のCUDAのサンプルごとSDKに同梱されていると思うけど。

>>127
ソースからコンパイルする目的は?

>>122
再帰的につかう変数をsharedにもってくると、断然速いだろう？

CUDAで再帰って使えるんですか？

CUDAで再帰って使えないと思うような根拠があるんですか?

ところで、CUDAって何て読むの?

キューダ? シーユーディーエー?

>>133
お好きにどうぞ。
一般的な英語の発音ルールに乗っ取るなら「カダ」になるのかも知れませんが、
barracudaからの連想で「クーダ」と発音する方がそれっぽいかも知れません。

s/乗っ取る/則る/

クーダの方が語感的にいいな

Kirk氏の講演を聴いたことがあるが
クーダと言っているように聞こえた

管

Windows版の開発もLinux版のようにコマンドラインだけでやれる方法はありませんか？

VisualStudioを使わなければOK!
cl.exeやnmake.exeの使い方は自分で調べてくれ。

祭りだ祭りだ
Japan CUDA カンファレンス　3/6（木）＠本郷
http://www.loopinc.jp/nvidia2008/

ブロックやグリッド数を2つ以上必要な場合ってどのような場合なのでしょうか？

GPUのプロセッサ数はどこでわかりますか？

>>142
GPUの性能を活かすためには、並列度を上げる必要があります。
逆に、ブロックやグリッドが1だとプロセッサを使い切れなくてパフォーマンスが出ません。

>>143
GPUの仕様か、CUDAのドキュメントに書いてあります。
アプリケーションから動的に取得することはできないようです。

ブロック数が1つだと、並列には実行されることはないのでしょうか？
プログラミングガイドを見ると、ブロックの中に複数のスレッドがあるのですが、
このスレッドというのは、CPUのスレッドのように並列には実行されないのでしょうか？

だから用いるThreadの個数を指定するの
例のソースコードなんかを読んだ方が良いよ

そう言えば次の廃エンドになるらしい9800GTXではCUDAもやっと2.0とかになるのか?

CUDAは範囲外のメモリを読み込もうとしてもエラーを返さないのですか？

>>148
APIはエラーを返すけど、デバイス側の関数は必ずしもエラーを返さないね。
下手すると落ちるけど。

blockIdxとthreadIdxは0から始まるのですか？

はい。例えばfunc<<<m, n>>>()と呼び出した場合は、
それぞれ0からm-1, 0からn-1の値をとります。
二次元型、三次元型の場合も同様です。

ブロック数を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);

プログラムにもよると思うけどありえる話だろうね

何スレッドが同時に動くかはGPUによっても違うので一概には言えませんが、
少なくとも8の倍数にはなるのでスレッド数50や100は効率が落ちます。
例えばスレッド数を64にしても50のときと所要時間は変わらないかも知れません。
又、デバイス側の関数呼び出しは完了していても演算自体は終わっていないかも知れません。
cudaThreadSynchronize()してから時間を計った方がいいかも知れません。

研究室のPCにLinux入れてCUDA動かしてますが、
ユーザ権限実行でもOSが死ぬプログラムを余裕で書けてしまいますね
これはCUDA使う以上、仕方ないことなんですよねきっと
バグのあるデバイスドライバ使ってるのと同じことと、あきらめるしか
ないのでしょうか

vista対応と自動回復が働いてくれることを渇望しましょう

そして危険なプログラムが走らないような検証系に進むのであった。
といってもDirectX動かす権限があればOS死ぬのは容易いような。

vista対応マダー？

vistaは要らない子。
*BSD対応マダーー？？

いやむしろTRON対応

*BSD対応は確かに欲しいな・・・。
っていうか、オプソにするくらいの気持ちじゃないと
この規格は普及せんような・・・。
将来的にnVIDIAのGPUが携帯電話などに使われて行った時に
なかなか面倒じゃないかなぁ。

NVIDIAはしばらく内部仕様公開しないだろうね。
ptxは読めてもその先の動作がわからないから辛いわ。
AMDには圧倒的に差をつけていそうだし。

GeForce9600出ましたな。
8800はStreamProcessorが128個×1.35GHzだったけど，今度は64個×1.625GHzみたい。
並列度が低い場合でも早くなるっていうこと？

だめっぽい
ttp://www.4gamer.net/games/047/G004716/20080220030/

>>157-158
ユーザースペースで動作させた場合の性能が気になりますね。

>>159-161
*BSDやgentoo等への対応も欲しいですね。

欲をいえばVista以外でのHybrid SLIのHybrid Power対応も w
自動切換えとかGeForce Boostとかは無理でも、コマンド一発でオンボとカードを切り替えられるようになるといいのですが。

doubleにはいったいいつ対応するのか

http://pc11.2ch.net/test/read.cgi/tech/1188374938/137

http://www.nvidia.com/object/cuda_learn_products.html
9600GT，CUDAも実行できなければ遅い。doubleが使える世代への場繋ぎでしょう。

>>167
NVIDIAは伝統的に新しいGPUでCUDAが使えるようになるまでに少なくともドライバを更新してくるから、
今のところ未だ、9600GTでCUDAが使えるようにならないとは言えない。

普通に使えるようになるでしょ。つかヘッダとか弄ったら使えたりして。

で、VC2008+cuda1.1の食べ合わせって駄目なのかよ。1.0はいけるらしいけど orz
http://forums.nvidia.com/index.php?showtopic=53494

Vistaのbeta版が３月の終わりごろに出るらしい

http://forums.nvidia.com/index.php?s=70658fb133a8b4a52524f6ec7f2f3408&showtopic=60431

>>170
�d

ずっと上の方で誰か質問しててまだ回答無いけど、
ゲフォ8400GSでCUDA動きますか？
WinXPでは動いたんだけど、Fedora Linuxでは
認識しなくてエミュレーションに成ってしまった。

何かと金の掛かるWinではCUDAやりたくないです。

>>172
一覧にはあるから使えるんじゃないかというレスがあった希ガス。
問題は、ドライバをインストールできるかどうかでそれは端末とBIOS次第かな。
ToolKitはFedora7用のがあるそうだから、逆にそれ以前のはダメかもしれない。
あーでも、エミュレーションでコンパイルできているなら大丈夫か。
コンパイルオプションを指定すれば、ちゃんとGPUコードのモジュールもできるんじゃないかな。

デバイス認識させるのは癖があるから要注意で。
Xが動く状況じゃないと、巧く認識できなくて苦労した記憶がある。
巧く認識できていれば/dev/nvidia0と/dev/nvidiactlができている筈なんで、
それがないならroot権限で正しくコンパイルできたGPUコードのサンプルを動かせば
ドライバの.koを使ってデバイスのエントリができる筈。
Linuxのデバイス周りは詳しくないんで、頓珍漢なことを書いているようなら失礼。

しかし、VisualStuido2005EEと組み合わせれば無料でできるのに「金の掛かる」とはまた意味不明な……
WindowsならGPUのプロファイリングツールも使えるし、それはそれで悪くない選択だと思うんだけどね。
# といいつつ、未だ自宅で環境整備してないけどw

8600GT　WIN2Kの俺に謝れ！！！

XPが無い…

>>175
Win2kで動かないのであれば、WinXP買うのも馬鹿馬鹿しいね。ゴメン、配慮が足りなかった。

XPが嫌いってだけのもあるんだけどなー
ムダに重たくて

>>177
テーマをクラシックにするか、システムプロパティでパフォーマンスを優先にすればWin2Kに較べて重くないと思いますよ。

構ってちゃんの相手するなよ・・・

そもそもWindowsが糞高いし、ソース付いて来ないし。
不自由過ぎてハックする気が萎える。

こんちわす。
.NET環境からCUDAとデータの受け渡しをする例、なんてのは
無いもんでしょうかね。既存のVB.NET、C#.NETの統計解析を
高速化致したく。

相関係数＝Σ(xj-xk)^2÷√Σ(xj-xjの平均)^2*(xk-xkの平均)

なんてのが多発し、j,kは1〜1000、それぞれ1万点ずつ、x1〜x1000の全ての
組み合わせで計算、という具合です。でも「総和を取る」とかはGPUは
苦手でしょうかね・・・。

総和は苦手なので、適当に分割して小計を求めてから最後に総計を求めることになると思う。

それはさておき、先ずはその相関係数の式自体を変形してみてはどうだろう。
分子も分母のそれぞれも平均値を使わない形にできるはずだが。
# ついでに言えば分母の最後は2乗が抜けている希ガス。
それをやると、平均を使う2パス処理ではなくΣxj, Σxj^2, Σxk, Σxk^2,...などを求める1パスですむ筈。

あ、肝腎のVB/C#からの呼び出し方はわかんね。CUDA側は"*.obj"になるからリンクさえできればなんとかなるのかな?
まぁ、計算量が充分大きいならCUDA側を"*.exe"にして起動するようにしても大差ない気もするけど。

どうもどすえ。自分なりにお勉強してみました。
.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」第三版を注文したので、もうちょっとお勉強してみるす。

だから、相関係数を求める式を変形しろって。
CUDA云々以前だぞ。

え、相関係数を求める式にはもう「平均」は無いすよ。
ごめんどこ怒られてるのかわかんね・・・

「相関係数」でぐぐれかす

お前ら両方とももういい加減にしろ

統計の基本:
相関係数は普通ピアソンの積率相関係数を指す。
これは、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回のスキャンで充分なので、各要素を保存する必要がない。

上のほうにもあるけど、総和の計算って実際どーすればいいんだ？
スレッドブロックごとに小計を求めて、各ブロックの結果をまとめるとかいう
方法になりそうなのは予想できるが、実際どんな風に書けばいいのかわからん。

具体的にはベクトルの長さが出したいんだが、各要素の２乗の和がうまく出せない。
ヒントでもいいので分かる方教えてください。

1つのブロックを総和の計算に割り当てる…とか
そんなことしか思い浮かばん

後で暇があったら実際のコードの集計部分だけ晒してみるか。

CUDAを使ったjpegの展開ライブラリはありませんか？

IDCTはそこそこ効率よさそうだけど前段のビットストリーム分解が不向きっぽいぞ。

>>191
そうしてくれると助かります。
よろしくお願いします。

デイビッド・カーク氏来日，CUDAカンファレンス2008開催

ttp://www.4gamer.net/games/032/G003263/20080306051/

やっとできたらしい日本語版のCUDAZONEとプログラミングガイドの日本語版。
ttp://www.nvidia.co.jp/object/cuda_develop_jp.html
しかし、機械翻訳そのままなのかな。余りに訳が酷くて原文読まないと意味が判らんとは……

CUT_EXIT(argc, argv)はなんでargc, argvを与えなくちゃいけないのでしょうか？

>>197
引き数にnopromptが指定されていない場合は
"Press ENTER to exit..."のメッセージを出して入力待ちにするためです。
詳しくは、cutil.h参照で。

floatであった変数をいくつかまとめてfloat4を使うようにしました。
そうしたところプログラムの実行速度が10%ほど落ちたのですが、float4よりfloatの方が
実行効率がよいのでしょうか？

よく知らないが、float4ってSIMD用の型じゃないの？

単にまとめただけじゃ速くならないと思う。一番いいのは-ptxして出力眺めることなんだけどね。

なんで、nvidiaのフォーラムには韓国語はあるのに日本語はないの？

流石にそれはnvidiaに聞いてくれ。需要が少ないと言うか、要望が少ないのだろ。
半島人と違って日本人は奥床しいからw

アメリカの地下鉄とかの案内板や自販機にハングルはあるけど
日本語はないとかいう話を思い出しました w

東大キャンパスであった、CUDAカンファレンス2008の動画とかは公開されないのでしょうかね？

ずいぶん前の話だが、20-26あたりで議論されてる問題って結局どんな風に
並列処理するのが最適なんだ？
総和の計算の話が出てたから思い出したのだが。
前に試行錯誤したが、高速化できなかったり、値がおかしくなったりして
あきらめてしまっていた。
どうぞ皆様知恵をわけてください。

この前のカンファレンスでも言われていたがチューニングにかける時間を
惜しんではいけないのだね。

総和みたいな計算は向いてないとも言ってたけどね

>>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配列への演算結果の格納が所要時間の大半を占めたために適当にチューニングして放置。

尚、予告した集計のロジックは、総和ではなく↑とも違う小計算出だったので条件が違うと言うことでパス。
総計は興味があるので、その内テストできたら改めて晒そうと思う。

いきなりですが質問です
xp64でVisualStudio2005と8800GTS(G92)がありCUDAドライバー、
ツールキット、SDK、と入れたんですが、結局ビルド方法とか
どのPDFに書いてあんだろ、というのがよくワカンネですよ。

VS2005のプロジェクト設定をx64に変えて、cutil32D.libを
拾ってきたZIPから解凍して入れたら、mandelbrotとか一応
動いたけど。他のmandelbrot描画ソフトより劇速なんで、
動いてはいるみたいです。

うーん、IDEは使ってないから判らないなぁw
pdfじゃなくて、ReleaseNoteか何かにビルド方法かいてなかったっけ?
後で見てみるわ。

>>206
確かに言っていた。まぁそれは言われなくとも自明なことなんだが。
ついでに、CUDAがあるからGPUに向いていない処理も容易に書けるようになってしまった
とも言っていたな。
だが、膨大なデータをホストに書き戻して、総和を計算し、その結果をまたデバイスに
送って使うということは出来れば避けたかった。
GPUが苦手な処理でもデバイス側でやらせるというのがいいのか、転送コストを覚悟の上で
ホストに戻して計算したほうがいいのか実際確かめるために試行錯誤してみている。

>>207
説明ありがとう。
おかげで24でかかれていたことが理解できたよ。
その考え方は思いつかなかったな。
俺ももうちょい考えてみていい感じにできたら晒そうかと思う。

ベクトルの総和を出すscanlarge.pdfを見てみると、
elementsが65536ならCPU−GPUでほぼ一緒、
100万elementsでGPUがCPUの5倍速い、ようですね。
なるほどGPU向きではないな・・・

HostがCore2Duo辺りなら未だいいんだけど、Woodcrest辺りだと厳しいんだよねw
超越関数を大量に使う演算だと、GPUが随分有利になるんだけど。
# SSEには超越関数がない→ベクタ化のためには関数テーブルが必要→メモリ消費→キャッシュミス

cudaでfloat4*float4ってできないの？
各要素ごとに4行に分けて書かないとだめ？

>>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}; //
--
こうなった。

FreeBSDのLinuxエミュでCUDAを使っているんだけど、
Linuxと比較してどれくらい速度が落ちているのかな?

ほぼネィティブと同程度ならこのまま使っていこうかと
思っているんだけど。

Windows上でコンパイルしているのですが、nmakeは使えないのでしょうか？
cygwinのmakeコマンドであればコンパイルできるようなのですが、nmakeを使うと
NMAKE : fatal error U1077: 'C:\CUDA\bin\nvcc.EXE' : リターン コード '0xc0000005'

Stop.
となってコンパイルできません。このエラーは何を意味しているのでしょうか？

float4ってsimdで計算してくれないの？

>>217
そもそもsimdなんてありませんが何か。つーか、>214で回答ついてるっしょ。

>>216
nvccが終了ステータスに5を積んでいるようですが、
ちゃんと動いた結果ならnmake側で無視することができると思います。
そもそもちゃんと動いてますか?

>>215
4種類ほどボード別に速度を測ったデータがあるからそれでも載せてみましょうか。

>>217
早速のご回答ありがとうございます。
Makefileの内容ですが、非常に初歩的な

test.exe: test.cu test_kernel.cu
    nvcc test.cu

です。直接、コマンドプロンプトから打ち込んだ場合、問題なくコンパイルできます。

>>219
レス番違いますぜ。
取り敢えず、nvccがエラーステータスを返しているかどうかと、
nmakeでそれを無視できないか調べてみてはいかがでしょう。

そういうわけで、各ボードの比較表を作ったので貼ってみる。
--前半
"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",

--後半
"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"

これってプログラミングガイドの最後にあるやつ？

CUT_EXITを実行せずにプログラムを終了した場合、どのような悪影響がありますか？

>>221
SharedMemoryって16KBしか割り当てられないのですか？？

幾ら何でも同じスレの中くらい検索してくれたっていいじゃないか(TT

>>223
いいえ、サンプルプログラムの出力です。

>>224
なーーんにも。どうしても気になるならcutil.hを眺めてください。

>>225
この表では、1block辺りのSharedMemoryのサイズを示しています。

Cのmallocの場合、NULLかどうかでメモリが確保できたか判別できますが、
cudaMallocでメモリが確保されたかどうかを調べる方法を教えてください

CUT_CHECK_ERROR

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.

ああ、あとデバッグモードじゃないとだめだよ

warpサイズって簡単に言うと何？

同一インストラクションが同時に走る、演算ユニットの集まり?

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の道険しいﾅﾘ

初歩的な質問かとは思いますが、各ストリームプロセッサにそれぞれ1つずつスレッドが割り当てられるのですか？

__syncthreads()ではブロック内でしか同期できませんが、全ブロックを同期したいときは
カーネルを抜けていったんCPU側に戻すしかないのでしょうか？

cudaMallocHostとCのmallocの違いを教えてください

>>235
確かカーネルは非同期で実行されたはずだから、
カーネルを連続で呼んだら同期されない・・・と思う。

<<<カーネルの呼び出し>>>
cudaThreadSynchronize();
<<<カーネルの呼び出し>>>
ならいけるんじゃないの？自信ないけど。

>>236
cudaMallocHostでメモリ確保すると、ページングなしスワッピングなし
物理メモリに連続して常駐、のメモリ領域が確保されるはずなのだ！
と思ってcudaMallocHost使ってます。mallocとmlockの合せ技？

Windows, VC2005なんですが、PTXを出すためのオプションはどこに
書けばいいんどすか( ・ω・`)
苦労してPhenomの6倍速いまで持ってきたんですが、さすがに
この先はアセンブラ見てみないと何をどうすればいいか分からんどす。

モノは信号処理用の畳み込み(4096)どす。

>>239
プロパティでカスタムビルドであることをチェックできない?
オプションは言うまでもなく -ptx

グラフィックカードを2枚挿した状態で、メインメモリからGPU側に転送する際に
どのようにしてGPUを指定すればよいのでしょうか？
カーネルから見られるのは、1枚のGPUだけという認識でよろしいのでしょうか？

>>241
cudaGetDeviceCount()からcudaChooseDevice()あたりでは。プログラミング
ガイドのD.1。二枚刺しﾃﾗｳﾗﾔﾏｼｽ

>>240
プロジェクトのプロパティのカスタムビルドステップ、等にもコマンドライン
指定が何もなかとです。次はVSディレクトリに潜って.cuのRuleファイルを
探してみるです。

RuntimeAPIとDriverAPIの違いを教えてください。

CUDA素晴らしすぎです。ようやくグリッドとブロックとスレッドと
globalとsharedとconstantを理解して、今「Xeonより100倍速いぜ！」
な速度を叩き出しています。もっとチューニングがんばってみよう。

>>243
RuntimeよりDriverの方が低レベルとして書いてありますけど、
似たようなの一杯有るし、違い良く分からないですよね。

恐らくは、Driverの方はC++で使うことを想定してない。
実際問題として、nvccはC++としてコンパイルするのでRuntimeAPIの方が使いやすいと思う。

>>244
segdmm のC800では120GFlops程度と言われていますが

4coreのXeonでも75GFlopsです。

どのくらい違いますか教えてください。

>>246
やってみたのが4096単位データ×128個に対して4096のフィルタで
コンボリューション掛ける、なんですが、Xeon2GHzの1スレッド
だと250M操作/秒位しか出来ないですよ。キャッシュミス連発？
GeForce8800ですと、28G操作/秒で動きます。

あ、そんで、Xeonでやるとこうだってのはもれじゃ無くて既存の
なので、弄れないし、作りはよく分からんのです。とりあえず
「このルーチン速くするライブラリ作れんか？CUDAとか言う奴で」
と言われてやってみたら100倍速いので怒りが有頂天になった所です。

でも冷静に考えると、CPU側もカリカリに書けばあと10倍位速くなる
可能性はありますわな。たぶんSSEとか使ってないし。
でもそこから10倍にするには、つまりもう何台かパソコン追加しないと
ならんわけなので、やっぱりCUDA可愛いよ可愛いよCUDAｗ

私のところじゃ、Xeonの3-10倍だなぁ。単体テストで。
# 他の演算と組み合わせると、1.5-2倍が限度（TT

「Xeon」てみなさんが言ってるのは、DPマシン8コアのこと？それとも1コアのこと？
1コアに負けるんじゃ確かにちょっと悲しいよね。

1スッドレ、って書いてるじゃん

つまりこういう暑苦しい感じが最強だと。
喧嘩してる場合じゃねぇ！Quadコア×デュアルCPU＋8800GTXの2枚刺し、
の友情で最強を目指そうじゃないかお前ら！夕日に向かって海まで走れ！
敵は巨大行列のLU分解だ！手ごわいぞO(N^3)！

グリッド数、ブロック数、スレッド数の決め方がよくわかりません。
例えば独立な300個のスレッドがあった場合、
一つのブロックに300のスレッドを割り当てるのがよいのでしょうか？
それともスレッド数は1つで300のブロックを作成するのがよいのでしょうか？
ワープ数などはどのように考慮すればよいのでしょうか？

>>253
あくまでも漏れの場合ですが、
ブロック数＝16。これは使っているカードの「Multi Processor」に合せて。
　違うグリッドに属するデータのやり取りは出来ない事に注意。
スレッド数＝512。目安で。
で、カーネル関数でループで処理を回す場合、最外では
for (i = 0; i < maxcalccount; i=i+スレッド数×ブロック数)
のようにします。これで8192単位での並列化になります。
ループ内部では、グローバルメモリからたとえば１ワード読むなら
sharedメモリ[スレッドID]
　＝グローバルメモリ[スレッド数×ブロック数＋ブロックID×スレッド数＋スレッドID]
となります。
floatを使う場合、スレッドが512なら、各スレッドで最大8ワードもてますが、
sharedメモリは使い切らないほうが良いようです。

ブロック数は、どうせ多く割り当ててもCUDA内部で直列に並べるだけだから
非同期で少しでもCPUと並列にしたい場合を除けば大目に割り当ててOK。
スレッド数についても多い分はどうせ別のワープに割り当てられるから多めでOK。
但し、同期を取る場合には多過ぎるとダメ。
手元のデバイス関数の場合、ブロック数*スレッド数は少なくとも1024か2048以上必要(8800GTXの実測で)。
これらを踏まえると、スレッド数が32ならブロック数は64以上、スレッド数が64ならブロック数は32以上くらいか。
ブロック数の上限は、実測しながら適当に調整するとして、大体1024を超えるといくつでも変わらないと思う。
# これも、具体的なテスト用のサンプル用意したいところだね。

>>188みたいな話が載っている本で使っている本があれば教えてほしいなあ。

入門書の理屈を素直になぞっているような人間には、普通のcpuとgpuとの差を出せなさそう。

>>256
>188は統計学を知らんと自分で導けないし、統計学学ぶ香具師がプログラマになるとも限らんしね。
Webでも知識があることを前提としているか、律儀に平均値との差を求める方式しか載ってなかったり。

私の場合は、電卓の統計機能が個別のデータと平均値の差を使わずにどうやって標準偏差を得るのか
不思議に思って調べて以来の知識かな。
ってことで、統計学を知らない私も>188みたいな話の解説があるなら見たい希ガス。

CUDAというかGeforceってSIMDなの？
SSEは確かに1つの命令で4つとかの足し算が行われるけど、
CUDAの場合、どれがそれに当てはまるの？

>>254>>255
早速のご回答ありがとうございます。
基本的にはスレッド数は目一杯であとはブロックすうを調整すると言うことでいいみたいですね。
ところで、処理する数がスレッド数×ブロック数で割り切れない場合、最後のループはどうするのがいいのでしょうか？
やはりifでやるしかないのでしょうか？

>>259
んにゃ、私はブロック数もスレッド数も実測で決定している。
で、割り切れない場合はループ数で制御するのがいいみたい。
# やっぱり説明するためにはさんぷるが必要かw

>>258
GPUの場合は、SIMDじゃなくてMIMDということになるのかな。
CUDAで考えると、Warp単位で同じインストラクションを走らせて、
複数のデータスレッドを処理するイメージ。

だから、条件分岐をCUDAで書いても、Warp単位で同じインストラクションなので
データスレッドごとに条件が違うと無駄が生じてしまう罠。
詳細は手ごろなサンプルがないので割愛w

>>260
やっぱり基本は実測なのですね。
ところでループ数で制御するというのはどういうことなのでしょうか？
処理する数/(スレッド数×ブロック数)=ループ数
だと思うのですが・・・

例えばこんな感じ。
__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"。これをプログラミングガイドでは「結合した」と訳しているけど……

>>262
ありがとうございます。
なるほど非常にうまい手ですね。
プログラミングガイドはなんか日本語訳がめちゃくちゃで非常に読みにくいですが、
がんばって解読したいと思います。

shared memoryを使ったら、global memoryを直接使うより遅くなってしまいました。
どうも、カーネルを実行する際のshared memoryの確保に時間がかかっているようなのですが、
カーネルを実行するたびに毎回確保するのではなく、常時確保するような方法はないでしょうか？

>>264
その現象は確認していないので不明です。
が、保持する方法は多分ないでしょう。
実は一つ問題があって、shared memoryは16バイト同時アクセスができないので
その点でglobal memoryより遅いのは確認済みです。
# ptx出力を見れば見当がつきますが。

>>261
もれは、CUDAに投入する前に、CPU側で後詰めゼロとかやって、
「CUDA内では一切条件判断は不要」を原則にしています。

>>264
global、sharedへのアクセスがcoalescedじゃないと、めがっさ遅くなります。
それと、sharedに確保した後で、ブロック内の各スレッドでデータを
shareしない場合には、そもそもsharedを使う理由無いです。レジスタ
に順番にグローバルから読んで処理すればいいです。

ブロック間の同期をとる方法を教えてください

「ブロック」は、そもそも同時にブロック001とブロック002が動いているか
どうかも分からないんで、同期出来ないんでは？
CPU側で、カーネル関数その１呼び出して、cudaThreadSyncronizeして、
カーネル関数その２呼び出す。ではないかと。

一応、__syncthreads()という手はありますよ。
但し、全ブロックが同時に動いているわけではない点には注意。
つまり、上の方に書いた分割方法で書いたような大きなブロック数の時には
恐らく巧くいかないかかなり待たされることになるかも知れず。

私の手元のロジックでは、ブロック数が64の時には巧いこと同期が取れています。
# C870, 8800GTX, GTS, GTでしか確認してないけど。

>>269
それって各ブロック内のスレッドの同期じゃないの？

あー、そうそう。失敬、その通りです。
>269は全て、「ブロック」を「スレッド」に読み替えてください。
# どっかでブロック数に切り替わってしまった……_/￣|○

と言うことで、やっぱりブロック間の同期は取れないのでした。

>>268
カーネル2はカーネル1が終了してから実行されるんじゃないの？

>>272
そこが良く分からんのです。
サンプルではcudaThreadSyncronizeでカーネル呼び出しを挟んでるんで、
一応もれも呼び出すことにしておるのですが。

>>273
今、挟んだのと挟んでないのを比べてみたけど、全く同じ結果だった。
ifがあるので、各ブロックはばらばらのタイミングで終わるはずなのだが。

前にも書いたかもしれないけれど、カーネル終了を明示的に知る手段はないので注意。
つまり、<<<>>>による呼び出しから復帰しても、GPUの処理が終了したわけではないので。
なので、次のような呼び出しをした場合、k1()とk2()は並列に動く可能性があるわけ。
--
k1<<<1, 8>>>();
k2<<<1, 8>>>();
--
終了を保証するには、cudaThreadSynchronize()を呼ぶかcudaMemcopy*()を呼ぶかすればOK。

カーネルのループの段数を増やしたら、スレッド数が512だとカーネルの実行ができなくなった。
256ならできる。257はだめ。なんでだろう・・・

1BlockのThread数の制限512ってのと関係があるのでは？

ちょっと舌足らずだったかもしれませんが、カーネル内のループの段数を増やしたところ、そのカーネルが実行できなくなってしまいました。
そこで、スレッド数を512から256に減らしてみたところうまくいきました。
スタックか何かがあふれてしまったのではないかと疑っているのですが・・・調べる方法がわからず行き詰まってしまいました。

実行できないってことは何かエラーが出るの？

使っているレジスタの総数が8192を超えたとか。
unrollしすぎだとか。

カーネルを2つ同時に実行できますか？
2枚のカードでそれぞれ別のタスクを割り当てて、並列で実行したいのですが。

ブロック数が少なければ、勝手に同時実行するな。

Maximum memory pitchって何？

VCのnmakeって使えないの？
cudaをコンパイルできるサンプルMakefileがあったら教えてください。

>>281
サンプルのMultiGPU、8800GTS＋8400で試したら動いたよ。
単にこれ、WINAPIのCreateThreadして中でカーネル関数呼んでるだけですな。
でWaitForMultipleObjectsでWindowsスレッドの終了待っている。

>>283
2次元配列でGPUのメモリを使うときの幅じゃないすか。
cudaMemcpy2Dとか使うときに気にするんだと思われ。
ごめん二次元配列使ってないので良く分からない。
8800GTSも8400も262144bytesです。

演算性能がボトルネックになっているのか、バンド幅がボトルネックになっているのか知る良い方法はないものか

>>286
演算抜きでcudaMemcpyのHostToDevice、DeviceToHostだけやって
時間を測って見る、そして演算入りの時間と比べる、でどうだろう。

__constant__ floatとconst floatの違いって何？
__constant__の方が速いの？実測してみたけど、全然違いが出ないのだけど。
ちなみに何故か定数を直接、式の中に書いたら唯一5％ほど遅くなった。

プログラミングガイドの5章が全然わかんねー
結局のところcoalescedって何よ
Figure 5-2の右側は何でダメなの？？？？

教えて偉い人

>>288
__constant__は全部で64Kバイトある定数格納用メモリ、ホストから
memcpyすることが出来る。const floatは8192個のレジスタを使って
処理される。なので、普通のコード内定数はconst、ホストからどさ
っと書き込みたい固定データは__constant__でいいと思い。

>>289
5-2の右側は、132=33×4、だから？　各スレッドは、
16×sizeof()×何か＋スレッドID、にアライメントされた
アドレスにアクセスしないとcoalescedしてくれないと。

>>290
coalescedを辞書で調べても「合体した」とかしか出てこないので、いまいち意味がよくわからないのですが、
結局0番スレッドは
a[0]
a[16]
a[32]
:
１番スレッドは
a[1]
a[17]
a[33]
:
以外のアドレスにアクセスしないようにすればよいと言うことでいいのでしょうか？

要は、配列のアクセスはarray[m * blockDim.x + threadIdx.x]のような形で
blockDim.xは16の倍数になればよろしって話なんじゃないかと。

cuMemGetInfoフリーメモリと合計メモリを取得したのだけど、４MBにしかならない。
使い方は
unsigned int f,m;
cuMemGetInfo(&f,&m)
でいいんだよね？

>>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。理由分かったら教えてください。

>>290
__constant__って書き換えられるの？

>>295
__constant__は定数メモリと考えればいい。
つまり、GPUからは書き換えられないのでCPUから書き換えると。
# constな変数は初期化はできるけど代入できないのと似てるかもしれない。

>>294>>293
>294のdebugでの数値は多分正しい値が出ているみたいだね。
>293の方は、もしかしたらエラーが返っていないかな?
私のところでも、整数にして201だったかのエラーが返ってきていたのだけれど、
その状態では値をセットしてくれないようだった。
つまり、もしかしたら>293の4MBとというのはfとmが未初期化で偶々入った不定値かも。
# 気になるなら、0で初期化してみればいい。

>>294>>291
coalescedなアクセスパターンにすることは、sharedだけでなくglobalでも重要なので
慣れておくことをお勧め。まぁ、例えば>262に書いたようにすればいいだけだけど。

>>295
cudaMemcpyToSymbol()でCPUからセット出来るよ！
速いし64KBあるんで結構便利です。

問題は、ベクタアクセスできる__global__と違って1データずつのアクセスになるから
逆に遅くなるケースもあることだな。

ややこしさはCellとあんまり変わらなそうだな。

>>299
とにかくメモリのコピーがヤヤコシイんだわ。ホスト側、GPUの
グローバルメモリ、GPUのチップ内メモリで転送しまくらないと
いかんで。やり方間違えると全然性能出ないし。
Cellはその辺どうなの？

256kしかないLSでやりくりするのが大変って聞いたな

両方試した私に言わせて貰えば、どっちもどっち。
確かにCBEは256KiBの壁がねぇ。GPUも64KiBの壁やcoalescedの沼があるけど。
超越関数を使える点ではCUDAが有利。ホストの性能でもPPEじゃ結構泣けるし。

PPEはひどいよな。ホステスにC2D使ってるけど、ifがいっぱいあるような場合はC2Dの方が速いしね。

>>298
それ実験してみたんだけど、変わらないみたい。
・global→shared(行列多数)と__constant__に置いた定行列で行列積
・global→shared(行列多数)とsharedの一部に置いた定行列で行列積
で、後者が1%遅いくらいだった。リードオンリーなだけで、コアとの距離や
所要クロック数はconstantもshared・レジスタも同じなのかもと？

CUDAで逆行列を求める方法を教えてください

page-lockedってどういう状態を表すのでしょうか？

>>304
-ptxでニモニックを出力してみれば違いが判るかと。
例えば、ld.global.v2.f32はあるけどld.const.v2.f32はないからld.const.f32が2回になってしまう。
coalescedなアクセスができるglobalは4クロックでアクセスできるからsharedやconstと変わらないわけで。
尤も、一旦そのパターンから外れるとglobalは数百クロックだそうだから途端に劇遅になるけど。

>>305
先ずは、逆行列を求めるプログラムをCで書いてみてください。
それの、ループの部分を分割する方向でGPU関数を作るのが第一歩になります。

つーかさぁ、一行質問する人達ってなんなの?
情報交換する気がなくて、単にくれくれの精神なんだったら勘弁して欲しいんだけど。
# なんか回答者が数人とそれより少し多いレス要員だけで持ってる気がするよ……

>>304
あーそうそう、書き忘れたけどレジスタアクセスは1クロックじゃなかったカナ。
確実にメモリよりも速いみたい。

nvccにはO2とかの最適化オプションは無いのでしょうか？

-Wallが使えないみたいね
たぶん無さそうな予感

一部のオプションはgccと共通ですよ。
例えば、-O3(恐らく-O2なども)や-pgは使えます。

>>299
どっちもどっちで、用途や慣れの問題だと思う。302がいうように、SFUの有無は
大きいが、それも用途しだい。
自分的には比較的情報がオープンなCellのほうがいろいろいじりやすいと思うが、
最近人気なくて心配。たぶんそのうちLarrabeeが全部かっさらってくんじゃないかな。

GeForce 9800 GX2で現バージョンのCUDAは動きますか？

ストリーム番号の同じカーネルやメモリコピーは順番に、違うものは並列に実行される？？の認識でいい？？？

テクスチャメモリの容量はどのようにすれば調べられるのでしょうか？
どこかに資料があるのでしょうか？それとも何かAPIがあるのでしょうか？

いつも適当に答えてるおじさんです。まことにもうしわけない。
>>313
もれ、8800GTS(G92)で動いてるから、×２で二個見えるんじゃないかと
予想しますけど、一応どっちも対応外だから。漢は度胸でひとつ人柱に。
>>314
ストリームが二つあっても、実行順は関係ないのではと予想。ストリーム
１の半分実行した後ストリーム２の半分実行、もありえるかと。
cudaStreamSynchronizeでストリームの実行完了を待たないといけないです。
>>315
テクスチャメモリはグローバルメモリのキャッシュですが、実サイズが
不明ですね。cudaGetDevicePropertiesでは「alignment=256Bytes」と
読めます。これ以下と以上で実行時間を調べてみると、キャッシュミス
してるかしてないかわかるのではないでしょうかと。

openSUSE 10.3でCUDAは動きますか？

動くらしい

次のバージョンっていつ頃出るの？
新しいカードも出ているわけだしそろそろ、SDKの方もアップデートしてほしいものなのだが

Vista64と9600,9800対応版が欲しいすね。
早く出してくれないとAMD/ATIに浮気したくなりますよ。

サンプル作ると言いつつなかなか暇が取れない私が来ましたよ。

>>313
ドライバが対応しているかどうかが問題。ドライバが認識さえすれば、動くとは思いますが。

>>314
ストリームのご利用は計画的に。ちなみに、CUDA1.0世代の(要はG8xかな)GPUだと
処理の隠蔽ができないらしく、全く速くならないという噂もあります。

>>317
>27-30
荒らしじゃないんだったら、スレくらい検索してほしいと思う今日子の頃。
# いや、荒らしなら検索如何に関わらずご遠慮願いますが。

>>319
取り敢えず、待つしか。下手すると夏の新製品までお預けの可能性も。

>>320
え〜、待っていれば倍精度も来るのに。

英wikiでは、9600GTでもcudaできることになっている
悪い子と漢は居ね゛ーがー

>>322
もれ9600GT買おうと思ってる。
ARCTICのS1Rev2が付く、という報告待ち・・・。ファンレスで使いたくて。

http://www.nvidia.com/object/cuda_learn_products.html

9800GX2 CUDA OKだとでてるよん

おー、更新されている。情報THX。
ドライバは変わっていないから、チップ依存なんですかね。

ML115祭りに参加してOpetron使用をポチりました。
で、nbody.exeを実行すると・・・・ 160〜170GFLOPSという値が。

うぅ・・速過ぎる。初めて見るそのパワーに感動ものですた。

忘れてた　GPUは8800GTですぅ。

そろそろまじめに実用的実装に入らないのかな？
圧縮解凍ソフトに組み込むとか
暗号化ソフトに組み込むとか
一番簡単なのは画像処理ソフトだろうけど

大真面目に実用的な実装で悪戦苦闘していますが何か。

画面の解像度によって使えるメモリ量は変わりますか？

CUDA使うアプリをインストーラで配るにはどうしたらいいもんだろうか？
というところが。Vista対応してないし。

CALならドライバ要らずなんだけどなぁ。
なんで別途ドライバが必要なんだろ。
つかVista版ドライバは3月の終わりごろじゃなかったのかよ！

XP用のビデオドライバをインストールできるはずだから、その状態でXP用のCUDAドライバ入れられれば動くんじゃないかな
と無責任なことを言ってみる

>>332
たぶん時差

ttp://japanese.engadget.com/2008/03/27/windows-vista-nvidia-28-8/
cudaなんかやってる場合じゃないよな

戦犯はマイクロソフトだろ･･･

黒歴史確実なOSだからスルーでいいよ

え、俺パフォーマンス100倍UPにちょー期待してんだけど

普通にopenGL使ったオープンソースのキット使ったほうがいいだろ

今月リリース予定のCUDA 2.0ベータ版でVistaサポートだってさ

CUDAを有効に使うのなら、gccとの相性から考えてもLinuxで使うだろ。JK

仮想メモリが使ってみたいんだもの。

ローカル変数はどこに格納されるのですか？
shared memoryでしょうか？

スタックフレームはないようなので、全てレジスタに格納されると思って宜しいかと。
ローカルに大きな配列取ろうとしたらどうなるのかはしらんなぁ。
ptx出力して読んでみたら?

>>341
商売にならんじゃないか。建前はどうでもいいんですよ。
売れるかどうかだけです。

「商売にできない」の間違いですね。

エンコソフトでCUDA使ったやつ売れれば儲かりそうね。
動作検証が非常に大変だろうけど。
Vista対応不可って所で今の所どうにもならないですなぁ。

漏れがいま一番熱望してるのは、XilinxがCUDAに対応して、
FPGAのコンパイル・シミュレーションが劇速になること。
ああ夢のようだ・・・

GPUで計算させているのだけど、その間のCPU使用率って100%なんだけど、これって正常なの？

正常です。
ブロック数がPE数より多い場合は待たされることがありますし、
メモリ転送や同期を取るときには当然待たされます。
実は待っている間もしっかりCPU時間を消費するのです。

従って、GPUとCPUを巧く連携させて高パフォーマンスを狙うには
いかに待たずに済ませるかが鍵になるので、Streamの使用は必須になってくるかと。

くそ速いglReadPixelsとして使えるかと思ったが
CUDAで扱うデータしかWrite Readできないのね

ぐすん

CUDAにもGLサポート関数があるから、もしかしたら連携できるんじゃないの?
サンプルの、simpleGL辺りになんかない?

PBO使え

GT200と99GX2はどっちが高性能でしか？

このスレの住人なら知っていますね、あの糞開発ツールのことを

・自分のプログラムのバグなのかコンパイラのバグなのかわからない
・他の仕事に応用できない糞開発ツールの独自世界を必死に学習している
・テキストエディタで書いたほうが効率的なのに糞UIツールを懸命に使っている
・糞開発ツールを批判すると「性格が悪いから糞ツールを批判するんだ」と言われる

糞だけど、政治的な理由で無理やり使わされているんですよね。
もう、あんな厨の作った糞ツールを我慢して使うのはやめましょう。

・糞開発ツールを部下に押し付ける上司の命令は無視しましょう。
　上司は糞開発ツールが使われる実績を作ることであの会社のごきげんをとっているのです。
・糞開発ツールを使わせる上司の下では働けません、と上司の上司に直訴しましょう。
・あの糞開発ツール提供会社には「おたくの糞開発ツールは話にならない」と突き放しましょう。
　バグレポートなどしてはいけません。改善要求などもってのほかです。
　あの会社はあなたたちのことをテスター/モルモットとしか思っていません。
・あの会議で「糞開発ツールを使ったら生産性がxx％アップしました」
　なんて話が出たら力強く机を叩き、会議室を出ましょう。
　あの人たちは糞開発ツールをマンセーすることで立場を確保しているのです。

糞な開発ツールを糞だと言える、そんな当たり前の環境をみんなの力で取り戻しましょう。

GX2はSLIと同じように2つのgridを使う必要があるのでしょうか？それとも、1つのgridで両方のGPUを使えるのでしょうか？

CUDAにSLIは関係ありません。
二枚挿しのことをSLIと呼ぶのは間違いです。
何故なら、SLI用コネクタを挿さなくてもCUDAではGPUを二つ使うことができるからです。

で、肝腎な9800GX2ですが予想ではGPUが二つ見える筈なので、一つの時のままだとダメな気がします。

2つのGPUを使うには、CPU側もマルチスレッドにする必要があるのですか？

Streamを使えばシングルスレッドでも何とかなるんじゃないかと思うので、是非とも試してみてください。
# 因みに、StreamはCUDA1.1からの機能なのでサンプルのmultiGPUでは使っていません。

ビデオカード2枚刺しの場合、2枚のビデオカード間のデータのやりとりはPCIe経由なのですか？

>>359
CUDAにはチップ間の転送なんてなかったと思いますが。

CUDA 2.0　in this month

http://forums.nvidia.com/index.php?showtopic=63956

Vistaもサポートするよ。

一方のVistaはMSに見捨てられた。

PGがゴシップ記事鵜呑みにするなよ・・・

Linux でCUDA使っていますが、pthreadを使って以下のようにスレッド内でcudaMemcpyを呼ぶと
invalid device pointerでcudaMemcpyに失敗します。なぜでしょうか。
スレッドを作らずに、mainから直接testCopy(NULL)と呼び出すと、問題ありません。

#include <pthread.h>
#include <cutil.h>
#define NUM 512
int *h_buf=NULL;
int *d_buf=NULL;
pthread_t thread;
void *testCopy(void* args) {
CUDA_SAFE_CALL(cudaMemcpy(d_buf, h_buf, sizeof(int)*NUM, cudaMemcpyHostToDevice));
CUT_CHECK_ERROR("Copy failed");
}
int main(int argc, char **argv) {
CUDA_SAFE_CALL(cudaMalloc((void**)&d_buf, sizeof(int)*NUM));
CUDA_SAFE_CALL(cudaMallocHost((void**)&h_buf, sizeof(int)*NUM));
for(int i = 0; i < NUM; i++) h_buf[i] = i;
pthread_create(&thread, NULL, &testCopy, NULL);
pthread_join(thread, NULL);
}

Linuxよく判ってないけど次の点をチェック。
・cudaMallocHost()はメモリのページロックをしてしまうので、malloc()にしたらどうなるか。
・そもそもスレッドを分ける意図は何か。Stream系APIでは事が足りないのか。

CUBLASの使用においても他スレッドで確保したメモリは扱えないので注意が必要です

>>365
mallocも、newも試したのですが、だめでした。
CUDAの制御をするスレッドと、GUIの面倒を見るスレッドを分けたかったのです。
>>366
スレッド、またいじゃだめなんですか。それは、CUDAの仕様なのでしょうか。

>>367
CUDAの面倒を見るスレッドでメモリ確保すればいいだけじゃん。

ttp://pc.watch.impress.co.jp/docs/2008/0416/nvidia.htm

NVIDIA必死だなｗ

AdobeのPremierとか、先に取り込んだとこの勝ちになるんじゃね？

おまえら2.0（Beta版）きたぞ

http://forums.nvidia.com/index.php?showtopic=64974&pid=363542&st=0&#entry363542

otu

CUDAプログラミングガイドに書かれているGPUの解説は
CUDAで使う場合限定の話ですか？CUDA対応GPUが行う処理全てについてですか？
例えば複数のグリッドで共有するメモリなど。

>>374
共通する点もあるとは思いますが、CUDAでは(利便性と引き換えに)GPUの利用に
制限が掛かっていることから類推して当てにならないと思った方がよさそうです。

夏ぐらいまでにCUDAで一本何か作ろうと思うけど
今から始めるなら2.0がいいの？というか何が違うんだ？

まずはEmuでつくれ

Visual Studio 2005じゃないとだめなのね

サンプルが2005じゃないとビルドできないですね

AMDの新しいbetaSDKって
Radeonじゃ動かんよね？

>>380
スレ違い。

>>379
Linux版はインストールされているgccがあれば大丈夫。
# cygwinのgccが使えるといいんだけどねぇ。

>>379
金がないってことか?
Express試用したら?

>>381
スレねーじゃん
ケチくせーこと言わないで教えろよ

GPGPUのスレあるだろ

>>381
逆にgccのバージョン違ってると上手く動かないけど。
個人的には*BSDで動いて欲しい。それも64bitで。

>>385
そんときは、違うバージョンのgccを入れて設定を変えておけばいいらしい。
なので、ドライバは兎も角nvcc(とemu)は*bsdでも動く可能性はあるかと。

あー、今確認のために見に行ったら2.0betaの案内がCUDAZoneにも出てますね。
CudaVisualProfilerも2.0betaが出ているらしい。

>>384
だってAMDのは知らないし〜

かけ算にかかるクロック数や足し算にかかるクロック数を知りたいのですが、
そのような資料はないものでしょうか？

ガイドブック。

>>387
基本的に、掛け算も足し算も同クロック。
問題は、そこに至る過程なのでptxを出力してループ内の行数を数えるとか
メモリアクセスの個数を数えるとか。
そもそもプロファイラを使いこなせとか。

SLI環境で、GPUから別のGPUに直接データを転送することはできるのでしょうか？
もしできるのでしたら、やり方を教えてください

>>390
>355-360

CUDAってｎVidia限定でしょ？
サンプル動かして感動したけど
せっかく作っても使える環境が限られると意味ないね

確かに、仮令MacやLinuxで使える場合があったとしても、Windowsの半数以上で使えないからね。

まぁ、使いたい人が使うからいいのよ。

スレッドをいくつか使うときに

スレッド0、１，２の計算が終わってから
スレッド4の計算を始めるってどうやって書けばいいの？

待ち同期どうやって作ればいいか解らない

GPU側関数で同期を取るのは__syncthreads()でできるけど、
恐らくそれではパフォーマンスが出ないと思われ。
スレッドは数個だけ動かしている積もりでもWarp単位でしか動かないので、
分割の仕方を見直した方がよさそう。

ついに、gtx 280がでるね。
第２世代型の統合シェーダーアーキテクチャってどんなんかな。
倍精度サポートとか、10秒で計算打ち止めへの対処とかいろいろありそうだけど。

自分的には、分岐への弱さの改良とか、ブロック間の同期機構とか、手を入れてくるんじゃないかと思う。
あとは、各種メモリの配分をどうしてくるか注目。shared memoryを倍増とかあるかな。

取らぬ狸のなんだが
GTX 280/260ではかなりの高速化が見込まれるがどう？

消費電力がやばすぎると思う

>>398
お前そんなネガ発言すると
このスレをN房の溜まり場にすっぞコラ？

400Wゲトー

IU鯖で使えるくらいになってほしいとこだよなあ

今頃になって説明の間違いに気づいた……

>188の説明において、途中に出てくる分散と共分散の式が全て(n-1)で割られてないや。
相関係数の段階では分母子に出てくるから消えてなくなるんで忘れていた。
スレ違いだけど間違いを放置するのも難なんで念の為。
# Excelの分散は分母が(n-1)で、共分散は分母がnなんて仕様なのに気づいて、ここを思い出した。

CUDAを使ったプログラムは専用ドライバが入っていない環境では動かないですか？
配布は無理がありますか

っ[knoppix for CUDA]

普通に使うWindowsソフトの高速化として使うものではないと

最新のnVIDIAドライバにだったら入ってるでしょ

でも手軽にCUDAテクノロジを使えるのソフトというのも面白いね。

>>406
最新のnVIDIAドライバにはCUDAが使える機能が入ってるってこと？

スレチなんですけど、ほかに質問できそうなところがなくて・・

nVIDIA のグラボで水平スパンやDualViewの設定をアプリから制御する API みたいなのはあるのでしょうか？

アプリケーション起動時に自動でそのあたりを設定したいのですが。

あると思うけど、しらね。NVIDIAに聞いたら?

CUDAは別途ドライバが必要
Geforx8以上ならドライバ入れれば動くけど
一般ユーザーにドライバ入れるなんて作業が簡単に出来るわけないので
ソフトとして配布するのは無理じゃないかな
専用のインストーラーでも作るなら別だろうけど

>>410
以前 別件でnvidiaジャパンに問い合わせてみたけどシカトされたからなあ。
とりあえず ELSAジャパンにメールしてみました。

なんで標準ドライバにCUDAのドライバ入れとか無いんだろう・・・・

>>412
あ、やっぱり?w
1000枚単位の客じゃないと相手にしないって噂は本当だったか＜NVIDIA
ELSAもどうだろ。最近連絡くれなくなったからなぁ。

NVIDIAジャパンは営業ばっかで技術スタッフ皆無だからな

>>411
そうなんすかー
>>413
ですよねー

配るならDirectXやOpenGLでやるしかないということか。

>>410
エルザから返事がきましたよ！

>基本的には制御不可能かと思われます。

泣ける。あとは nVIDIA のデベロッパーサイトのリンク。号泣。
1000枚とはいわんけど、30枚くらいは買ってるんだがなあ。

だから、最初からドライバの入ってる
knoppix for CUDA と一緒に配布すれば
良いじゃん。

ドライバ内の公開向けには作ってない制御APIを、
公開しろとか騒いでも無駄だと思うぞ。
ドライバのバージョンでコロコロ変わってもおかしくないもんだし。

GPUはin-orderですかそれともout-of-orderですか？

ミクロではインオーダ
マクロではアウトオブオーダ

CUDAで見えるレベルだったらOoOと思っといてよい。

速攻でミスった。

× CUDAで見えるレベルだったらOoOと思っといてよい。
○ 〜キホンインオーダ。ただしコンパイラはそれなりの最適化を施す。

>>417

今のバージョンで使えるかどうかはわからんけど、NVControlpanel（nvcpl.dll）のAPIはある。

ttp://developer.download.nvidia.com/SDK/9.5/Samples/DEMOS/common/src/NvCpl/docs/NVControlPanel_API.pdf

年代からして、Vistaはおそらく無理。

CUDAのサイトがリニューアルしてるぞ。
賛否両論あるとおもうが、とりあえず整理されている。

ROPユニットとストリームプロセッサの違いって何？

>>424
見た見た。

そして、例によってCUDA-Enabled ProductsにGTX280/260が載っていない罠。

GTX280 なんすかこれｗ

Vipperより糞じゃんｗ終わってるな

>>413
>>416
最新のドライバから統合するようになったみたいだよ
ttp://www.nvidia.co.uk/object/winxp_177.35_whql_uk.html
提供されてるのが200シリーズ向けだけなので現在確かめようがないけど

ん、いや統合はvistaだけなんだろうか・・・？
ttp://www.nvidia.co.uk/object/winvista_x86_177.35_whql_uk.html
こっちにははっきり
Adds support for CUDA? Technology.
と記述されてる

PhysXとCUDAは同時に使えますか？

ゲロビディア終わったなｗ

だれかgtx280でcudaやったやつおらんのか。

ゲロビディア氏ね

>>432
やってみました。確かに速いが

>>434
確かに早いが？。。。その先が気になる。
すまないが、他のボードとの比較とかの情報があるととても嬉しい。

発熱が凄いってこったろう

GPUとしては今はラデのほうが良いみたいね。
Brook+ってどうなのよ？

実用にはまだちょっと厳しい。
F@HはBrook+で書いてるし使えないこともないってレベル。
今のところOpenGLやDirectXと協調出来ないからそこも問題かな。

これから暑い夏を迎えて、発熱が凄いボードはどう評価されていくか興味深いね。

洞爺湖サミットでNvidia名指しで批判されるらしい
だから焦っていろいろアピールしてるらしい
環境もっとも悪い製品を作ってる会社の代表格て
声明が盛り込まれる予定

CUDA 2.0 Beta 2
http://forums.nvidia.com/index.php?showtopic=70171

cudavideodecodeなんて今更何に使うんだと思ったけどDXVAがないOS向けか

倍精度計算したいどの変数使うの？

すいません。助けてください。
VC2005でサンプルをビルドしようとすると以下のエラーが出ます。

Visual Studio configuration file '(null)' could not be found for installation at 'C:/Program Files (x86)/Microsoft Visual Studio 8/VC/bin'

コンフィグファイルが必要なのかと思い、C:/Program Files (x86)/Microsoft Visual Studio 8/VC/binの下にnvcc.profileを
置いてみたのですが結果は変わりませんでした。

どうしたらいいですか。

ＯＳはwindows xp 64で、CUDAのtoolkitとSDKは1.1です。

>>444
cuda_build_rule.zip入れた？

倍精度で計算ってどうやって書けばいいの？

大学の課題でで使うことになったのでマニュアル読んでみたけど、日本語訳がファッキンなのは仕様ですか？
冗談抜きで英語の方が判りやすかったｗ

>>446
授業ページおしえてください

>>447
うちの大学の教授が俺用に30秒で考えて出した課題なので、特に授業ページはありません。 ごめんね。

>>446
授業ページおしえろｺﾞﾗｧｱｱｱｱｱｱ

>>446
NVIDIAジャパンには、まともな技術者がいないので仕方ありません。
尤も、営業にもまともなのがいるかどうか些か疑問ではありますが。

Gerovidia Japanには屑営業しかいねーじゃん
実際CUDAとか国内研究で利用していて
すごそうなのは、Gerovidia本社から来るし

cudaのおかげで英語力が向上しました

cudaのおかげで彼女ができました

double float はいつ???

いつって、2.0β使えばすぐできるんじゃないの? エミュで。

CUDAが.NETから使えるようになってるらしいので興味持ったのですが（今現在動作環境なし）
C#とかVSとかとの親和性とか、ネイティブと比べて処理性能どれくらいかとか、試した人いたら教えてください

CUDA.NET
http://www.gass-ltd.co.il/products/cuda.net/

ランタイムだけじゃね

アセンブリ読み込めば後はどの言語でも使えてしまうんですかね
PowerShellやIronPython（pyCUDAというのもあるみたいですが）から使えたりしたら面白そう、面白そうなだけですけど

超初心者ですが質問お願いします。

自分のパソコンにグラボが2枚刺さっているのですが
CUDAで使うグラボを選択するにはどうしたらいいんですか？

めっちゃ初歩的な質問ですが回答お願いします。

サンプルにあるだろ

さわりだけだけど日本語資料みっけた
ttp://www.easize.jp/download/Workshop20071003.pdf

>>461
今月のインターフェイス８月号にも似たような資料があったYO!

誰かGpuCV使ったやついねーか…
makeできねーん、OpenCVで聞けばいいのかこっちなのかわかんね

>>462
環境書かけよくず

>>462
釣られてやったけどVS2005、VS2008
SuseLinux RedhatES 5.0 Ubuntu 8.04 CentOS 5.1で
ビルドできたんだがｗ

ビルドすらできないPGってカスもいいとこだろｗ

うわぁぁぁぁファッキンなのは日本語訳だけかと思っていたらSDKもだったぁぁぁぁ

誰だcutil_math.hの float / float[2,3,4]の演算子オーバーロード書いた馬鹿は
俺の3時間返しやがれ
inline __host__ __device__ float2 operator/(float s, float2 a)
{
　　float inv = 1.0f / s;
　　return a * inv;
}

ひょっとしてcuda SDKのcommon/inc/以下のヘッダーってわりと危険だったりする？

>>463-464
ちょっｗｗｗｗｗ

環境はCentOS 5.19 64bit
GpuCV 0.4.2 driverはcudaのやつ174.55

gpucv,resources,sugoitoolsをダウンロードして

ttp://picoforge.int-evry.fr/cgi-bin/twiki/view/Gpucv/Web/WebHomeDownloadResources#SugoiTools
を参考にファイルコピー

cp -Rf bin\gnu\ ..\resources\bin\gnu\
cp -Rf lib\gnu\ ..\resources\lib\gnu\
は　sugoitools/bin/gnu,libが無いので実行できず無視して↓のみやった
cp -f include\SugoiTools\*.h ..\resources\include\SugoiTools\
cp -f include\SugoiTools\*.inl ..\resources\include\SugoiTools\
cp -f include\SugoiTracer\*.h ..\resources\include\SugoiTracer\

しかしgpucvに入って./createSolutions.shで

./createSolutions.sh
Generating GNU makefiles:
...GPUCVHardware
...GPUCVTexture
...GPUCVCore
...GPUCV
...GPUCVConsole
...GPUCVSimpleApp
を読み込めません: そのようなファイルやディレクトリはありません
を読み込めません: そのようなファイルやディレクトリはありません
sed: ./projects/gnu/example/Makefileを読み込めません: そのようなファイルやディレクトリはありません

無視してmakeしても

#make
==== Building CUDA plugin ====
==== Building GpuCV ====
==== Building GPUCVHardware ====
GLBuffer.cpp
<コマンドライン>:1:20: 警告: missing whitespace after the macro name
/usr/include/GL/glxew.h:150: error: ‘GLulong’ has not been declared
make[2]: *** [../../../Debug/gnu//GPUCVHardware/GLBuffer.o] エラー 1
make[1]: *** [GPUCVHardware] エラー 2
make: *** [SUB_DIR_ALL] エラー 2

>>465
cutilはサンプル集だと思ってた方がいい。

カーネル内で二次元配列を使うと

Advisory: Cannot tell what pointer points to, assuming global memory space
というwarningが出るのですが、どのようにすれば出なくなるのでしょうか？

>>469
「カーネル内」ってどこのこと?
いっそ再現する最小限のソースを貼ってくれた方が話が早そうな希ガス。

>>467
必要なもの
premake(これはビルドするかパッケ拾ってくるのじゃいいな
OpenCV
libglew1.5-dev

SugoiToolsとかいうふざけた名前のライブラリSVNから盗んでくる
ちなみにこのSugoiBakaToolsを作ったやつはUnixとLinuxの.soを作る流儀を
知らん白雉なので許してやれ

svn co https://sugoitools.svn.sourceforge.net/svnroot/sugoitools/trunk sugoitools
cd sugoitools
premake --file Premake.lua --target gnu --os linux
make

これで./libにlibSugoiTools.soが生成される。

次に、GPUCVをSVNから盗んできてくれ

cd gpucv
premake --file Premake.lua --target gnu --os linux
cp ../sugoitools/lib/gnu/*.so ./lib
ln -s ../sugoitools/include/gnu/SugoiTools SugoiTools
ln -s ../sugoitools/include/gnu/SugoiTracer SugoiTracer

make

後は必要なライブラリ入れるだけでうまくいく
いかなかったら>>470にゴルァしてくれたらまた何かかくぉ？

>>471
うひょーPremakeの使い方を始めて知ったお！

でもGPUCVTextureのコンパイルで止まる…

>必要なもの
>premake(これはビルドするかパッケ拾ってくるのじゃいいな
>OpenCV
>libglew1.5-dev

premake,OpenCVはインストしました。libglew1.5-devがCentOSで
rpmがなかったので
ttp://chihara.naist.jp/people/STAFF/imura/computer/OpenGL/glew/disp_content
を参考に導入しました

>svn co https://sugoitools.svn.sourceforge.net/svnroot/sugoitools/trunk sugoitools
>cd sugoitools
>premake --file Premake.lua --target gnu --os linux
>make

>これで./libにlibSugoiTools.soが生成される。

>次に、GPUCVをSVNから盗んできてくれ

>cd gpucv
>premake --file Premake.lua --target gnu --os linux


ここまで完璧です！

>cp ../sugoitools/lib/gnu/*.so ./lib
コピー先は../resources/lib/gnuでしょうか、
でもgpucv/lib/gnuのフォルダにGPUCVHardwared.soができてるのですが。。。
うーん、両方(/resourcesと/gpucv/lib)試しましたがうまくいきません。。。

>ln -s ../sugoitools/include/gnu/SugoiTools SugoiTools
>ln -s ../sugoitools/include/gnu/SugoiTracer SugoiTracer
コピー先は../resources/include/SugoiTools等ですよね。。
こっちは/gpucv/Sugoitoolsで試してません
コピー先は
ttp://picoforge.int-evry.fr/cgi-bin/twiki/view/Gpucv/Web/WebHomeDownloadResources#SugoiTools
を参考に決めました。。

でmakeすると
# make
==== Building CUDA plugin ====
==== Building GpuCV ====
==== Building GPUCVHardware ====
..中略..
Linking GPUCVHardware
==== Building GPUCVTexture ====
DataDsc_GLTex.cpp
..中略..
TextureGrp.cpp
<コマンドライン>:1:19: 警告: missing whitespace after the macro name
../../../src/lib/GPUCVTexture/TextureGrp.cpp:100: error: prototype for ‘bool TextureGrp::AddTextures(DataContainer**, unsigned int)’ does not match any in class ‘TextureGrp’
../../../include/GPUCVTexture/TextureGrp.h:118: error: candidate is: bool TextureGrp::AddTextures(DataContainer**, size_t)
make[2]: *** [../../../Debug/gnu//GPUCVTexture/TextureGrp.o] エラー 1
make[1]: *** [GPUCVTexture] エラー 2
make: *** [SUB_DIR_ALL] エラー 2

うわああああああああああんごｒｒｒｒｒっるうううううううううううああああああ

>>473
GPUCVの馬鹿どもはMakefileすら満足にかけないみたいだから
あれだけど

premake間違ってるからちょっと修正
premake --file Premake.lua --target gnu --os linux --cuda

あとコンパイル関係のログは一番上見て
○○.h No sucha file or directoryって出てるはず
きちんとログは最初から理解していきましょう。

以下の場所にGPUCVxxxx.makだかって糞Makefile入ってるから
ここで○○.hが足りないって言われたら　-Iでパス足してやって
vim ./project/gnu/lib/

それでうまくいくはず、この作業はちなみに6回ぐらい繰り返すので
GPUCVプロジェクト市ねって500回唱えるのがいい

sugoitoolsってなんだよｗ馬鹿にしてんのか？
と思ったらマジで日本語の「凄い」から命名してんだな
ttp://sugoiaware.free.fr/

>>475
どっかのインターフェイスカードかとおもったよｗｗｗ

確かにｗ

>>475
カナダにはSUGOIっていうスポーツウェアブランドもある
それくらいで驚いちゃ駄目だｗ

Windowsなんだけど、出来上がったexeって、
他のPCで動かすためには、exe以外に必要なものって何？

cudart.dllが無いって言われるのは、CUDA対応カードじゃないから？

cudart.dllをsystem32に放り込むかpath通せ。

.cuファイルって外部のincludeファイルは読み込めないのでしょうか
template.cuに
#include <cv.h>

と1行書き加え、（OpenCVという画像処理用のライブラリです）

Makefileに
# OpenCVのためにパスを通す
NVCCFLAGS+=`pkg-config opencv --cflags`
LIB+=`pkg-config opencv --libs`

でMakeすると
/usr/lib/gcc/x86_64-redhat-linux/4.1.2/include/mmintrin.h(49): error:
identifier "__builtin_ia32_emms" is undefined

/usr/lib/gcc/x86_64-redhat-linux/4.1.2/include/mmintrin.h(62): error:
identifier "__builtin_ia32_vec_init_v2si" is undefined


....................というのがダラダラ続く…

Error limit reached.
100 errors detected in the compilation of "/tmp/tmpxft_000010b9_00000000-4_template.cpp1.ii".
Compilation terminated.
make: *** [obj/release/template.cu_o] error 255

環境は
OS:CentOS 5.2 64bit
CUDA driver: NVIDIA Driver for Redhat Enterprise Linux 5.x with CUDA Support (174.55)
cuda toolkit:CUDA Toolkit version 1.1 for Redhat Enterprise Linux 5.x
CUDA SDK: CUDA SDK version 1.1 for Linux
です。よろしくお願いします

>>481
普通になんでもインクルードできるよ。
だからこそ、cv.h経由か何かで読み込んでも解釈できないmmintrin.hまでインクルードしてしまっているわけで。
取り敢えず、cv.hが何をインクルードしているのか順に見てみたら?

どうみても、Intrinsicの関数が無いと言ってるだろ

>>479
ドライバ入れてもらわないとダメじゃないのか

普段Radeonを使ってるんだけど、CUDAを使うならやっぱり取り替えなきゃダメかな？
もう一枚追加して普段はRadeonでCUDA用にだけGeforceを使えたらと思ったんだけど…

最近CUDACUDA言ってるやつ多いけど、ShやBrookGPUも忘れないでね。

>>485
ドライバがバッティングしないなら、その方が楽だね。GPU一枚だと、CUDAで暴走したときに画面が崩れる場合がある。
# ついでに言えば、Radeonで3Dアプリを動かしていてもCUDAの処理速度への影響が出にくいと言うメリットもある。

SIGGRAPH2008でいろいろやるらしーよ
民主化だってさ
http://www.siggraph.org/s2008/attendees/techtalks/sessions.php

GPGPUやってるやつらの論文
精度甘くてうそばっかりなんだけど

クイックソートレベルで嘘あるって
なんなんだろう

SIGRAPHも悪の片棒担いでるし
国内の正会員も歯切れ悪い胡散臭い爺多いし
あいつら本当に計算機科学全般に害悪な
ことばっかりしはじめたなぁ

>>489
具体的に晒してください

>>490
さらせねーよｗ俺社会的に終わる
それだけはヤダ

>>491
つまり、いいがかりなんですね

CPUに比べて○○倍ってよく見るけどCPUの方は最適化してなかったり１コアしか使ってなかったりやけに古いCPU使ってたり
数字をよく見せるために胡散臭いのはあるな

>>487
ドライバの干渉が一番の不安なんだよなあ。
Geforce2枚刺して用途別にするのが一番いいんだろうけど完全に環境を変えなきゃいけないし…

>>493
海外のも見てみてよ
査読ありの論文でレベル的ほとんど査読無しと
変わらないヒドイ質の論文多いから

GPUへのフェッチ性能とかそれのオーバヘッド
0扱いとかいう凄まじい論文あるぞ

>>493
結構同意
CPUの数十倍早くなりました、とあるが
CPU，GPU間の転送時間を含んでいませんとかアホかと

>>489
SIGGRAPHはもう正しく論文を精査できてない。
で、年度によってはとりあえず載せちゃえってスタンス。
それでも多過ぎてSIGGRAPH ASIAやっちゃえって事に。

あのーここは学会加入してないとだめなのでしょうか
皆さん頭のよさそうな話ばかりで困ります

同意する点もあるけれど、言いがかりに近いと思う点もある。

転送時間を含まないケースでも、CUDAに関して言えば転送時間を隠蔽できる可能性もあるし
初回だけで中間に転送がいらないケースもあると思う。

実際のところ、意味があるかはさておき単純なロジックを組んでみるとちゃんと理想値に近い時間で
処理できるのは事実だし、最新でないCPUでもGPUを使って高速化するという運用はありだと思うのだけど。

まぁ、一般論で片付けられる問題ではないから個別にレスするのは控えるけどね。

>>499

>転送時間を含まないケースでも、CUDAに関して言えば転送時間を隠蔽できる可能性もあるし
>初回だけで中間に転送がいらないケースもあると思う。
CUDA内部の処理と
CUDA-バス-キャッシュ-CPU間の話どっちを
しているわけなの？厳密に答えてくれないかな？思うじゃなくて

そこ厳密に計測すると汎用的なアルゴリズムになるとせいぜい40GFぐらいしか
出てない。8CoreのXeon２台MPIするよりも全然遅くなっちゃうんだけど

ついでに、WoodcrestXeon3GHzで8800GTを使った場合の実運用での処理時間について一言。
行列演算みたいなGPU向きの演算と違ってGPUには不利な演算なのだけど、
単体演算ではCPUのみに対してGPUを使った場合は約2倍の処理能力が得られた。
尤も、ファイル処理なども含めた実運用状態での処理能力比は1.3倍。
まぁこんなもんでしょってことで。要は、如何にGPU向きの演算に落とし込むかだね。

スレチだしそろそろ自重しようぜと過疎スレで言ってみる

>>501
それは何コア使って1.3なの？
SSEも入れて、TBL使ってMathLib使った場合と
比較して1.3倍？

>>500
「汎用的なアルゴリズム」で40GFlops出て、8coreXeon2台のMPIと勝負できるならいい方なんでない?
8coreXeonでかりかりにチューニングしたら、GPUなんて使わない方が速くて当然だと思う。
そういうCPUと勝負するなら、GPU単体の性能じゃなくてCPU+GPUでCPUより「何割」速くなるかでしょ。

あー、書き忘れてた。>501に書いたのは2core*2CPUのシステムで、CPUのみとCPU+GPUの処理能力比ね。
CPUでも1coreだけ使うよりは4core使った方が当然4倍近く速いのだけど、そこにGPUを足すことで
更に1.3倍になったということ。1coreだけと較べてどのくらい速くなったかは失念している。
1coreのみを1とすると、2core2CPUで3倍、それにGPUを足して4倍位じゃなかったカナ。

具体的な論文名出せよ、そうしないから話がややこしくなる

>>482

>取り敢えず、cv.hが何をインクルードしているのか順に見てみたら?
すみません、こういう時どう対処していいのかわからなくて。
includeしていく順番を探せば何か見えてくるのでしょうか

cv.h - cxcore.h - cxtypes.h - emmintrin.h - xmmintrin.h - mmintrin.h

とつながっていました。mmintrin.hはMMX？xmmintrin.hはSSEでしょうか
私の直感だとこのくらいの解釈が限界です…

includeファイルを順にコメントアウトしてみましたがやはり通らず…
gccではコンパイルを通っているのにnvccではダメということは
MMX？が使えないようになっているのでしょうか

__builtin_ia32ほげほげが何者なのかさっぱりです…

ググるとWinXPではOpenCVが使えているっぽいので
XPのnvccではMMXが使えてLinuxのは使えないのでしょうか…

何か足りてないんだろうね。

コンパイルできるようになりました。

cxtypes.hの中でSSE2がオンになっている場所
#if defined WIN64 && defined EM64T && (defined _MSC_VER || defined CV_ICC) \
|| defined __SSE2__ || defined _MM_SHUFFLE2
#include <emmintrin.h>
#define CV_SSE2 1
#else
#define CV_SSE2 0
#endif

を見つけました。見てみるとWIN64のときだけオンになるみたいで…
ここで
#include <emmintrin.h>
#define CVSSE2 1
をコメントアウトして
#define CVSSE2 0
と書いたらコンパイルできるようになりました。
Vistaではこのようなことをしなくても
コンパイルできたのでは32bitからだったのでしょうか…

>>509
お疲れさん。そう言えば、Windowsでは64ビットでコンパイルできないってレスがあったからその辺りも関係するかもね。

SIGGRAPHではないがCanny Edge Detection on NVIDIA CUDAなんてのを読むとなかなか笑える。

・Matlab並に速いアセンブラで最適化されたOpenCVと比較したよ！
　~~~~~~~~~~~~~~~pugya-
・使ったCPUはIntel Core2 CPU 6600 @ 2.40 GHzだよ！
　　　　　　　　　　　　　　　　　　　＾ EなのかQなのかはっきりしろと。
　　　　　　　　　　　　　　　　　　　　ま、所詮OpenCVだから結果に影響ないけどｗ
・GPUはGeForce 8800 GTX, 768 MB、OCなんてしてないよ！
　　　　　　　　　　　　　　　　　　　　　　　~~~~~~~~~~~~~~~~~~~~研究レベルでOCを考慮するッ
　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　 そこにシビれる！あこがれるゥ！

研究なら並列アルゴリズムだけ書いておけばいいんじゃね

opencvってアセンブラで最適化されてたっけ？

>>512
移植性考えてまったくされてないよ
デフォルトではシングルスレッドだからマルチコアも使われてない

今後GPUはえらい成長遂げるのに対しCPUはもう脚詰まりだから
いきなり8coreXeon2台とGPU一個が張り合えるという事実は大したことだ
後はプログラムの処理内容を最適化させれば良い

ハードgf98GX2 *2
ソフトwinXP VS2005
ドライバ　cuda2.0β

上記の環境でC++とcudaを混在させることはできるのでしょうか？
キャプチャーカードのSDKがC++なので・・・

>>515
ドライバはCUDAじゃねぇぞ。

CUDAの開発では、*.cuのファイルがnvccでコンパイルしてデバイスモジュールとホストモジュールが出力される。
その後リンクするのはVSのリンカになるので、この段階でVSのオブジェクトモジュールとリンクできる。
但し、nvccはベースがgccなのでextern "C"を使うなどの工夫が必要かも知れず。

GCC使う以上、素直にLinuxで開発すれば良いのに。
なんでワザワザ苦労してWin糞使うのか？

大昔のように、
プログラム開発者＝ハード＆ソフト環境を自由に整えられる人
じゃないんだよ。

nvccってclのバージョンでエラー吐くから
Windowsだとclベースじゃねーの。

>516
ありがとうございます。自身で試行錯誤してみます。
>517
MFCを使ってwindowsアプリケーションを作るのにwinの方がいいかと思いまして

対応カードのドライバがXP〜しか対応してないみたいだけど、Win2kでプログラミングしてる人はいますか？
余ったPCでCUDAをやってみようと思ってるんですが>>173あたりを見ても対応して無さそうで…

余ってないPCに二枚挿した方がよくね?

余ったPCにLinux載せちゃうのがよくね?

余ったPCを俺にくれればよくね？

いつのまにかにノート向けも対応リストに載ってるのね

2.0 北っぽいよ
ttp://www.nvidia.com/object/cuda_get.html

>>515
.NET for CUDA つうのがあるみたいなんだわ
突撃してみてほしいです。

日本版でもBeta2取れたけど相変わらずアナウンスはないな

SDK2.0入れて見たんだが、threadMigrationってサンプルプログラムが
"cuCtxDestroy","cuCtxPopCurrent","cuCtxPushCurrent"が定義されてないってエラー吐いてる。
他のサンプルはコンパイル通るので、インストールが悪いのかSDKが悪いのか分からん。

バージョンアップの観察はいいから何か作れよお前らｗｗｗ

作ってるよ〜
仕事だから詳細語れないけど。

作ってるけど研究なので言えない
けど上手くいかない
ああああああ

はぁ…

研究って学生の卒論ですか？
それとももっと高いレベル？

自宅研究員…

>>533
CUDAをどう使うかと、論文のレベルは関係ないだろｗ

http://gpu4vision.icg.tugraz.at/index.php?content=downloads.php
CUDAを使った領域選択

CUDAのビデオ講義、リンクきれちゃってるね
せっかくiPod nano買ったから英語の勉強に聞きたかったのに…

２日でCUDAマスターしますた
すごい簡単だなこれ
なんか質問あれば聞いて良いよ

>>538
100万桁のπは何秒くらいで計算できますか？
1995年当時のスパコンで5秒らしいですが超えられますか？

>>539
やってみないとわかりません

CUDAでソートするアルゴリズム思いついた　うひょ

だれか多倍長をCUDAで効率的に実現する方法教えてくださいー

>>541
http://www.nvidia.co.jp/object/cuda_sample_data-parallel_jp.html

>>543
バイオニックソートってやつと同じｗｗｗ
車輪の再発明かｗ

当方 Mac なんだけど、CUDA 2.0 インスコしてサンプル make したのは
いいんだけど実行しようとすると

dyld: Library not loaded: @rpath/libcudart.dylib
Referenced from: /Developer/CUDA/bin/darwin/release/scalarProd
Reason: image not found
Trace/BPT trap

って出て先に進めない。どういうことなの。。。。。。
UNIX 詳しい人教えて下さい

並列に計算して計算結果を１つの変数に合計するってのをやってるんだけど
10回に1回くらい計算結果が狂うのはなぜだ

>>546
Ａｔｏｍｉｃ関数使ってみるとか・・・後、parallel reductionを読んでみると
良いかも

>>546
そういうのはCUDAの最も苦手とするところだ。
全スレッドで同期を取って、代表1スレッドが合計するのが手っ取り早いが遅い。
全スレッドで同期を取って、代表nスレッドがmスレッド分合計してからnスレッド分を合計するのが無難か。
或いは、n個になった時点でCPUに転送してしまう方がいいかもしれない。

転送が遅いんだよなあ

ストリームを使えば殆ど隠蔽できるよ。

>>548
ｎ個分の計算結果を別メモリに退避しておいて1個のスレッドでそれを順次合計するってこと？

>>550
kwsk

>>551
そのままＣＰＵ側へ転送してＣＰＵで合計しろってことでしょ？

すんません、日本語版ドキュメントが公開されたってあちらこちらで書いてあるので
公式サイト探したんですが、なんかないっぽいんですが、これってもしかして削除されたの？

日本語のマニュアルが必要なほど内容濃くないぞｗ

グローバルメモリを使っても速度に限界を感じたんで、
テクスチャメモリを使ってみたいんですよね

で、そのために日本語ドキュメントも読んでみたかったんすよ

テクスチャメモリはグローバルメモリの代わりにはならんと思うが。
共有メモリも使い難いし、定数メモリはデバイスから書けないし。
# 書けないのはテクスチャメモリもそうだけど。
アクセスパターンを見直したほうが医院で内科医?

vs2008には対応するんだろうか

2.0出たけど対応してなくてがっかりした

>>559
俺は結局2005入れた。

>>549
もれのやってるのでは転送に3msec、演算に400msecくらいなので
全然オーバーヘッドにはなってないすわ。演算の負荷小さ杉なんでは。
>>556
書き換え不要な定数行列をまとめて__constant__に置いて見たすが
ほんの3%くらいしか変わらなかったすわ。

両方とも、もれのコードがだめな可能性ももちろん有るけど。大有りッスけど。

CUDAじゃなきゃ困るって用途がいまいち思いつかない

それは、CUDAじゃなく直接GPUを扱うほうがいいということか、AMDのStreamナントカでもいいということか、なんだんだ？

>CUDAじゃなきゃ困るって用途がいまいち思いつかない
なら使わなきゃいい

100万回以上回るループとか、
何千回単位の二重・三重ループが有るならCUDAｻｲｯｺｫｩﾝｷﾞﾓｯﾁﾞｲｲｲｨｲ

>>564
バーカwwwwwwwwwww
頭悪いなお前

>>565
確かに
初体験のあの気持ち・・・忘れられない・・・

気持ちいのは分かるけど普段使わないようなものだしね
既にあるものを自力でCUDA対応させるのも気力が沸かないしね
倍精度浮動少数が扱えないしねｗｗｗ

>>568
CUDA 2.0は？単精度に比べてかなり速度が落ちるらしいけど。

ハードの問題だから無理だよ
やってみた

>>570
GT200世代じゃないとハードの問題で使えないのか・・・

二重ループは兎も角、三重ループとなるとCUDAは苦手だと思うが。
一重は並列にしても、二重ループが残ってしまう。
最近のIntelCPUはループが無茶苦茶高速だから、WoodcrestでOpenMPでも使われたら太刀打ちできなくなる。

CUDAをループの自動並列化だと思っていらっしゃるｗ

ところで2.0正式版は皆さん安定してる？
うちは、テクスチャmemoryがなんだか変。
エミュで正しく動いているのに実機だと挙動がおかしい。

CUDAでクラスが使えないのが痛いな
木構造系のアルゴリズムは並列処理にかなり向いてるし応用範囲も広いのにな
CUDAで無理やりやろうとすると無駄な処理をわざとさせないといけなくなるし
プログラムがむちゃくちゃ汚くて見てられない

PyCUDAかCUDA.NETあたりを使ってぜひ感想を聞かせて

CudaArrayに、48ビットや24ビットのRGBのデータを入れて、テクスチャにバインドできている方います？
うちではうまくいかないんですよね。

CUDAってCPUモードがあるけどドライバが無い環境だと自動的になるの？

エミュレーションモードの話かな?
自動的に切り替わるほど融通は利かないよ。

dim3 threads(100,1);
method<<<1, threads>>>();
これはいけるんだけど

dim3 threads(100,100);
method<<<1, threads>>>();
ってやると一回も呼ばれないんだけど
何か勘違いしてる
threadIdx.xとyで２次元的に呼び出せるんじゃないの？

スレッド総数は512まで(詳細はdeviceQueryを実行するべし)。
従って、100*100は拙い。

追記:
その条件だけなら、dim3 threads(ThreadsOfBlock, ThreadsOfBlock)にして
dim3 blocks(100 / ThreadsOfBlock, 100 / ThreadsOfBlock)を追加して
method<<<blocks, therads>>>()するのが定番かな。
勿論、methods<<<100, 100>>>()でもいいけど効率は若干落ちることになりそう。

method<<<dim3(100,100),1>>>();
にしたらいけた
もしやブロックって並列処理じゃないの？

GPUのプロセッサからあぶれた分は時間軸方向に並列になります。

whetstoneとかのベンチマークをCUDAで動かしたいんだけど、
とってきたソースをCUDA環境で動かすには書き換えないとダメなの？

sharedメモリなんだけど制限とかあるの？

method<<<dim3(1000,1000),4, 1000*1000*4>>>
とか

ガウシアンぼかし３ｘ３を５００ｘ８００のフルカラー画像で約０．３秒だった
8400GSですけど
こんなもん？もう少し早いのを期待したんだが

そのアルゴリズムはCPUでやるとどのくらいかかった？

ブロックの分け方をいじったら０．１５秒になった
CPUで同じの組むのめんどいです
エミュレーションモードってどうやってやるの？

>>589
＞エミュレーションモードってどうやってやるの？
nvccにオプション指定するだけだよ。
あくまでもエミュレーションだから余計に遅くなるけど。

>>587
サンプルで似たようなのなかった?
フィルタ類は色々あったと思うから眺めてみるといいと思う。

>>586
あんたもdeviceQueryを実行する必要がありそうだ。
sharedにそんなに取ると、物理量を大幅に超えるから巧く動いたとしても無茶苦茶遅くなるぞ。

>>585
whetstoneなんて、並列演算に向かないと思うのだけど。
あー、繰り返しを並列にすればいいか。それだったら移植もそれほど難しくはない。

デバイス上にRGB(unsigned char)の画像配列を確保してある状態で
これをOpenGLのテクスチャにホストを介さずにバインドして利用出来る？

すいません質問です。
CUDAでFFTやってるサンプルコードとかどこかにないでしょうか。
探しても見つからなかったんです。

>>592
simpleCUFFT違う？

>>590
「並列演算に向かない処理を並列にすればよくってよ？」
これ最高よね

>>592
そのくらい自分で考えれ、そんなに難しい事じゃなじゃん

>>593
＞これ最高よね
???

>>592
CUFFT使うだけならnvccも要らない。

>>590
誰かが「このアルゴリズムは並列化には向かないうんぬん」と言った
アルゴリズム×128本を同時に実行してしまうとか最高よね、という
意味なのではないかな。確かに上司の驚愕を呼ぶね。

俺は並列化できそうな新たなアルゴリズムを考案しろという意味かと思った

PIの計算だって本来並列には向かない

CUDAでデバッグDLLが付属してないようなのだけど
例外処理ってどうやればいいので？
try..catchとか使えるの？

無理。

多体シミュやりたいなーと思ってたら，本屋で見かけたGPU Gems3に載ってた
でもそれだけの為に買うのもなー・・・３はいつ原版が公開されるんだろー

なんかDirectX11で並列演算に対応するのと
11世代のOpenGLの仕様にOpenCLっていうCUDAの類似品が実装されるらしい
たぶんCUDAはこのまま消えていく

GPGPUの可能性を示してくれただけで十分だよ

カーネルのネストって出来るの？
__global__ void a(){
...
}

__global__ void b(){
a<<<dim3(100,100),1>>>();
}

void main(){
b<<<dim3(100,100),1>>>();
}

みたいな

なんか根本的にひどい勘違いしてなくね？

>>604
GPU内部からカーネルを発行することは不可能

__device__ でプログラミングガイドを検索汁。

>>602
CUDAの内部にOpenCLが含まれる構成だからCUDA

処理に時間がかかるとGPUが完全停止してタイムアウトでドライバレベルのエラー出すんだけど
一回そのせいで画面全体がぐちゃぐちゃになってWindowsが操作不能にまでなったぞ
非同期実行じゃないんかい
どうすんのこれ？

>>609
それは質問なのか？愚痴なのか？
質問だとしたら・・・分かるな？

>>609

OSがVistaなら下の資料を参考にレジストリ弄ってみるとか。

WDDM によるタイムアウトの検出と GPU の回復
ttp://www.microsoft.com/japan/whdc/device/display/wddm_timeout.mspx

192コアとか216コアとか240コアとか

使いづらいですよ！漏れの弱い頭がパンクしそうです！
お願いだから隠された力を覚醒させて256コア版出してくだしあ！！！１１

いや一回計算方法実装すれば後はそれを使いまわすだけだろ

>>612
CUDAをVer2にするんだ。デバイス情報取得APIでコア数が判るようになっているぞ。
# つーか、256コア版の歩留まりが悪くてサブプロセッサ単位で減らして対応しているんじゃないの?w
# PS3のCBEが7SSEなのはそういう事情だそうだし。

CUDAってなんて読むの？くーだ？

んだ。

cubaがキューバなんだから
cudaはクーダだろ・・・常識的に考えて・・・

キュ〜(＞＜)〜だ

8400GSだと、h264エンコで実速出ないねぇ
12fpsがやっとだよ
もう少し速いかと思ってたんだけどな

ＣＵＤＡを使いたいと思っている初心者です
macのxcodeでもできますか？？
imacで８８００ＧＳです

>>619
よりによって、CUDAが動く最底辺の方のGPUを使わなくても……

>>620
NVIDIAの公式サイトが全てなので、そこを読んで判断してください。
つーか、xcodeってなに?

>>620
５万払って雪豹もらってください。
動くレベルじゃないって話だけどね。

初心者です。
CUDAサンプルを動かしてみて感じた事なのですが、
HLSL,GLSL,Cg言語それらを使わずに、使った時のような絵が出せるのでしょうか？
CUDAもGPU上で計算しているみたいなので・・・

それじゃあ俺も初心者です

じゃあおれも

>>623
サンプルを見たのなら判ると思うけど、OpenGLはほぼそのまま使えるようですよ。

http://en.wikipedia.org/wiki/OpenCL#Example
これ見るとCUDAにそっくりだな
CUDAをちょっと修正するだけでOpenCLに対応出来そうだけど

>626
OpenGLは使えるのはわかってます。
CUDAでピクセル単位の計算できるのかが知りたいのです（汗
Cg言語を使ってバーテックスシェーダで計算していた処理を
CUDAで実現いてみたら、うまくいったのですが、
フラグメントシェーダでやってた処理が実現できなくて・・・
そもそもできるのかどうか・・・
ってところが知りたいのです。

そのためだけにCUDAは作られました

つかサンプルにSobelFilterってのがあるだろあれみれ

>>630
SobelFilter見てみました。
見落としてました。
ありがとうございます。

日本の公式フォーラムできた
ttp://forum.nvidia.co.jp/

なんとも直訳のような回答ばかりw

直訳なんかしたらまともな日本語になってるわけないじゃないか

あの回答がまともな日本語だと思うのか?

ああ俺の中で日本語ドキュメントの思い出が醜化されていたようだ

悪くない

え、俺これ読んですごいwktkしてるんだけど。
サンプル投稿みたいな質問したら、気が向いたらNVIDIAが答えてくれるかも知れないって事でしょ？

正直、ここで聞いた方がましだと思う漏れもいる。

投稿者：　NVIDIA CUDA Team
どんな人　専門家
自信　　　自信あり

ちょっと面白いｗ

少し読んでみたが日本語のあのドキュメントよりはずっと読みやすいよｗ
あれは酷すぎた

それでもどう考えても英語のほうが読みやすいけどな。

アトミック処理に放射線の危険はありません。:-)

質問です。
CUDAを使って、GPU上で計算している部分の一部を、
FBOとCg言語を使って、オフスクリーンで計算させることってできますか？

もちろん、無駄なことはわかっているのですが・・・

なんでCUDAのサンプルってコンソールアプリばっかりなんだろう

コンソールアプリじゃないときは
CUT_DEVICE_INIT(argc,argv);
CUT_EXIT(argc, argv);
のargc,argvってとりあえず 0 と nullとか渡しとけばいいの？

>>644
日本語でどうぞ。

>>645
そもそもCUT_DEVICE_INITはMultiGPU環境でGPUの番号を指定したりするためにあるようなもの。
0とNULLでもいいけど、敢えて使う必要もない。
CUT_EXITに至っては、プログラム終了時にプロンプトを出すためにあるようなもんだ。
# 詳細は、cutil.hを読め。

>>647
ありがとうございます

>>645
ウィンドウ出してグラフィックだすようなサンプルも多いが。

普通にイメージクラスでCUDAを実装してDLL化して使ってるが
DLLだからコンソールだろうがWindowだろうが使えるぞ
でもお前らにはやらない

>>694
頭のほうしか見てなかった...

ここでロングパス！

ttp://pepper.is.sci.toho-u.ac.jp/index.php?%A5%CE%A1%BC%A5%C8%2FCUDA%2FProgrammingGuide2

今日からcudaプログラミングを始めようとしている超初心者です
nvidiaのGTX280のピーク性能は933GFlopsだって歌われているのですが、
しかし240のコアで1296MHｚで動作しているのなら、311GFlopsになるはずですよね？
この3倍の数値の差はどこからくるのでしょうか？

あ、もしかして1サイクルで３つオペランドを消費するような命令があるってことで
しょうかね？

maddならmul+addだから2倍なんだけど、3倍なんてあったかな? とこの前から思っている。

http://journal.mycom.co.jp/special/2006/g80/007.html
madd + mul で3の気がする

ベクタ計算じゃないので最内ループの記述は楽々
それでいて300GFLOPS台の性能がでるなんて夢のようですね

巧く嵌まればね。書くのは楽だけど、チューニングが大変なのは変わらないわよ。

ローコストにCUDAプログラミングやチューニングのコツを掴むのに適したマシンってありますか？

いま持ってるノーパソのGPUがnvidiaならよかったんだが、そうじゃないので
安くCUDAできるマシンを買おうか検討中。「試し」なので自作とか高価なのは避けたい

参考になる話があったら聞かせてくださいまし

予算を教えなはれ。
14万でMacBookとか？
後は、Nvidiaチップ内臓の自作で10万切ることも可能
今デスクトップPC不所持で5万の予算だったら無理。

意味のある試しプログラミングができればいくら安くてもいい
中古ノートのオンボロで構わないと思っている

奮発しても10万といったところですかね

私は一世代前のCore2Duoで8800GT入れているけど9万ほどだったかな。
今だったら、45nmのCore2Duoでメモリ1GB積んで9600GTでも入れれば結構楽しめると思う。
DosPara辺りのゲーム用PCが丁度そんな感じのスペックじゃないかな。
# 都合いいことに、「ゲーム用」はVistaじゃないことが多いしね。

>>661
10万と書いたが、MacBookの14万というのはちと食指をそそられるな…

当方unixやlinuxはそこそこ扱えるが、Macは十年以上触ったことがないし
開発環境とかどうなっとるんでしょうか

あ、そうそう
今のノートに買い替えてから使ってないけどモニタとキーボードはあります
モニタったって今時CRTのSXGAですけどね

>>661>>663
ありがとうございます。参考にします

ちょっとDosPara見てきたけど、XP足しても8万くらいで作れるようね。後は余裕次第かな。
この手のBTOの常で、キーボードとマウスは嫌でもついてくるけど。

それは別として、MacBookで色々苦労したいのならそれはそれでありだと思う。
私なら、ミニタワーで安く済ませた分で、1280x1024を越える大きさの液晶モニタでも買うけどね。

>>663
やっぱりMacは今でも苦労が多いのかw でもまあひと通り調べてはみよう。
ドスパラのBTOも、後で自分でチェックしてみることにします。

bootcampでlinux入れちゃえばOKジャマイカ？
と思ってググッたら酷い、お勧めできない。

GTX260は単体で3万円切っているので、工夫すれば10万でデスクトップ組めるかも？？

別にWinノートPCを持っているなら、ミニタワーはLinuxでもいいんじゃね？
CUDAをWinで使うとnvccがgccベースだからC++関連の内部関数がめんどくさいよ。

まあ先に言った通り自作は手控えたいです。トラブルを楽しむほどの余裕が今はないので。
ボード一枚突っ込むくらいですかね。自分の手でやってもいいやというのは。

>>670
VS2008があるので、できればそっちと連携させながらやりたい
（ガワとか作るのはC#が便利なので…）と思っていたんですが、難しいですかね。
cudaが基本gccだというのは判っているので、まあ何ならIPCで繋げばいいやくらいに
テキトーなことを考えていたわけですが。

VC向けのプロジェクトテンプレートあったよね

>>672
ああ、ありましたね

これから風呂入って寝てしまうので、今日はこのへんで。
皆様どうも、いろいろご親切にありがとうございました。

一応、明日以降もときどきこのスレをチェックするようにします。

グラボだけ変えればいいじゃん

PCIEマシンを持ってないなら
E1200+9400GTあたりで組めば3万以内で組める

ちなみにうちのE1200は３Gで動いてたけど
E7200に変えたのでサブにした
３Gあればメインマシンとしても十分だと思うけど
ただマザーはそれなりにOC耐性の高いものが要るので
どうしてもミドルクラスの1万5000円前後になるけどね

>>675
だいたいそのへんで検討中です。
Windows XPをインストールするとケチっても5万くらいにはなっちゃうようですが
（自宅PCは98SE→Vistaノートへポーンと飛ばしたので、XPを持ってない）ま、それはそれで。

強力な浮動小数演算ができる装置を手に入れても、物理屋、ごく一部の化学屋、工学屋の超上位層ぐらいしか使い道が無いぽ。

と無知な輩がわめいております

実際ない

GPU Gemsの3とかに沢山載ってるよ。
暗号化解析とかパケットフィルタリングとか。

パケットフィルタの為に２００W近く常時負荷掛けるなんて無理だ

＞物理屋、ごく一部の化学屋、工学屋の超上位層ぐらいしか
もともとゲームやCGでさんざん使い倒されているデバイスなのに何を言うのか

CUDAは、ということであれば、たとえばこれを土台にファイナンス系のミドルウェアが作られたら
そっち系の需要が一気に開花するかもしれない
今んとこCUDAのサイトで紹介されている事例はいささか高尚すぐる気がするけどな

どっかの銀行でやってるんじゃないか?

銀行ってイメージ的に扱うトランザクションは半端なく多そうだけど、それがSIMDではなさそうな。
必要なのはPOWERに乗ってるような十進演算器じゃないの？

そう言えば、NVIDIAの営業が得意気に「銀行系には1000台単位でお買い上げいただいてます」とか言ってたな。

それ演算用ではなくて表示用だったりして。

DoubleFloatのみで対決したら、最新GPU1機 vs. 最新のクアッドコアCPUのどっちが勝つと思う？
CPUはamd64と、SSE等使用時（誤差を考えればむしろこっちと比較するべき？）の両方で予想してほしい。

CPU

用途による。
考えるのが面倒くさければCPU。

メモリアクセスのペナルティがあるから、その観点で比べてもしょうがないよ

銀行で使うとしたら、勘定系ではないだろ。
そうではなくて、商品開発やシミュレーションなど、1円2円ずれてもいいような業務

銀行とかトラフィックが確かにすさまじいけど
システム改変するコストもすさまじいぞ
数十年に一回やれるかどうかだろ
それに負荷１００％で２４時間なんてカードが耐えられるとは思えない

常時負荷100%という状況になっていること自体設計ミス
60〜70%が適正

>>671
マーケティングの人に直接聞いた話では、そのうちそのへんは改善されると思うよ。
あ、機密事項と言われてるので具体的には言えない。

最近銀行システムの開発で、6000人のSE集めた超プロジェクト失敗したものねえ。まあ当然だが。
SEが6000人だからねえ。プログラマはもっと多いとかもう想像つかない。

TMPGencのCUDA対応版が出たんでインストールしたんだけど
CUDAの項目にチェックできないのは何故・・？

ドライバは178.24でグラボがASUSの8800GTS(640MB)

スマソ自己解決
g80はダメなんだってねOTL

>>698
ｲｷﾛ。

G80はストリーム系のAPIが使えないからねぇ。

>>696
JRとかの鉄道や、電力といったインフラ系はもっと大きい。
しかし、大きいが故にPJ失敗しまくってる。

人数を増やせば増やす程、集めた人材の質は低下する。
そして頭脳労働の場合、一番質の低い人のレベルに
足並みを揃えなきゃいけなくなるからなぁ。

しかし戦中・戦後に一気に作ったシステムが老朽化して、
銀行どころでなく大規模な改修がどれもこれも必要なのだが。

#mixiで見掛けたよ＞団子の中の人

>>662
HP ML115サーバ機に、GF9400GTあたり刺して、
Linux入れたら？ 1CD-Linuxの knoppix for CUDA
なら、最初からCUDA環境が構築済みで、サンプル
も憑いて来るし。

慶應義塾大学泰岡（やすおか）顕治研究室 Yasuoka Laboratory
http://www.yasuoka.mech.keio.ac.jp/cuda/

個人的にはGF8200なM/BのオンボでCUDA走れば、
裸M/BのCUDAクラスタ組もうかと思ってるが、
CPUやメモリの値段を考えると、ML115の方が
安上がりなんだよな。

>>702
これはビックリ！こんな激安サーバがあるなんて知らなかった…
激安なのにPCI-Expressとかついてて（x16必須な）nvidiaのグラボもちゃんと動く、
ということでゲームの人達にも人気があると…ふむふむ。

ところで素のML115はメモリ512Mなのだけれど
上記研究室のページによるとknoppix for CUDAの推奨動作環境はメモリ2G以上、とある
ML115を使う場合、安いやつを別に買ってきて刺し換えればよろし、ということですね？
（ML115もhp直販だとメモリ増設オプションはECCつきの高いやつしかない…）

ML115が16k、9400GTが9k、2Gメモリも安いのは3k未満、で30kを切りますな。
個人的にサーバ機もAMD64もknoppixも使ったことがないので、
それらの組み合わせとなると微妙に不安だ（笑）が、いずれにせよこの値段は魅力的

大変参考になりました。ありがとうございます。

デモ機で借りたTeslaC1060使っているんだけど、ホストCPUがAMDのPhenom。
Xeonに較べて遅い遅い。普段使っているXeonに8800GTの組み合わせの方が早いって何さw

つまんない質問だけどGTX2xxの人は電源いくらよ。
＋150Wくらいはマージンとったほうがいいと思うよな？よな？

なんに対して+150?
GTX280ボード単体での消費電力は236W、GTX260でも180Wクラス消費するからね。
ついでに言えば、補助電源用コネクタもGTX280は6ピン+8ピンの特殊コネクタが必要だし。

>>706
システム全体で。500W電源以上推奨って言ってるけどじゃあ500Wで安定するかっていうと
信用できねー

無理。

とすると、マシン一式組んで貸し出してもらうのがベストだよな？
よし参考になった。

8800GTなら100Wだし、補助電源コネクタも6ピンだけで済むよ。

うん、俺も8800GTまでなら550Wで余裕といえるラインかなと思っている。

GTX280を使うのなら、700Wクラスの電源が欲しいところだね。

CUDAは8800以上のクラスで無ければ意味ない。
8500とかはとりあえず走るだけでパフォーマンスは全然駄目。

8400GSとかになると額面性能でもCore 2シングルコア以下だな。

　　そだ　 |------､｀⌒ｰ--､
　　れが　 |ﾊ{{　}} ）））ヽ､l l ハ
　　が　　　|､{ ﾊリﾉﾉﾉﾉﾉﾉ）､ l l
　　い　　　|ヽヽｰ､彡彡ﾉﾉﾉ}　　に
　　い　　　|ヾヾヾヾヾヽ彡彡} 　や
　　!! 　　 /:.:.:.ヾヾヾヾヽ彡彡} l　っ
＼＿＿／{ l ii | l|} ﾊ､ヾ} ﾐ彡ﾄ
彡ｼ ,ｪ､､､ヾ{{ヽ} l|l ｨｪ=ﾘ､シ} |l
lミ{　ﾞｲｼﾓ'ﾃ､ミヽ}ｼｨ=ﾗ'ｧ､ }ﾐ}} l
ヾﾐ　　 ￣~'ｨ''': |ﾞ:ｰ.￣ 　 lﾉ/l | |
ヾヾ　　 " 　: :　!､ 　｀　　lｲﾉ l| |
　>lﾞ､　　　　ｰ､,'ｿ 　 　 /.|}､ l| |
:.lヽ ヽ 　　ｰ_ ‐-‐ｧ'　 /::ﾉl ﾄ､
:.:.:.:＼ヽ 　 　 二"　 /::// /:.:.l:.:.
:.:.:.:.:.::ヽ:＼　　　 　/::://:.:,':.:..:l:.:.
;.;.;.;.;;.:.:.:.＼｀ｰ-- '" //:.:.:;l:.:.:.:l:.:

サブノートPCでCUDA動くようにならんかな。
通勤、出張の途中でいぢってみたい。こんな
時でもないと、仕事に直結しないプログラム
組んでる暇無いからなぁ。

>>703
AMDの場合、メモリコントローラがCPUに内蔵なので、
ECCでもnon-ECCでも使える。安い通常のnon-ECCメモリ
1GBx2枚買ってくればOK。ML115はNTT-Xで買えば、
13800円(送料込)。

つ N10J

>>716
回答ありがとうございます。

>>716
俺はサブノートでソースは書いてるよ。動作確認は自宅に戻ってからだけどね。
一発で動けば気持ちいいもんだ。

質問です

OpenCLが出たらCUDAはお払い箱ですか？

>>719
エミュは動いている?

>>720
いいえ、画像処理だけがCUDAの使い道ではありません。

Apple主導の言語処理系って流行らんだろ。
GPU版Objective-Cだと思え。


ちなみにNVIDIAから補助もらってる俺は仕事につながるって言えるのかな？

>>722
その仕事、こっちにくれw
情報少なくて、参ってるんだ。

>>722
ObjectiveCは言語仕様からしてクソだったから流行らなかった。
それだけです。

>>723
メールサポートだけもらってるけどマニュアル落として自分でやったほうが早いしなぁ

俺のほうこそ各ptx命令のレイテンシ・スループットの資料欲しいんだけど。
Intelはそういうのまめに出してくれるから助かるんだが

YellowBoxだっけ？
WindowsでもMacでも動くアプリケーションが動くフレームワークとか
大風呂敷広げてあれ結局どうなったっけ？

MicrosoftはDX11があるからOpenCLの標準化なんて破談する可能性大
Appleのフレームワークは地雷ばかりで困る。

>>725
なんだ、ないのか。NVIDIAの日本法人は、ろくに情報持ってないっぽいんだよね。

>>725
ptxは中間言語だろ？

>>726
先入観が身を滅ぼすだろう。

質問スレッドなので、唐突に質問するわけですが、ごきげんよう

CUDAのSDKに付いてくる
Programming Guide Version2.0の60ページ目の真ん中あたり

For devices of compute capability 1.x,
the warp size is 32 and the number of banks
is 16 (see Section 5.1);
a shared memory request for a warp is split into one request
for the first half of the warp and one request for the second half of the warp.
As a consequence, there can be no bank conflict between a thread belonging to the first
half of a warp and a thread belonging to the second half of the same warp.

が分からない。
何が分からないのかというと、これはShared Memoryの最適なアクセスに関する記述なんだけど、
ワープの中に並列実行できるスレッドが32個あるというのにshared memoryのバンク数は16個しかない。
普通に考えたら２つのスレッドが同時に１つのbankにアクセスするわけで、
思いっきりバンク競合するはずよね？
でも、この記述はバンク競合が起こらないって自信を持って記述されているわけよ
nVidiaの人教えてちょ

Half Warp(つまり16スレッド)ずつスケジューリングされるんじゃなかったかな
だからバンク競合は起きない
nVidiaの人では無いが

なら何でWarp=16スレッドとしないんだろう…というのが俺の疑問

中の人いわく
命令レイテンシ隠蔽のためにクロック毎にインタリーブしてるだけだから細かいことは気にすんな

なるほど〜
ワープの正体は16並列と見つけたり
ってことですな！

GeForce 9400MってCUDA使える？

2.1でサポートできるように頑張ってるけど間に合わないかもしんないって言ってた

今月中に何かしら動きが・・うわなにをする
くぁwせｄｒｆｔｇｙふじこｌｐ；「’」

個人的には1.3世代の1スロット厚のGPUボードが欲しいのだけれど……
# 出ますと言ってた奴はその後連絡寄越さないしなぁ。

コードの実行時動的生成（分岐除去とかパラメータの定数化とかってレベルで）って
CUDAではいまんとこ無理なんだよな？
Larrabeeが出たらそういう最適化できる部分はXbyak使おうかなと思ってるんだが


っていうか、SPMDじゃないプログラミングモデルまだー？

>>730-732
に補足。

各SPは最大2issue同時実行なんだけどデコーダは半速。
1SPあたり4スレッドでインターリーブして同じオペレーションを実行するとちょうど命令供給が間に合う構造だな。


1warp＝
16にすると、デコーダは等速か、半速×2にしないといけない。
デコーダの負荷を抑えたかったんじゃないの？

CUDAはじめようと思って調べ始めたんだが、
7xxxシリーズはなんで切り捨てられたのか・・・

今週末に9600GT買いに行かなきゃ

どうせなら260だか280あたりにしといたら

電源とかケースとかの敷居高くない？

>>740
切り捨てられたんじゃなくて、始めから想定されてない。

ttp://journal.mycom.co.jp/column/graphics/index.html
この連載のはじめの方のGPUの進化を追うと、少しは判るかも知れない。
で、どうせならQuadroFX3700をお勧めします。8800GTとほぼ同一仕様でお値段10倍w

CUDA-Zなんて便利なものがありました。
forum.nvidia.co.jp

それよりレイトレベンチマークのほう、Ｂｉｏ１００％が作ったのか！
SuperDepthとかカニミソとかが蘇ってきたぜ

スレ違い済まん

>>745
http://noridon.seesaa.net/

>>745-746
まだ生きてたのか！
PC-98では大変お世話になりました。

そしてブログを読んでみたら、超わかりやすい！
coalescedの意味とか、8/29のエントリみたいなメモリアクセスが遅い理由とかよく分からなかったんだよ。助かった。

vista sp1にCUDAをインストールしたいんだけど
ドライバ：○
↓
tool kit：×インストールが終了しない。。

で上手くインストール出来ないんですが、誰かしりません？
強制終了したらアンインストールの項目にtool kitの項目があるのにアンインストールするとerror:5001で失敗しやがるし。。。

最悪

TMpegEncのCUDA対応は4フィルタだけで今のところあまり効果がないみたい
AviUtilのCUDA対応フィルタもパフォーマンス出ないという理由で公開停止
今からでもチャンスありますかね?

作りたいなら是非作ってくれ

>>748
管理者権限でやっている?
後柱ね。

なんかLinux向けのドライバにCUDA2.1入ってるらしーよ
ttp://www.nvidia.com/object/linux_display_ia32_180.06.html

>>738
http://www.nvidia.com/docs/IO/47904/VolumeII.pdf
スライド87

>>752
ついに来てしまったか>>736

ラジオシティできるソフトってありますか?できればソース付きで...

Cg勉強しようと思って調べてたら、CUDAってのもあるんだな。
それぞれできる事って、具体的に何が違うの？

とりあえず7600GTしか持ってないんで、CUDAは使えないんだが、
Cg勉強するぐらいなら、CUDA勉強したほうが圧倒的に良いなら
対応グラボ買おうと思うんだが

そういう何処にでも載っていることすら調べられないのならCgにすれば。

CUDA 2.1 beta
ttp://forums.nvidia.com/index.php?showtopic=82290

VC++9とDX10インターオペラビリティがやっと

cuda sdkのサンプルを実行するとtest failedと出て実行できないんですけど。
環境はos xp 64, quadro FX 4600です。
先ほどnvidiaからドライバとツールとSDKをダウンロードして
インストールしました。ドライバは更新されています。
visual studio 2005も入れました。

Teslaを使っているのですが、電源コードを抜く以外の方法で、装置を再起動
する方法はないでしょうか。
バグのあるコードを何度も実行した結果、cudaMalloc()が返ってこない
状態になっています。

たわけた質問だと思いますが、お許しください。
NVIDIA製のカードが入っていないPC上で、
nvemulate.exeを利用してCUDAを使用する事は可能なのでしょうか？
実際の処理に使うのではなく、CUDAプログラミングの練習に使うのが主です。

>>760

Sシリーズならホストを再起動するだけで復活しませんか？

やっぱ大学くらいしかまだ使ってないのかね

最近発売された、GeForce9300、9400を積んだＭＢ、
少し前のGeForce8200、8300を積んだＭＢでも実用ではないですが、
CUDAのプログラミングをして走らせる事は可能なのでしょうか？
誰もmGPUでCUDAを使っていないので…

みんな何の計算させてるの？

株価予測をリアルタイムに

株価の予測はできんだろ。アホか。

>>767
阿呆丸出し乙

株価の予測ができたって言ってるのは、数年前の慶応が出してた論文ぐらいじゃねーの？

>>768
予想と予測は明確に違うんだぜ？

論点がづれてるー髪もづれてるー

>> 762
shutdown -> 電源切断 -> 電源投入の手順を踏むと、復活しました。
ただのrebootで良いかどうかは試していません。

>>772
色色と掲示板の使い方を間違っているw
で、reboot試してないなら報告しなくていいから。

>>771
髪はずれないと思うぞ、髪は。

>>765
私の所では、最近はFFTWの代わりにCUFFTでFFTを計算させている。

>>764
NVIDIAのサイトのCUDA ZONEでリストアップされていれば、使える。

>>763
んなこたーない。

--device-emulationでは正しく動くけれども、実機では動かないときには
ソースコードをにらむしかないのでしょうか。

nvcc --device-debug (-G) というオプションがあったので、これをつけて
コンパイルすると、ptxas が Parsing error を出して失敗します。

http://forums.nvidia.com/index.php?showtopic=66291
の会話を見ると、--device-debug は今年5月の段階ではまだ使えなかった
らしく、私の場合と現象が似ているので、以前としてまだ使えないままか
と思ったのです。

>>769
そりゃ、予想は出来る罠
ただ、外乱はいつも不明だし、確定解は得られない。
つまり、最尤推定しかできないし、当然推定結果が大ハズレってことも、
初めから推定理論に謳われてる

株価予想が正確になればみんなそれを信じて買うようになるでしょ
予測自体が株価に影響を与えだして本来の予測とは違う値動きを始める
そして的中率は下がる
つまり一定以上の正確な予測を行うことは不可能なのだ

ここには、当たり前の簡単なことを、必死に難しく言おうとしてる
能無しがたくさん居るようだねｗ

１つの.cuの中で実装しているglobal関数の個数によって、
Kernel呼び出しのターンアラウンドタイムが変わるという奇妙な現象に遭遇してます。
特にKernelで処理するデータが少ない時に顕著になります。
関数の数を５〜１０個で変えてみると、ターンアラウンドタイムは
最悪値で80μsec、最良で30μsecでした。
この値は
timer.start();
for (int i=0;i<100;i++) test_kernel<<grid,thread>>(test);
cudaThreadSynchronize();
timer.end();
みたいな書き方で調べてます。

9個目、10個目あたりで底があるようなのですが
こういう現象について、合理的な説明はありますか？
僕にはさっぱり見当がつかないのであります。

何かを勘違いしている

>>778
再現できるソースを貼ってくれたら解析するじょ。

#include <windows.h>
#include <stdio.h>
#include <cuda_runtime.h>
//Round a / b to nearest higher integer value
int cuda_iDivUp(int a, int b) {return (a + (b - 1)) / b;}
#define BLOCK_DIM ( 32)
template <unsigned int loops>
__global__ void testcuuuuKernel(float* d_h0, unsigned int size)
{
unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < size) {
float d = d_h0[i];
for (int j = 0; j < loops; j++) {d -= j * 0.1; d += 0.9;}
d_h0[i] =d ;
}
}

void dummy() {
dim3 block(BLOCK_DIM, 1, 1); dim3 grid(1, 1, 1);
testcuuuuKernel<4><<<grid, block>>>(NULL, 0);
//testcuuuuKernel<5><<<grid, block>>>(NULL, 0);
//testcuuuuKernel<6><<<grid, block>>>(NULL, 0);
}
int main(int argc, char* argv[]) {
cudaDeviceProp deviceProp;
cudaGetDeviceProperties(&deviceProp, 0);
cudaSetDevice(0);
#define DATA_SIZE ( 100)
unsigned int byte_size = DATA_SIZE * sizeof(float);

float* data = new float[DATA_SIZE];
for (int i = 0; i < DATA_SIZE;i++) {data[i] = i;}
float* d_data; cudaMalloc((void **)&d_data, byte_size );
cudaMemcpy(d_data, data, byte_size, cudaMemcpyHostToDevice);
LARGE_INTEGER nFreq, nBefore, nAfter; //TIMER初期化
DWORD dwTime;
memset(&nFreq, 0x00, sizeof nFreq);
memset(&nBefore, 0x00, sizeof nBefore);
memset(&nAfter, 0x00, sizeof nAfter);
dwTime = 0;
QueryPerformanceFrequency(&nFreq);
#define LOOPNUM 100
dim3 block(BLOCK_DIM, 1, 1);
dim3 grid(cuda_iDivUp(DATA_SIZE, block.x), 1, 1);
for (int k = 0; k < 10; k++) { //試行の試行
//start!
QueryPerformanceCounter(&nBefore);
for (int i=0; i<LOOPNUM;i++) { testcuuuuKernel<3><<<grid, block>>>(d_data, DATA_SIZE); }
cudaError err=cudaThreadSynchronize();
//stop!!
QueryPerformanceCounter(&nAfter);
cudaMemcpy(data,d_data,byte_size,cudaMemcpyDeviceToHost) ;
dwTime = (DWORD)((nAfter.QuadPart-nBefore.QuadPart) * 1000000 / nFreq.QuadPart / LOOPNUM);
printf("%d usec for %d times kernel launch\n", dwTime, LOOPNUM);
Sleep(400); //ちょっと待つ
}
cudaFree(d_data); delete [] data; getchar(); return 0; }

再現できるコードを書いてみました。
Dummyという関数でテンプレート展開されている__global__関数の数を調整してみてください。
ちなみに使っているチップはGTX260です

>>783
面倒だから動かしてもじっくり読んでもいないのだけれど、
カーネル関数はGPUに都度転送することになるから
一回の呼び出し粒度が小さいと転送コストが目立つことになるよ。
その位だと、恐らくは起動コストも無視できないからもっと処理させるべき。
つーか、カーネル呼び出し(<<<>>>)をループで包んだらそりゃ遅いって。

>>784
もちろんそれは分かるのですが、カーネル呼び出しの処理の内容は、<<<>>>の中の次元数に束縛される
傾向にあると思います。
データ構造などが異なる色々な処理を連携して実行する場合は、カーネル内の分岐では限界があるように思われます。
なので、１回当たりのカーネルのレイテンシを正確に把握しておきたいわけです。

＞データ構造などが異なる色々な処理を連携して実行する場合は、カーネル内の分岐では限界があるように思われます。
何にも判ってないと思われ。

っていうか、分岐ってさ、プレディケートで全部実行するんだよな？

if (cond) { //ここの条件は要素ごとに変わる
　　funcA();
} else {
　　funcB();
}

だったら、funcAとfuncBをインライン展開して全部プレディケートつき実行する感じだと思ってるんだが。

困ったことに、団子に同意。

条件分岐したら負けかなと思ってる by GPU

そいえばCUDAって１つのカーネルのサイズが制限されてない？
でかいやつがまったく動かなくて苦労したんだけど

どの位かは知らんが、そりゃぁ制限はあるだろうねぇ。

64Kのコンスタントメモリがあるじゃん。
コンスタントメモリは自分自身では中身の入れ替えは不可能。
あとはわかるよな？

cudaで自己書き換えプログラムってできますか？

GPUのカーネルコード自身で書き換えるって意味なら無理。
PTXのバイナリコードを動的生成とかなら何かやれば可能かもしれない。

その辺の資料を中の人に要求したら

「機密事項ですのでお答えできません」

大学の研究室にCUDAプログラミング用のコンピュータが導入された！と喜んでいたら、
HP ML115 + げふぉ8400GSカードだった…orz 学習用仕様で萎えた…

まぁ、学部生なら十分だろ。
どうせ大した論文も書かないくせに。

>>796
お前みたいなの、必ず湧くよなｗ
人を馬鹿にしたら、自分が偉くなるとでも思ってんの？

でも俺もそこまでは言わないにしても
学習用仕様で萎えたってセリフは贅沢だと思うよ。

値段を見て萎えたって意味だろうが性能的には十分な気がするんだけど

そもそもGPGPGUはなんなんだ？8800はVGA用途だろ

>>797
この程度で反応するなよｗ
社会でやっていけないぞ

>>795
それなんて俺
CUDA使ってみたくて買い換えを考えていたけど
PCI Express x16バス搭載で1万だったのでML115と
玄人志向の8400GSカード買ってきてCUDA環境を手に入れたぜ
メインで使ってるマシンより性能がよくてこのままメインになりそうな予感

事前に断っておくと俺は批判されるとチンコがたつよ

自宅でCUDA使ってる人のコンピュータスペックってどんなもん？

>>804
c2d e6320

>>804
PenD920 GeForce9600GT

>>804
Core i7 920 + GeForce GTX 280

>>804
Xeon + GeForce GTX280

>>799
いや、新技術半可通の俺が察するにGeF8000系は
CUDA対応ハードの底辺だから文句言ってるんだよ。
研究室で導入する(自腹じゃない)んだからもっと良いのよこせってことだろ？

ぶっちゃけ研究といっても今の時期じゃ全体的に未成熟だし
そのスペックでさえ限界動作なんぞせんだろと思って贅沢だなと。

>>800
そんなことはない。GPGPUは対応ハードならグラボ1枚でも2枚でも一応使えるんだよ。
実例は押さえてないんでなんともいえないが
競合さえおきなきゃRadeon(VGA)・GeF8800以上(GPGPU専用)ということも出来るはず。

用途と価格の兼ね合いもあるだろうけど、
とりあえず試したいとかコストパフォーマンス的に有利な部分もあるかもしれん。

4gamerの記事が参考になるよ。
http://www.4gamer.net/games/050/G005004/20080615001/

>>809
底辺どころかメインストリーム

>>811
ラインナップ上は確かGeF8世代からの対応だから一応底辺だろう。
>>795の萎える理由に対しての推測だからこの説明でいいんじゃね？

もっとも現状ではSP数の差による性能差がそこまで決定的じゃないから
2x0とかQuadroを期待してるんだったら贅沢なんじゃねって話。

GT2x0:性能は兎も角、(電源などの)要求仕様が一般的でない。
QuadroFX:同程度の性能のGeforceの数倍の価格と言う段階で、論外。
# なんて書くと、「アキバ的発想云々」って言われるんだけどねw
もしこれらが理由なら、見識不足と言わざるを得ない。
単に、>795はML115が気に入らないんだろw

まぁ、今時なら8400GSでも8600GTでも9600GTでも大して値段変わらんと思うが。

ML115の電源じゃパフォーマンスレンジのGPUは無理ね

ML115の電源はサーバー向けだけあってそれなりに良い電源だよ
電源投入後の全開爆音は凄いけど
ちなみに容量は370W

それじゃ精々8600GTSか9600GT止まり。8400GSは案外無難な選択じゃないのかな。

ちょっと誰か日本の公式フォーラムのcudaggさんの日本語をなんとかしてあげて

Ubuntuの話なら、一応意味は判るでしょ。
つーか、ここでフォローしても意味ないしw

>>813
GT2x0で一般的でないとか言ってたらCUDA自体一般的じゃないだろjk
ML115が不服と言うのでも十分贅沢……。俺が行ってた某理系大学なんか(ry

>>819
GT280は200Wの電源とそれに見合う通風環境が必要になるってこった。
この際、CUDAが一般的かどうかが問題じゃない。

>>820
CUDAスレでの話だからマシン云々というよりは
単に出たばっかのGT2x0系じゃなかったことに萎えてるのかと思っただけじゃん。

そもそも研究室とかだったら要求仕様が一般的でないなんてことは些細なことだ。
例えばOracle高価だから絶対使いませんなんて真似しないだろ。
そういうところは必要あれば用意するだろうよ。
GT2x0載せるとしたらそりゃML115じゃ無理だろうけど、それに載せろとは誰も言ってないし。
一般的云々言われなきゃ俺だってCUDAが一般的かどうかなぞ持ち出さんわ。

グラボの方に萎えてると思ったもうひとつの理由が、
CUDAプログラミング用ということだから
マシン自体のスペックはそれほど高くなくても良いはずなんだ。
電源さえ足りてればな。深読みしすぎた俺が悪かったよ……。

そこいらでレスをつけてるやつらの中に、8000系では一部使えない命令がある件を指摘するやつが居ないのは何でだぜ？
お前ら実は妄想だらけでなんもしてないんじゃないの？

8400GSならCompute Capability 1.1だから
GT200で追加された命令でもなければ使えるわけだが。

>>823
1.0世代は8800GTX,GTS,ULTRAに積まれたG80くらいで、8600GTや8400GSに積まれたG84やG86は
1.1世代だと言う知識もなしにここ見てたの? 他人を妄想だらけなんて指摘してられる状況じゃないじゃんw

CUDAを単純に使いたいやつは、

HP ML115 (http://nttxstore.jp/_II_HP12591344)
と
G98の8400GS買えばおｋだよ。

グラボ早い順に並べるとどんな感じ？
やっぱりSP数×クロックなの？

仕様を限定せずに一般的に言うなら、そりゃそうだ。他にどんな要素が来ると?
特定の応用でメモリ転送が重視されるのならメモリ帯域やバス仕様にも依存するし、
メモリ量が多い方が高速なアルゴリズムを採用できる応用ならメモリ量にも依存するわけだけど。

あー、1.0世代だとストリームが使えないと言う大きな欠点はあるね。

チップ外は考えてなかったφ（．．）ﾌﾑﾌﾑ…

後からSPが減ってクロックが上がった製品が出るじゃない？
だいたい値段は一緒だけど、劇的に性能が違うのかなと…

同じ名前のSP減ったのは、恐らく歩留まりの関係で在庫処分しているだけだと思う。
使い方にも拠るけど、クロックが利く応用なら結構変わるかもね。
# SP数が利く応用なら逆に遅くなりかねない。

やっぱみんなPDEとか立てたり解けたりできるの？

そうだね
SDEとかね

>>795
あぅっ、それはウチかも知れない。

まぁどーせ学生の作るプログラムなんて、レポート見てると
自分で作ってるのは15%ぐらいで、残りの人は友人のコピペ
改変か、ググって見付けたページのコピペ改変だからな。
予算の都合ってやつで、1人1台用意するならML115+GF8400
に成ってしまう。許せ。
おっ、これは?!と思うプログラムが出てきたら、GTX280で
走らせてやるから頑がってくれ。

>>816
PCI-Ex補助電源コネクタ無いので、GF8600GT
か、GF9500GTまでが無難なところ。

texture<float, 1, cudaReadModeElementType> tex;
のテンプレートで
> error C2018: 文字 '0x40' は認識できません。
ってエラーが大量に出るんだが、みんな出ないですか。
オレッスカ
textureのテンプレート使わなければでないのですが。

CUDA 2.1 + VC++ Express 2005

>>835
PDFからコピペしてない？

>>836
一回行コピペしてエラー出たので消して書いたのですが、
ファイルに変なコードが残るとかありますかね。
でもその行消すとエラーなしで、手書きで書いたら、
↓のようになります。

1>.cu(54) : error C2018: 文字 '0x40' は認識できません。
...
1>.cu(54) : error C2018: 文字 '0x40' は認識できません。
1>.cu(54) : error C2065: 'COMPILER' : 定義されていない識別子です。
1>.cu(54) : error C2146: 構文エラー : ';' が、識別子 'ERROR' の前に必要です。
1>.cu(54) : error C2065: 'ERROR' : 定義されていない識別子です。
1>.cu(54) : error C2143: 構文エラー : ';' が '<template-id>' の前にありません。
1>.cu(54) : error C2146: 構文エラー : ';' が、識別子 'tex' の前に必要です。
1>.cu(54) : error C2275: 'texture<T,dim,__formal>' : この型は演算子として使用できません
1> with
1> [
1> T=float,
1> dim=1,
1> __formal=cudaReadModeElementType
1> ]
1>nv_som_gpu.cu(54) : error C2065: 'tex' : 定義されていない識別子です。

texture<float1, 1, cudaReadModeElementType> *tex = new texture<float1, 1, cudaReadModeElementType>;
ならOKなので、スタックにtextureを置くとダメな条件があるのでしょうか。

kernelに引数で渡せないので
スタックにおいたらダメみたいでした。
解決。

まだ、初めてみようかと考えている初心者以前のものです。
現状の自作数値解析プログラムはMPIで各ノードの各コアにプロセスを振っております（ノード内もMPIで並列）。
これの自然な拡張としては、MPIの各プロセスがGPGPUを使うという形になるかと思います。
しかし、マルチCPUのノードの場合、ひとつのGPUを共有することになります。
一つのプロセスがGPUのシェーダのうち1/4だけ使って、4プロセスから同時にGPUを使うなんてことは可能なのでしょうか？

>>839
テクスチャはコンパイルするとただの数値みたいになるので、
グローバル変数としてそのまま使う以外のことはほとんど不可能、らしいよ。

関数にパラメータとしてを渡すのも、ポインタを得るのも無理。

>>840
CUDAを使う場合、4プロセスから同時に使うことも可能。
但し、GPUをどう割り振るかはCUDAドライバの御心次第。
厳密に制御するには、GPU1枚ごとに担当スレッドを設けることになる。
# その場合、プロセス間通信でJOB型にするか1プロセスだけでGPUを占有するかは設計次第。

それはいいけど、プロセス全部分けるとマルチスレッドに較べて効率落ちないかい?

>>842
ありがとうございます。
やはりリソースの競合が起きるようですね。

マルチスレッドは単純に覚えることが増えるのでやってませんw
スパコンのマニュアルでMPIを覚えたので、WSクラスタでもそのままの手法を持ってきてます。
一応、ノード内では共有メモリを使って通信するようなオプションでmpichをインスコしてるので多少はマシでしょうという感覚です。

NVIDIA、PhysX/CUDAを活用する「パワーパック」の第2弾を提供開始
http://journal.mycom.co.jp/news/2008/12/03/001/index.html

0x40でググったら
>文字 '0x40' は認識できません。 原因： ソースコード中に全角空白 ' ' が使われています。

まさか、な

こんな感じで網羅的に点を回転させるプログラムを書いているのですが、
Thread数が320を越えるあたりでrotate関数に全ての引数が渡らなくなってしまいます。
その時でも計算量が少ないからか、roll_axisのz成分などは、引数として機能しています。
roll_axisのx,y成分も適当な値に変えると引数として働くようになります。
rotate(&coord, &pitch_axis,ANGLE);をコメントアウトすれば,roll_axisのx,yはThread数が多くても引数として渡ります。
計算量が多くなると(特に三角関数？)起きる気がするのですが、何が回避する方法はあるのでしょうか？

--ptxas-options=-vはUsed 42 registers, 56+28 bytes lmem, 2080+32 bytes smem, 8024 bytes cmem[0], 88 bytes cmem[1]て出ます。
-deviceemuでは正常に動作します。

>>846のソースコードです。
####kernel(Thread１つが回転させる点１つに対応)####
__device__ runDevice〜〜の一部
float4 yaw_axis = make_float4(0, 0, -1, 0);
for(iyaw = 0; iyaw < limY; iyaw++){
float4 pitch_axis = make_float4(-sin(radian*iyaw), cos(radian*iyaw), 0, 0);
for(ipitch = 0; ipitch < limP; ipitch++){
float4 roll_axis = make_float4(cos(radian * iyaw) * cos(radian * ipitch), sin(radian * iyaw) * cos(radian * ipitch), sin(radian * ipitch), 0);
for(iroll = 0; iroll < limR; iroll++){
if(tid == iroll + ipitch * limR + iyaw * limR *limP)
g_mem[tid] = coord.x; //回転できているか確かめる。

rotate(&coord, &roll_axis, ANGLE); //回転させる点の座標と、回転軸と角度を与える。
}
rotate(&coord, &pitch_axis, ANGLE);
}
rotate(&coord, &yaw_axis, ANGLE);
}

__device__ void rotate(float4 *coord, float4 *axis, int angle){
coord->x = axis->x;
coord->y = axis->y; //とりあえず現段階では引数が渡るか確かめてるだけ
coord->z = axis->z;
}

誰か、すげぇ簡単単純だけど、真理を付いてる様な神サンプル晒してよ

>>848
どんなの? SDKのサンプルやbioのblog辺りじゃお気に召さない?
具体的なテーマがあって、実装が難しくなさそうなら作ってもいいけどね。

ATI Streamから来ました。
スイマセン、場違いな場所に来てしまった。
AもまだなのにCなんて出来ません

共産主義者の書いたマンガによると、Aより簡単らしいぞ。

最近の若い連中はしょっぱながCだからな

cudaってさ
C-daにしておけばギャグっぽくてよかったんじゃね？

それならCarracudaの方がｗ

teslaの話題が少なくて絶望した！

えー、QuadroFX5600のアナログ回路をとっ外しただけの代物の、何を語れと言うのさ。
聞いてくれたら答えるけど。
あ、4桁シリーズは白根。どうせ事情は一緒だと思うけど。

>>855-856
このスレの住人は　GPGPUの用途にQuadroFX使ってるの・・・？
ていうかそもそもみんな何使っているの？

ML115最強伝説

Tesla D870 が至高

>>857
このスレともう一つのCUDAスレを見ると、QuadroFXのメリットが書かれていると思うが。
# 普通、使うわけないだろって。
8800GTX、QuadroFX5600、Tesla C870の価格を見る限り、Teslaは未だましだと思う。

>>859
D870ってNVIDIA謹製ミニタワーにC870を二枚入れた代物だっけ?
あれだったら同じ筐体のQuadroPlexの方が潰しが利くと思うのだけど。

GeForce GTX 280 極上

９６００GT買ったらCUDAでエラー出まくりで問い合わせたら
CUDAのようなメーカー付属のドライバで対応してない機能は保障外とのこと
CUDAをやるために買ったのにCUDAが動くことを保障してないなんて
はっきり言って詐欺だよ

DELLにVisual Basicで組んだプログラムでエラー出るんだけど？
って聞いてるようなもんだな。

>>862
nvidiaのドライバだとどうなんだ？
マジに教えてくれ

>>863
CUDA SDKのテストプログラムでエラーが出まくる
実行する度に計算結果が違ってくる

>>862
どういう意味？NVIDIAのドライバ以外にCUDAが動くの？
付属のCDのドライバ以外はNVIDIAの公式ドライバでも保障しないと言われた

>>862
一体どこの糞メーカのPCなんだ?

>>865
>付属のCDのドライバ以外はNVIDIAの公式ドライバでも保障しないと言われた
それが普通ですよ。販売メーカとしては、自分のところでテストした
ドライバ以外は保証してないよ。

だから公式ドライバ入れる時とかもよく言われるでしょ。
トラブっても自己責任で、ってね。

まあ当たり前なのは分かるんだけど、そこんとこ盲点だよ。
付属のバージョンだと根本的にCUDA機能が入ってないからね。
必然的にアップデートしないといけないのに保障外でしょ。
CUDAは素晴らしいとか大々的に宣伝してるけど動くかどうかは運次第という対応だからね。
これが世間に知れたらCUDAなんて誰も相手にしなくなると思うけど。

tesla売るためだろ

そうだろうね。TESLAは唯一保障された製品だからCUDAやりたい奴はTESLA買えってことなんだろう。
一般人には何の恩恵もない誇大広告だったというわけですか。

メーカー純正フロッピーディスク以外使わない奴なんていなかっただろ。
みんなそうやって保証外の物を使ってきているんだ。今更何を言っているんだか。

で、そんなふざけた事言うのはどこなんだ？

どこかの誰かが同様にCUDAが使えないとなると自分がCUDAを使う意味がない。
だからCUDAは必要なくなるので別に何も問題ない。
誰もが使えないプログラムなんて作っても時間の無駄でしょ。
保障がないってのはそういう事よ。個人的な恨みとかじゃなく。

ノートPCなんてNVIDIAが公式ドライバを提供しないんだぜ。
メーカーも提供しないから、アマチュアの作ったものを探してくるしかないっていう。

あー、WindowsでCUDA試してないな、うちのMac

保証が無いからどうだとかいうのは自作板でやってくれ。

>>857
犬板にも書いたけど、JetWay HA05-GTのオンボッボで動きますた。
USBメモリ起動で、薄いCUDA nodeの出来上がりです。
オンボッボももっと沢山SPU載る日はいつかなぁ。

【GPGPU】CUDA/ATI　STREAM　速度・画質検証スレ
http://pc11.2ch.net/test/read.cgi/jisaku/1228764782/
　

これからCUDAを使おうと思っているのですが、visualC++でうまくビルドできません。どなたか解決方法を教えてください。
環境はXPx64Edition,QuadroFX570でx64版のバージョン2.0のドライバとToolkitとSDKを入れました。
VisualC++は2005Expressです。

すみません、879ですが、ビルドできないのはSDKの中にあるprojectsのsimpleTemplateです。
webで調べてみると皆さん問題なく実行できるようなのですが。。。

どんなエラーが出るとかの情報はないの？

失礼いたしました。下のようなエラーがでます。
nvcc fatal : Visual Studio configuration file '(null)' could not be found for installation at 'C:/Program Files (x86)/Microsoft Visual Studio 8/VC/bin/../..'
>>444さんと同じような問題と思われます。VisualC++のコンパイラとの相性が悪いのでしょうか。詳しい方、よろしくお願いいたします。

こんな連中ばっかりだし、クマには生きづらい世の中になったぜ

C++もろくに使えない人がいきなりCUDAに手を出すのもどうか

なんか64bit環境でもVCのclは32bit版使わないとだめみたいなことCUDAフォーラム（英語のほう）に書いてあったよーな

nvccから呼ばれるほうのclのことね＞32bit版

>>882
パスが通ってないんじゃね？

PATHは通っていると思います。消してみると、
　　nvcc fatal : Cannot find compiler 'cl.exe' in PATH
となりました。platformSDKのclにすると↓のようになりました。
　　nvcc fatal : nvcc cannot find a supported cl version. Only MSVC 7.1 and MSVC 8.0 are supported
結局、CUDAtoolkitを32bitにしましたところOKとなりました。885さん、886さん、887さんありがとうございました。

最適化してregisterの数を減らしたいんですが、
なんか心がけることとかコツってありますか？

ptx出力を読め。以上。

なんでわからないのでご教授を・・・・
VS 2005にカスタムウィザードを組み込んでパスも通して空の.cuファイルに書き込んだのに
1>LINK : fatal error LNK1181: 入力ファイル '.\Debug\sample.obj' を開けません。
と出るのですが・・・・何か環境設定間違っているのでしょうか？

http://sourceforge.net/projects/cudavswizard
どうぞ

>>892
それ使ってるんですが・・・readmeを嫁ということですね

ptx出力がどうなるように最適化していけばいいんですか？

>>894
>889なら、registerの数を減らしたいんだろ。ptxの出力を読んで少なくなっていればOKだw

>>862
NVIDIA、モバイルGPUドライバを配布開始
〜ノートでPhysXなどCUDAアプリケーションが利用可能に

12月18日(現地時間)配布開始

　米NVIDIAは18日(現地時間)、モバイルGPU用の汎用ドライバを同社ホームページ上で配布開始した。
対応GPUはGeForce 8M/9M、およびQuadro NVS 100M/300Mシリーズ。

　従来、NVIDIA製モバイルGPUのドライバは、ノートPC本体のメーカーが提供していたが、
今回よりNVIDIAが汎用ドライバを提供することになった。対応OSはWindows XP/Vista。

　バージョンはデスクトップ版の180代に近い179.28 BETAで、CUDAに対応。これにより、PhysXや、
CUDA対応の動画編集ソフト、Photoshop CS4などでアクセラレーションが効くようになる。
ただし、PhysXについてはビデオメモリが256MB以上必要となる。

　なお、ソニーVAIO、レノボThinkPad、デルVostro/Latitudeシリーズ、およびHybrid SLI構成の
製品については対象外となっている。

　現バージョンはベータ版で、WHQL準拠の正式版は2009年初頭の公開を予定している。

http://pc.watch.impress.co.jp/docs/2008/1219/nvidia2.htm

なんでvaioは対象外ですか？いじめですか？？

vaioとか一部のノートはリコールされてるから

それ以前にノートだとメーカーが色々やってるから
一般的にドライバ更新さえノートのメーカー対応待ちだろ。

シェーダプロセッサってエラーがあってもゲームでは無視されて
認識出来ないレベルで画面が崩れるだけだから
一見正常に動いてると思っててもＣＵＤＡは動かない人が多いかも
ドライバが対応してるとか以前にハードが用件を満たしてない

どこでエラーになっているのかデバッグしろよ

ちょいちょいCUDAをいじり始めてるんですが、OpenCLの規格が決まったことで、
CUDAの仕様も大幅に変わってしまうんですかね。

MacBookProなんかだと自前でドライバ突っ込めばCUDA使えたんだけど，
VaioやThinkPadだとどうですか？
使っている人がいれば情報求む．

5秒制限の条件とか回避の仕方とかがよくわからないのですが
教えてもらえないでしょうか？

ThinkPad T61はOpenGL関係以外は使えた。

OpenCLはAPIのみだからCUDAとは別にライブラリが用意されるだけだと思うよ
こっちは専用コンパイラも必要ない

>>904
Gridを何回も起動するしかないのでは。だから、途中結果をメモリにのこしとかないといけない。
面倒だよね。

>>904
Xを使わない、グラフィック機能を使わなければいけるはず。

…何のためのグラフィックカードやねんって感じだが

>>907>>908
やっぱりそれぐらいしかないですか。
現状でグラフィック機能止めて使うのは無理なので
Gridの分割を考えてみます。

返答ありがとうございます。

Vistaだとレジストリ弄ればもうちょっと長くできるよーみたいなこと言ってたような、言ってなかったような

使うグリッド数を制限すればなんとかならなくもないんだが、効率は落ちるんだよね。
つーか、効率重視ならもう一枚挿すしかない希ガス。

こんなの教えて貰った
http://compview.titech.ac.jp/Members/endot/adv-app-hpc/2008schedule

どのみちCUDA実行中は画面の描画が止まるから割り込み入れる工夫は必要でしょ

描画要求の方がGridよりも優先されているようにしか見えないんだけど。
止まってくれたらどんだけ効率が改善することかw

>>912
面白そうだけど、東工大は遠いなぁ。ワープアSEに新幹線は高い。
わがまま言うと、神戸でやって欲しい。

大岡山の東工大なら歩いて行けるぜ！

>>915
六甲台か？
あの急勾配の坂は登りたくないな。

>>912
これは紹介レベルだから、たいした内容ではないよ

CUDAをクアッドコアで動くように出来ればいいのに

>>919
どういう意味?

CUDAをQuadCoreCPUで動かすことなら何の問題もなくできるのだけど。

っOpenCL

VIPPERの諸君　立ち上がれ！
以前韓国を打ち破ったのは誰か知らない、小学生や、他サイトの管理人がVIPを馬鹿にしている！

全員集結せよ。そして全員で打ち倒すのだ！

もし時間がない人は、この文章をそのまま他のスレに貼ってほしい！

対策本部
http://takeshima.2ch.net/test/read.cgi/news4vip/1231149782/

今こそ　Vipperの意地を見せつけるのだ！

nc = 100*100
bs = 50
dim3 dimBlock(bs,bs)
dim3 dimGrid(sqrt(nc)/dimBlock.x,sqrt(nc)/dimBlock.y)
kernel<<<dimGrid, dimBlock>>>(idata, odata, sqrt(nc)

__global__ void kernel(float* idata, float*, odata, int nc)
{
index=blockIdx.x * blockDim.x + threadIdx.x +
+(blockIdx.y * blockDim.y + threadIdx.y) * nc
}
この時のイメージは、Gird:2x2、Block:50x50でよいのでしょうか?
それとこのままグローバルメモリで計算するのはできるのですが、一旦
シェアードメモリに退避して計算してグローバルメモリに戻す方法が
サンプルを見てもうまくいきません。どういう感じになるのでしょうか?

いつも疑問に思いつつ使っているのですがグローバルメモリ使うときも
共有メモリ使うときも
float* data;
CUDA_SAFE_CALL(cudaMalloc((void**)&data, sizeof(float) * 100)
CUDA_SAFE_CALL(cudaMemcpy(data1, h_data, sizeof(float) * 100, cudaMemcpyHostToDevice))
になるのでしょうか?これはグローバルメモリにとっているのですよね。
共有メモリの話は、kernelのほうで使用するかどうかですよね?


コンスタントメモリ使うときだけどうも
CUDA_SAFE_CALL(cudaMemcpyToSymbol(data, h_data, sizeof(float) * 100))
としてホストからコピーするんですかね?

>>924
cudaMemcpyToSymbolは定数メモリの他、グローバルメモリでも使えるよ。
要は、device空間で定義されたメモリと言うニュアンスでしかない。
逆に言えば、cudaMemcpyToSymbolを使わないとhost側のメモリ空間に割り付けられたアドレスを利用しようとしてしまう。
# この辺が判り難いのよね。

共有メモリはグリッド実行中しか使えないからhost側からは為す術がない。

>>923
どのサンプル見てどうやったらどう巧くいかないのか具体的に。
少なくとも、>923のような処理なら共有メモリの使いようがないんだけど。

warp単位32スレッドで同じ命令を行って、
実際に一度に動くのは8スレッド

こういう考え方でいいのですか？

>>925
結局現状では、cudaMemcpyXXXの違いは理解できそうにないです。
まあそれはよいとして(全然良くないですが)
>>923
のidataは、実は、一次元配列なんです。なのにカーネルには2次元で利用しよう
としているんです。
float host_idata[nc]
float* idata
cudaMalloc((void**) &idata, sizeof(float) * nc)
cudaMemcpy(idata, host_idata, sizeof(float) * nc, cudaMemcpyHostToDevice)
こんな感じのグローバル変数がidata。
としたとき、ncとbsの設定値によってはうまく処理できないことがあるのです。
これが謎です。わかる方いませんか?もちろんncはsqrtで整数になる値、設定する
値は、割り切れる値です。nc=1000*1000,bs=500で処理結果がおかしくなっています。
kernelの読んだあとエラーでinvalid configuration argument.になってます。

それと、共有メモリのほうですがこれは分かりました。
kernel<<<dimGrid, dimBlock, sizeof(float) * nc>>>(idata, odata, sqrt(nc))
第3引数の共有メモリのサイズを定義しないとだめだった。ちゃんと説明書を
読んでいなかったというレベルでした。これは解決です。

>>927
ブロックは512x512x64と出ているけどいちお3次元は分からないので512x512
としてxで512使い切るとどうもyではもう1しかとれないようです。
ncは、そのままだとエラーが出ないレベルは、bs=20つまり20x20=400ブロック
使用。21,22は割り切れないので設定できない。みたいです。
とすると一次元配列を使用するのにわざわざ、x,yの次元のブロックを使いきる
テクニックは必要なさそうですね。と推察できると思います。結果からですが
ただ、これがバグであるなら次のCUDAのバージョンアップで使えるようになる
かもですね。

>>928
次元とは別に1ブロックあたり512スレッドが上限じゃなかったっけ？
どっかにきちんと書いてあったはずだからバグではないよ。

私は1次元でやってしまうことが多いなぁ。
で、スレッド数についてはdeviceQuery参照で。

スレッドの最大数はAPIで取得できる。全部512だと思う。
ブロックの最大は65535だから65535*512以上は1次元のグリッドでは扱えない。
あと最大は512だけど、192か256くらいのほうがパフォーマンスが出る。

CUDA 2.1 Release
ttp://forums.nvidia.com/index.php?showtopic=85832

情報サンクス。早速、Linux版ゲットした。

9600GTでVideoDecodeのサンプル動かそうとしたら正常に動かなくて、
結局SDK2.0+178.24まで戻さないとダメだった。

2.1にもソースは付いてるけど、2.0のままみたいだね。

>>931
1 warp = 32 threadで、GeForce 8/9が24warp/block、GT200で32warp/blockが最大だから
768thread/blockと1024thread/blockが最大なんだけどね。本来は。
CUDAのドライバ側が512でリミッタかけてるんだ。罠としか言いようが無い。
逆に言うとCUDAを経由しなきゃ目いっぱい使えるかもね。

ただ、スレッドインターリーブすると１スレッドあたりで使えるレジスタ本数が減っちゃうんだよね。
メモリレイテンシを隠蔽するならスレッドを目いっぱい使ったほうがいいし
逆に一時変数を何度も再利用する場合は、thread/blockを減らして１スレッドあたりの仕事量を増やしたほうがいい。

CUDAってシーユーディーエーって読んでいいの？
それともクダとか、なんか読み方あったりする？

>>936
クーダ

ネイティブの人は「クーダ」って発音してたよ。
Cellは「くた」だから注意な。

ありがとう！明日から会社でクゥゥゥウーダ！！って叫びまくるわ

きゅーだ　って発音してる俺は異端ですかね？

>940
そんな人いたんだ。

英語的に　クー　よりも　キュー　だと思わね？

>>941
異端だな

どっちでもいいんじゃね

ちなみにnudeはニュードとヌードの両方の発音が有

キュもクも似てるし、まあ仮に間違って言ったとしても気づかれないんじゃない？

ホムペにはクーダと嫁
と書かれてたはず

"The cuda and my wife"

たしか英語版のfaqには
kuh-da
と書いてあったはず

2.1はJITサポートがミソ？

2.1をつついた人に質問
2.1ではカーネルに渡す引数でポインタ配列は使えますか?

共有メモリからグローバルメモリに連続領域をコピーするにはどうすればいいでしょうか
カーネルから呼び出せるmemcpyみたいなのがあれば教えてください

cudaMemcpyが長いようで、the launch timed out and was terminated.が
でてしまいす。これなんとかするほうほうありましたっけ?

>>952
転送量を減らす。

>>951
cudaMemcpy

>>951
スレッド内で連続的にメモリにアクセスするようなコードは望ましくない。
coalscedにアクセスするとなると、自分でコードを書くしかない。
共有メモリから転送を始める前に、同期を取ることも忘れずに。

>>952
参考までに、どんなハードウェアの組み合わせでどれくらい転送しようとしている?
8800GT辺りだと転送できるだけ全部転送してもタイムアウトしないのだけれど。

>>954
カーネル実行後の結果データコピーでエラーになっているのですが
sizeof(float)*1000バイトを100回連続でコピーしています。
100個とも別のアドレスに結果をコピーしたいためそうしてます。やはり
いっぺんにメインメモリにいったんコピーしてからやったほうがいいので
しょうか?

>>955
sizeof(float)*10000バイトを100回連続コピーでした。

>>954
＞ >>952
＞ どんなハードウェアの組み合わせ

>>956
同一データ部分を100箇所にコピー？
独立の100箇所をコピー？