【GPGPU】くだすれCUDAスレ【NVIDIA】

このスレッドは、他のスレッドでは書き込めない超低レベル、
もしくは質問者自身何が何だが分からない質問を勇気を持って書き込むスレッドです。
CUDA使いが優しくコメントを返しますが、
お礼はCUDAの布教と初心者の救済をお願いします。

CUDA・HomePage
http://www.nvidia.com/cuda

関連スレ
【GPGPU】NVIDIA CUDA質問スレッド
http://pc11.2ch.net/test/read.cgi/tech/1190008468/
GPUで汎用コンピューティングを行うスレ
http://pc11.2ch.net/test/read.cgi/tech/1167989627/
GPGPU#2
http://pc11.2ch.net/test/read.cgi/tech/1188374938/

関連サイト

CUDAに触れてみる
http://chihara.naist.jp/people/STAFF/imura/computer/OpenGL/cuda1/disp_content

cudaさわってみた
http://gpgpu.jp/article/61432191.html

CUDA のインストール
http://blog.goo.ne.jp/sdpaninf/e/9533f75438b670a174af345f4a33bd51

NVIDIAの「GeForce 8800 GT(G92)」と次に控える64-bit GPUアーキテクチャ
http://pc.watch.impress.co.jp/docs/2007/1031/kaigai398.htm

CUDAを使う
http://tech.ckme.co.jp/cuda.shtml

NVIDIA CUDAを弄ってみた その2
http://dvd-r.sblo.jp/article/10422960.html

CUDAベンチ
http://wataco.air-nifty.com/syacho/2008/02/cuda_2044.html

KNOPPIX for CUDA Tutorials
http://www.yasuoka.mech.keio.ac.jp/cuda/

>>2
もうちょっとましなサイトはないの？
ブログってチラシの裏と大して変わらないのが多いのだが

>>3
すみません、ぐぐって適当に載せたのですが、良いサイトがあれば教えてください。

みんなGPGPUでどんなプログラム書いてる？

>>1
スレを分ける訳は? あっちのCUDAスレだけでいい希ガス。

teslaってどうなの？使ってる人いる？

TeslaC870はQuadroFX5600からアナログ回路周辺を取り除いた代物と思っていい。
性能的には、VRAMは広くて使い出があるけどそれ以外は8800GTXと大差ない。
副作用として、8800GTXなどだとオンボードGPUを切り離してしまうタイプのPCでは、
GraphicChipと認識されないC870は便利と言えば便利。

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

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

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

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

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

AMDも来たね
http://ati.amd.com/technology/streamcomputing/sdkdwnld.html

nVIDIAといいATiといい、Vistaの開発者をバカにしてるのか

1次元テクスチャって通常のfloatの1次元配列のように扱えるの？

sharedメモリはカードによって使える容量は違うのですか？

>>13
1block辺りの容量が決まっているから、結果的にGPUによって違う場合があることになりますな。

CUDAでトリップ検索したら早いの
１２８あれば１２８倍早そうな気がするんだけど
すぐもれちゃいそうで怖いんだけど

↑早いの？です

>>15
(比較的)単機能の(必ずしも)早くない演算器が(現在の最上位機種でも)256個あるだけなので、
単純に256倍速くなるわけではありません。
実際、4コアXeonよりも速いプログラムを作るのに苦労しているくらいですから。

AIも256倍賢くなるのかな？

コンパイルすると
Parsing error near '[': syntax error
というエラーが出ます。何がいけないのでしょうか？

[の近くの構文がいけないのでしょう

>>17
なるほど夢見すぎでした�d

CUDAのプログラムを配布する場合、何を含めればよいのですか？

>>22
NVIDIAの対応ドライバとCUDAToolKitをインストールしてもらう前提で、実行モジュールだけ。
ToolKitのDLLが同梱できるかは不明。

nvccの-Oオプションて数字はいくつまであるの？gccと同じ-O3まで？

gccにそのまま渡すだけだと思うから、きっと3。

CUDAコードは最適化されないのですか？

されますよ。もう、徹底的に。
# まぁ、それほどドラスティックではないけれど。

cudaの実行速度を最適化したい場合、Oオプション以外にどのようなオプションがありますか？

頑張って自分で最適化してください。

-O3 -use_fast_math くらいかな。私が使っているのは。
GPU側のロジックは、実は-Oの有無では変わらないかも知れないと思う。

nvccのオプションで-O3を指定した場合、-Xcompilerで改めて指定する必要はないのですか？

ブロック数はどのように決めればよいのですか？
できるだけ多くするのがよいのでしょうか？

>>32
多くすれば良いと言う物でもないみたい
分割数を多くすることによるオーバーヘッドの方が大きくなるのか
良く分からないんだけど

いくつぐらいが適当ですか？

もう一つのCUDAスレにヒントがあるので参考にどうぞ。

単精度を2つ使って倍精度の計算をする様なことはできませんか？

実は現行CUDAにも一部倍精度演算があります。
が、まともに使うにはCUDAv2を待ちましょう。

今日やっと、CUDA_Profiler_1.1.txtの存在に気付いてProfilerをまともに使えるようになった。
環境変数で設定ファイル名を指定して、設定ファイルで出力項目を指定するなんてややこし過ぎる……
しかも、timestamp以外に4つまでの制限なんて、なんて半端な。
つーことで、こんな感じ。
-- cuda_profile.setting
export CUDA_PROFILE=1
export CUDA_PROFILE_LOG=cuda_profile.csv
export CUDA_PROFILE_CSV=1
export CUDA_PROFILE_CONFIG=cuda_profile.config
-- cuda_profile.config
timestamp
gld_incoherent
gld_coherent
# gst_incoherent
# gst_coherent

# local_load
# local_store
# branch
# divergent_branch

instructinos
warp_serialize
# cta_launched
--
# もうCUDAスレのどっちがどっちだか判らなくなっているのは内緒。

CUDAと心中しても大丈夫でしょうか？

心中する覚悟をするのは勝手ですが、実際に心中されたら大迷惑です。

CUDA 2.0(Beta)

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

OpenGL系がダメダメだね。Vista。
ただ、nbody.exeはなぜか動いている。

> OpenGL系がダメダメだね。
付属のバイナリはダメだが。

ソースからビルドし直せば、動くぞ。

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

188 名前：デフォルトの名無しさん[sage] 投稿日：2008/04/21(月) 19:17:35
Linuxはここだ
http://forums.nvidia.com/index.php?showtopic=65067

どうでもいいけど明後日の仕事が増えた(謎

質問スレッドっていきてるの？

たった今回答して来ましたが何か。

Visual C++のバージョンはチェックしているみたいで、合っていないとコンパイルできませんが、
gccはバージョンが合ってないとコンパイルできないのでしょうか？
gccの最新版を入れても大丈夫なのでしょうか？

>>47
http://pc11.2ch.net/test/read.cgi/tech/1190008468/386

ローカル変数は、どのメモリに格納されるのでしょうか？
sharedメモリでしょうか？

レジスター

2つのGPUを同時に使いたい場合、
for (int i=0;i<2;i++)
{
    CUDA_SAFE_CALL(cudaSetDevice(i));
    somekernel<<< 64, 512>>>( );
}
のような形で書けばよいのでしょうか？
それともこれだと、一つ目のカーネルが終わらないと2つめのカーネルの実行が開始されないのでしょうか？

上記のように書いているのですが、実行速度が2倍になりません。

>>51
スレッドを分けるか、CUDA1.1で追加されたストリーム機能を使う必要がありそうです。
後者の場合、こんな流れでしょうかね。
--
for (int ic = 0; ic < 2; ++ic) {
cudaStreamCreate(& stream[ic]);
cudaSetDeviece(ic);
kernel<<<64, 512, 0, stream[ic]>>>();
}
for (int ic = 0; ic < 2; ++ic) {
cudaStreamSynchronize(stream[ic]);
postProcess(ic);
}
cudaThreadSynchronize();

>>52
ありがとうございます。
上記のようにやったのですが、
for (int ic = 0; ic < 2; ++ic) {
の部分を
for (int ic = 0; ic < 1; ++ic) {
に書き換えると、実行時間がちょうど半分になるので
やはり並列で実行されていないようです。
上記コード以外になにか付け加える必要があるのでしょうか。

>>53
CUDA_SAFE_CALLマクロは毎回cudaThreadSynchronize()を呼ぶから要注意。

>>52さんの書かれたとおり、

for (int ic = 0; ic < 2; ++ic) {
cudaStreamCreate(& stream[ic]);
cudaSetDevice(ic);
kernel<<<64, 512, 0, stream[ic]>>>();
}
for (int ic = 0; ic < 2; ++ic) {
cudaStreamSynchronize(stream[ic]);
}
cudaThreadSynchronize();

のようにしているのですが、逐次実行になってしまいます。

どこかにサンプルコードとかはないでしょうか？

cudaSetDevice()するとそのCPUスレッドで使うGPUデバイスが決定されるので、
それ以降の変更はできない模様。cudaSetDevice()しないでAPIを使うと、
無条件にDevice0を使うことにされてしまう模様。

CPUスレッドを分けるしかなさそうね。

NVIDIAのForumに行けばサンプルも転がっていそうだけれど……

あー、書くの忘れた。CudaProgramingGuide(1.1)の4.5.2.2参照で。

ストリームだけではなんともならないということでしょうか？

ならないっぽいですよ。ドキュメントを読む限り。
2.0betaのドキュメントも書き換わってないからそのままの悪寒。

geforce9以降もCUDA対応してるんでしょうか？

>>61
勿論。

>>60
そうですか、そうなると結構面倒ですね。
ところで、スレッド分けで2つのGPUを使う場合、pthreadなどでも良いのでしょうか？
それとも、サンプルにあるような方法のみが認められているのでしょうか？

>>63
ちっとはcutilの中身も見ようよ。cutStartThread()は、Linux版ではpthread_create()を呼んでいる。

GeForce GTX 280で倍精度を使うと、処理速度はどの程度落ちますか？

単精度の半分ぐらいでしょ。JK

いいえ、12分の1です。

GTX280　78GFLOPS
GTX260　62GFLOPS

うーむ。GPU2枚使うのって実はコツとかいる？
ホスト側でOpenMP使って2スレッド走らせて、それぞれのスレッドに別のGPUを割り当てて
るんだが、答えとしては正しいものが返ってくるんだが速くならない。
むしろ、ホスト1スレッドで1つのGPUで計算させたほうがいくらか速いくらい。
きちんとGPU2枚使ってそうだというのは確認できたんだが速くならない理由がわからない。
なんだかGPU0の処理が終わってからGPU1の処理が始まるとかやってそうな予感がしている。
GPUの切り替えの分だけちょっとだけ遅くなるというオチじゃないかと。。。
サンプルのmultigpuとか読むとホスト側のスレッド生成法は違うけど同じようなことやって
るんだよなー。
2枚使って速くなったって人いたらどういう風にやったか教えてくれませんか？
かれこれ2週間くらい悩み中だ。

>>69
pthreadでは速くなった？
開発環境は？

pthreadって要するにホスト側スレッド生成にcutStartThread()を使うってことだよね？
それは実はまだやっていないんだ。これからやってみようと思う。
とりあえずスレッド走らせればいいんだと思っていたから、使用経験のあるOpenMPを使った。
もしかして、cutStartThread()で生成したスレッドを使用して各GPUを使うっていうことが
暗黙の条件だったりしないよね。
cutStartThread()で速くなった人いますかね？
開発環境はVisualStudio2005。

うちでは速くなってるけど・・・
こちらはLinuxだけど、もしソースが載せられるならうｐしてみて。

>>72
ソースは諸事情のため色々と削ったりしたもの。
ちなみにCore2Duo＋GPU×2なのでホスト側スレッド数＝GPU数な環境。

float *h_data = 0;
h_data = (float*)malloc(mem_size);
(h_dataの中身をこしらえる)
cudaGetDeviceCount(&num_gpus);
int num_cpus = omp_get_num_procs();
omp_set_num_threads(num_cpus);
#pragma omp parallel
{
unsigned int cpu_thread_id = omp_get_thread_num();
unsigned int num_cpu_threads = omp_get_num_threads();
unsigned int element_per_kernel = element / num_cpu_threads;
float *sub_h_data = h_data + cpu_thread_id * element_per_kernel;
int gpu_id = 0;
CUDA_SAFE_CALL(cudaSetDevice(cpu_thread_id % num_gpus));
CUDA_SAFE_CALL(cudaGetDevice(&gpu_id));
printf("CPU thread %d (of %d) uses CUDA device %d\n", cpu_thread_id, num_cpu_threads, gpu_id);

unsigned int mem_per_kernel = mem_size / num_cpu_threads;
float *d_data = 0;
CUDA_SAFE_CALL(cudaMalloc((void**) &d_data, mem_per_kernel));

dim3 dimBlock(num_thread);
dim3 dimGrid(num_block/num_cpu_threads);

for(i=0; i<T; i++){
#pragma omp single
{
total = 0;
for (j=0; j<element; j++) total += h_data[j];
(h_dataの値を色々いじる)
}
CUDA_SAFE_CALL(cudaMemcpy(d_data, sub_h_data, mem_per_kernel, cudaMemcpyHostToDevice));
calculate<<< dimGrid, dimBlock >>>(d_data, total);
CUDA_SAFE_CALL(cudaMemcpy(sub_h_data, d_data, mem_per_kernel, cudaMemcpyDeviceToHost));
}
}

諸事情によりomp singleを使っている場所があるんだがこいつが悪さをしているとかあるのかな。

質問です（全く素人です）

OpenGLを使えばGPGPUのプログラミングが出来ると聞いたのですが、
かたやCUDAなどの専用の言語が出てきているのはどうしてでしょうか？

だったらcとか使わずfortran66でも塚っテロ

>>76
いえ、ですから、OpenGLではGPUの使えない機能もあるのかなぁ、と思いまして。

GPUで計算中でも画面は通常通り表示されますか？

>>78
ハードについて勉強するよろし

CUDA実行中に、時たま怪しげな影が飛んだりフリーズしたりします。

Folding@HomeのGPU版をしている人居ます？
PS3の280GTXは6倍以上をこなしているようですが。

ベータ版が公開されているが時間がなくてまだ何もしていない（困った。）

>>81
マルチすんな。つーか、PS3の280GTXってなんだよpgr

>>69
OpenMPの使い方が根本的に間違っていると思う

>>83
サンプルのcudaOpenMPを参考にしてやったんだけどなぁ。
これじゃどうだめなのかぜひとも教えてくれ。
ちなみにOpenMPの使い方はサンプルにあわせている。
ただsingle構文はサンプルでは使っていないがな。

VisualStudio2005はExpress？だったらOpenMPディレクティブは機能しないよ

nvidiaドライバを削除する方法ってどこかにのってないか＠CentOS5

新しいグラボ買ったんだけどドライバの削除がわからないウンコー

grid, threadの意味がわからない
アニメに例えて教えろ

今amazonみたらGPUGems3の日本語版の予約を受け付けてるじゃｎ

ノート（OS:Vista）で動かそうとしても動きません
CUDAに対応したGPUを搭載しておりますし（8700M GT）
VC2005Expressをインストールしたのですが、コンパイルすると
nvcc fatal: nvcc cannot find a supported cl version.
Only MSVC 7.1 or 8.0.
と表示してしまいます。まったくわからなくて困ってます。
よろしければお教えください

2005以外のバージョンのVCも入れてる？
nvccが参照しているVCが2005じゃないように見える。

いれておりません。
インストールしたのは
VC2005EXとVC2005SP、VC2005SP Update for Vista
だけなのですが。

>>85
VisualStudio2005Professionalを使っています。
とりあえずcuda使わないで純粋にCでOpenMP使ったら高速化されているのでOpenMP自体は動いている模様。

ttp://vipprog.net/wiki/CUDA.html

>>93
なにこの糞サイト

質問です

PhysXがGeforceでも使えるということを聞いたのですが、
PhysX SDKとCUDAやOpenCLというのは同じレイヤーでの話なのでしょうか？

>>84
pthreadでは高速化された？

>>88
アノ値段じゃ学生はp-子するから売れん罠

Gem3はGPGPUに限って言うならあんまり役に立たなかった

>>86
--uninstall

>>99
ありがトマト
ところでSobelFilter見ててShareMemの使い方がわからないんだが
フォーラムに質問と返答あったが内容わからん＋英語で全然わからん
誰か全訳してくれ

http://forums.nvidia.com/index.php?showtopic=52612&pid=286191&mode=threaded&show=&st=&#entry286191

誰か誰か！

>>86
訳すのはめんどいからパスだけど、ポイント絞って質問してくれたら回答するよ。

>>86
C:\WINDOWS\system32\drivers
から目的のドライバを手で削除すればよい。

>>102
神ktkr
ポイント絞るからちっと待ってて…

問題は、自宅が規制に巻き込まれていてなかなか書き込めない辺り。
取り敢えず、注意点を列挙しておく。
・共有メモリを確保するのはglobalFunc<<<blocks, threads, sharedMemorySize, streamNo>>>(parameters)の
　三番目のパラメータでサイズが指定されたときだけ。
・共有メモリは一回の<<<>>>の呼び出しの間だけしか有効じゃない。
# つまり、次の回には残っていない。
・共有メモリをハンドリングするには、extern __shared__ anyType * nameで宣言するだけ。
# つまり、コンパイラは型のマッチングやサイズのチェックをしないので自分で管理しないといけない。
・共有メモリはblock間で独立、block内ではthread数に関わらず共有。
# つまり、実際のデバイスにそぐわないthread数を指定した場合はCUDA側で同期処理が入るのか、遅くなってしまう。
・あるthreadが共有メモリに書いた後、別threadが読む前には__syncthreads()で同期を取らないといけない。
# ある意味当然なんだけど、その所為で遅くなるのも事実。

あー、ついでにメモリの違いを簡単に。
※グローバルメモリ
・読み書きできる。coalescedにアクセスできれば結構速い。消えない。広い。
・ホスト側スレッドごとに独立している。ホスト側から見ると、毎回同じアドレスになるのでどのくらい使えるか判りにくい。
※共有メモリ
・読み書きできる。遅くない。
・呼び出しごとに消えてしまう。余り広く取れない。事実上同期を取る必要があって使い難い。
※定数メモリ
・速い。消えない。そこそこ広い。
・例えばfloat2を読み込むインストラクションがないので実はグローバルメモリからfloat2を読むより遅くなる場合もある。
・デバイス側から書き込めない。ホスト側スレッドごとに独立している。複数スレッドからCUDAを使うと毎回転送しなおすのか?
※レジスタ
・読み書きできる。速い。厳密に型チェックされる。つーか、型ごとに違うインストラクションが使われるからptxファイルで追える。
・呼び出しごとに消える。他のメモリに較べれば狭い。
# ローカルな配列は宣言したことないからよく判らん。

NVIDIAの仕様書見てもcoalescedの意味がいまいちわからないのだけど、どういうこと？
どこかわかりやすく解説しているサイトない？

カーネル内で__shared__つけて配列を宣言するのと、つけないで配列を宣言するのでは何が違うの？

>>106
GPUはWarp単位で同じインストラクションが走るから、要は16人17脚みたいに考えればいい。
メモリアクセスを16人17脚によるパン喰い競争みたいに考えると、自分のパンが目の前にある状態がcoalesced。
そのとき、2,3人パンを喰う必要がなくてもスルーするだけだから問題ない。
処が、二人のパンが入れ違っていたらそこで入れ替える間、みんなが待たされることになるって感じ。
# 判っている人には判るけど、判っていない人には判らない説明だなw

>>107
共有メモリを使うかどうか違うだけだと思うけど。ptx出力させて較べてみたら?

>>108
その説明、いただいてもいいですか？

>>109
本にするならもっと書かせてくれw
Vipのwikiに載せるなら是非やってくれ
金取って講習するのに使うのなら分け前よこせw

>>105
その通りですシェアードメモリとブロック数が理解できない

1. プロック数
dim3 blocks = dim3(iw/(4*BlockWidth)+(0!=iw%(4*BlockWidth)),

ih/threads.y+(0!=ih%threads.y));

なぜblocks.xはiw/threads.x+(0!=iw%threads.x)じゃなくて
上の式になるのか。

2. シェアードメモリ
int SharedPitch = ~0x3f&(4*(BlockWidth+2*Radius)+0x3f);

int sharedMem = SharedPitch*(threads.y+2*Radius);

SharedPitchはなぜ上の計算になるのか。
0x03fの意味、4*の意味、BlockWidth+2*Radiusの意味が理解できない

とりあえずSharedメモリの使い方はどこを調べればわかるんだ！ｳﾝｺ！

>>111
どんな数字を入れるとどんな結果になるか、Excelでも使って計算してみたら?
0x3f使うのは64の倍数にするためでしょ。

>>112
うーんそのなんで６４にするのかがわからないのよ

関係ないかもしれんがシェアードメモリを使ってないSobelFIilterも
SobelTex<<ih, 384>>>でなぜ384かわからないお
こっちは１行１グリッドにして、１スレッド１ピクセルなんだと思うが
なんでスレッド数を画像の横幅iwにしないで384にするんだぁ
スレッド数が384だと計算が速い理由でもあるのかお！

単純に、warp数の適当な倍数になるからってだけじゃなかろか。

GPUコードではmemcpyは使えないの？

>>115
デバイス側で、デバイス間のコピーをしたいってことなら、自前で書くしかないんじゃないかな。
でも多分、そこがボトルネックになると思う。

global memoryからshared memoryへのコピーの時間と
shared memoryからglobal memoryへのコピーの時間が
異なるのは何でなんだぜ？
よく分からない…

>>117
グローバルメモリへの書き込みはcoalscedでも遅いと思う。
そうでないなら、ptxファイル見てみないとなんとも。
プロファイラを使うともう少し様子が掴めるかも知れない。

そだ、プロファイラの使い方って、日本語で書かれたものがWeb上で見つからないんだよね。
誰か、まとめてない?

>>113
指定できるスレッド数の上限は合計512までだからだろ
もしくは、スレッド数を多くすると使用するレジスタ数がパンクするから

>>115
自前で作成したmemcpy関数（サンプル付き）
記憶で書いているのでデバッグは自分でよろしく
注意：sizeは4の倍数のみ

__device__ void memcpy1D(long* p_dst , const long* p_src , unsigned int size)
{
const long* p_end = p_src + (size >> 2);

p_src += threadIdx.x;
p_dst += threadIdx.x;

while (p_src < p_end)
{
*p_dst = *p_src;

p_src += blockDim.x;
p_dst += blockDim.x;
}
}

__device__ structHogeHoge g_data;

__global__ void sample(void)
{
__shared__ structHogeHoge s_data;

memcpy1D((long*)&s_data , (long*)&g_data , sizeof(s_data));
__syncthreads();
}

NVIDIAのサイトから、lameをCUDA化したサンプルコードをダウンロードして、コンパイルしてみたのだけど、
かえって遅くなるのだけど、速くなった人いる？
CPUはAthlon2.0GHzでGPUは8800GTXです。
コンパイルはサイトの指示通り、USE_GPU_HPFを有効にしてかつ、そのCPUパートはコメントアウトして実行しないようにしています。

>>121
nVidia「全然速くならないから誰か代わりにやってくれ。速くても賞金出すのは北米在住者のみな＾＾」
というコンテストだから当然

来週締切なのに今のところまともに投稿してるのが1チームという惨状

賞金はともかく、NVIDIAの書いたコードだから、勉強になると思ったのですが

>>123
甘いな。ハードの設計している連中とドライバを作っている連中と
CUDAを作っている連中とそれの応用を作っている連中が勝手にやっているのがNVIDIAだ。

http://3.14.by/en/md5
http://forum.insidepro.com/viewtopic.php?t=2051

ttp://www.insidepro.com/eng/egb.shtml

MD5 Crack on CUDA で腕試しが流行っているみたいだよ。
日本も頑張らなきゃ。

だれかフォートラン仕様にしてくれ

>>122
2チーム目が来たね
2倍ちょっとって、これならquad core使った方がよくね？

GTX280でCUDAすると素人考えで単純にSP数増えてる分だけ
速くなりそうな気がしてしまうんだが実際はどーなの？
Warpとかいろいろ考えることもありそうなんだけど。。。

>>127

英語読めていないよ。こんな簡単な読みも出来ないとなると
本当に困る。

Teamの人数が２人だよ。
登録が200を超え、提出者が20組　トップは６回目のUploadだ。
提出者が20組だから結構な盛況だと考えるよ。

NVIDIA、GPGPUに関する説明会を開催 - 今年中のFortranのサポートを計画 (2) CUDA2.0正式版は近日中にリリース予定
http://journal.mycom.co.jp/articles/2008/07/23/gpgpu/001.html

2.0とマルチCPUサポートはそろそろ、マルチGPU・多言語・デバッガ・プロファイラサポートは多分おそらくきっとそのうちだってさ

cudaを覚えはじめるのにぴったりな本ないかなぁ？

CUDAに特化した本は未だないだろ。

>>131
本家プログラミングガイド
変な日本語訳のプログラミングガイドもある。

>>132
出たら買うから　出してｗ

じゃ、書くからNVIDIAと出版社相手の交渉は任せた。

何か並列処理の本のおまけみたいな形じゃないと
正直売れないだろうなぁ

>>135
たぶん、猫でもできる・・・・とかと同じように、ＨＰたててそこで解説したら出版社から来ると思う。
わざわざＨＰスペース借りてーとかでなくても、ＷＩＫＩでもいいと思われ。
交渉とかは向こうが書いてくれっていうだけだと思う基本的に。
んで、ワードなり、一太郎なりで原稿書けばＯＫじゃね？
画面キャプチャなどの画像は、まとめて管理したほうがいいぽいよ。
なんか、ワードとかで埋め込んじゃったりすると、逆に手間かかるらしい。

NVIDIAの日本語サイトがようやく
更新されたｗ

gpuはcpuより劇早なんだから
使いたいよねえ

パスワード解析とかにも使えそう

速くねーよ。

CPUをセダンに例えるならGPUは貨物車だな。
荷物を大量に扱う場合だけは効率がいい。

それはそうとパスワードクラックなんて浮動小数演算能力が全く役に立たないものの代表格だな。

そうだね
並列処理すればいいじゃない？

そういや去年CUDA使ってそんなソフトが出てたな

>>141
MD5はやっていますが？
他に現在手に入るクラック専用マシンでCUDA以上のC/P
があればお話が成り立ちますがね。

>>144
それやってるのは整数論理算術演算じゃん
乗算すら浮動小数じゃないよ。


つーかMD5なんてすでにハッシュアルゴリズムとしては死んでるんだが。
ワンタイムパスワードとしてならまだまだ現役でやれるだろうけど
性質上瞬時にクラックできないとクラックの意味もなくなる。


ああFPGA焼いた方がいいんじゃない？

今各社が力入れてるのはAESなどの128ビットブロック暗号だな。

つーか個人でGPUで遊ぶ分にはアリかもしれんけど
でかいシステム組めるTeslaになると別に高CPでもなんでもないんだよね。
特に整数性能に関してはPS3でクラスタくんだ方がまだ良いくらい。

GTX280含め現行GPUの整数性能はまだまだ残念な印象。

パスワードクラックほどの極小のプログラムでも
Xeon＞PS3だって自分のページで言ってなかったっけ？

暗号関連ならCellはAESみたいなブロック暗号はそれなりに速いよ。
例のvpermもどき命令で1コア16並列でエンコード／デコードできるからね。
ただAESアクセラレーションはWestmereでの強化点にもなるのでこの先はわからん。


まあいずれにしてもCUDAはいろいろ残念だな。
ATiとのGPUシェア争い無視してでも汎用性能をとるか来たるLarrabeeとガチるかの瀬戸際だし。

>性質上瞬時にクラックできないとクラックの意味もなくなる。
笑った。連れの言い分と全く一緒だからな。
流れているデータをリアルタイムで解読できなければ暗号解読ではない
というような考え。ある意味では正しいが、ネットのデータは記録可能
という点を考慮していないし大穴になる。

AESアクセラレーションハもう６年も前にある技術だし今さら。

しかし認証段階で破られたらAESは即死に体だよ。実際解読ソフト
使っているのでww

別にCUDAは整数性能ウリにしてないし
このスレってムチャを前提にケチつけてるだけで結局何も出来ない無能者の集いだな・・・

並列演算の応用として暗号解読の話題を語ってるだけだと思うんだが・・・

>>151
それがわかってないといわれてる原因なんでしょ…

そうなんだ

　弱点だった倍精度もGTX280で改良された。しかし何かイマイチなんだよなぁ。
消費電力あたりの性能がたいした事ないからかなぁ

CUDAでレイトレーシングがいました。　ＣＰＵの１５倍パフォーマンスが出ているようです。
http://noridon.seesaa.net/

SSEも使ってない2コアの速度と比べて15倍じゃ今一かも

http://www.nicovideo.jp/watch/sm4095447
http://www.nicovideo.jp/watch/sm4095398
レイトレーシング

>>154
このスレ的には正常進化じゃね？

だれか4850にCUDA載せてくれ

Core2Duo(3GHz)×2 = 最大 1.89FPS
CUDA(GTX 280) = 最大52.48FPS

27.7倍高速出ているようです。

http://noridon.seesaa.net/

それatiのgpuじゃ実行できないよね？

CUDAですから

OpenCLはCUDAを採用したと聞いたが…本当なん？

OpenCommonLisp

4 x 2.5 GHz の 3 fps と比べても 17倍
これは大したものだ

IntelC++ と比べてどうなんかな？

うごかねえｗ

>>162
んなわけない
既にMacOSXにもCUDAがあるのにわざわざ10.6でOpenCL採用を謳うかよｗ

インテルは０からスタートするわけだからどのぐらいの速度でIntel製GPU普及するのかは見ものだね。　NVDIAのGPUはすでにゲームやCAD、３Dグラフィックの用途にすでにかなり普及しているからね。今後はAMDやIntelのQuadcore CPUでCUDAが使えるようになるしね。

SPがPentium-Mベースだから割と簡単に追いつけるんじゃない？

>>168
そのとおりだけど、intelには事実上CPUと抱き合わせ、という裏技があるからなあ。
今回はかなり本気っぽいし、微妙かも。

質問です。

libshをコンパイルしようとしたら以下のようなエラーが出ました。
何が足りないのでしょうか？

ShSwizzleImpl.hpp: In member function 'void SH::ShSwizzle::copy(const SH::ShSwizzle&, bool)':
ShSwizzleImpl.hpp:191: error: 'memcpy' was not declared in this scope
ShSwizzleImpl.hpp: In member function 'bool SH::ShSwizzle::operator==(const SH::ShSwizzle&) const':
ShSwizzleImpl.hpp:251: error: 'memcmp' was not declared in this scope

ｱﾊﾊ

夏だな

俺は夏かどうか確認するために書き込んだわけではありません。
馬鹿正直にくだらない質問に答えてくれる人を待っているだけですよ。
質問に答える気がない人間は不要。
スレ違いだと言わざるを得ませんね。

夏だね

そんな失礼な態度のヤシに親切に答えてくれる人がどれだけいるか
待ってるより自分の頭で英語の意味を理解した方がずっと早いと思うがｗ


>error: 'memcpy' was not declared in this scope
>error: 'memcmp' was not declared in this scope
普通の頭の持ち主なら必要なincludeファイルがincludeされてないと解釈するだろうな

>>176
一つ言わせてもらう。

お前らに　「聞いてやってる」　んだよ。

聞いてほしくてたまらないんだろ？ニートで暇だから。
お前らが答えないなら別のやつに聞くだけ。
聞く相手はいくらでもいるんだよ。

>>176
ああ、ちなみにもう解決したから。自力で。
お前らみたいに大した知識もないくせにもったいぶった無能に聞いたのが間違いだったな。

ちなみに、ソースに手を加えずに解決した。
解決方法なんて書かないよ。
同じ問題に出会ったやつは　ググレカスｗ

>>178
で、どこがCUDAに関係してたの?

凄いのが来たなぁ
釣りであると信じたいが天然厨房っぽいな

いや、質問主はまだ困ったままだろｗ
質問横取り煽り厨だろ

>>160
俺もやってみた
4コアCPUが必死に動いても8800GTに25倍くらいの差付けられた
レンダリングの革命やでぇ〜

GPGPUレンダはまだ無いから実用性のあるの早く作って売れば金になるんじゃね
今の外部レンダは綺麗だけど時間かかって静止画までしか実用性無いからな
お前らちゃんすだぞ
絶対金になるだろうから、多分作ってるところはもう取り組んでると思うけどね
ていうか俺が早く出してほしいんだよ

gemsもそうだが、なんでOpenGLやGPU関連の書籍というのべらぼうに高いんだ？

>>184
需要がない割りに書くのが大変だから。

OpenGLは需要あると思うが。

>>184
訳してる会社が（ｒｙ
「誰が買うんだこれ！？」って思う値段の本は大抵原著は安かったりする
酷いのは原著の2倍以上

OpenCLはIntel製GPUでも使えるのだろうか。

>>188
Larrabeeのこと？

>>188
使えるでしょ
自社が参加してるワーキンググループが発表するんだから

ttp://www.gpureview.com/nvidias-g92-and-g94-parts-are-reportedly-failing-article-721.html
うーむ……

http://jp.youtube.com/watch?v=AgoEri9UX7o
GeForceでスパコン終了
一時間が30秒に

ゲームをやってた助手が「あ、どうも」、わざとらしくてﾜﾗﾀｗ

こんちにわ。LinuxでCUDAのプログラムを作成している者です。
早速質問ですがMPIを用いて、複数のクラスタ（PC５台、そのうち１台が親機で４台が子機）で
CUDAプログラムの動作できた方はいませんでしょうか？
プログラムの内容として、子機で演算させ、親機では適当なメッセージを出すのみという、
並列演算とか関係なく、単にMPIでCUDAを動かすテストプログラムです。

下記のようなMakefileでコンパイルを行っています。
ROOTBINDIR = ./
ROOTDIR = /home/usr1/NFS/NVIDIA_CUDA_SDK/
# Add source files here
EXECUTABLE := mpicuda_matrixMul
# Cuda source files (compiled with cudacc)
CUFILES := mpicuda_matrixMul.cu
# C/C++ source files (compiled with gcc / c++)
CCFILES := \
matrixMul_gold.cpp

SMVERSIONFLAGS := -arch sm_11
############################################################
# Rules and targets
INCLUDES = -I/home/usr1/mpich-1.2.7p1/include -I/home/usr1/NFS/NVIDIA_CUDA_SDK/common/inc
USRLIB = -L/home/usr1/mpich-1.2.7p1/lib -L/home/usr1/NFS/NVIDIA_CUDA_SDK/lib -lmpich -lcutil
include ../../common/common.mk

続きです。
上記のようなMakefileをNVIDIAのCUDA関連のBBSで見つけたので参考に作成しました。
しかしコンパイル時に下記のようなエラーが出てきます、、

/usr/bin/ld: cannot find -lcutil
collect2:　ldはステータス１で終了しました
make:　***[linux/release/mpicuda_matrixMul] エラー１

上記のエラー内容を調べたところ、libcutil.aというライブラリがリンクされていない、
という内容と思われるのですが、Makefileには正しい場所にリンクさせてあると考えています。
開発環境として、OSがはCentOS５、親機のGPUがGeForce7900、子機が8600です。
（親機でCUDAの演算は行わないので7900）

それでは、どなたかよろしくお願いします。

>>195
cutilはmakeした?
した積もりなら、ちゃんとライブラリファイルはできている?
OpenGLがインストールされていない環境だと、恐らく失敗している筈。
その場合、OpenGLをインストールするかOpenGL絡みのエラー修理辺りをコメントアウトすればmakeできるはず。

MPIでCUDAを試したことは無いけどGLUT，GREWがインストールしてないと
>>195のようなエラーが出たことがある

なあ、あのへんてこ日本語プログラミングガイドのリンクが消えているような気がするんだけど、
前からこんなんだっけ？

いいことだw

http://jp.youtube.com/watch?v=AgoEri9UX7o
530万ドルのスーパーコンピュータ「CalcUA」で一時間
一台デスクトップPCで30秒

http://www.atmarkit.co.jp/news/200807/23/nv.html
ベルギーのアントワープ大学では、それまで使っていたAMDベースの256ノードのクラスタサーバ「CalcUA」の性能を、8GPUを使ったデスクトップPCシステム「FASTRA」が上回った。「CalcUAは530万ドルのスーパーコンピュータ、FASTRAは7000ドルのデスクトップだ」（グプタ氏）。
NVIDIAによればGPUを使った同様のクラスタサーバは、NCSA、イリノイ大学、ノースカロライナ大学、マックスプランク研究所など、すでに十数の組織で使われているという。

>>196 さん 、>>197 さん
お返事ありがとうございます。
cutilのライブラリは確認済みです。CUDAインストール時に作成される、
〜/NVIDIA_CUDA_SDK/lib の中にlibcutil.aがありました。

OpenGLは前からインストールしており、またそのプログラムは動くので、
GLUT等はインストールされているはずなのですが。。

/home/usr1/NFS/NVIDIA_CUDA_SDK/と〜/NVIDIA_CUDA_SDK/libが違うって落ちじゃないよな。

自分からテスト専門です、って宣言してる派遣テスターって何なの？

将来プログラマとかSEになりたい、とかならわかるけど。
向上心ないよね、頑固だし。
そういう派遣テスターって、仕様書は読めない、
テスト仕様書も作れない、テストプログラムも作れない
やれることは「テキトーにプログラムを触る」ことだけ。

俺は派遣プだけどさ、こういう派遣テスターがいると
派遣全体がバカにされるんだよ。
テスト専門派遣なんて氏んで欲しいよ、まったく。

今日も正社員の人が派遣テスターに仕様書を読んで
テスト仕様書を作ってください、って説教してたよ。
その派遣は頑固に「何故、仕様書が必要なんですか？」って
反論してたから、きっとテスト専門派遣テスターだな。
仕様書も読まず、テスト仕様書も作らず、ただテキトーに
プログラム触るだけで給料もらおうなんて頭おかしいんじゃねーの？

あ〜あ、あの派遣テスターが3ヵ月後に切られるまで、
仕様書も読まねーでテキトーにテストしたバグ票がまわってくんのかよ。
そんな糞なもん、読んで処理する派遣プの身にもなってくれよ。
うわ〜、しかもそいつが切られる3ヵ月以内に中間納品あるじゃねーか！
テスト仕様書もなしにテキトーにテストして納品か。
中間納品後にソッコウクレームでデスマ必至だな。俺の休みも返上かよ。

派遣専門テスターさんよ、少しは向上心持てよ！
頑固な性格直して仕様書読めよ！テスト仕様書作れよ！

うちの客先はテスト専門部署があるけど
環境構築とかテストプログラム作ったり
海外のテスト専門会社へ行って研究したりしてるぞ

バカの作ったものをテストするなんて最悪な仕事の1つなんだぜ

以前入社してすぐやめた会社でテスターをやらされたが、
そこの会社、3Dの知識ないやつが、OpenGLで懸命に2Dの描画プログラム作ってた

まさにバカの作ったものでテストするのがバカらしくなって会社を辞めた

判ったからマ痛に行け。

>そこの会社、3Dの知識ないやつが、OpenGLで懸命に2Dの描画プログラム作ってた
こういうバカ、たまにいるよな
2Dだけ描画すればいい仕様なのにOpenGLしか知らないからとりあえずOpenGLで書いちゃうという
すぐに辞めさせないとプロジェクトが悲惨な結果になる

>>208
AutoCADは2DもOpenGLで描画していますが。

>>209
>207

3Dも描画するならいいだろ
2Dしか描画しないのにOpenGLはバカ

GDIで描画したらLinuxでWineが必要になっちゃうじゃん

いや、2Dだけ描くのにOpenGL使うようなバカはWindowsだけだろ

このスレでうだうだ3D話を続ける奴って頭弱いの?
それとも只の構ってちゃん?
しかも切っ掛けはコピペだって辺りが痛過ぎるよね。

きっかけはコピペでのバカの作ったものをバカにされるべき
GPU使っているはずなのにBITBLTしかしてないとかー
もーバカ杉

BITBLTの後ろ三つはベーコン・レタス・トマトの略だが、前のBITは何の略なんだ？

マルチプラットフォームで2Dを描くのに仕事で使えるレベルの表現力と速度を兼ね備えたライブラリって存在するの？
あればどんなものか教えてください。有償無償問わず。

>>217
OpenGL

マルチプラットフォームの仕事なんて普通ないだろ

プロならプラットフォームに最適化されたライブラリを選択しろよ

OpenGL専門学校の学生さんにマジレスするなよ

>>219
>マルチプラットフォームの仕事なんて普通ないだろ
？

>>217
http://d.hatena.ne.jp/keyword/RenderWare

>>221
これから増えるだろう
お前は近視か？

間違った。>>219だ。

>>219
openofficeはいったい何なんだ？マルチプラットフォームのソフトウェアじゃないのか？

お前は仕事でopenofficeにかかわったのか？
違うだろ
派遣先のプラットフォームだろ

お前の会社がマルチプラットフォームの製品を開発してんのか？
違うだろ
派遣先のプラットフォームだろ

現実見ようぜ

>>226
まず、仮定からおかしい。
・マルチプラットフォームのソフトウェアはopenofficeだけではない。
・>>225の仕事の内容を断定する根拠が見当たらない。
・なぜ無数にある可能性から>>226の仮定を導き出したのか不明。

説明をお願いします。

>>227
>・マルチプラットフォームのソフトウェアはopenofficeだけではない。
まず仮定からおかしい。
実際に仕事でマルチプラットフォームを開発しているのかを問題にしている。
>>225がopenofficeを開発したというのなら謝る。おみそれした。

>・>>225の仕事の内容を断定する根拠が見当たらない。
まず自分で開発してそうにないopenofficeを出してくるあたり、
ただのバカ＝派遣ってこと

>・なぜ無数にある可能性から>>226の仮定を導き出したのか不明。
派遣に可能性なんてない

荒らしは放置の方向で。

219のほうがバカだと思います。
マルチぷラットフォームの仕事はあると思います。

だからお前はマルチプラットフォームの仕事をしたことあるのかと

サターンとプレステで同じゲームを同時に作ったぞ
同時に作っただけだけど

>>231
Firefoxはマルチプラットフォームじゃないのかね。

マルチプラットフォームの要件はオープンソースではかなり多い。
オープンソース系の企業ならそういう仕事していてもおかしくないね。

日本語読めないの？
聞かれてるのは「お前は仕事でマルチプラットフォームをやったことがあるのか？」だぞ

>>235
話のすりかえ。
俺自身は学生時代にFirefoxの開発に携わったが、
今の話題は「マルチプラットフォームの仕事があるかないか」。
お前が間違っていることに気づけ。

そもそも、ここの住人が開発に携わったことがあるかどうかなんて誰も興味を持っていないし、
信憑性もない。
もっと有益は議論をすべきだと思うけどね。

スレ違い　お前らネットイナゴは佃煮にされて死んでしまえ

派遣先企業で派遣先企業のプラットフォームに合わせたものを実装するのがおまえら派遣プログラマの仕事だろ
派遣先正社員からマルチプラットフォーム対応なんて指示されてないだろ
くだらないこと考えて余計な工数使うなよ

>派遣先企業で派遣先企業のプラットフォームに合わせたものを実装するのがおまえら派遣プログラマの仕事だろ
派遣先がパッケージベンダってこともあるだろう
なにムキになってるの？

>>239
何故派遣にこだわる？

ID出すとしょうもない素人以下の奴がいろんなスレに粘着してた

>>242
この板はID出ないじゃん

なんだみんな派遣だったのか

派遣に異常なコンプレックスを持っている人がいるようですが、
池沼の方なので基本放置でおねがいします。

夏の間だけ2chもアカウント認証式にすればいいと思う。

正社員にコンプレックスでしょ？
派遣のどこにコンプレックスをもてと

ひどい脱線ぶりだな。

派遣野郎はスレ違い

>>232
それありそう。PS3 と Xbox360 とか。どう似せるかに力が注がれる・・・

ここはなんのスレなんだ

くだすれ

CUDAはマルチプラットフォーム。

アホかい。

━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━
■ 今週の名言
──────────────────────────────────────
　　　　　　　　　「HPCは“ハイ・プロダクティビティ・コンピューティング”」
　　　　　　日本SGIのHPC・サービス事業本部で本部長を務める田坂隆明執行役員）

CUDAってSIMD並列といいうことは、過去のスパコン用のコード
移植したら激速ってことでOK?

>>256
SIMDじゃないよ。仮にそうだとしても、そんなに単純な話じゃない。

SIMDだよｗ

SIMDじゃないなら、何なのか問いたい。

SIMDだったりMIMDだったり

実行はSIMDだけど
コードはスカラだから面倒臭いよ

NVIDIAはSIMTと言っているね。SIMD＋αくらい。
オンボメモリはSX8i並だからやっぱ速いんじゃねーの。

理屈はいいからなんか作れよ
＞＞過去のスパコン用のコード移植したら激速ってことでOK?
OKの一言で終わることをなにぐだぐだやってんだか

要するに、「移植できるもんならしてみろ」ってことだろｗ

だね。単純に移植できるとは思わない方がいい。

たしかにブロック切り分けの所もう少しスマートにならないのかと思うのだが
あんなもんデバイスの種類で最適化した組み合わせを自動算出すりゃいいだろ

>>266
CUDA2でストリームプロセッサ数なんかも取得できるようになったから、
是非とも自動算出関数を作ってみてくれ。

いや作るのは簡単なんだけどCUDAってクラスも使えないしラップ出来ないからだめじゃん

あいや出来るのか
使うのなら作るぞ

CUDAに興味があっていろいろ調べてるんだけど
http://slashdot.jp/hardware/comments.pl?sid=421471&cid=1432609
こんなにメモリアクセスがきついのか…
へたれの俺はCPUで頑張ろうって気になってきた
このあたりの条件が和らぐのはどれくらい後だろう？

それでもビデオカードは一世代速いメモリを積んでいるから、レイテンシを誤魔化せれば、CPUより多少はいいんじゃないかな。
CPUもGPUも、演算能力よりメモリのスピードがネックなんだよねえ。

>>270
そうそう、グローバルメモリはコヒーレントな読み出しなら4クロックなのに
ランダムアクセスすると100倍のクロックを喰ってしまう。
スラッシュドットの記事の書き方はちょっと言葉が足りてなくて、
実際にはいかにメモリをコヒーレントな読み書きで済ませるかが鍵ということになる。

例えば、CUFFTの2Dを1000回回すのと1Dバッチと自前の転置を2回ずつ動かすのとでは、
64x64辺りだと余り変わらないのに128x128になると途端に2Dの方が速くなる。
プロファイラで調べると、CUFFTの内部で用意されている転置関数はコヒーレントな読み書きしかしていない。
一方、自前の転置はコヒーレントな読み出しとインコヒーレントな書き出しになっている。
つまり、1000回呼び出すコストをインコヒーレントな書き出しコストが上回ってしまうということだ。

CPUより遅くなるということはない
電気代の割りに効果が少ないというのはあるｗ

TMPGEncで軽い処理はCPUより遅い。
VRAMとのメモリ転送あるし。

メインメモリと共有するタイプのGPUで処理した方が速い事も状況によってはあるのかな

CUDAでは有り得ないからスレ違い。GPGPUスレにでもどうぞ。

すいません、超低級な質問です。

dim3 threads(128, 1, 1);
dim3 grid(128, 1, 1);

hoge_kernel <<< grid, threads >>> (d_ptr, 128);

とかで関数を呼び出したんですが、ホスト側のスレッド数が128個生成
されるみたいです。これって、GPU内にスレッドが出来るんだと思って
たんですが、違うんでしょうか？

ちなみに、-deviceemu はつけてません。コンソールに以下のように
出るので、GPUにはいってると思います。

Using device 0: GeForce 8600 GT

>>277
＞ホスト側のスレッド数が128個生成されるみたいです。
そう判断した根拠は?

それとは別になるが、128x128が妥当かどうかは検証が必要だと思う。

質問です。

__device__ void hoge_kernel(void)
{
}


extern "C" __global__ void hoge(void)
{
dim3 threads(16, 1, 1);
dim3 grid(16, 1, 1);

hoge_kernel <<< grid, threads >>> ();
}


上をコンパイルすると、
「error: call can not be configured」
となるんですが、これはどういうことなんでしょうか。

__global__な関数から__device__の関数を<<<>>>で呼ぶことはできない。

ちょっと質問させてくらさい。。

８ストリーミングプロセッサ（SE)単位で構成されている
ストリーミングマルチプロセッサ（SM)でのことなのですが、
グローバル（ローカル）メモリ上へのデータのロード、
ストア命令は100サイクル以上かかってしまいますよね。

その間、8つのストリーミングプロセッサはおお休みしている
のでしょうか？

それとも、他のWarpのインストラクションが割り当てられたり
するのでしょうか。そうであるといいのですが、
そうするとSharedメモリがあとからのWARPにのっとられてしまって
まずそうですよね???

レジスタ的には32bit×2048本あるようなので(280の場合)
OKそうですが、Sharedメモリは16KBytesしかないし。

実際LD、STに160サイクルぐらいくったら、
160命令ぐらい無駄にしちゃって効率をかなり落としてる気がするので、
なんかやです。

24WARPぐらい、バッファリングしていて、割り当てられるWARPを
out of oderで実行するという記事も、どっかで見たような
気がするのですがSharedメモリがじゃまして無理があるような・・・

なんか良い方法があるのかな???

っていうか、なにか根本的な部分で俺の勘違い？

・グローバルメモリアクセスは、最大400(?)クロック掛かるが、最短では4クロックで済む。
# そのためには、coalescedにアクセスできるように工夫する必要がある。
・各ストリーミングプロセッサは、独立して動作する。Sharedメモリも同様。

例えば、行列の転置のような処理の場合、普通に書くとcoalescedに読んでincoherentに書かざるを得ない。
# 或いはその逆か。
そこで、CUFFT内で行なっている転置処理では、(プロファイルで見る限り)一旦共有メモリにおいて同期を取ることで、
読み書き共にcoalescedアクセスを維持しているようだ。

>>282
早速のレス、さんくす。

なるほど、最短４クロックですか。coalescedにしてもレイテンシー
だからだめかなって思ったけど、よくよく考えると、DRAMへの直接
アクセスに数百クロックってのはおかしいことに気づいた。

各ストリーミングプロセッサは、独立だけど１インストラクション単位で
同期してるんじゃないんかな（んなことぁない？）

・・・っておもったけどブランチがあるから、TPXレベルから見ると独立か???

最初見たときSharedメモリも他のSPのregも０レイテンシで使える
ようだったんで、演算と独立にグローバルメモリとのロード、ストアができん
１６K程度のSharedメモリ、あまり意味ないじゃんとかおもったけど、
そんなことなさそだね。

なかなか面白そうなチップだ。
分岐粒度が32って実際、汎用でどうなんだろ。
ベクトルつかったことないんでわからんが、
適当にセルオートマタでも乗っけて遊んでみまつ。

CUDA-Zなんて便利なものがありました。
forum.nvidia.co.jp

>>284
kwsk

ttp://cuda-z.sourceforge.net/

あげてすまん。

CUDAのためにGTX280つきのPCかったのだが
ほかにもう一枚グラボ買わないといけんのかな？

あとPCIバスしか残ってないのだが・・・
LINUXで使う予定。

>>287
別に一枚でも大丈夫。二枚挿しても(開発には)却って面倒だったり。
まぁ、二枚あれば表示に足引っ張られる心配なくなるから安定はするけどね。

>>288
そうなのか、安心した。

ありが�d

質問です。

float にてゴリゴリ計算して、結果を返すプログラムを書いてみたのですが、
普通のアルゴリズム、CUDA+GPU、CUDA+CPU(deviceemu) の2つで比較して
みたところ、思ったより差が大きいのですが、こんなもんですか？

もしくは、機種依存があったりするんでしょうか？

>>290

すいません、間違えました。
「2つで比較して」->「3つで比較して」

何の差?
演算精度なら、そりゃぁあるさ、floatなんだもの。

ホストは、デフォルトではfloatをfloatで計算してくれないからな。

SSEを使うなりで、ホストがちゃんとfloatで計算すれば、結果は一致するんジャマイカ？

>>293
一致する保証はないよ。CUDAのドキュメントにもあるけれど、超越関数はGPU内部の組み込み版を使うと若干誤差が残る。
いずれにしてもfloatの想定の範囲内だから、実用上は問題にならないけどね。

x87 SSE CellB.E. CUDA の浮動小数サポートの対比表みたいなのがCUDAのマニュアルにあったな。
確かに完全じゃない。
糞まじめに準拠してるのはx87くらいだ

>>292-295
なるほど。出てきた結果が、CPU上の double で計算した結果と、
GPUの float で計算した結果が、最大1%程度違ったから、正直
驚きました。普段doubleしか使って無くて、誤差なんかほとんど
気にしなくてよかったので。

誤差が蓄積してくようなタイプのアルゴリズムではないと思っていた
だけに、少し驚きました。

Cで使っていた自作ライブラリは、
nvccでコンパイルし直さなきゃダメなの？

食べてみたけどおいしかったよ。
一つどうぞ、あーんして

飾りだ、食うなボケ。

僕のもうないです
>>299から貰ってね

>>297
ホスト(CPU)上で実行させるものならVC++と一緒にリンクできる。
（cppIntegrationサンプル参照）

GPU上で実行させたいならそれなりにいじくらないと無理。

-deviceemuでは動くプログラムが、GPUを使うと誤動作します。
__device__関数に引数が一部しか渡らなくなります(float4のz成分だけしか渡せない)
ループ回数を極端に減らすと改善されるのですが、これはregisterメモリがパンクしてしまったということでしょうか？

ソースを見ないとなんとも言えませんが、ループ回数とregisterは依存関係にはありません。
nvcc -ptxでptxファイルを出力して眺めてみては如何でしょう。

ptxファイルですか。。。
正直どこをどう見たらいいのかわからないので敬遠していましたが、
遂に避けられないとこまで来てしまいましたかねえ。

見るべきポイントとかありましたら、よろしければ教えてください。

ptxもネイティブコードじゃなしに中間コードの断片だからな。
レジスタが何本割り当てられるかとかそういった情報すら持ってないから困る。

その点Larrabeeは16本+16本+αとかレジスタ本数が決まってて
コード生成時点で割り当てが決まってしまう。
結果的に静的な最適化がしやすい。

良し悪し。

ﾘｮｳﾃﾆ●ヽ( ・∀・)ﾉ● CUDA！

.oとかにしてgccでリンクすることってできますか？

>>307
何を? どこで? なんで質問もろくにできないの?
ちなみに、エスパー募集ならお門違い。

>>307

普通にできますよ。
-l/usr/local/cuda/lib -lcudart
等のオプションは当然自分で追加する必要があります。

Windowsで、とかいう話だったりしてね。

.objじゃなくて.oって書いてあるんだからそれはないんじゃない？

それだったらわざわざ聞くかねぇ。まぁいいか。

【GPGPU】CUDA/ATI　STREAM　速度・画質検証スレ
http://pc11.2ch.net/test/read.cgi/jisaku/1228764782/
　

これからCUDAを始めようといろいろサイトを巡回しているのですが、
グリッドサイズとブロックサイズはどのようにして決めたらいいのでしょうか?
同期をとるかとらないかで変わると思いますが、とりあえずとらない場合でお願いします。

>>314
BlockSize(スレッド数)は32の倍数で余り大きくない辺り。GridSize(ブロック数)は充分大きく。

>>315
ありがとうございます。
もう一つ質問させていただきたいのですが、たとえば、512x512のようにブロックサイズ、グリッドサイズ
ともに最大値未満の場合、<16, 32, 1>のように複数の次元に分けるのと、<512, 1, 1>のように1つの次元
にまとめるのとどちらがいいのでしょうか?
多分通常はブロックサイズは各次元の最大値<512, 512, 64>をオーバーすることが多いと思うので、こちらは
分けることになると思いますが、グリッドサイズの最大値は<65536, 65536, 1>もあるので、
次元を分けなくてもいいときは分けない方がいいのでしょうか?

threadサイズが実際の分散処理の単位になる
つまりthread数だけ並列演算が行われる
例えばthread(1)にしてblock(10)とかだと順次処理と変わらない
ただしthreadは256までしか出来ないのでblockという単位が容易されてる
blockは単にプログラムがやりやすいように用意されただけのもので
実際に分散処理には影響しない

>>316
言いたいことがよく判らんが、演算量が512x512ならブロックを8x8にしてグリッドを64x64でいいんでない?
或いはxyの区別が重要でないならブロックを64、グリッドを4096とか。
私は後者の方針でptxファイルの出力とプロファイルの結果を睨みながら決定することが多いけど。
# 横着するときは前者だなぁw

>>316
単に内部ループが
for(x=0;x<?;x++){}
_global_
か
for(x=0;x<?;x++){
for(y=0;y<?;y++){
_global_
}
}
になるかの違い

元の位置を特定する為にglobal関数でインデックスを計算することになるから
もしxやyのどちらかで収まるなら１つだけ使うのが良い
global関数内部の計算コストが一番ネックになる部分だからね
インデックスの計算部分が一番邪魔で削れるだけ削った方が良い

>>317,319
ありがとうございます。ようやく喉のつっかえが取れました。

>>318
試しに昔書いた画像処理のコードを移植しようとしていたのですが、入力によって
いろいろサイズが変わるので、512x512で固定するのはちょっとやりづらいかなと…
ただ、ここまでの話からするとchannel(RGB)*height*widthの三次元配列に画像をつっこんで、
grid = <height, channel, 1>
block = <32, width / 32, 1>
で処理を回せそうですね。画像の幅をを32の倍数に合わせる必要はありそうですが。

>>320
ありがとうございます。実際に使うときは上のような感じにブロックとグリッドの数を決めようかなと
思いますが、どうでしょうか?

heightは超えてもいいの？
範囲は内部でインデックスが範囲かどうか判定するよりは
端数を含めたメモリを確保して必要無い部分も計算させてしまった方が速いよ

>>322
そこまで大きな画像を入れるかどうかは分かりませんが、確かに高さが65536を
超える場合は分割してやる必要がありますね。

>端数を含めたメモリを確保して必要無い部分も計算させてしまった方が速いよ
ということは、配列にコピーするときに幅を合わせてコピーするのではなく、後ろに空白部分を
まとめた方がいいということでしょうか?

#これって多次元配列を引数に渡すのってどうするんだろう…

いや、最終的には全て一次元の配列でそれが適切な間隔でアクセスされるように並んでいるのが一番いいことになる。
だから、画像なんかの場合だと座標値に意味がないなら詰め込んで構わない。
実際には、コンボリューションフィルターなどで隣の画素を参照しないといけないことが多いだろうけど、
その場合はどうせ共有メモリを使ってアクセスを減らすなどの工夫が必要だからどの途32の倍数に拘る必要はない気がする。

論文用のならともかく普通の画像処理で画像サイズが固定されてることはまずないし
サイズも共有メモリじゃとても足りない
共有メモリを使うならCPUサイドであらかじめ画像を分割して渡すという方法しかないけど効率悪いよ
毎回隣接した画素情報も含めないといけないから

>>324
カーネルサイズにもよりますが、sharedメモリはちょっとサイズ的にきつそうですね…
constantメモリならブロックサイズを調整すれば何とかなるかも?

>>325
単なる趣味グラムなので、ぶっちゃけ全部globalメモリからとってきてもいいような気がしたのですが、
逆にそれだとCPUで処理するよりも遅くなりそうですね

遅くなりそうというかどんだけがんばっても最終的にはglobalメモリを使うのが一番早いという結論になるよ
まあそれがGPGPUの一番のネックなんだけどね
CPUサイドのメモリを共用するようにいずれはなるんだろうけど今の限界はそこ

一番(手っ取り)早いのは同意するが、一番速いとは限らない。
それこそ、思考放棄じゃないかな?

詰まってます、助けて・・・
私はVS2005で開発しようとカスタムウィザードを入れ、ツールキットとSDKをインスコし
パスも通したのにLNK1181エラーでるんです・・・存在しないオブジェクトファイルを指定するのはなぜでしょう？
ライブラリパスがみつからないってどういうことなんでしょうか？

"Program" "Files"と認識してるとか

Windowsってなんで重要なディレクトリ名にスペースなんか入れるんだろう。
MSって頭悪いの？

低脳を排除するためかもねｗ

プログラマお断り。ずっと消費者で居てくだちい。

デスクトップやマイドキュメントのスペースや半角カナフォルダはなくなったが、Program Filesはそのままだな。

>>331
その昔１６ビットＯＳが主流だったころＭ＄は８文字ファイル名しか扱えなかった
そしてUnix系ＯＳは既にロングファイルネームに対応していた
そこで登場したのが32bit Window 95だった
長いファイル名が使えるんだぞ、凄いんだぞと新発明でもしたかのように宣伝していた
そして自慢げにフォルダ名を無駄に長くした

VistaはUNIXerには優しくなったな。
上位エディション限定だけどSUAも使えるし

CUDAをさ、VS2005で使えてる人っているの？
つか、なんで英語のマニュアルしかないんだよ。わかんねーよ

むしろ2005以外でどうやって使うんだ？
2.1βで2008対応したけど

>>338
おお、それはいいこと聞いた
やっと2005アンインストールできる

>>338
どこか日本語で環境設定を教えてくれるサイト知ってたら教えてくれ・・・

環境設定なんか何も要らんがな。スクリプト走らせるだけ。
bashのバージョンによってはウォーニング出るけど、ちゃんと動く。
スクリプト猿にも成れないとは、どんだけゆとり脳なんだよ?スポンジなの?

どのみちあの程度の英語が読めない人に優れたプログラミング書くのは無理だと思うんだ。
俺の私的な業務の手伝いしてくれるなら日本法人さんに代わりにチュートリアルとかの要求
してあげてもいいんだけど、職権乱用か。

>>340 ここで解らなかったらもう諦めたら？
http://chihara.naist.jp/people/STAFF/imura/computer/OpenGL/cuda1/disp_content

どうでもいいけど、↑のURLは「disp_content」を削らないとエラーになった。

>>340
サンプルプロジェクトを開いて設定を見て勉強するのがいいよ

英語読めないのにプログラミング言語してる男の人って・・・

英語に優越感もってるやつって、さもしいな

>>345
可愛いもんだよねｖ

英語なんて誰でも１ヶ月もあればある程度は読めるようになるよ
出来ないと思い込んでやろうとしないだけで

vs2005の環境で拡張子cuの場合、定義参照はできるのですが、
インテリセンスが機能しません。freeとかvcのランタイムも同様に
機能しないのですが、やり方分かる方いますか?
ツール->オプション->テキストエディタ->ファイル拡張子にcu c++として
登録するとできるのは上記だけです

いっそ、nvccの拡張機能以外は全部cppでやったら?

それがスマートだね

質問です。
VC++2005で開発する際に、コマンドプロンプトにデバイス名も含め一切文字を表示させたくないのですが、どのようにすればよいでしょうか？

>>352
長いパス名が邪魔なだけなら set PROMPT=$G とか？

>>352
CUTILを使うと標準出力があるって話なら、リダイレクトするかCUTILを使わなければいいと思うのだが。

プロンプト出したくないって話なら_tmainじゃね

>>353-355
回答ありがとうございます。

>>353
質問があいまいですみません。
パス名が邪魔、ってわけではないです。

>>354
おそらくこの指摘にあてはまるのだと思います。
＞リダイレクト
リダイレクトしても、コマンドプロンプトに表示が出てしまいます。
＞CUTIL
cutil.hを使わずに、cutil.hで提供されている関数を使わない方法、もしくは、代替可能なヘッダーファイルはありますか？

>>355
_tmainを使ってしまうとWindowsでしか使えないソースコードになってしまうので、別の方法があれば教えてください。

自己レスですｗ

誤：cutil.hで提供されている関数を使わない方法
正：cutil.hで提供されている関数を使う方法

の間違いです。

system("cls");
とか

>>356
> _tmainを使ってしまうとWindowsでしか使えないソースコードになってしまうので、別の方法があれば教えてください。
Windows かどうかで #ifdef するくらいは許されるんじゃないの？

プロジェクトの設定でWindowsアプリにすればいいだけでしょ

>>_tmainを使ってしまうとWindowsでしか使えないソースコードになってしまうので、別の方法があれば教えてください。
mainだろうがWindows依存部分のコードが発生するのは当たり前でしょ
CUDAの部分とC++の部分とWindows依存部分とファイルを分離するのが鉄則だよ
複数の環境で使う場合は#ifdefプラグマを使って分離したWindows依存部分の
#includeを切り替えるようにするだけ

コンパイラのオプションで指定できるでしょ
環境ごとにmakefile書くだけじゃね

>>361
#ifdefはいつからpragmaになりましたか?

>>363
#ifdefディレクティブと言いたかったんじゃないの?
で、使い慣れない言葉なんでつい間違えたと。
句読点もちゃんと使えないようだし、日本語に慣れていない三国人なんでしょ。

cudaのMDコードどこかに追ってませんか？

たった1行でここまで意味不明なのも凄いな。

×追ってませんか？
○オッドアイいませんか？

>>365氏はcudaで書かれたMD(分子動力学）コードどこかに落ちてませんか
と聞いてるのかと。当方も知らないので答えられませんが。
Fortranで書いたものを全部はアレなのでGPU上で実行したいサブルーチンだけ
Ｃに変えてCUDAで動かしたいのですが、そんな例とかは落ちてないですかね。
mainルーチンその他関係ないところまで全部Ｃに移植するのが嫌ってだけな
んですが。あ、当方はIntelFortran使用。
当方まだCUDA触りたて、試しにSTREAM BMTのtriadだけ手で適当に書いて
GeForce9600GTで40GB/s弱（効率7割弱）のメモリバンド幅。あ、じゃもうでき
るじゃんとか勝手に思い、Fortranのコードに挑もうとしてあえなく止まってますorz

つttp://www.easize.jp/product/fortrantocudamenu.html

>>369
ありがとうございます。m(_ _)m
多次元配列はどのみちGPU上では一次元化するつもりでしたがなるほど。
参考にさせていただきます。

>>369
そのページ、ポインタの扱いがメタメタだな
わかってるなら別にいいが

うん、ぶっちゃけ「"cuda fortran"でぐぐって一番上に来たリンク貼り付けただけ」なんだ、すまない

何か普段使うものをCUDAに移植しようと思いつつ適当な物が見当たらない

ぶっちゃけ、俺もそうなんだよね。おまけで付いて来たサンプル書き換えて
動かしてFLOPSベンチしてるだけ。
CUDA動かす環境をコスパ良く構築するには?と考えて色々やってみたが、
構築した環境で動かすモノって、結局サンプルの改造ばっか。まんまと
nVidiaの販促戦略に乗せられたぜw

zip圧縮とかjpg圧縮とかを移植したらライセンスの関係はどうなるの？

圧縮アルゴリズムに関しては問題ないんじゃね？

CUDAプログラミングモデルの概要
ttp://http.download.nvidia.com/developer/cuda/jp/CUDA_Programming_Model_jp.pdf
CUDAプログラミングの基本（パート I - ソフトウェアスタックとメモリ管理)
ttp://http.download.nvidia.com/developer/cuda/jp/CUDA_Programming_Basics_PartI_jp.pdf
CUDAプログラミングの基本（パート II - カーネル)
ttp://http.download.nvidia.com/developer/cuda/jp/CUDA_Programming_Basics_PartII_jp.pdf

デバイス側に確保したメモリにホスト側から
cudaMemcpyのように一括でコピーするのではなく、
一部分だけコピーする、又は書き換える良い方法がありましたら教えてください。

・cudaMemcpy()で一部分だけコピーする。
・cudaMemset()で（ｒｙ
・そういうカーネル関数を用意して呼び出す。

一部分だけってのはたいてい順次処理になるからCPUにやらせたほうが有利だよ

CUDAのプログラミングでZIPアーカイブのパスワード解析とか、早くなりませんかねぇ。
エンコード/デコードに使えるんだから、どうなんでしょう？

CUDAのfortranサポート予定ってGPUカーネルの部分をfortranライクに
書けるようになるって理解でいいのかな？2.0からって見かけたけど、それっぽい記述が全く無い…

>>379-380
ありがとうございました。
参考にさせていただきます。

>>381
総当りか辞書型か知らんが演算能力よりI/Oの問題だろ普通。

>>381
早くなるよ

>>384
ファイル数が多い場合はI/Oも問題になるかもしれませんが、
ある程度のサイズであればWindowsでもメモリに読み込むので、
それほど問題にならないのでは？
毎回物理的に読みに行くわけではないし。

ファイルI/OじゃなくメモリI/Oなのでは？

>>387
お前は何を言っているんだ

>>388
pciバスつかうんだろ？
立派なioとおもうが



つかクーダってほんとマイナーなんだね。

そりゃNVIDIA限定だからだろ？グラボを選ぶのなら汎用的なアプリは作れない。
CUDAよりもっといろんなグラボに対応しているOpenCLへのつなぎに過ぎないよ。

ZIPのパスワード解析にZIPファイル全体へのアクセスが必要だと勘違いしてる馬鹿ｗ

>>381-391
ああ、俺がぼけてた。なんか知らんけど自分でもなんであんなレスつけたんだろ orz
わけわからない流れにしちまって、すまん。

正しくはこうだな。↓
パス解析では総当りにしても辞書にしてもforなどのループ速度が求められるだけで、
CUDAによる計算は意味があるのかい？と思った。
MD5とかデータから導き出す数値ならば本当の意味での解析だから意味ありそうだけど。

こんな記事もあるしな

GPUを利用した無線LANのパスワードクラッキングが主流に
http://itpro.nikkeibp.co.jp/article/MAG/20090123/323378/

>>393の言ってるのはまさに本当の意味での解析ってやつだしな。
WPA/WPA2-PSKからの復元は総当り/辞書比較と違ってGPGPUにする意味がある。

Catalystだけど、トリップ検索は随分速くなるみたい
ttp://sourceforge.jp/projects/naniya/wiki/Radeon

ホスト側でconstantメモリを動的に確保できないのでしょうか？
方法がありましたら教えてください。

>>396
目的が良く分からないのでなんとも言えませんが、
どうせ64KBしかないのだから静的に64KB確保しておくのではだめなのでしょうか。

>>396
あるけど、なんでprograming guide読まないの？？

>>397
静的に確保しておいてもよかったんですが，綺麗にかけるならとw
>>398
一度目を通したつもりでしたが，素通りしていたようです。
もう一度読み直してきますm(_ _)m

誰かZIPの暗号解析ツール作って＞＜
それか解析のループ部分だけうｐしてくれたら、カーネルは俺が書きます。

総当りするだけならどっかのオープンソース実装の
パスワードチェックと復号ルーチンとってくればいいだろ

ただでさえ消費電力がデカイのにパスワード解析に何週間も動かしてられっかｗ

超解像のような考え方は昔からあるようだな

フリーソフトImageD2
http://www.tiu.ac.jp/~zohzemi/imgD2/index.html

これは時間軸方向にも参照するものだ。
NVIDIAのMOTIONDSPと同じ考え方だな。
西川善司の大画面☆マニア 第113回 超解像
http://av.watch.impress.co.jp/docs/20090120/dg113.htm

この方法は数フレームを参照することで低解像な映像のブレから情報量をアップさせるようだが、
この方法だとシーンチェンジやめまぐるしく動く映像では逆効果でめちゃくちゃな映像になるのではないか？
（これは上記のフリーソフトの別ページでも注意点として載っていた）
でも東芝などの日本の各社がやってる1フレームだけで行う超解像はそもそも無理がある。だから不自然な画質になったり、情報量が逆に消失したりする。
Lanczosなどでそのままアプコンしたほうがずっと情報量あるし自然な画質だ。比較してみれば一発で分かる。
http://plusd.itmedia.co.jp/lifestyle/articles/0812/24/news031_2.html
ここの元画像を720×480にし、それをAVIUTLなどでLanczosでフルHDにしたもののほうがずっと綺麗。
超解像は単純な処理だから柵とか崩れてるし、文字も駄目だし、元からあった情報を処理によって消しちゃう副作用のほうが強い。
超解像、超解像と目新しくいって盛り上げようとしたいのは分かるが、こんなのはまやかしだよ。

http://www1.axfc.net/uploader/Li/so/24824.zip&key=pass
比較用画像もうｐしておいた

一方、数フレームでやる方法は計算が大変だが、シーンチェンジや盛大な動きの問題さえクリアすればかなり使えそうではある。
MOTIONDSPや↓はそのあたりちゃんとクリアしているのだろうか？
http://www.flashbackj.com/red_giant/instant_hd_advanced/

10年ぐらい前にカーネギメロン大学で勉強してた頃に超解像のプログラム作れっていう課題を出されたことがある

ちょー解像は判ったからマルチすんなや。

>>403
スレ違いじゃない？
GPGPU 向きな処理なのは確かだし、そもそも Cell なり SpursEngine は GPU じゃないし。
画像処理スレとかあったと思うよ？

アメリカのメロン大学は一応名門だけど日本校はどうなんだ？

>>404
なにその高レベルな課題
俺の大学時代と雲泥の差

超解像の原理が未だに理解出来ません
複数の映像フレームの同じポイントの色の出現頻度に一番高い色を適用するってことですか？

は？今の超解像って、時間軸の補完までやってるのか？

MPEGの原理は主にフレーム間差分をとってJPEG圧縮なんで、それを改めてチェックしたところで・・・

>>407
神戸のハーバーランドの近くにあるあれか？
兵庫県の財政圧迫してるらしいよ。
学生少なくて撤退の噂もあります

同じビルの同じ階に神戸電子専門学校の学校法人が経営する大学院がある。
何を勘違いしたのか兵庫県立大学も情報科学の大学院を神戸市内に置いてる。

神戸という都市に何を求めてるのか、理解不能である。

たしかに古くからの工業都市で組み込みソフト屋の数がそれなりにいるのはわかる。
最大の勘違いは、日本のITドカタは職場を離れて大学院に通えるほど裕福ではないことだ。

俺はてっきり、ラーメン大学とか、洗車大学と同じノリで、
神戸電子専門学校が作ったヒト寄せパンダだと思ってた。
ってか別法人なのか?
学生少なくて、って言う前に、派遣のワーキングプア、
デジタル土方に成る為にわざわざ苦労して大学通う馬鹿
が神戸のどこに居るんだよっ?




俺が居るよ…orz

兵庫県もネギメロンなんて誘致するくらいなら神戸大に寄附講座でも作ったほうがよかったんじゃね？
現状兵庫県内で優秀な学生がいるんだし。


おっと、「学生」の枠で括ると近くの中高一貫校のほうがよっぽど・・・

情報系の大学院は多すぎるよな
旧帝でも定員割れする世の中なのに

このうえ西和彦が秋葉に大学院作るんだから馬鹿としか

>>411
誰がMPEGの話をしてるの?
誰も君の知識を披露してくれなんて頼んでないよ?

圧縮で壊れるったって圧縮率によるし、位相が生き残ってりゃ理論的には解像度は上げられる。
圧縮率によっては壊れた分が回復するだけかも知れないけどな。

誰がMPEGの話をしてるの?
誰も君の知識を披露してくれなんて頼んでないよ?

それを言うなら、なんでいきなり超解像の話になったんだろ？

メロン日本校のカリキュラムを見たけど、ウンコすぎて話にならねぇｗ
こりゃ学生もこねーわ

>>416-417
だれも連投してくれなんて頼んでないよ？

GPGPU向けの画像・動画関連のソフトってどれも速さばかり求めて品質は二の次だな
ウンコなエンコーダだとRGB/YUV変換で腐る
一番腐るのは量子化だろうけどな

Geforceは何気にGTXシリーズで初めて64bitに対応してるけど
演算装置は１個しかないｗ
並列処理が得意な分野でfloatだけで済むようなものはほぼ無いだろｗ

>>420
連投じゃなくて別人
2ch歴長いんだからそれくらい分かれ

甘いな。floatでも速ければ使い物になる用途は結構あるもんだよ。
例えば、近似計算なんかはfloatで近似させてからdoubleで更に近づけることもできるしね。

そういや西和彦もこっちの出身だな。確か甲陽から和田大
だから、落ちこぼれか。そもそも博士余りで高学歴ワープアPD
が社会問題に成ってる今日この頃に、大学院新設とかもうね
アボガドバナナ…。
鴨葱メロンと言えば、金出教授もこっちの出身だったな。
沢山カメラ使って超解像みたいな論文も有ったような無かったような。

>>425
お前、自分の考えまとめるの下手だな。

まあ、CUDAは英文のドキュメントが読めてある程度知能がないと
使えないからな。

スレ違いの話で盛り上がるのも分かるな。バカばっかりだしw

CUDA版のTripcode Explorerができたみたいですね。最適化に期待します。
http://tripper.kousaku.in/
http://download.kousaku.in/trip/Tripcode-Explorer-CUDA-test1.zip

ようこそ、バーボンハウスへ。
このテキーラはサービスだから、まず飲んで落ち着いて欲しい。

うん、「絶対に動かない」んだ。済まない。
仏の顔もって言うしね、謝って許してもらおうとも思っていない。

でも、このネタプログラムを見たとき、君は、きっと言葉では言い表せない
「ときめき」みたいなものを感じてくれたと思う。
殺伐とした世の中で、そういう気持ちを忘れないで欲しい、そう思って
５分ででっちあげたんだ。

じゃあ、注文を聞こうか。





--------------------
1M超のバイナリファイルに何が詰まってるか疑問な人は、テキストエディタで開いてみればいいよ＞＜

CUDAはコアレスと分岐の扱いを把握すれば、やりたいことは大体クリア出来ると思われ。

結局プレディケートなんだよね。
SMあたりの命令デコーダは1基だからSP毎に別々のフローを実行することができない。

Larrabee（Ct）なら分岐は容易に表現できる。

>>431
プレディケードって？？

えーとね、たとえば
if (cond) {
　　funcA();
} else {
　　funcB();
}
なんてコードがあるとしよう。
普通のCPU向けのCだと、 cond の条件にしたがって、funcA()を実行するブロック、
あるいはfuncB()を実行するブロックにジャンプする。
すなわち命令ポインタを操作してコードを飛ぶ。

CUDAにおいてはシェーダマルチプロセッサ一つにたいし、命令デコーダは1つしかない。
にもかかわらず、condは要素ごとに変わるわけで、条件分岐先はSPごとにバラバラになる可能性がある。
んで、そこで使うのがプレディケートなわけだけど、簡単にいえば、ifとelseの両方を通るようにする。
funcAとfuncBをインライン展開して、条件ビットで選択的に実行するコードに展開する。
んで、各要素に対して、実行するか実行しない（あるいは実行結果を反映しない）かを選択的に行うわけだ。
並列度の高いプロセッサではよく使われる方法だ。

んで、こっからはこのアプローチの弱点。
問題はif-elseを何重にも組み合わせたり、switch文を多用する場合、総当たりにかかる計算時間量が
並列化によるパフォーマンスメリットを相殺し、逆に遅くなるケースもある。
並列処理を諦めて素直に要素ごとに逐次処理をさせてくれたほうがかえって効率がいいかもしれない。

しかしCUDAってそのへんの融通がきかないんだよね。基本的に【並列処理しか記述できない】から。
正確には逐次処理は専用のプリミティブなんかを使って限定的に逐次処理はやれるけど記述面では
かえって面倒になる。
GPUで不得意な処理はCPUでやれってアプローチだからそのへんの融通を利かせる気は無いらしい。

CUDAで困るのはその点のほかに、並列数を途中で変えられないこともあるよね。
一度ホストに処理を戻すと遅くなりかねないし、共有メモリが失われてしまうし。
私の関わっているプロジェクトでは演算処理が中心なので、ある程度融通が利いてくれないとね。

前にも有ったけど、条件分岐したら負け。
Crayだってそうだったじゃん。

CPUだとforループが多重になる部分をGPUに
丸投げすればいいんでしょ？
3項演算子程度は実行して欲しいけど

> forループが多重になる部分

重要なのは回数であって多重かどうかはあまり関係ないです。
32の倍数であることは有る程度重要かな

一番重要なのは依存関係がないこと。
たとえばループを逆順で実行しても結果が等価であったりとかね。

240 SP（30 SM）を使えるとすれば、フルに使うには、最低960程度の演算が並列実行可能である必要があります。
ただし、全部のSMでまったく同じ処理をする必要はないです。

>一番重要なのは依存関係がないこと。
そうだね。そのためiCotの時も関数型言語の並列化
に七転八倒してたね。
今ならerlangで良いんじゃね? CUDAスレでこんなこと
言うのも何だけど、Cライクな手続き型言語だとどうして
もすぐに壁にぶつかってしまって、スケーラビリティが
出せない。もしくは出そうとするとプログラマの負担が
重過ぎる。
個人的には今更lispやprologに復活されたくはないけど。

>>438
そう、わかったよ
じゃあ俺様がCUDA用erlang処理系書いてやる

ループを並列処理に展開するのって自動化出来そうだけど

>>424
おいおい調子に乗って嘘つくな。
「簡単にできるぜ！」っとか鼻高々なのはいいけどそんなのないよ。

>>441
いや数値計算なら反復法とか1次連立方程式の陰解法で使えるだろ

また自信満々な人の嘘つき合戦ですか？

なんだよ微分方程式、というか積分使うのか。
積分を近似計算といっていいのか？
実用内ではあると思うけど、それもfloat(7)からdouble(16)だろ。何回ループするつもりなんだよ。

積分を近似計算と言ってはいけない理由がわからん。

数値計算なら兎も角、シミュレーション関係なら大いに有り得る話だな。
自分が知っていることが全てではないと認めることができれば世界は広がるのに。

他人を否定することでしか自分を正当化できない、ということか。

>>447
数値計算と近似計算の違いを教えてくれませんか？

>>448
1+1=2となるのは、数値計算ではあるが近似計算ではない。

CUDA 2.1 Notebook Drivers for Windows
ttp://forums.nvidia.com/index.php?showtopic=91157

βだけども

>>446
とすると、貴方の世界ではシミュレーションとは近似計算をしてるってことですか？

物理にしろ確率モデルにしろコンピュータシミュレーションは近似計算だろ

整数演算は近似計算ではありません。

approximationの意味分かってるの

なんかめんどくさそうな人がいっぱいいますね。
どうでもいいですが、早いところストリーム用のプログラミング手法を確立して、ストリームに特有な技法を紹介する本を出してくださいな。

どうして煽りを入れる人ほど知識が足りないんだろう？

>>456
知識が足りないから煽るじゃないですかね？
そういう低脳な人は「煽ることしか出来ないよね」っとも言いますけど。
そんなことどうでもいいんで「ストリームんグ・プログラミング技法」とかいうブログを早く作ってくださいな。

「approximationの意味分かってるの」とかめんどくさそうな人との議論とかいかめしい顔した人が言う哲学に興味はないんで。

↑
低脳な人の煽りの見本

458 名前： デフォルトの名無しさん [sage] 投稿日： 2009/03/06(金) 17:14:51
↑
低脳な人の煽りの見本

プログラム板なんだから少しはプログラム出せって

↑
低脳な人の煽りの見本

↑
低脳な人の煽りの見本
↓

>>448
1+1=2となるのは、数値計算ではあるが近似計算ではない。

VIPのまとめWikiに書いたこともあるが、今はネタがない。

あーそうそう、x16バスに繋がった8800GTの方がx8バスに繋がったGTX280よりも速かったってこと位か。
但し、転送量が多目の用途だからだとは思うが。

VIPPERプログラミングスレの派生なのかここ？

大丈夫、私はvipには書いていないw
でも何故かまとめWikiには複数投稿している罠。

自分が知らない事はWeb見て知ったかぶらないで馬鹿げたレスする暇で
amazonで本の一冊でも買えばいいのに。

ところでCUDAで性能だすためのまとまった日本語の文書ないかな？

そもそもCUDAに関して有用な日本語資料がなくね？
公式でさえ日本語マニュアルはあんなだったし。

大丈夫、英語資料すらろくなのないから。

やはり本家のドキュメントにあたるしかないのか。
めんどくせー。環境の開発もいいけどドキュメントの整備も力入れてほしいわ。

逆に日本語ドキュメントがあっても大して意味無いよ。
IntelのプログラミングマニュアルなんていまだにPentium 4のことしか書いてないぞ。
日本法人仕事しなさすぎる。

CUDAを勉強するより前に英語アレルギーを克服したほうが何かと良くなるかも。

英語アレルギーってなに？

そう言いたくなるくらい、英語から目を背ける人は世の中に意外と多い。

俺の場合英語と日本語だと読むスピードが１０倍〜１００倍違うorz.

母国語でないと読むスピードが遅いだけじゃなく小さなとこで思い
違いがでてきて結局後からまた参照したりして嫌だ。

技術的文書に機械翻訳はどの程度通用するんだろ。奇想天外な訳になってしまうのかな。連投ｽﾏｿ

英文学読めって言ってるんじゃないんだし、書ける・話せるも別問題。
技術ドキュメントの英語なんて、有る程度形式ばった言い回ししかやらないので
単語を摘み出すだけでも回数を重ねればそこそこ意味はわかるようになると思う。

慣れてくれば技術系ニュースサイトとかも読んでみたり。

だんごさんはどーゆーサイトみてますか
スレ違い申し訳ない

Intelの開発者ブログとかRSSに入れてる

団子はGPGPU嫌いなんじゃなかったの？

逆に、好きな奴いるのか？
非生産的で変態だけど性能のために仕方なく使う類のモノだろ

おれもGPGPUなんて嫌いだな。
開発したことあるけどPSシリーズも大嫌い。

x86でいうMMX/SSEって、分岐が除去できるとか直列方向のパフォーマンスメリットがあった。
GPGPUって並列方向のスループットありきで、ホスト側のコードでお膳立てしてやらないといけない。

>非生産的で変態だけど
それってintelアーキのことじゃん。昔からずっと言われ続けてることだが。
mc68kやsparc,mipsの方がよっぽど素直に書ける。

けど市場規模のために仕方無く使わされてる。

アセンブラはmc68とx86しかやったことないけど、
mc68はかきやすかったな〜。

欲を言えば16本すべて汎用レジスタだったらよかったんだけどｗ

>>484
あー？
メモリアドレッシングモードが貧弱すぎるんだけどー？
まじうけるー？
ﾊﾟﾈｪっすよ

インテルに慣れきってるとそう思うかもね。
どうせ団子はインテル一筋なんだろ？w

> mc68とx86
その時代だと8086だろ。「x86」って基本的に32ビット以降のことを言うと思うんだけど。
32ビットだとぜんぜん自由度が違うっしょ。セグメントなんて使わなくていいし。
4GBの論理メモリ空間をリニアにアドレッシングできるし。
んで、案の定ローエンドサーバだけにとどまらずHPCもx86に惨敗して虫の息じゃないか
古くからあるRISCなんて。
MIPSも組み込みに逃げたけどARMに食われたね。

それはともかくSSE・MMXも経験ない男の人がCUDAなんて・・・




さて、CUDAの話なんだけど、基本的に最小の演算単位は32ビット×32のSIMDで
メモリロード・ストアも、各要素ごとに計算してscattering/gathering機構付きの
ロード・ストアユニットで、
このへんはCUDAのアーキテクチャマニュアルにも載ってる通り。

従来SIMDって基本的に連続的に並べないと性能出ないけど、
CUDAは動的にベクトルを再構成することで、一気に柔軟性が向上した。

逆にこの強力なロード・ストアユニットを載せたせいで、連続したデータに対する
ロードストアの効率が悪くなってね。
一時変数をどっかに置いとこうとした場合にも、32要素ごとにバラバラにアドレスを計算する
scattering/gathering機構つきのロード・ストアユニットに通す羽目になる。
これじゃエネルギー効率的にもよくないでしょ。

んで、レジスタにそのまま保持すればいいじゃないってことで、それで
1つのシェーダコアあたりのレジスタファイルが、32KBとか64KBみたいな巨大なことになってる。
それにしても一般のCPUのL1キャッシュよりレイテンシの大きいレジスタファイルって一体・・・

団子の脳みそがx86のアーキテクチャで凝り固まってて、現代風のプログラミングパラダイムについて来れないってことだろ。
あと１０年もすればおまえの持ってる小手先業などは博物館の展示資料でしかないし、おまえの能書きなど頑固オヤジの戯言同じなるだろう。
インテルのブログで洗脳されまくっちゃうのもいいけど、アーキテクチャマニュアル云々よりも団子が頭の切り替えをできるかどうかのほうが問題なんじゃないの？

現代風のプログラミングパラダイムって何だ？

斜め上をいく愚言に感謝する。

しかしながらscattering/gatheringによる柔軟なアクセスはSIMDの新時代を切り拓くものだ。
実際Intelも2〜3年先のSIMD拡張では256ビット、512ビットと幅が広くなってるため、
AoS/SoAの変換をいかに効率よくこなすかがテーマになってくる。
（ちなみにLarrabeeにはscatter/gather命令そのものを導入する）

このへんはむしろCISC的なプロセッサの美学だと思うがね。
AltiVecとかCellのSPEでなら何十命令かかる命令を1命令でこなす。
1クロックサイクルスループットでこなせない命令を実装しないのがRISCだろ。
モダンなCPUではパイプラインの前半部分のほうがALU自体よりもコストがかかるしまってるから
それで処理単位がリッチなCISCのほうが効率がよくなってるわけさ。
このへんは　http://gimpo.2ch.net/test/read.cgi/i4004/1220728356/76あたりと同意見

しかしさ、16要素とか32要素とか、全部バラバラのアドレスだとしてみ？
とてもワーストケースで要素数分だけメモリアクセスが必要だぜ。
RISCの守備範囲じゃねーよ

んで、個人的にCUDAの問題は、scatter/gatherスカラ命令を備えないことなんだよね。
常に32並列単位で演算しないといけない。それで小回りがきかない。
スカラレジスタでアドレス指定するベクトル単位のロード・ストアと
scatter/gather
Larrabeeあたりがまさにこれをやってるわけだが。

> あと１０年もすればおまえの持ってる小手先業などは博物館の展示資料でしかないし、おまえの能書きなど頑固オヤジの戯言同じなるだろう。

残念だが俺は流行りものの言語・フレームワークには目がない。
Ruby On Railsとか大好きだし。むしろ高級言語をより効率的に使うためにマシン語レベルで理解する必要があるんだよ。
たとえばさ、LLって性能的にはネイティブマシン語より遅いから、LL向けのJITコンパイラ書きたいとするじゃん。
どうしてもアセンブラの知識は必要なんだよね。もちろん業務じゃないよ。
ということでプロ高級言語er、趣味マシン語er
それでARM語もx86語もそれなりにたしなんでおきたいわけ。

○んで、個人的にCUDAの問題は、スカラ命令を備えないことなんだよね。

頑固オヤジの戯言ごとと同じになるだろう。

ニート相手に５行も書くの面倒だから誤字脱字なおすのも面倒だよな。
「CISC的」とかいう概念がもう古いパラダイムってこと。
おまえみたいな純粋な「消費者」の戯言などどうでもいいけど、ストリームなのに128/256bits単位とか全く鼻糞だろ。
ストリーム演算やってるのに、「スカラ演算もやりたい！」「アドレッシング！」という考え自体を改めたほうがいいと思うけどね。

どうでもいいけど、ストリーミング・プログラミングの小技を集めたブログをはよ作ってよ。
C#だとスニペットというんだったか？そういうイディオム集みたいのでもいいから。

Intelの中の人のブログって言っても、本当に自社製品のプログラミングがらみの話題って
月に1回出るかどうかのレベルだぜ
次期Windowsの話題だったり、XMLやLLなんかのWebまわりの技術がどうこうだったり。
中の人の興味のあることが書いてあるって感じだけど、頭の悪い技術系ゴシップサイト
よりはよっぽど為になる。さすが半導体総合メーカーだわって思うわ。

NVIDIAのニュースも購読してたけど本当に自社製品向けのコンピュータグラフィックスのノウハウとか
グラフィックよりの物理演算が中心で、そっち方面はそんなに深入りする気はないので読む価値なしと。
（そっち方面で食ってる人ごめんなさいね）

>>493
> どうでもいいけど、ストリーミング・プログラミングの小技を集めたブログをはよ作ってよ。
> C#だとスニペットというんだったか？そういうイディオム集みたいのでもいいから。
ﾌﾟﾌﾟﾌﾟﾌﾟﾌﾟ

ソースコード例文をいんたーねっつで検索してきてコピペをするのが
プログラミングだと思ってる人はそう言うのに本質を求めるよね。
いや、いいんだけどね。
俺とて業務では最高級の言語から低級言語で書かれたライブラリを使わせてもらってる立場だし。

ちなみにマルチコアとかSIMDを使いこなして最適化コード書いたりできる人間は稀少性があるから
長い目でみれば食いっぱぐれしないよ。

今でこそ団塊COBOLerの後釜需要があったりするくらいだし
（徐々にJavaや.NETに置き換わってるので将来性を考えれば微妙だが）

自動並列化ランタイム環境使えばいいとか言うだろ？
そう言う考えの三流プログラマは食いっぱぐれる。間違いなく。
じゃあその並列化ランタイムは誰が書くんだと。書きもしないのに沸いてくるのかと。
最近流行のJavaScriptのJIT部分のコードでも見てみればいい。各CPU用のバイトコードの山だ。

その点、覚えさせれば小学生でも出来るような、コードをコピペして貼り合わせる能力なんて誰が評価するんだよ。
知識が無いと難しい作業こそ高い市場価値がある。

CUDAはまだ市場として育ってないがな。とがってる分、苦手なことが多すぎて。

俺の団子が火を吹くぜ！

っていうか、電子の移動度の限界とか云々でクロックが上がらないのでフリーランチ終焉、
SIMDやマルチコアを明示的に使いこなさないと性能出ませんよ
これ以上1スレッドの負荷の重たいソフト書くなよ、なんて、何年も前から言われてることなのに
「価値がなくなる」だとか何を妄言はいてるんだか。
10年後に100GHｚとか200GHzとかいくのかよ。
数十コアとか数百コアになって最適化屋の需要拡大することはあっても、縮小することなんてねーよ

要するにSIMD・マルチコア使いは10年先もナウい。ﾊﾟﾈｪ

最適化できる奴は別にたくさん要らないよなぁ・・・
結局ライブラリ作って終わりだし。
そういうライブラリがなかったり高額だったら、誰も使わないからあまり流行らないわけで、どんどん忘れ去れていく技術なだけだしなぁ・・・
GPUとは関係ないけど、MSの提唱してる技術とかかなり不発が多くて流行らずに忘れ去れてるの多いでしょ。
（スカラの）マルチコアとライバル関係だけど、運が悪いとGPU（ストリーム）の方が流行らずに終わってしまうことだってある。PCってのはそういう世界だったよな。
どうでもいいけど人柱がんばってよ

せいぜいコードコピペで済む単発案件こなしてなよ
希少価値のある技術には見えないがね。

どっちかというとコピペプログラミングこそ自動化できそうだけどなぁ
お絵かきツールだけでプログラムのフロー書くASTERIAみたいなツールも出てきてるし

既にゲーム業界では下っ端レベルからそういう技術が要求されるようになってるけどね
PS3とか360やってるところなら半ば強制だぜ
脳天気でいられるのは高級言語屋とローエンド組み込みCPUソフト技術者くらい

CUDAは流石に今のポジション以上の普及はないと思うよ
「汎用」ってものをわかってない。

GPGPUの【GP】に関してならLarrabeeに食われるだろうね。
たとえば普通のCを使うとして、たとえばtime.hすら使えないのがCellのSPEなら
CUDAはそれ以前の問題だし

スニペットとかコピペってのは、結局コードのモジュール化ってことでしょ。
オブジェクト指向による再利用促進とも言うけど、それは時代の流れって言うよりもう当たり前じゃないのか？
IDEとか便利だし、かゆいところは自分でコード書けばいいんじゃないか。
今の時代、３０分で作れるのに一からメモ帳作る奴はよっぽどバカでしょ。

ああ抜けてた。

コピペって簡単に言うけど、典型コードの再利用なわけでだからこそメモ帳アプリが３０分で作れる威力があるんだけど。

そういえば、ム板でコテ名乗ってるのは団子ぐらいしかいないよね？他にいるの？

>>504
コピペの単純工程をやるプログラマもいれば
ライブラリを書くプログラマもいるわけで

法律事務所のアルバイトと弁護士くらいの格差は出てくるかもね
いや、既に出来てるか

>>505
使い回しでメモ帳に30分ってかかりすぎだろ。3分でやれよ。
テキストコントロール配置してファイル読み書き機能付けるだけで終わりだろ
IDEの雛形だけでほぼ完成なんだからさ

それともGREP機能でも搭載するのか？

30秒だろ

#include <stdlib.h>
int main(void) { system("notepad.exe"); return 0; }


再発明する価値もない。

無いものを作る、あるいは既にあるものをより良くすることに知的労働の価値があるわけで
劣化コピーの再発明で金とるなど馬鹿の所業だろ。

>>509
ワロタｗ

30分で作れる程度のエディタなんて誰も使いたくないな

なんでおまえらはそのうちいい情報を提供してくれそうな人を叩くんだよ

いい情報を提供するのが自分じゃないと気がすまないからさ。
そのために全体が遅延しても問題なし。

CUDAは既存の一握りのプログラムの再発明のためデバイス・言語処理系だろ。
性能はともかく効率CUDAでできることは普通のCPUでもできる。
より高いスループットを得るためにこそある。
プログラミング対象を選ぶし、性能を出すには工夫がいる。

テキストエディタの話じゃないけど、生産性を言い訳にして自分で創意工夫が出来ない奴には不向き。

,,・´∀｀・,,）っ-○◎●　に嫉妬してるだけじゃね？

まぁ、団子は必ずしも間違ってはいないからな。
CUDAに未来はないかもしれないけれど、OpenCLはAMDも担いでいるからもう少し生き延びるだろうし。

OpenCL（笑）

なんかの魔法の言語のように思ってないか？
OpenCLは「GPU版Java」じゃない。
共通化されてるのは言語の基本仕様の部分だけで、細かいところは処理系依存。

んでもって、CUDAやCAL/Brook+のプログラミングの敷居を高くしてるのは言語処理系じゃなくて
少ないスクラッチパッドメモリとレイテンシの大きいメモリと
やたら小回りが利かないベクタ演算ユニット、その他諸々のGPUのパイプライン・・・
要するにシェーダコアの構成そのものにあるのであって、それが解消されない限り
CPUを置き換えて普及していくことなどあり得ない。

普通のCPUと同じ定番言語のC/C++言語をまがりなりにもサポートしてるのに
業界の評価のお寒いCellを見れば、課題は言語じゃなくて汎用プロセッサとしての
柔軟性にあることくらいわかるだろ？

その意味、OpenCLを効率良く実行できるのはよりCPUに近いLarrabeeだと思うよ。
というか本質的にOpenCLなんて要らない。
どうせCellなんかと同じくハード専用にカリカリにチューニングしなきゃいけないんだし。

>>518
世の中それほどぎりぎりのチューニングまではしないけどちょっとは速く走って欲しいなんて用途が結構あるのよ。
で、私自身はOpenCLはAMDが必死こいてアピールしているだけで実際には普及しないと思っているのよね。
どうせLarrabee出て来る頃にはCtも来ているだろうから、NVIDIAもAMDも青息吐息でしょ。

まぁ、CUDAスレなんだからLarrabeeの待つ未来を語るのは程々にしましょ。

期待してなんか無いよ。
Cellと同じくニッチ市場を食い合うだけ。

ゲーム屋の意見としては、SPUの数とメモリが倍あったらCellも悪くないと思う。
あとはメモリのバンド幅か。
柔軟性もあったら嬉しいけどね（整数や分岐とか）。

>期待してなんか無いよ。
おっと、だんごさんの悪口はそこまでだ

>やたら小回りが利かないベクタ演算ユニット、その他諸々のGPUのパイプライン・・・
Crayだってそうだったじゃん。Personal CrayとしてCUDAは良く出来てると
思うけど。
メモリの不自由な階層は何とかしてくれ、と思うけど。Cray同様、IPも持って
一般I/Oも出来て欲しい。
あと出張先でデモ出来るように、CUDAの動くnVidia GPU載ったサブノート
が出てくれないと…。学会発表しようにも、デスクトップ担いで持参しなきゃ
ならんってのは勘弁。

つ[Asus N10]

つ[新Mac Book]

いや、でも、アカデミック畑の人の求める特化型プロセッサって一般のニーズとかけ離れてると思うよ。
CellやGRAPE-DRでワードやエクセルが動くかっつーの。

当たり前だけどアカデミック色の薄いアプリケーションって書く人少ないのよね。
サンプル探しにCUDA-Zone逝っても「なんとか論文ps.gz」みたいなのしかないし

ここにアカデミック色の殆どないアプリケーションを書いている人が居るんだが、
残念なことに特定用途向けだし契約の都合もあるんで公開できないんだわさ。

アカデミック色って例えば何？
ブラックホールのシミュレーションとか？

俺も書いてたよ
NVIDIAの営業さんじきじきに頼まれたがめんどくさくなった

>>528
そういえばGRAPEのコミュニティではCUDAはやたら受けが良いらしいね。

>>523
モバイルCUDA環境が欲しくてN10jc買った
性能は
./nbody -benchmarkで16.472GFLOP/s
./nbodyでタイトルバーにでるやつだと80GFLOP/sくらい

割と出るんだね

大学時代にやった熱力学シミュレーションのレポートをまた引っ張り出してきてCUDAで実装してみるかな。
Rubyで書いたらアホみたいに遅くてC++で書き直した覚えがある。

Rubyで書いてCより性能でればいいのにね。無理言うなって感じだが

アルゴリズムが悪いんじゃないの。

まさに「グリッド」（格子点）だよ。
アホみたいに並列化しないと性能出ないCUDAには向いた問題

RubyはCでかかれたインタプリタであって、
1語句ごとにループ・switch文で処理を行う以上
それ自体の致命的な遅さはどうしようもない。

YARVとかJRubyなら多少速いかも知れんが
本家はまだJIT以前の問題だし。
Matz氏はXbyak見て「いずれは考えなきゃいけない」的なこと言ってたんだけどね。

団子の中の人って、大学逝ってたんだ。

>>531
意外とやるな。電池で動いてそれなら上出来だと思う。
ARM+DSPでは桁違いに負けてると思う。しかし、所詮
ネトブクに毛が生えただけなのに、ThinkPad Xシリーズ
より重いのか。
Linux対応はどない? EeePCのLinux対応はすこぶる良
かったから期待してるのだが。

>>537
CentOS5.2はおk
サウンドは自分でドライバ当てる必要あり
無線LANは認識してる
（ドライバ入れてないから使えるかどうかは不明
あとはカメラと指紋認証が使えないくらい

他の鳥は試してないからわからん

BIOSでHT切れないのが気持ち悪い

>>537
金を気にしないならネットブックは辞めたほうがいい。
1024x600は割と不便。

EeeUbuntuなら、最初からEeePC向けのカメラやBluetoothの設定ユーティリティが
インストール済みだが。

>Matz氏はXbyak見て「いずれは考えなきゃいけない」的なこと言ってたんだけどね。

いつ？

>YARVとかJRubyなら多少速いかも知れんが
>本家はまだJIT以前の問題だし。

YARVはすでにRuby本家だけど？

>>540
http://www.rubyist.net/~matz/20070116.html#p04

参考になるかもしれない、じゃん

>>541
リンク先読んだが、Xbyakじゃなくて「Gecko 3.0にはJIT付きJavaScriptエンジンが添付されるということだが」が、将来の参考になるという風にしか読めないんだが・・・
2007年の時点なら、Matz氏がRuby用のJITについて参考にするという文脈なら、XbyakじゃなくてYARVのJITが暗黙でしょ。

別に"へるみエンジン"を検討してるなんて言ってないが
「JIT」としか言ってねーよ

JITじゃなくて、「「いずれは考えなきゃいけない」的」と「参考になるかもしれない」は違うだろって話でしょ？
あと、
>YARVはすでにRuby本家だけど？
についてはノーコメントのなの？

YARVはJIT実装があったろ？
あれこそ亜流だけど

>>545
ちなみにYARVとか鬼車のJITは環境非依存の中間コードに変換するだけであって
CPUネイティブじゃないよ。
んで更にそのバイトコードをインタプリタで動かしてる。

ネイティブコードのJITに言及したのは↓だけ

> _ [言語] IA32(x86)JITアセンブラ Xbyak

ま、Rubyが動かせそうなGPUはLarrabeeが最初で最後だろうな

いつJITの実装の話になったんだ。
話そらすのが上手いなｗｗ

そもそもMatz氏はXbyakについて「「いずれは考えなきゃいけない」的」な事は言ってないので（参考にするのはGecko 3.0の方）、>>541以降のお団子さんのコメントは見当違い。

Xbyakを採用するなんて俺は言ってないし君が勝手に勘違いしただけでしょ

もともとはRubyがC++よりクソ遅いって当たり前の話だろ。
スクリプト言語が静的コンパイル言語を超えられる訳がない
それだけのことよ

団子、いい加減にしろ。最近のお前はオカシイぞ。
形式言語より、日本語勉強し直せ。マジで。
コミュ力無さ過ぎ。

自分が思考短絡してるのを棚に上げて他人を避難するｳﾞｧｶがいると聞いて

自己紹介、乙。
そんなヴァカ呼んでないから、「避難」してこい。

テンプレ入りか
> Matz氏はXbyak見て「いずれは考えなきゃいけない」的なこと言ってたんだけどね。

コテ団子の相手はするな。キチガイになっちまうぞ！

>>553
自分の技術力をいくら上げても、無責任な発言ばかりしていると誰も君のことを信用しなくなるよ。気をつけたほうがいいと思う。

「JIT」について話してるのに
一番近くにある単語「Xbyak」を「検討」ということにしたがる思考短絡ぶりがゆとり脳

Matz氏はXbyak見て（JITの仕組みを）「いずれは考えなきゃいけない」的なこと言ってたんだけどね。

これでいいかな？

温度分布の立体グラフをExcelでプロットしたいんだが、なんかいい方法ある？
俺もゆとりだからCSVで吐き出して読み出すとか原始的な方法しか思いつかない

隔離スレなのか、ここはw

>>560
Excelなんかを使いたいなら、csvでいいんでない?

つーか、団子もそれに噛み付く奴も自分の言葉が足りてないことに気付けよ。

ここはグダスレじゃないぽ

くだをまくスレです

どのスレでもゆとり脳の団子が来ると荒れる。
そして人がいなくなる。

まだゆとりがどうのこうの言ってる時代錯誤な奴がいるのか

おまえはヒマになると２ｃｈ開いてるだろ？ｗ

お前は○○だろ

↑↑自分がそうだから他人も同じだと思っている奴の決まり文句

○○な>>568

Vista x64
Device 0: "GeForce 9600M GT"
4096 bodies, total time for 100 iterations: 663.110 ms
= 2.530 billion interactions per second
= 50.602 GFLOP/s at 20 flops per interaction

>>567
暇じゃなくても開いてるわボケ

忙しいときほど2ch開いちゃう、ふしぎっ

■後藤弘茂のWeekly海外ニュース■
KhronosがGDCでGPUやCell B.E.をサポートするOpenCLのデモを公開
http://pc.watch.impress.co.jp/docs/2009/0330/kaigai497.htm

素人質問で恐縮ですが……
Tesla C870を手に入れたのでCUDAで画像処理をしようとしているのですが、
CUDAでテクスチャフィルタリングユニットの機能を使うにはどうすればいいですか？
○○の○ページを嫁！で構いませんので、教えて下さい。

tex.filterMode = cudaFilterModePoint;

＞Tesla C870を手に入れたのでCUDAで画像処理をしようとしているのですが、
あー、8800GTXからアナログ回路を減らしてメモリを増やした、最早今となっては1万円ちょっとで買える
8800GTと数割程度しか能力の変わらない癖に値段は10倍以上と言う代物ですね。

テクスチャ関係は私はやってないからお役に立てませんがw

MV探すのに16x16のSADをCUDAで計算してるんだけど、なんでこんなに遅いんですか？

組み方が悪いんでしょ。

SADするのに、組み方どうこうとかあるんですか？
テクスチャ使ってるのに、なんかキャッシュミス多い感じだし。。。

>>575
ありがとうございます。
cudaFilterModePointでググったら、それらしいものが見つかりました。
ttp://forum.nvidia.co.jp/EokpControl?&tid=10733&event=QE0004
これから勉強します。

>>579
コードも晒さず、自分の無知を曝け出し、文句だけ言うなんて、馬鹿なの?

なんでこう沸点低いの？馬鹿なの？

ION採用ミニデスクトップAcer AspireRevo、オンライン予約開始
http://japanese.engadget.com/2009/04/14/ion-acer-aspirerevo/

Mac用の2.1ってツールちゃんと入ってる？

誰かN10JでCUDA使ってる人いる？
N10Jにtool kitインスコしようとすると失敗するんだけど。。。

今、ブロック数を増やして並列度をあげてみるといったことを
作った行列の積の計算にあててみようと思ったんだが

http://tech.ckme.co.jp/cuda_diff.shtml
に書いてるブロックを複数使った場合の問題は、カーネル内でブロック間の同期を
とる方法が存在しない点である。そのため、下記のプログラムでは、1回計算するたびに、
カーネルを終了し同期をとっている。

というのは1回毎の計算をホストにコピーしてやりたい回数分ループさせるというので
いいのかな？

>>586
いちいちホストにデータ転送してたら時間もったいないでしょ？

つか参考にしてるページ見たけど、かなり酷いコードなんだが。。。
>>586が何をしたいかが具体的に判らないから、アドバイスしづらい。

１ブロックの最大スレッド数を使った計算じゃ、GPUの処理速度がCPUに対して上回らなかったので
ブロック数を増やして計算しようと思ったんですが、１ブロック制限に到達した時、どうやって次のブロックに
移動すんのかが、記述の仕方がかなりよくわからないんです。

dim3 grid(16, 1, 1);
dim3 threads( 512, 1, 1);

testKernel<<< grid, threads, mem_size*2+sizeof( float)*2 >>>( d_idata, d_odata);

カーネルのほうの計算にこの値を元に何か記述すればいいとはわかってるんですが・・・
何か参考になるとこありませんか？

>>588
大いに勘違いしている希ガス。
先ず第一に、>586のサイトは参考にならない。
第二に、スレッド数は必ずしも多いほど速いと言う訳ではないし、共有メモリは使わないで済むなら使わない方がいい。
第三に、行列の積の計算なら、NVIDIAのプログラミングガイドにそれなりのサンプルがある。

>>589
レスサンクス、ガイドとＳＤＫもう一回見てきます

CUDAスレって何でこんなに勢いが弱いの？

ぶっちゃけ2年後位には廃れてると思うからやる気がしない
日本語資料少ないし

.netでもやってる方がつぶしがきく

そうか、ＧＰＧＰＵだと他にまともな環境はないだろ

GPUにこだわる意味がないっていう

みんなcellで思い知っただろ？
そういうことだ。

.NETかGPGPUか選べる立場なら前者でいいんでない？
宗教上の理由でGPUの中でしか選択できない人がいるのももちろん知っております

強いて言えばOpenCLか？
http://www.nvidia.com/object/io_1240224603372.html

ドトネトなんてLinuxで動かないじゃん。
*BSDでも動かない。糞。

とにかくGCCで動くようにしろよ。話はそれからだ。

.NETはmonaで動くだろ

ﾓﾅｰ

.NETはMONOで動くが、GCCで.NETアプリってコンパイルできたっけ？

CUDAかC#かって、ベクトルが全然別ですがな

>>601
Mono入れたらmcsってコンパイラが使えるようになるはずだが。。。
貴殿はGCCに入ってないという理由でPerlやPHPをも嫌うのですか？

問題はなぜこのスレは勢いがないのかってことだ

CUDAとOpenCLの認識の仕方として、
抽象レイヤ的にこんな感じかな？？

APP
--------
C/C++
--------
OpenCL
--------
CUDA
--------
driver and runtime

>>603
見た目簡単そうに見えて実は使いづらくて、本質はCellよりも更に応用分野は厳しいからね。

「CPUの数十倍とか言ってたけど全然遅いじゃん！」で、使い方を理解しないままみんな匙j投げる
いや、使い方がわかったところで、その正しい使い方が、本質的に目的のアプリケーション向きじゃなかったり。

そうそう、その演算だけに絞れば確かに速いんだけど、アプリケーション全体で見るとXeonに勝てなかったりね。
ボードメーカ側も自覚しているらしく、私の客先でのCUDA開発は2チップGPUボード4枚挿しするところまでいってしまっているし。

どうせララビーも期待外れに終るさ

nv社員乙ｗ

Larrabeeは、たかがx86、されどx86だな

Atomに毛が生えたような小規模なx86コアが数十コアあったら何が出来る？
汎用プロセッサとしては程度が知れてる分、逆に落胆しようがない。
良くも悪くも身の丈以上の期待はされてないからな。

流れをぶった切るが
GeforceはCELLより変態的な構造って認識でおk？

餅は餅屋

色々調べて見たけど結局CUDAのsuper piはまだ出てないんだな
CPUとGPUの比較が出来ると思ったのに

むしろスーチーパイがもっとリアルにぬるぬる動けば…

＞スッチーのπ
まで読んだ





とりあえずPTXの自己コンパイルは最低限だろ

どっかの営業さんが言うにさ
「たとえCPUより速くなくとも、CPUでやってる仕事を肩代わりしてやることが
出来るだけでも使う価値があるんじゃないでしょうか」

いや、それのお膳立てのためにCPU時間食うから本末転倒なのよ。

＞「たとえCPUより速くなくとも、CPUでやってる仕事を肩代わりしてやることが
＞出来るだけでも使う価値があるんじゃないでしょうか」

そういうことを臆面もなく語る営業マンを一人知っているんだけどw
同一人物と考えてよさそうだな。

肩代わりしてやるなんて大それたことを無理に言い張るから、おかしくなるんだよね。
重要なのは、ＣＰＵとＧＰＵとが各々の得意分野を担当し、住み分けをすることだろう。
ＧＰＵは汎用計算に向いていないのだから無理にＧＰＵを使わずＣＰＵを使えばいいし、
３Ｄゲームや科学技術計算などＧＰＵの方が効率的な計算でＧＰＵを使えばいいんだよ。

いわゆるアインシュタインとタイピストの喩えだ。
アインシュタインが優れた物理学論文を清書してもらうためにタイピストを雇ったら、
なんとまあそのタイピストよりアインシュタインの方がタイプが速かったとしよう。
じゃあ、そのタイピストを解雇すべきか？答えは否だ。タイプはタイピストに任せ、
アインシュタインは少しでも長い時間、優れた物理理論を考え出すことに費やすべきだ。
それが最も効率がいい。

>>615
営業ならだいたい同じこと言うんじゃねぇか？

>>616
>答えは否だ。
いや、答えは科研費の額によるだろうｗ

>>616
タイピストに指示だすのに、タイプするのと同じような時間がかかるから問題なんじゃね？

清書する段階で更に考えても無駄だろう

CPUを管理職、GPUを部下に例えてみよう。
CPUからGPUへの指示の中身が足りなかったりすると
CPU-GPU間のやりとりが増えてしまい遅くなる。
CPUから指示する内容がGPUの能力を超えると
なかなか結果が返ってこない。

逆にCPUの能力が低いとGPUへの指示や対応が遅くなる。
GPUの仕事に信用がおけないとCPU側でのチェックが
必要となり負荷となる。
GPUはCPUほど守備範囲は広くないし経験も少ない。

よいCPUやGPUを入手するには予算が必要である。
またCPU、GPUを動かし続けるには経費がかかる。（電気代、冷却設備）
あまり負荷をかけるとうるさくなったり、たまに壊れたりする。
overclockによる故障は保証の範囲外であることに注意。

GPUのIPコアが強化されれば良いんだが。
それをしようとして、intelに待ったを掛けられたんかな。

チップセットのバスライセンスと何の関係が？

たとえ話にすると細部の理解が必要ないから生半可な知識でも初心者が騙せて優越感に浸れてうめぇw

っていつも思う。

CPUとGPUは、お互いに交わる方向で
じきに差は無くなるんだろ
GPUいらねでおわりじゃねーの

昔GPGPUスレに書いたが、別のパラメータで同じコンテキストの処理をするようなときに
高速に処理できるのがGPUの利点。
別々のコンテキストが必要なら、丁度良いマルチプロセシングの環境を探しなさい。

どちらも歩み寄っているようだけど、ゲーム屋的には、現行世代機的なトランジスタ
バランスのマシンが次世代にも欲しいところ。

>>625
似たようなコストで作れるならな。
見当付いてるなら特許とって今すぐ始めるのがいいぞ。

GTX280って28SMじゃなかったか？

260のほうか
280は30か

>>619
常識的に考えてそんなことは起こらんだろ。
仮にタイピストに指示だすのにタイプするのと同じくらい時間がかかるなら、
それはこの喩えが適用できないケースだってだけの話だろう。

ＧＰＵに命令だすのにＧＰＵで計算するのと同じくらいＣＰＵ時間がかかるなら、
そりゃＧＰＵを使うのが不適切なケースだってだけのこと。

>>620
いや、清書してもらうのは既に考え出した理論であって、アインシュタインは
タイピストが清書してる間に次の理論を考えるんだよ。

たとえ話にすると細部の理解が必要ないから生半可な知識でも初心者が騙せて優越感に浸れてうめぇw

っていつも思う。

>>631に全面的に同意

>>624
>>631

>>625
ジョンカーマックが昔言ってたわな。それ。
ま、今は宇宙大好きっ子になっちまったが。

>>630
GPUを使うのが不適切なケースばっかりなのが問題。

アインシュタインは一人しかいないけど、CPUとGPUがアインシュタインとタイピストのような関係なら、
CPU増やせばいいんじゃね？無理にGPUにしなくても。

なんにも出来ないのになんでも出来ますと宣伝してるから問題なわけで。

そらあんた、ドラッグレーサーをそれなりに走らせるためには適当なコースとそこまで運ぶためのトランスポーターと
燃料などの消耗品が必要になりますがな。

>>638
>631

>>637 団子
自己紹介、乙。

だんごやさんだよ
だんごせんもんてんだよ

AviUtlがCUDAに対応するのを待つか…

某フィルタでシェーダで書いたより遥かに遅くて駄目だしされたような

今日、CUDASDK入れてみた。CudaSetup-2.1とNVIDIA_SDK10_CUDA_2.10の入れる順番なのか
環境変数でコンパイルエラー、何度かやっているうちになんとか、サンプルが起動できるようになった。
SDKのサンプルはどこにインストールしているんだ　アホか C:\に持ってきた。
サンプルへのパスを追加してやっとコンパイル、起動できた。

>>643
2.1のサンプルって意味不明なとこおかれるよね。
VistaのUAC対策かと勝手に思ってるけど。

NVIDIAは昔っから何でもそう。
ドライバも一旦C:\に展開してからインスコしてくださりやがる。
まぁ、GCCやそれ用のライブラリが、スペースの入ったパスを
嫌うからかも試練。

2.0はちゃんとProgram Files配下に置かれてたよ。

>>635
ＧＰＵはゲームや科学技術計算では実績をあげているので、
不適切なケースばかりではないだろう。

>>636
アインシュタインという不世出の天才物理学者と
タイピストという（当時は）いくらでもいた職業を
比較しているのが、この喩えの肝なんだよ。

ＣＰＵのコアを１個増やすより、ＧＰＵのＳＰを１個
（１０個でもいい）増やす方が、ずっと簡単だろう。

別にCPUコアなんて年間何億個でも量産できるだろ

タイピストが何人もいても意味無いだろｗ

こんなコア橋の下に捨てますよ!

Compute Capability 1.3 の GeForce って、GTX だけ?
GTS とかはだめなの?
ファンがうるさいのはやだな〜

GTSは9800シリーズのリネーム

初歩的なこと聞くけど、これってグラボ一台でもできるよね？
表示用と別にCUDA専用のグラボって必要？

>>653
その程度が分からないと厳しいかと思いますよ。
分からなくても、やってみて駄目だったら買い足すってことで問題ないと思うけど。

>>653
ttp://ja.wikipedia.org/wiki/CUDA#.E3.83.8F.E3.83.BC.E3.83.89.E3.82.A6.E3.82.A7.E3.82.A2

>>655
赤くなっている。。。
GTS250持っているからできると思ったんだけど、二台必要なのかな？
GPUGRIDに参加しようかと思ったらドライバ入れてるのにCUDA対応のデバイスが見つかりませんっていわれるし・・・・。

1台でも使える。
ただ処理中に画面が固まってOSが制御不能になることがある。

>>657
うーん、CUDAがちゃんと入ったかどうか確認する方法ってありますか？

PyCUDAなんてあるのか、おもしろそう

RuCUDAが必要だな

>>656
参加したいGPUGRIDがどんなのか分からんが、
もし、倍精度浮動小数点の演算が必要なものなら、
GTX200シリーズじゃないと無理。

ちなみに、モニタがつながっているかPhysX指定がされてないと、
CUDAでデバイス列挙されないぽい。

Py損とかルビィとか手続き型スクリプト言語は向かんだろ。
ocamlとか、Earlangとかの関数型言語をGPGPU対応に
した方が御利益は大きいんじゃね?

並列計算の場合、副作用とか、計算の依存関係が有ると
性能出ないんで…。

Earlang（笑）

統計解析ソフト「R」用のパッケージ「gputools」：
ttp://cran.r-project.org/web/packages/gputools

これのWindows版バイナリを作ってくださるネ申はいらっしゃいませんでしょうか。　m(ﾟ-ﾟ；)ｶﾐｻﾏ…

>>660
こんなのはあるみたいですが
http://ruby-opencl.rubyforge.org/

GPUの計算部分で
レジスタを多く使っちゃうようにコンパイラで最適化されちゃうんだけど
部分的に無効にする方法はありませんか？

volatile

>>667
？volatileは最適化から外すだけで、レジスタには適応されるっしょ

つか>>666 のレジスタ使ったら嬉しくない事ってのが想像できない。。。

>>648
ＳＰだって年間何億個でも量産できるけど。

>>649
誰も、ＧＰＵを何個も用意しろとは言ってない。

>>668　volatile __shared__

>666の動機によっては__shared__では何の解決にもならないような。
確認していないけれど、恐らくレジスタを使い回さずに消費しまくる方が速いんだろうねぇ。

SP2＋未公開パッチで7RC以上に軽くなってるよ

Windows Vista SP3 Part1
http://pc12.2ch.net/test/read.cgi/win/1241571715/225

ｺﾞﾊﾞｰｸ

SSE　４コアフルに使ったら
最上位品でも大差ないｗ

http://code.google.com/p/thrust/
>Thrust is a CUDA library of parallel algorithms with an interface resembling the C++ Standard Template Library (STL).

きた！STLきた！これで勝つる！　　


ないない

brookみたいだな。

仮想マシン上でCUDAのインストールに成功した方はいらっしゃいますか？
当方、ホストOS：Vista、ゲストOS：Ubuntu8.04です。
仮想マシンであるUbuntu上で、NVIDIAドライバ: NVIDIA-Linux-x86-180.22-pkg1.runを起動してみました。
すると、「 You do not appear to have an NVIDIA GPU supported by the 180.22 NVIDIA Linux graphics driver installed in this system」とエラーがでました。
GPUは、GeForce 9800 GTです。どなたか、アドバイスお願いします。

仮想マシンは無理

>>675
合計なんかCUDAでやって早くなるのか？

分割統治法は並列化の基本だな

合計求めるのは苦労したなぁ。
結局、分割数(128とか256とか)置きに足していって、その結果はPCで足した記憶がある。

とりあえず公式の３つをインストールしたのですがTMPGEncで確認できませんみたいなことを言われました。
インストールするだけではcudaの恩恵を与れないのでしょうか?

すいません、直ぐ解決しましたorz
ｸﾀﾞがちゃんと動いているか確認したいのですが方法はありますでしょうか？

>>684
SDKをインストールしたのなら、サンプルもインストールしてビルドしてみよう。サンプルが動けば、大丈夫。
# 特に、deviceQueryは便利。

こんにちは。CUDA初心者です。質問があります。
Visual C++ 2008、CUDA tool kit ver2.1、CUDA SDK ver2.1
で、サンプルのsimpleCUBLASをビルドすると、
1>LINK : fatal error LNK1181: 入力ファイル 'cutil32D.lib' を開けません。
と出ます。
そこで、CUDA SDKのlibを調べたところ、
cutil64D.libがあり、32のほうはありませんでした。
この場合、どうすればビルドできるのでしょうか？

リリースバージョンをリンクする。

リリース構成でビルドしたところ、今度は、
1>LINK : fatal error LNK1181: 入力ファイル 'cutil32.lib' を開けません。
と出ました。
CUDA SDKには、64があり、32はありません。
32と64の違いって一体何なのでしょうか・・・

パスが通ってないんだろ

OSが64bitだと、32bitのCUDAライブラリはインストールされなかったような。

ご回答ありがとうございます。
パスが通っていないということですが、
「パスを通す」について、詳しく説明していただけませんか？
知識不足で申し訳ありません；；

環境を名に使ってるかによるが、
Visual Studioだとプロジェクトのプロパティからインクルードするファイルがあるディレクトリのパスと、
libがあるディレクトリのパスをついかする
linuxだとコンパイラのオプションに追加する

詳しくはぐぐれ

何度も回答していただきありがとうございます。
リンカの追加ライブラリを調べたところ、
ちゃんと、SDKのcommon/libが指定されてました。
しかし、この中には、cutil32.libではなく、64があります。
ということは、690さんのおっしゃるとおり、
32bitのCUDAライブラリはインストールされなかったということなのでしょうか？
もしそうでしたら、サンプルプログラムは64bitに対応してないが、
自分でプログラムを作る分には、上記のようなエラーはでないということでしょうか？

サンプルのリンカ設定を編集して64bitのライブラリをリンクすればいいんじゃね？

アクティブソリューションプラットフォームとプラットフォームをWin32からx64へ変更したところ、
エラーがなくなりました。
そのかわり、
1>------ ビルドのスキップ: プロジェクト: simpleCUBLAS ------
1>
========== ビルド: 0 正常終了または最新の状態、0 失敗、1 スキップ ==========
とスキップしてしまいました。
何が原因なのでしょうか・・・

threadIdx.xがうまく値を返さなく困っています。
最小のプログラミングだと　ちゃんとした値を確認できるんですけど、
規模のあるプログラム書いた物では、threadIdx.xをみると最大でも1000以内の数値が40000を超えていたりします。

かなりラフな書き方していて、グローバル変数使いまくってるのが意見ないのでしょうか？
__device__　 kouzoutai[2000];//グローバル変数
とか宣言しまくって搭載メモリーを超えちゃってるかもしれませんが、その時は明確なエラーとか出ますか？

エラーは基本的に出ない

>>696
threadIdxはプログラムで変更できないと思うが。
どうやってthreadIdx.xを「みる|確認できる」としたのか知らんが、その手段を確認すべきでは?
ptx出力を眺めればthreadIdx.xをコピーした先で壊してたりするのが分かるんじゃない?
グローバル変数にたった2000要素くらいなんてことないと思うが、巨大な構造体でも置いてるの?
あーそれから、メモリをオーバフローしてもコンパイルエラーも実行時エラーも出ないどころか、
突然システムごとフリーズするかもしれないから覚悟しておいてね。

>>501
　　　？
　　∧∧
　 (´･ω)
　_｜⊃／(＿＿＿
／ ヽ_(＿＿＿_／
￣￣￣￣￣￣￣

寝た子を起こしたなw

>>699
　　！
　 ∧∧
　(･ω･ )
　_| ⊃／(＿＿
／ ヽ-(＿＿_／
￣￣￣￣￣￣

　
　　あ"
　 ∧∧
　(◎ω◎)
　_| う／(＿＿
／ ヽ-(＿＿_／
￣￣￣￣￣￣

>>701
　　　　　ｵﾊﾖｰ!!
　　　　∧∧ ∩
　　　 (`･ω･)/
　　　⊂　　ノ
　　　　(つノ
　　　　 (ﾉ
　＿＿_／(＿＿_
／　　(＿＿＿／
￣￣￣￣￣￣

質問。
グリッド・ブロック・スレッドのカーネル側での計算への応用方法が
よくわからないのですが・・・・特に動きです。

const unsigned int tid =blockIdx.x * blockDim.x + threadIdx.x;
がスレッドＩＤを示しているというのはわかるのですが・・・

for(int x=0;x<height;x++)
for(int y=0;y<width;y++){
　　　　　　　　sum=sum+tex2D(tex, x, y)*((float)(cos((M_PI*((x*width+y)+0.5)*tid)/number)));

}
outputdata[tid]=sum;

の中で、tidはどういう働きをしてるんでしょうか？
特に回収する配列であるoutputdataの動きが知りたいです。
わかる方ご教授お願いします・・・

>>703
　 ∧,,∧
　( `･ω･)　ｳｰﾑ…ここは？
　/　∽ |
　しー-Ｊ

>>704
何が判らんのか判らん。
tidが個別のデータスレッドの番号を指すように働くことが判っているなら、

その番号でoutputdataにアクセスしていることくらい判りそうなもんだが。

仮に、blockDimが4でgridDimが2だとしたら、blockIdxは0か1、threadIdxは0から3になるわけで。
つまりはtidは0から7になるわけだな。

つーか、VIPのWikiにも情報があるよ。
http://vipprog.net/wiki/%E3%83%97%E3%83%AD%E3%82%B0%E3%83%A9%E3%83%9F%E3%83%B3%E3%82%B0%E8%A8%80%E8%AA%9E/CUDA.html

>>730
カウンター3連打って言うけどスローで見るとカウンター2発に追撃1発って感じ
まあそれでも凄いって言うか、何この異次元映像ｗ
もはや芸術的とか圧倒的とかそういうレベルを超越していて笑いしか出てこないわけだが？ｗｗｗ
しかもメディナは立ち上がるしｗｗ不死身かｗｗｗｗ
この試合はボクシング以外の何か別の競技だわｗｗｗｗｗｗｗｗ

誤爆すまそ
俺が代わりに言っといてやったぜ

CUBLASに加えてCULAPACKとか用意して欲しい
LU分解もできない環境なんて

GPGPUはDirectX11で花開くんだろ。
エンコード高速化対応お待ちしております。

DirectX 11講座（中編）
非ゲーマーのアナタのパソコンも
DirectX 11演算シェーダーでスーパーコンピューターに変身する！
http://game.watch.impress.co.jp/docs/series/3dcg/20090327_79998.html
GPGPUの場合は、組んだGPGPUプログラムがATIかNVIDIAのどちらかでしか動かない、
踏み絵的な「ATI対NVIDIA」なのだ。
これではGPGPUの進化、発展、普及が望めない。
そこでGPGPUの標準化プラットフォームが強く望まれるようになった。
このGPGPU標準プラットフォームを、Windows環境下で提供しようと出てきたのがDirectX 11ということになる。
DirectX 11は、このGPGPUの標準プラットフォームとして「DirectX Copmute Shader」（演算シェーダー）を提供する。


演算シェーダーでなにができるのか
http://game.watch.impress.co.jp/img/gmw/docs/079/998/html/3dd1103.jpg.html

>>710
エンコードって意外に並列処理できる箇所少ないんだけど。。。CUDAでどのベンダも対応してこないのはそのため。
Compute Shader出てきても状況は変わらないと気が

それマルチ

今日本屋見てきたけど
CUDAの入門書とかって無いね。
世の中にはあるのかな。

__device__ int particle_position[30][30][50][30];
__device__ int particle_position_num[30][30][50];
　：
とか
__device__ __constant__ float K_dWpress;
　：
な感じで、結構たくさんグローバル変数をデバイス側に作ってるんだが、
これって、CUDA的にやっていい手法？てか数次元配列ってカーネル関数内で普通のcっぽくアクセスできる？

とりあえずメモリは足りるみたいだから、デバイスホスト間でのデータ転送は基本的にしなくて、
結果のみホスト側に持ってきたいんだけど。

コンパイルエラーはないけど、カーネル起動させたあとエラーが出て、
一瞬画面が消える。

追記
ちなみに、Emuモードだと問題なく終了できます。

>>714
CUDAはメモリ管理なんてしないから、デバイス側の変数がどれだけあるのか自分で把握すること。
一瞬画面が消えるのは表示用とCUDA用を同じGPUでやっている所為だと思われ。
もっと酷いときはフリーズするから要注意で。

CUDA初心者です。

サンプルをいじくっているのですが、いくつかわからないことがあり質問させてください。

�@テクスチャメモリ
キャッシュが効くからグローバルメモリよりも高速との事ですが、
時間を計ってみるとグローバルメモリと変わらない・・・。
どのような用途で使用するのでしょうか？

�A異なるカーネル関数の同時実行
ブロック、スレッド数を決めてカーネル関数を実行しますが、
cudaThreadSyncrinize関数やメモリコピー関数を呼ばなければ
同期待ちはしないですよね？
であれば、カーネル関数を2つ書いたときに同時実行するのでしょうか？
マニュアルを見る限りではできないようですが、会社の先輩が出来るというのでホントかなと。

お分かりの方宜しくお願いいたします。

俺も初心者だが。

1. 線形補完できるよ！ 速度面で違わないなら、それ以上はどう違うのか知らん。

2. なんか俺の動かしてる感触だと同時実行してる気がする。 少なくともCPUとGPUで同期はしていないよ。

>>718

サンクス。

�@は線形補間機能があるのは存じていますが、
補間せず生データを使用したいものでして・・・。

�Aは自分でも本家CUDAZONEで調べてみましたが、どうもできないっぽいです。

http://forums.nvidia.com/index.php?showtopic=84740&hl=cudaStreamSynchronize

このスレの前にもありましたね。
よく見てませんでした、すみません。
cudaStream〜って何のためにあるんだか。

CUDAを始めようかと考えていますが、下記の様な処理はCUDAで速くなりますか？

function hoge(const b:string):boolean;
begin
result:=true;
end;

>>720
すみません、間違えて書き込んでしまいました。
改めて、
typedef struct _table_t {
float x,float y, float z;
} table_t;

typedef sturct _vec {
float vector[4];
} vec;

typedef union _vtemp_t {
vec V[3], table_t data[4];
} vtemp_t;

table_t DATA[100];(DATAにいろいろと情報をが入っている)
vtemp_t temp;
temp.data[0] = DATA[15];
temp.data[1] = DATA[43];
temp.data[2] = DATA[11];
temp.data[3] = DATA[80];
(tempのメンバdataにＤＡＴＡの値をランダムに代入)
メンバdataがもつ、x､y､z情報をメンバVにx､y､z毎にまとめたいです。例えばこんな感じに。
V[0] = {data[0].x, data[1].x, data[2].x, data[3].x};
V[1] = {data[0].y, data[1].y, data[2].y, data[3].y};
V[2] = {data[0].z, data[1].z, data[2].z, data[3].z};
CellとかSSEならSIMDを使って出来ると思いますが、CUDAでも速くなるなら、
CUDAをやってみようかと思っています。

それ計算じゃなくてただのメモリ転送じゃん。 CUDAでやることじゃないよ。

Mac でCUDAをやろうとしています。
/Developer/CUDA/lib/にあるライブラリをライブラリのサーチパスに含めるには
どの環境変数を設定すればよいのでしょうか？
LD_LIBRARY_PATHに追加してみたのですが、サーチしてくれません

>>724
DYLD_LIBRARY_PATH

>>725
早速の回答ありがとうございます。
DYLD_LIBRARY_PATH に追加してみたのですがやはりうまくいきません。
ちなみに-Lで指定すればうまくいきます。

Coalesced、Non-Coalescedって簡単にいうとどういうこと？

>>108

>>728
よくわからん

ttp://noridon.seesaa.net/article/105559613.html

大前提として、ブロック内でデバイスメモリにアクセスするアドレスが連続していたら、コアレス。
んで対応アーキテクチャによってスレッド単位でアドレスが連続している必要があったりする。

まぁ簡単に言うと、バス幅有効活用できまっせ適な使い方をイメージすればいいかと。
なんでそのバス幅でアクセスできるアラインを考える事になる。

>>726
実行時じゃなくて、ビルド時のリンクのこと？
DYLD_LIBRARY_PATHは関係ないから-Lでやってね。

>>726
LIBRARY_PATHでできなかったっけ？

配列で最初と最後の変数だけ別の処理をしたいのですが、
どのように書けばよいのでしょうか？
従来のプログラムの場合、配列がa[N]だったばあい、
for(i=1;i<N-1;i++)と書けばよいのですが、
CUDAのカーネルで処理する場合どのようにすればよいのでしょうか？

>>734
とりあえず、現状のカーネル書いてくれれば教えやすいのに

NVIDIAが配布しているということで、PhysXの質問をさせて下さい。

キネマティックなアクターが物体と接触した時に受ける力の取得方法が分かりません。
ご存知の方がいらっしゃいましたらご教示下さい。

>>735
例えば、今は

__global__ void test(int thread_size, int one_size, float *res_d)
{
    int bx = blockIdx.x;
    int tx = threadIdx.x;
    int pos = bx * thread_size + tx;
    if (pos==0)
        res_d[pos] = cos(res_d[pos]);
    else
        res_d[pos] = sin(res_d[pos]);
}

このように、if文で0とN-1の場合だけ別の処理をするようにしているのですが、if文は遅くなるので、避けたいのですが。

具体的にやりたいのは、3次元の偏微分方程式を解く際の境界条件部分の計算です。
この場合各面、各辺、各頂点合わせてif文が20個以上になってしまいます。

何か良い方法はないでしょうか？

>>737
最初と最後だけ別の処理なら、そこはCPUに任せて、
残りをGPUにやらせればいいと思うけど、それじゃだめなん？
posを+1して、並列数をN-2にすればいけそうだと思うけど。

あと、そのソースだと N-1 の処理が別になってない気がするのと、
posがN以上かどうかのチェックが無いのが気になる。
（ソースは単に省略しただけで、チェックはそもそも必要ないように組んでるのかも知れないが）

>>737
一般的にアセンブリレベルで分岐をなくしたいときは、こんなテクニックを使ったりする。

// i == Nのときだけsinをcosにする

float s = sin(なんとか);
floac c = cos(なんとか);
int b = i == N; // true が (int)1に変換されることを期待する
return (c * i) + (s * (1 - i));

もちろん3行目は実装依存なので、プロセッサのマニュアル見ながらいちいち確認はした方がいいんだが、
だいたいのプロセッサで使えるテクニック。

>>737
それだと全てのパターンでsinとcosの両方が計算されるから

float s = sin(なんとか + b * PI / 2);

にすれば？

ごめん安価ミスった

>>740は>>739へのレスです

CUDAにはsincosねーの？
また三項演算子でいいんじゃねーの？

>>737
あらかじめ、特殊条件について位相を補正しときゃいいんじゃないの？と思うのは俺だけ？

>>736
PhysXスレがゲーム製作技術にあるからそっち行け

>>744
誘導ありがとうございます

>>742
あるよ。但し、ストリームプロセッサに一個しか超越関数演算機がないから使い方間違うと遅くなるけど。

>>746
超越関数演算機なんてハード的に存在したっけ？倍精度演算機じゃなくて？

>>747
SFU(Super Function Unit)がある。
ただし、4SPで共有で、>>747のいうように8SPで共有するのは倍精度演算器。

sin()、cos()、tan()は確かに遅いけど、
sinとcosについては__sinf()と__cosf()を使えば、
精度は悪くなるけど1op/clockでできる。
除算以外の算術演算が8op/clockでできることを考えたらそれでも遅いことには変わりないけど。

超越関数がレイテンシ1なんて凄いね
SFUだけ100倍のクロック数とか
してるのかなぁ

http://www.khronos.org/opencl/
OpenCL 1.0 リリース

First OpenCL demo on a GPU
http://www.youtube.com/watch?v=r1sN1ELJfNo

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

ドライバ一式は登録しないといけないようだ

>>748
SFUなんてあったのか、しらなかった。DPみたいに実行はSPと並行してできるのかな？

やっとFortran対応の話出てきたけど、有料なんかね
PGI and NVIDIA Team To Deliver CUDA Fortran Compiler
ttp://eu.st.com/stonline/stappl/cms/press/news/year2009/t2399.htm

ていうかいつのまにかに2.3のベータ出てたのか

GTX 285でようやくおれのようなへっぽこでも普通にcpuを凌駕するコードが書けるように
なったようだね。

後は（値段はともかく）消費電力が下がってくれれば、大ブレイクしそう。

３００Wなんてクアッドコアマシンを４台くらいクラスタリングしたような消費電力だからな

>おれのようなへっぽこでも普通にcpuを凌駕する
だったらCPU要らんがな。
むしろCPU<->GPU間の通信がボトルネックなので、CPU無くして
GPUに直接I/O繋げられるようにしろや！

>>758
それなんていうTesla?

>>759
それTesｌaでもない。

CUDAで乱数を効率的に生成する方法を教えてください

>>761
コアごとに違う種から乱数生成すればいいんじゃね？

質問なんですが、SDKに付いてたサンプルをmakeしようとしたところ
以下のようなエラーが出てうまくいきません、何か足りないものがあるのでしょうか

/usr/bin/ld: cannot find -lXi
collect2: ld はステータス 1 で終了しました
make[1]: *** [../../bin/linux/release/nbody] エラー 1

OSはfedora10です。色々初心者なのでどうかお願いします。

libXiというパッケージをインストールできないかね
俺もfedoraはよく記憶にないんだが

>>764
libXiというのがインストールされていないということだったのですか
早速インストールしたいと思います。どうもありがとうございました。

teslaとgeforceって何が違うの？

OpenGL向けかDirectX向けか

>>767
それquadroとgeforceでしょ。

>>766
teslaはhpc専用ボードで、ディスプレイの出力ポートがない。
NVIDIAがボードの製造まで一貫して行ってて、テストもしっかり通している…らしい。
あとメモリが4GB乗ってる。

CUDAというよりハードウェアの話なんですが質問させてください

TESLA C1060をMM3500に刺したいんですが、この状態だとBIOSが拝めません
ビープ音「−・・（長短短）」って音がするんでAwardで調べたらビデオアダプタ不良らしい
マザーボードだけを他のに変えたら正しく起動するので、電源やTESLAが悪いわけではなさそうです
TESLAなしのMM3500だけならUbuntuのBoot CDからMemtestが走るとこまで行きました

スレ違いどころか板違いだ
判ってるなら該当板いけよ

>>769
買ったとこに聞いた方が良いかと...

>>772
ありがとう　まずはVIAにメールしてみます

・GeForce
NVIDIAはパーツ選択にはノータッチ。各社さんが勝手に安い部品を積んでいるから動作しなくても当然。
プロならそんなアキバ的発想はやめて、QuadroFXかTeslaを使ってください。
・QuadroFX
NVIDIAの厳しいテスト基準に合格した高品質の出力を保証しているから決して安くありません。
プロにも安心してお使いいただけます。
・Tesla
QuadroFX以上に品質に重点を置いてテストを行なっています。HPCにはこれ以外の選択肢はありえません。

>>768
OpenGL向けもDirectX向けもないよ。>773にもある通り、NVIDIAがテストしているかどうかの違い。
なんせ、GeForceの製造販売を行なっているELSAの営業の前で堂々と>773みたいなことを仰ってる。
10倍の価格差も、当然ということなのでしょう。

そりゃカノプーが手を引くわけだよな
どうしたってGeForceは価格競争になっちまう

不特定多数の客向けのプログラムでGPGPUを使うコードを書くやつはアホウ、ってことだな。

全てアホウと決め付けるのはねぇ
　中には超やばいものを開発するかもよ（世界は広い）

特にロシア東欧は要注意！

>>777
アホウと言っているのが、nVIDIAの営業ってことだよ。

「阿呆」とは言わなかったけれど、「動かなくても当然ですから」くらいのことは言ってたなぁ。

今月末に CUDA2.3 が出るとさ。

PGI compiler 9.0 は2週間体験できるから、Fortran で単純な行列積の
コードを書いて試してみたけど、Core2 Quad + GeForce9800GTX で
CUDA 2.0 という古い組み合わせのせいか，
!$acc region の設定の効果が全く現れなかった。

pgf95 -ta=nvidia:cc11 matrix.f
というコマンドでいいと思うんだけど。

CUDA用の姫野ベンチをGTX285で実行しようとしたのですけど動きませんでした。
何が原因かわかりません、どなたかわかる方いらっしゃいましたらお願いします。

原因はわかりますが、何をお願いされているのかわかりません。

その原因について詳しく教えて頂きたいのですが

>>783
原因なら簡単ですよ。
あんたが人間じゃなくて人だからです。

動くかどうかも保障してないものが普及するとか思ってるのか？

>>785
>779のことを言っているのなら、QuadroFXとTeslaがあるから大丈夫だと思っている
頭に蛆が涌いているとしか思えないNVIDIAの営業に言ってあげてください。

無駄だけどw

visual profilerについて質問です
プロファイラの項目GPUTimeを見ると全体で3000くらいなのですが、
プログラム中で cutCreateTimer( &timer) を使って計測すると 18ms〜60msのバラけた値が出ます
これってどちらが正しい値なのでしょうか？

>>781
バイナリのCUDA SDK1.1だからだろ

>>787
cutCreateTimer()の実装を見てみた?

Fixtarsの資料面白いね、祝日に更新ご苦労さんだわ
この調子で本家資料全部訳してくれないかなｗ
ttp://gpu.fixstars.com/index.php/CUDA%E3%83%97%E3%83%AD%E3%82%B0%E3%83%A9%E3%83%9F%E3%83%B3%E3%82%B0TIPS

2.3来たね、やっと64bit版で32bitコンパイルできるようになった

Mac版の2.3.0ドライバーインストールしてみたら、一回目libcuda.dylibだけ入らなかったことに気付かなくて、全然見当違いの所で原因探してて苦労したよ。
一発で入るようにしてください、nVidia様

ttp://www.nv-info.com/cuda_for_highschool/

20年前の高校生は参加しちゃダメかな?

Teslaプレゼントしろよ
相変わらず一貫してないな

CUDAで動くトリップ検索ツール頂ける方いらっしゃいませんか

ﾉｼ

【レポート】中高生がCUDAにチャレンジ! NVIDIAがGPUコンピューティングの体験イベント開催 | パソコン | マイコミジャーナル
http://journal.mycom.co.jp/articles/2009/07/27/cuda/index.html

>>795
あれ？まだ誰も作ってなかったの？

あと6年早ければ・・・

>>798
らしいです

じゃぁ俺が作るか。
まずはCUDA無いバージョンのソースがあるURL教えてｗ

まあ、待て屋。

>>802
それはラデ用のしかなかったような…

CUDA無しを所望のようだからこれで良くね？
CPU用のコードも当然あるよ。

>>804
nVIDIAのVGAしか持ってないんだorz

>>805は>>795か？
俺は>>801に向けて書いたんだが・・・

>>806
あ、はい。
俺は>>805であり、>>795です。

ID出てないと面倒だね

>>807
トリップ使ってけれw

>>808
おｋ

>>807です

Larrabee向けのソースコードなら多分あるんだがｗｗｗｗｗ

まだ手元にある定額給付金で
ラデ買ってくるかな…

それも一つの選択肢

かつてTripcode ExplorerのCUDA版があったらしいね
でも、今はもう消えてる…

ギコハハハ

>>812
がはは！それはネタだからーw

>>814
そうなんだｗｗｗｗｗｗ

ふふふ

あらあらふふふ

これでいいじゃん
http://tmkk.hp.infoseek.co.jp/tripper/

３００Wも消費する物体を数週間動かすとは恐ろしいｗ

shared memoryって,CPUで言うL1キャッシュみたいなもん？？

>>818
これって使うとどれくらい寿命減るの？
作った人って絶対に高負荷かけて石を壊そうとしてるよね・・・

>>821
なにいってるの・・・
まあたしかにファンがクソだと熱で死ぬ可能性はあるが

>>820
大分違う。つーか、CPUとは考え方を変えないと使いこなせないよ。

>>820
CPUで言うレジストリみたいなもんだ

>>818
for Macだ…と

Windows用に直せるだろこんなん。（と言ってやらない）

Vistaだと困ったことに、描画に支障ない程度に負荷調整しないとAeroが死ぬｗｗｗ

だれかトリップ生成のアルゴリズム説明して？
趣味でCUDA処理のコード書いてみたいからー

>>818のコード読めって話だけど、ハッシュとか暗号系はアルゴリズム知ってないと、時間かかるんでお願いします。

309 名前：◆cZfSunOs.U []： 2009/06/19(金) 11:46:10 ID:MLNb4KfK0 (17)
生キー指定以外の12桁以上「##〜」形式も将来の拡張用にしておきましょう．
で，まとめ:

if (length $handle_pass >= 12)
{
my $mark = substr($handle_pass, 0, 1);
if ($mark eq '#' || $mark eq '$')
{
if ($handle_pass =~ m|^#([[:xdigit:]]{16})([./0-9A-Za-z]{0,2})$|)
{
$GB->{TRIPSTRING} = substr(crypt(pack('H*', $1), "$2.."), -10);
}
else
{
# 将来の拡張用
$GB->{TRIPSTRING} = '???';
}
}
else
{
use Digest::SHA1 qw(sha1_base64);
$GB->{TRIPSTRING} = substr(sha1_base64($handle_pass), 0, 12);
$GB->{TRIPSTRING} =~ tr/+/./;
}
}
else
{
# 従来形式
}

インデント消えたorz

Kernel用のCRYPTを自前実装しないといけないのか・・

C++用のコード晒してみる。これは実際に某ツールで使ってる。

char* trip_sha1(const char* key) {
SHA1Context sha1;
SHA1Reset(&sha1);
SHA1Input(&sha1, reinterpret_cast<const uint8_t*>(key), std::strlen(key));
uint8_t digest[SHA1HashSize];
SHA1Result(&sha1, digest);

static char cryptresult[13]; /* encrypted result */
static const char base64mod[] = /* 0..63 => ascii-64 */
"ABCDEFGHIJKLMNOPQRSTUVWXYZabcdefghijklmnopqrstuvwxyz0123456789./";
cryptresult[0] = base64mod[digest[0] >> 2];
cryptresult[1] = base64mod[(digest[0] << 4 | digest[1] >> 4) & 0x3F];
cryptresult[2] = base64mod[(digest[1] << 2 | digest[2] >> 6) & 0x3F];
cryptresult[3] = base64mod[digest[2] & 0x3F];
cryptresult[4] = base64mod[digest[3] >> 2];
cryptresult[5] = base64mod[(digest[3] << 4 | digest[4] >> 4) & 0x3F];
cryptresult[6] = base64mod[(digest[4] << 2 | digest[5] >> 6) & 0x3F];
cryptresult[7] = base64mod[digest[5] & 0x3F];
cryptresult[8] = base64mod[digest[6] >> 2];
cryptresult[9] = base64mod[(digest[6] << 4 | digest[7] >> 4) & 0x3F];
cryptresult[10] = base64mod[(digest[7] << 2 | digest[8] >> 6) & 0x3F];
cryptresult[11] = base64mod[digest[8] & 0x3F];
cryptresult[12] = 0;
return cryptresult;
}

ｓｈ１のCUDA版くらいありそうだけど

ご両人、コードさんくす〜
実はC住人なんで（汗

蛇足だが関数内のstaticの使い方が気持ち悪いです（笑

さー昔しらべたSHA-1の資料でもひっぱってくるかー

オリジナルのcrypt(3)でもchar配列はstaticでとってるんだよな
たしかにキモイ

std::stringとか使ってみ？
CUDAに移植できへんで？

笑えるくらいきれいなワンツーｗ

ごば

CUDAでrandは使えない？

>>838
使えないよ。自分で用意しないと。

CUDA上でのSHA1の実装というとこのあたりか？
http://forums.nvidia.com/index.php?showtopic=102349

r0,g0,b0,r1,g1,b1,r2,....

というようなグローバルメモリにあるchar型の配列に対して、GPU側からuchar4型変数で

{r0,g0,b0,r1} {r1,g1,b1,r2},....

というようにアクセスすることは可能でしょうか？

y*4+x

あーこれじゃ駄目だった
ごめんなさい、忘れてください

だめだー
SHA1と総当り攻撃じゃ現実的じゃない。。。
4桁適合ならまだ使える感じorz

誰かRADEON HD48xx向けのCUDAドライバ開発してないの？

>>845
どちらかというと
AMD Accelerated Video Transcoding（AVT）

では？

いや、CUDA互換機能つけて
CUDA対応ソフトを使えるようにならんかなーと。

>>847
OpenCLじゃだめですか？

ラッパー作ってくれるならいいけど・・・
既存のソフトでCUDAにしか対応してないやつとか
CUDA以外の開発環境が整わなさすぎて、見切り発車でCUDA使いたいなとか
あるんで。

CUDAってOpenCL取り込んでたような……気のせい？

CUDAは環境。

OpenCLは言語・フレームワークで、CUDA向けのプログラムを記述出来る物のひとつ。
従来の「CUDA言語」は "C for CUDA" という扱いになった。

x64でCUDAをしようと思ったんだけど，vitual studioでプロジェクトが作れない．
CUDA_VS_Wizard_W64.2.0を入れて，
CUDA64WinAppはテンプレートとして表示されるようになったが，
選択してもうまく生成されない．
CUDAWinApp(32bit)の方はうまく生成される．

64bitで開発してる方や，
症状の原因がわかりそうな方，ご教授願う．

おっと，自己解決．
スレ汚しスマソ．

>>853
いや、教えてよ（笑）

共有メモリのBank Conflictがいまいちわからないのですが、
同じ配列のアドレス(添え字)に複数のスレッドが同時にアクセスしなければいよいのでしょうか？

最近のマザボってPCIEx16が普通に2,3個あるよな
ところで違う世代のGeForce9800、GTX280とかを積んだ場合
CUDAはどちらのGPUも使ってくれるの？

>>856
CUDAはGPUチップを一つだけしか使えない。
どれを使うかは、cudaSetDevice()で指定しないといけない。
サンプルの、MultiGPUを参照のこと。

ということは，
最近のグラボ一枚にGPUが2つ積んであるヤツでも，
GPUを指定してやらないといけないのか．

そそ、しかも1スレッド1GPUに拘束される。

>>858
未確認だけど、SLIの場合はドライバで1つのGPUに見せかけるらしい。
2チップを積んでいるカードは内部でSLI接続されているらしいし、
1デバイスの場合と同じでもいいかもしれない。

SLIとCUDAは同時使用できないとかじゃないのか？

http://forum.nvidia.co.jp/EokpControl?&tid=11178&event=HE0004
* SLIグループ内の各GPUは個別に列挙されるようになったため、
計算アプリケーションはSLIがグラフィックス向けに有効になっている
場合でも、マルチGPUの性能を活用することができます。

と、リリースノートにはあるね

>>862
2.3からはSLIの場合でもCUDAからは個別に見えるように変更なのか。

SLIでゲームの中でPhysXとかCUDAを使う場合とかは特定のチップに負荷が集中して
そのせいで利用効率が落ちたりしないのかな？

プログラマが生理整頓するか，
CUDAさんよろしくお願いしますするか，どっちがいいかってことか．
でも，一度の演算に対して，1コアにスレッドが重複しない限りは，
無駄に複数のGPUに処理を分配する必要性はないよな．
演算の命令ごとにGPUを切り替えるなら，効果はありか．
SLIの場合って，GPUのRAMは共有されるのか？

>>859
GPUはGPU内の大量のコアがそれぞれ1スレッドずつ受け取る仕組みじゃないのか？
1スレッド1GPUって並列演算ってレベルじゃねーぞ．

>>864
「スレッドを受け取る」っていう意味がわからんのだが。
CPUの１スレッドがGPU１個使って、そのGPUが複数スレッドを実行する。

>>865
CPUのスレッドなのね．
スレッドのことをGPUのハードウェアスレッドかと勘違いしてた･･･orz

CUDAがでてきた2005年当時は、シングルコアの限界も見えてきて、それなら

有り余るグラフィックボードの性能を、3Ｄレンダリングや、フォトショップフィルタなど、ゲーム以外にも有効活用と歌ってたけど

あれよあれよという間に、インテルCPUがマルチコア化してきて、そっちの方が上記の用途に、大きな効果があり、メインになってきて、

結局グラフィックボードは、結局またGPUはゲーム以外に用なしになったなｗ

U-BASICってバージョンアップしないの？

>>867
コンシューマーはそうかもしれんけど、エンタープライズはまた少し違うんじゃない？じゃなきゃNVIDIAもTESLAなんて作らんよ。

ああシールの貼り替えで10倍の金を取るために始めたアレか

>>870
本気でそう思ってるなら偏狭すぎる。少しは頭使えよ

>>869
むしろ、nVIDIAはTESLA以外作りようが無かったんでしょ。
だってCPU作れないんだもん。

NVIDIA は、2009年8月17日、業界標準の線形代数ライブラリ、LAPACK を、超並列処理の NVIDIA CUDA 対応 GPU に最適化した実装「CULA」のベータ版が EM Photonics 社からリリースされたと発表した。
LAPACK ルーチンは何百万人もの開発者が活用し、数値物理学、構造力学、電子設計オートメーションなどさまざまな問題の解決に役立てられている。
CULA により、ワークステーションやデータセンターで使われている NVIDIA Tesla GPU の力を活用し、クアッドコア CPU 1個の場合でパフォーマンスを最大10倍に引き上げることが可能になるとされる。
EM Photonics には、CULA Basic、CULA Premium、CULA Commercial からなる CULA tools という製品ファミリがある。CULA ライブラリとは、LAPACK ルーチンの実装の1つで、GPU の能力を活用して高速に処理が行えるようにしたもの。
LAPACK というのは線形代数でよく使われる関数の集合体で、科学分野や工学分野で活用されている。科学分野や工学分野の問題は線形モデルで近似できることが多いため、線形代数ルーチンで解が得られる。
CULA は、NVIDIA GPU が持つ超並列 CUDA アーキテクチャを活用し、LAPACK ルーチンの多くを高速で処理できる。

>>872
それ以外作りようがなかったからって…事前事業でやってるんじゃないんだぞ？「しょうがなくTESLA作る」とかありえないんだよ。純粋に、戦略的なマーケティングの結果だろ。

>>874
× 事前
◯ 慈善

戦略的なマーケティングの結果、マザボは作れてもCPUは作れないNVIDIAは仕方なしに、
TESLAをでっち上げたわけですね、判ります。

CPUは作れないってどういう意味で言ってるんだろ
x86互換プロセッサなんて作れるところはいくらでもある
ただライセンスの問題なダケ
あとはマーケットな

x86互換プロセッサぐらいなら大学の授業でFPGAに実装したことあるぞ

互換なんてダレでも作れるべ
性能はムリだが

DXCSがくるからCUDAはもういらないかもー

作るだけならVIAにだってできる

↑↑お前出来ないだろｗ

NVIDIAってCUDA捨てたん？
ここ1年ぐらい見る限りすげーやる気なさそうなんだけど。

捨ててないけどやることないだけだろ

今はOptiXに期待

http://news.developer.nvidia.com/2009/08/nexus-visualstudio-based-gpu-development.html

おまいら，
9月30日以降に期待ですよ．

CUDAで最適化したいプログラムに条件分岐満載で
実際の演算部分がほとんどなくてDate並列化できそうもないときって
あきらめるしかないのかな･･･。

http://www.atmarkit.co.jp/fcoding/articles/parallel/03/para03a.html

Fixstarsのエンジニアからの有り難い言葉

条件分岐でシーケンスを発生させる事そのものが目的でなけりゃ、
何か手はあるんでないの？
何しようとしてるか知らんけど、今使ってるアルゴリズムに以外に
GPU向きの手がないか調べてみたらいいんでない？

やっぱりx86コアが100個も200個もあるほうがいいんじゃね？

>>889
なるほど
もっとソースを読んで見るよ。

あとCUDAについてのオススメの文献あるかな？

>>891
GPU Gems3。訳本はアホみたいに高いが、英語でよければネットで公開されてる。途中までだけど。

>>890
そりゃそうかもしれんが、コスト見合わんだろ。

x86のほうが性能マシな演算用途こそ、GPUが何百台あってもコスト的に見合わないんじゃないかと

>>892
thx
でかい書店いって立ち読みして検討するかな。

>>894
そうとも限らない。
32bit単精度浮動小数点を扱えるプロセッサが数百個あるわけで、
プログラムの仕方によっては高速化する。
まあ向き不向きがあるのは事実だし、どっちが良いとは決め付けられない。

向き不向きでいえばGPUに向いてない（というか使うことすらできない）演算の方が圧倒的に多い。

円周率の計算とかに使えればいいじゃん。

なにそれ誰が得するの？

円周率を100万桁まで暗記してる人が、あれ？次なんだったっけって言って調べられる

GPUが今後発展するなら、GPUに向いたアルゴリズムが開発されてくるでしょう。
でも、普通のCPUが高速化してくるので、結局特定の分野しか使えないのかもね。まさにハードウェアアクセラレータ。

全く依存性がなく並列に計算できるものに向いてるんだっけ？
あんまり計算用途には向いてない気もするんだけど。

実際高速化されてるものもあるんだし、得している奴はいるだろ

粒子法では，GPUかなり美味しいです＾＾

画像や文字のパターンマッチングにもそれなりに使えるよ。

実は一般ユーザーが恩恵を受ける事はほとんどない事実。。。

ATI臭くなってきたぞ，このスレ

ATIのGPUでCUDAが使えるならともかく使えないのに、
なぜかCUDAスレでATIの話題がされる不思議。

ATOM３３０搭載PCでCUDAなグラボを付けたら幸せになりますか？

>>908
気持ち次第です。

NVIDIAへの財務次第です

あきらかに落ち目でしょNVIDIA

落ち目(笑)

クアッドコアＣＰＵにＳＳＥを加えたらＧＰＵなんてまだまだごみレベルでしょう
本気で速度を比較した人がいないからＧＰＵが早いなんて嘘に騙されてるけど

ぷぷぷぷぷぷぷぷ




GPGPUがCPUに及ばないのとはまったく別次元の問題だ

8コアのCoreシリーズが普及価格帯に落ちてくればGPGPUなんてゴミだろう。

自分が井の中の蛙なのを知らずに吠えている可哀想な人がいますね。
「本気で速度を比較した人がいない」なんて本気で言っちゃってるんだもの。

>>915
8コアのCoreシリーズとか、比較するものによっては現行のCellでも勝てるんじゃね？

>>917
スレ違い。CELLの話をしたいのならCELLのスレに逝っとくれ。
つーか、CELLが勝てると言いたいのかCELLに勝てると言いたいのかよく判らんぞ。
# いや、CELLが何かに勝てるとは思えないのだが。

例えばリアルタイム3Dグラフィックスをやるなら
Core7iだろうがなんだろうが、Vooboo3にすら勝てないだろ
当たり前だけど

Vooboo...

邪教も随分かわいらしくなっちまうな。

適材適所ってことが分からない奴って、意外と多いんだな……

AMDはCPUとGPUの処理の切替えをするデモをしてたな
あぁいうのが理想的な気がする
適材適所という意味でも

個人的にはAMDはGPGPUの開発に遅れてる気がする。
ラデ使いに言わせれば、”どうせ使わない”らしいんだが。

適材適所だね。
GPUをフル活用できるようになれればかなり強そうだ。
レジスタとかシェーダーメモリとか。
CUDAのプログラムを書いていて未だに複数のGPUを使い分ける方法が分からない・・・。

GPGPUやCellは人間で言うならアスペルガー症候群みたいなもの。
得意な分野きわめて優秀な劣等生であっても決して優等生じゃない。
周りがお膳立てをしてやってようやく実力を発揮できるかもしれないが
大半の人は評価される前に匙を投げてしまう。

俺もアスペです。ごめんなさい＞＜

>>923
＞CUDAのプログラムを書いていて未だに複数のGPUを使い分ける方法が分からない・・・。
>57を理解した上で、MultiGPUだかのサンプル参照。
cutilにも同じようなラッパ関数があるからそれを見てもいいね。

>>925
もどって読んでみたら面白そうなことやってるね。
CudaOpenMPなんてのもあるみたいだし、OpenMPも勉強してみようかな。
分かりやすい本が出ることを期待してみるけど、ドキュメントとかも読んでみる。
参考にします。

>>925
> >57を理解した上で、MultiGPUだかのサンプル参照。
> cutilにも同じようなラッパ関数があるからそれを見てもいいね。

たぶん>>923はそれくらい理解してると思われ。もっと根本的な事かと

そういえば東工大のセンセが10月くらいにCUDAの本出すとか何とかどっかで見たような

東工大の学食は安くて上手いよな

近くに行くと利用してしまう

そういや学食って学外の一般人にも利用可なのか？

そりゃ当たり前だろ
大学は外部の人も出入り自由だし、
金払って食べるわけだからな。
企業の人事の人とかもよく学食で食ってるよ。

そういやシリコンバレーに居た頃は
暇な時は、勝手にスタンフォード大の講義聞きに行ってたなｗ
俺の出身大は日本のＦラン大なのにｗ

>>930
>>931
全然当たり前じゃない。
本来学生証の提示が必要な学食もたくさんある。
学校の福利厚生で学食が運営されている場合はそういうのが多い。
ただ、食券の自販機が普及してから、提示のタイミング自体がなくなってるから
利用できてしまうわけだが。

あと、外部の人も出入り自由とかはキャンパスによるから気をつけろよ。
出入りできてしまうと言うだけで、本来は自由じゃないところも多いから。

大昔、御茶ノ水にあった明大の師弟食堂ってのによく行ってたが、
土方のおっちゃんも作業着のまま普通に飯食ってたｗ

>>932
授業料払わず聞き放題っておいしいなｗ
そういや、スティーブジョブズもスタンフォード大に在籍してないのに
忍び込んで講義聞いてたんだっけ？

>>934
ああみえて二部に通う学生だったりするんだぜ。

20年前通ってた神楽坂の理科大の学食は糞不味かったなぁ．．．
飯が何故か塩素臭いし。
そういや、学食で秋山仁をよく見掛けたよ．．．

>>933
具体的にどの大学が出入り不自由なんだ？
そう言うところがあるというなら、具体例の一つぐらい挙げられるよな。

>>938
俺933じゃないけど、都内のビル系の大学とかそういう所多くね？
あと別にそこ突っかかる所じゃないだろう。 牛乳飲めよ。
あとスレ違いだよ俺たち。 牛乳板行こうよ。

女子大はマークされたりしないの？

慶應は入る気しなかった。
あのへん食うところ無さ杉。
結局MKビルの地下で昼飯食った。

弁当作ってくれる嫁がいないのが俺らの弱さ
財布を握る恐妻がいないのが俺らの強さ

ネタを振っておいてなんだが、
お前らどんだけ学食好きなんだよwww

学食が好きなのではない。安くて腹一杯になるのが好きなんだ！

ここは食うだのスレだろ！

>>945
ごめん。
素直に感心した。

あまりに感心したから今日、明治大学の師弟食堂いってくるわ。
今名前変わって、この時期昼しかやってないらしいけど。

CUDA2.3って何が変わったの？
誰か日本語で教えてください

ヴァージョンが変わった

おしいな。

版の付番が変わった

って言わないと。

マジレスすると日本の公式フォーラム見ろ>>862

CUDAのコードをどの開発環境で書いてますか？
Netbeansやeclipseを使いと思っていますが、makefileでのコンパイルなどは
簡単にできますが、CUDA独特のキーワード(dim3など)などからくるエディター上の
エラー表示を消すことができず、赤いエラー表示がエディターにあふれてしまいます。
Netbeansやeclipse環境でエラー表示を消す良い方法はないでしょうか？

四の五の言わずにVisualStudioにしなさい．
そのうちデバッガツールも公開してくれるから．

それeclipseスレ行った方がいいと思う

VS2008が一番使いやすいね
有志で作ってるウィザードが優秀だし

ちなみにNVのCUDAプログラミングイベントでもVS2008使ってた

nvccのコンパイルオプションで、
プラットフォームはどこに設定すればいい？x64とかx86とか。

リンクオプションにはあるんだが、
どうやらコンパイルの時点でx64用オブジェクトしか出てないから
x86用バイナリをリンクしようとするとエラーになる。

Visual Studio 2008のプロジェクトのプロパティページには
「CUDA Build Rule v2.3.0」ってのが出てるんだが、
プラットフォームの設定が見当たらない。

>>955
-m32 でどう？

>>956
帰宅したら試してみる。
さんくす

VSの有志ウィザード使ってやってる者なんだが･･･

初めてアトミックオペレーティングのオプション付けてプログラミングしてるんだけど，
ビルドするときの
リンクに入ってからの時間が異様に長いんだがこんなもんなのか？

>>958
そんなもんだろ
コンパイルミスしてんじゃないかとか考えちゃうけど、気長に待つしかなさそう。

VSウィザードが一番やりやすいし

初心者の質問で恐縮なんだけども，

float4とかのベクタ型の中身をホストで見たいって思ったら，
別にfloat変数をデバイスに作って，そこにベクタ型の要素の値をぶち込んで，
ホストにコピーしないと見れないのか？

ベクタ型って面倒なのかなぁ・・・

>>960
ホスト側でもベクタ型変数が使えるし、メモリ配置とアラインが合ってれば別に型は気にする必要無しです。

OpenCLやるならなるべくベクタ長16で
書いておいた方が良いんじゃね。

明らかにララビー想定した仕様だよね、これ。

>>961
たしかにホスト側でもベクタが使えますね．
どうもです．

CUDA+thrustで麻雀の思考ルーチン作った。
整数演算だけどなかなか速い。
GTX295搭載PC買うかな・・・

9800GTを１万で買ったがファンがうるさすぎ　ファンレスのヤツが欲しいな

ファンレスのは自前で冷却手段を用意してやるのが前提っぽいから9600GTの低消費電力版でも買えば?

CUDAプログラミングするなら素直にnvccを直接叩いたほうが良いね。
VSウィザードは正しくコンパイルできるものと、エラーが出るものがあるしね。

Nvidiaが公式でもうちょっと凄いの作ってくれたらいいな。
デバックとかもできる奴。

普通make使うよ

make使ったら負けかな

プップッ

>>969
makeを使ったら負けかなっと思っていた時期もありました。
いい思い出です。

ｆｍ、makeが一番よさそうだな。

>>967
VSアドオンのデバッガはβ出たよ（登録制）
ttp://developer.nvidia.com/object/nexus_features.html

Nvidia looking for new use for CUDA---Antivirus!!
ttp://vr-zone.com/forums/492420/nvidia-looking-for-new-use-for-cuda--antivirus-.html

CUDAを勉強できる本はありますか？
プログラムについてはもちろん、技術についての解説本などもあれば教えてください

>>975
ない、そろそろ偉い人が出すはずなんだけど、出てない。
ツールキットに付いてくるプログラミングガイドよむといいよ。

【GPGPU】くだすれCUDAスレ part2【NVIDIA】
http://pc12.2ch.net/test/read.cgi/tech/1254997777/

ベクタ型に関する質問をさせてくり．
ベクタ型(int4 ex;)とポインタ(int *p;)を作ったとする．
ポインタを
p = &ex.x;
と初期化したあとで，
ポインタをインクリメントすることで，ポインタで
ex.y
ex.z...
へアクセスできますかね？

>>976
ですよね…どおりで探しても見つからないわけだ…

自己解決しますた．

あぁ，ミスった．
980は978に対するレスです．
スレ汚しまくってスマン．

どう解決したか書かないから酷いスレ汚しだな