08/03/22 11:34:17
関連サイト
CUDAに触れてみる
URLリンク(chihara.naist.jp)
cudaさわってみた
URLリンク(gpgpu.jp)
CUDA のインストール
URLリンク(blog.goo.ne.jp)
NVIDIAの「GeForce 8800 GT(G92)」と次に控える64-bit GPUアーキテクチャ
URLリンク(pc.watch.impress.co.jp)
CUDAを使う
URLリンク(tech.ckme.co.jp)
NVIDIA CUDAを弄ってみた その2
URLリンク(dvd-r.sblo.jp)
CUDAベンチ
URLリンク(wataco.air-nifty.com)
KNOPPIX for CUDA Tutorials
URLリンク(www.yasuoka.mech.keio.ac.jp)
3:デフォルトの名無しさん
08/03/22 16:56:39
>>2
もうちょっとましなサイトはないの?
ブログってチラシの裏と大して変わらないのが多いのだが
4:1
08/03/22 17:21:30
>>3
すみません、ぐぐって適当に載せたのですが、良いサイトがあれば教えてください。
5:デフォルトの名無しさん
08/03/23 17:13:40
みんなGPGPUでどんなプログラム書いてる?
6:デフォルトの名無しさん
08/03/24 10:38:43
>>1
スレを分ける訳は? あっちのCUDAスレだけでいい希ガス。
7:デフォルトの名無しさん
08/03/24 14:45:08
teslaってどうなの?使ってる人いる?
8:デフォルトの名無しさん
08/03/24 14:51:29
TeslaC870はQuadroFX5600からアナログ回路周辺を取り除いた代物と思っていい。
性能的には、VRAMは広くて使い出があるけどそれ以外は8800GTXと大差ない。
副作用として、8800GTXなどだとオンボードGPUを切り離してしまうタイプのPCでは、
GraphicChipと認識されないC870は便利と言えば便利。
9:デフォルトの名無しさん
08/03/25 01:28:36
このスレの住人なら知っていますね、あの糞開発ツールのことを
・自分のプログラムのバグなのかコンパイラのバグなのかわからない
・他の仕事に応用できない糞開発ツールの独自世界を必死に学習している
・テキストエディタで書いたほうが効率的なのに糞UIツールを懸命に使っている
・糞開発ツールを批判すると「性格が悪いから糞ツールを批判するんだ」と言われる
糞だけど、政治的な理由で無理やり使わされているんですよね。
もう、あんな厨の作った糞ツールを我慢して使うのはやめましょう。
・糞開発ツールを部下に押し付ける上司の命令は無視しましょう。
上司は糞開発ツールが使われる実績を作ることであの会社のごきげんをとっているのです。
・糞開発ツールを使わせる上司の下では働けません、と上司の上司に直訴しましょう。
・あの糞開発ツール提供会社には「おたくの糞開発ツールは話にならない」と突き放しましょう。
バグレポートなどしてはいけません。改善要求などもってのほかです。
あの会社はあなたたちのことをテスター/モルモットとしか思っていません。
・あの会議で「糞開発ツールを使ったら生産性がxx%アップしました」
なんて話が出たら力強く机を叩き、会議室を出ましょう。
あの人たちは糞開発ツールをマンセーすることで立場を確保しているのです。
糞な開発ツールを糞だと言える、そんな当たり前の環境をみんなの力で取り戻しましょう。
10:デフォルトの名無しさん
08/03/25 01:55:19
AMDも来たね
URLリンク(ati.amd.com)
11:デフォルトの名無しさん
08/03/25 07:56:36
nVIDIAといいATiといい、Vistaの開発者をバカにしてるのか
12:デフォルトの名無しさん
08/03/26 08:33:46
1次元テクスチャって通常のfloatの1次元配列のように扱えるの?
13:デフォルトの名無しさん
08/03/31 13:04:11
sharedメモリはカードによって使える容量は違うのですか?
14:デフォルトの名無しさん
08/03/31 13:12:12
>>13
1block辺りの容量が決まっているから、結果的にGPUによって違う場合があることになりますな。
15:デフォルトの名無しさん
08/04/01 08:47:11
CUDAでトリップ検索したら早いの
128あれば128倍早そうな気がするんだけど
すぐもれちゃいそうで怖いんだけど
16:デフォルトの名無しさん
08/04/01 08:48:12
↑早いの?です
17:デフォルトの名無しさん
08/04/01 09:57:13
>>15
(比較的)単機能の(必ずしも)早くない演算器が(現在の最上位機種でも)256個あるだけなので、
単純に256倍速くなるわけではありません。
実際、4コアXeonよりも速いプログラムを作るのに苦労しているくらいですから。
18:デフォルトの名無しさん
08/04/03 16:05:16
AIも256倍賢くなるのかな?
19:デフォルトの名無しさん
08/04/04 15:15:42
コンパイルすると
Parsing error near '[': syntax error
というエラーが出ます。何がいけないのでしょうか?
20:デフォルトの名無しさん
08/04/04 15:28:37
[の近くの構文がいけないのでしょう
21:デフォルトの名無しさん
08/04/04 17:26:04
>>17
なるほど夢見すぎでした㌧
22:デフォルトの名無しさん
08/04/06 09:35:45
CUDAのプログラムを配布する場合、何を含めればよいのですか?
23:デフォルトの名無しさん
08/04/06 09:52:28
>>22
NVIDIAの対応ドライバとCUDAToolKitをインストールしてもらう前提で、実行モジュールだけ。
ToolKitのDLLが同梱できるかは不明。
24:デフォルトの名無しさん
08/04/06 12:39:02
nvccの-Oオプションて数字はいくつまであるの?gccと同じ-O3まで?
25:デフォルトの名無しさん
08/04/06 15:28:30
gccにそのまま渡すだけだと思うから、きっと3。
26:デフォルトの名無しさん
08/04/06 16:11:04
CUDAコードは最適化されないのですか?
27:デフォルトの名無しさん
08/04/06 16:14:09
されますよ。もう、徹底的に。
# まぁ、それほどドラスティックではないけれど。
28:デフォルトの名無しさん
08/04/07 15:26:08
cudaの実行速度を最適化したい場合、Oオプション以外にどのようなオプションがありますか?
29:デフォルトの名無しさん
08/04/07 16:20:09
頑張って自分で最適化してください。
30:デフォルトの名無しさん
08/04/07 16:22:20
-O3 -use_fast_math くらいかな。私が使っているのは。
GPU側のロジックは、実は-Oの有無では変わらないかも知れないと思う。
31:デフォルトの名無しさん
08/04/09 07:43:09
nvccのオプションで-O3を指定した場合、-Xcompilerで改めて指定する必要はないのですか?
32:デフォルトの名無しさん
08/04/09 16:50:10
ブロック数はどのように決めればよいのですか?
できるだけ多くするのがよいのでしょうか?
33:デフォルトの名無しさん
08/04/09 22:01:54
>>32
多くすれば良いと言う物でもないみたい
分割数を多くすることによるオーバーヘッドの方が大きくなるのか
良く分からないんだけど
34:デフォルトの名無しさん
08/04/10 17:46:35
いくつぐらいが適当ですか?
35:デフォルトの名無しさん
08/04/11 01:13:42
もう一つのCUDAスレにヒントがあるので参考にどうぞ。
36:デフォルトの名無しさん
08/04/15 06:31:17
単精度を2つ使って倍精度の計算をする様なことはできませんか?
37:デフォルトの名無しさん
08/04/15 08:45:23
実は現行CUDAにも一部倍精度演算があります。
が、まともに使うにはCUDAv2を待ちましょう。
38:デフォルトの名無しさん
08/04/16 00:59:13
今日やっと、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スレのどっちがどっちだか判らなくなっているのは内緒。
39:デフォルトの名無しさん
08/04/16 19:29:17
CUDAと心中しても大丈夫でしょうか?
40:デフォルトの名無しさん
08/04/17 00:46:36
心中する覚悟をするのは勝手ですが、実際に心中されたら大迷惑です。
41:デフォルトの名無しさん
08/04/17 14:43:17
CUDA 2.0(Beta)
URLリンク(forums.nvidia.com)
42:デフォルトの名無しさん
08/04/19 13:25:57
OpenGL系がダメダメだね。Vista。
ただ、nbody.exeはなぜか動いている。
43:デフォルトの名無しさん
08/04/19 15:53:33
> OpenGL系がダメダメだね。
付属のバイナリはダメだが。
ソースからビルドし直せば、動くぞ。
44:デフォルトの名無しさん
08/04/22 00:19:42
GPGPU#2
スレリンク(tech板)
188 名前:デフォルトの名無しさん[sage] 投稿日:2008/04/21(月) 19:17:35
Linuxはここだ
URLリンク(forums.nvidia.com)
どうでもいいけど明後日の仕事が増えた(謎
45:デフォルトの名無しさん
08/04/23 23:13:45
質問スレッドっていきてるの?
46:デフォルトの名無しさん
08/04/23 23:17:54
たった今回答して来ましたが何か。
47:デフォルトの名無しさん
08/04/29 19:05:47
Visual C++のバージョンはチェックしているみたいで、合っていないとコンパイルできませんが、
gccはバージョンが合ってないとコンパイルできないのでしょうか?
gccの最新版を入れても大丈夫なのでしょうか?
48:デフォルトの名無しさん
08/04/29 19:26:11
>>47
スレリンク(tech板:386番)
49:デフォルトの名無しさん
08/05/04 01:23:12
ローカル変数は、どのメモリに格納されるのでしょうか?
sharedメモリでしょうか?
50:デフォルトの名無しさん
08/05/04 08:40:58
レジスター
51:デフォルトの名無しさん
08/05/09 16:59:41
2つのGPUを同時に使いたい場合、
for (int i=0;i<2;i++)
{
CUDA_SAFE_CALL(cudaSetDevice(i));
somekernel<<< 64, 512>>>( );
}
のような形で書けばよいのでしょうか?
それともこれだと、一つ目のカーネルが終わらないと2つめのカーネルの実行が開始されないのでしょうか?
上記のように書いているのですが、実行速度が2倍になりません。
52:デフォルトの名無しさん
08/05/09 23:36:57
>>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();
53:51
08/05/10 08:54:52
>>52
ありがとうございます。
上記のようにやったのですが、
for (int ic = 0; ic < 2; ++ic) {
の部分を
for (int ic = 0; ic < 1; ++ic) {
に書き換えると、実行時間がちょうど半分になるので
やはり並列で実行されていないようです。
上記コード以外になにか付け加える必要があるのでしょうか。
54:デフォルトの名無しさん
08/05/10 09:22:20
>>53
CUDA_SAFE_CALLマクロは毎回cudaThreadSynchronize()を呼ぶから要注意。
55:51
08/05/10 09:29:14
>>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();
のようにしているのですが、逐次実行になってしまいます。
56:51
08/05/10 09:30:35
どこかにサンプルコードとかはないでしょうか?
57:デフォルトの名無しさん
08/05/10 11:34:03
cudaSetDevice()するとそのCPUスレッドで使うGPUデバイスが決定されるので、
それ以降の変更はできない模様。cudaSetDevice()しないでAPIを使うと、
無条件にDevice0を使うことにされてしまう模様。
CPUスレッドを分けるしかなさそうね。
NVIDIAのForumに行けばサンプルも転がっていそうだけれど……
58:57
08/05/10 11:35:55
あー、書くの忘れた。CudaProgramingGuide(1.1)の4.5.2.2参照で。
59:51
08/05/10 12:35:11
ストリームだけではなんともならないということでしょうか?
60:57
08/05/10 13:12:58
ならないっぽいですよ。ドキュメントを読む限り。
2.0betaのドキュメントも書き換わってないからそのままの悪寒。
61:デフォルトの名無しさん
08/05/10 18:52:45
geforce9以降もCUDA対応してるんでしょうか?
62:デフォルトの名無しさん
08/05/10 20:00:46
>>61
勿論。
63:51
08/05/11 19:48:55
>>60
そうですか、そうなると結構面倒ですね。
ところで、スレッド分けで2つのGPUを使う場合、pthreadなどでも良いのでしょうか?
それとも、サンプルにあるような方法のみが認められているのでしょうか?
64:デフォルトの名無しさん
08/05/11 20:10:22
>>63
ちっとはcutilの中身も見ようよ。cutStartThread()は、Linux版ではpthread_create()を呼んでいる。
65:デフォルトの名無しさん
08/06/17 19:19:34
GeForce GTX 280で倍精度を使うと、処理速度はどの程度落ちますか?
66:デフォルトの名無しさん
08/06/17 22:48:50
単精度の半分ぐらいでしょ。JK
67:デフォルトの名無しさん
08/06/18 00:54:10
いいえ、12分の1です。
68:デフォルトの名無しさん
08/06/18 05:03:23
GTX280 78GFLOPS
GTX260 62GFLOPS
69:デフォルトの名無しさん
08/06/20 13:29:07
うーむ。GPU2枚使うのって実はコツとかいる?
ホスト側でOpenMP使って2スレッド走らせて、それぞれのスレッドに別のGPUを割り当てて
るんだが、答えとしては正しいものが返ってくるんだが速くならない。
むしろ、ホスト1スレッドで1つのGPUで計算させたほうがいくらか速いくらい。
きちんとGPU2枚使ってそうだというのは確認できたんだが速くならない理由がわからない。
なんだかGPU0の処理が終わってからGPU1の処理が始まるとかやってそうな予感がしている。
GPUの切り替えの分だけちょっとだけ遅くなるというオチじゃないかと。。。
サンプルのmultigpuとか読むとホスト側のスレッド生成法は違うけど同じようなことやって
るんだよなー。
2枚使って速くなったって人いたらどういう風にやったか教えてくれませんか?
かれこれ2週間くらい悩み中だ。
70:デフォルトの名無しさん
08/06/20 15:38:55
>>69
pthreadでは速くなった?
開発環境は?
71:69
08/06/20 17:34:03
pthreadって要するにホスト側スレッド生成にcutStartThread()を使うってことだよね?
それは実はまだやっていないんだ。これからやってみようと思う。
とりあえずスレッド走らせればいいんだと思っていたから、使用経験のあるOpenMPを使った。
もしかして、cutStartThread()で生成したスレッドを使用して各GPUを使うっていうことが
暗黙の条件だったりしないよね。
cutStartThread()で速くなった人いますかね?
開発環境はVisualStudio2005。
72:デフォルトの名無しさん
08/06/20 17:40:20
うちでは速くなってるけど・・・
こちらはLinuxだけど、もしソースが載せられるならうpしてみて。
73:69
08/06/20 18:41:08
>>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);
74:69
08/06/20 18:42:06
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を使っている場所があるんだがこいつが悪さをしているとかあるのかな。
75:デフォルトの名無しさん
08/06/27 02:51:33
質問です(全く素人です)
OpenGLを使えばGPGPUのプログラミングが出来ると聞いたのですが、
かたやCUDAなどの専用の言語が出てきているのはどうしてでしょうか?
76:デフォルトの名無しさん
08/06/27 03:13:39
だったらcとか使わずfortran66でも塚っテロ
77:デフォルトの名無しさん
08/06/27 12:11:54
>>76
いえ、ですから、OpenGLではGPUの使えない機能もあるのかなぁ、と思いまして。
78:デフォルトの名無しさん
08/06/27 16:55:21
GPUで計算中でも画面は通常通り表示されますか?
79:デフォルトの名無しさん
08/06/27 22:56:10
>>78
ハードについて勉強するよろし
80:デフォルトの名無しさん
08/06/28 00:16:42
CUDA実行中に、時たま怪しげな影が飛んだりフリーズしたりします。
81:デフォルトの名無しさん
08/06/29 11:59:05
Folding@HomeのGPU版をしている人居ます?
PS3の280GTXは6倍以上をこなしているようですが。
ベータ版が公開されているが時間がなくてまだ何もしていない(困った。)
82:デフォルトの名無しさん
08/06/29 12:04:52
>>81
マルチすんな。つーか、PS3の280GTXってなんだよpgr
83:デフォルトの名無しさん
08/06/30 07:04:03
>>69
OpenMPの使い方が根本的に間違っていると思う
84:69
08/07/01 22:08:51
>>83
サンプルのcudaOpenMPを参考にしてやったんだけどなぁ。
これじゃどうだめなのかぜひとも教えてくれ。
ちなみにOpenMPの使い方はサンプルにあわせている。
ただsingle構文はサンプルでは使っていないがな。
85:デフォルトの名無しさん
08/07/02 04:05:38
VisualStudio2005はExpress?だったらOpenMPディレクティブは機能しないよ
86:デフォルトの名無しさん
08/07/02 15:57:17
nvidiaドライバを削除する方法ってどこかにのってないか@CentOS5
新しいグラボ買ったんだけどドライバの削除がわからないウンコー
87:デフォルトの名無しさん
08/07/02 16:15:29
grid, threadの意味がわからない
アニメに例えて教えろ
88:デフォルトの名無しさん
08/07/02 21:30:09
今amazonみたらGPUGems3の日本語版の予約を受け付けてるじゃn
89:デフォルトの名無しさん
08/07/03 01:44:22
ノート(OS:Vista)で動かそうとしても動きません
CUDAに対応したGPUを搭載しておりますし(8700M GT)
VC2005Expressをインストールしたのですが、コンパイルすると
nvcc fatal: nvcc cannot find a supported cl version.
Only MSVC 7.1 or 8.0.
と表示してしまいます。まったくわからなくて困ってます。
よろしければお教えください
90:デフォルトの名無しさん
08/07/03 01:47:35
2005以外のバージョンのVCも入れてる?
nvccが参照しているVCが2005じゃないように見える。
91:デフォルトの名無しさん
08/07/03 10:56:29
いれておりません。
インストールしたのは
VC2005EXとVC2005SP、VC2005SP Update for Vista
だけなのですが。
92:69
08/07/03 13:12:53
>>85
VisualStudio2005Professionalを使っています。
とりあえずcuda使わないで純粋にCでOpenMP使ったら高速化されているのでOpenMP自体は動いている模様。
93:デフォルトの名無しさん
08/07/04 21:25:01
URLリンク(vipprog.net)
94:デフォルトの名無しさん
08/07/08 02:11:24
>>93
なにこの糞サイト
95:デフォルトの名無しさん
08/07/08 02:42:51
質問です
PhysXがGeforceでも使えるということを聞いたのですが、
PhysX SDKとCUDAやOpenCLというのは同じレイヤーでの話なのでしょうか?
96:デフォルトの名無しさん
08/07/08 05:08:31
>>84
pthreadでは高速化された?
97:デフォルトの名無しさん
08/07/08 09:16:41
>>88
アノ値段じゃ学生はp-子するから売れん罠
98:デフォルトの名無しさん
08/07/08 11:01:42
Gem3はGPGPUに限って言うならあんまり役に立たなかった
99:デフォルトの名無しさん
08/07/08 19:02:12
>>86
--uninstall
100:86
08/07/09 13:42:42
>>99
ありがトマト
ところでSobelFilter見ててShareMemの使い方がわからないんだが
フォーラムに質問と返答あったが内容わからん+英語で全然わからん
誰か全訳してくれ
101:86
08/07/09 13:44:25
URLリンク(forums.nvidia.com)
誰か誰か!
102:デフォルトの名無しさん
08/07/10 12:22:28
>>86
訳すのはめんどいからパスだけど、ポイント絞って質問してくれたら回答するよ。
103:デフォルトの名無しさん
08/07/10 12:28:53
>>86
C:\WINDOWS\system32\drivers
から目的のドライバを手で削除すればよい。
104:デフォルトの名無しさん
08/07/13 21:12:38
>>102
神ktkr
ポイント絞るからちっと待ってて…
105:102
08/07/14 11:56:25
問題は、自宅が規制に巻き込まれていてなかなか書き込めない辺り。
取り敢えず、注意点を列挙しておく。
・共有メモリを確保するのはglobalFunc<<<blocks, threads, sharedMemorySize, streamNo>>>(parameters)の
三番目のパラメータでサイズが指定されたときだけ。
・共有メモリは一回の<<<>>>の呼び出しの間だけしか有効じゃない。
# つまり、次の回には残っていない。
・共有メモリをハンドリングするには、extern __shared__ anyType * nameで宣言するだけ。
# つまり、コンパイラは型のマッチングやサイズのチェックをしないので自分で管理しないといけない。
・共有メモリはblock間で独立、block内ではthread数に関わらず共有。
# つまり、実際のデバイスにそぐわないthread数を指定した場合はCUDA側で同期処理が入るのか、遅くなってしまう。
・あるthreadが共有メモリに書いた後、別threadが読む前には__syncthreads()で同期を取らないといけない。
# ある意味当然なんだけど、その所為で遅くなるのも事実。
あー、ついでにメモリの違いを簡単に。
※グローバルメモリ
・読み書きできる。coalescedにアクセスできれば結構速い。消えない。広い。
・ホスト側スレッドごとに独立している。ホスト側から見ると、毎回同じアドレスになるのでどのくらい使えるか判りにくい。
※共有メモリ
・読み書きできる。遅くない。
・呼び出しごとに消えてしまう。余り広く取れない。事実上同期を取る必要があって使い難い。
※定数メモリ
・速い。消えない。そこそこ広い。
・例えばfloat2を読み込むインストラクションがないので実はグローバルメモリからfloat2を読むより遅くなる場合もある。
・デバイス側から書き込めない。ホスト側スレッドごとに独立している。複数スレッドからCUDAを使うと毎回転送しなおすのか?
※レジスタ
・読み書きできる。速い。厳密に型チェックされる。つーか、型ごとに違うインストラクションが使われるからptxファイルで追える。
・呼び出しごとに消える。他のメモリに較べれば狭い。
# ローカルな配列は宣言したことないからよく判らん。
106:デフォルトの名無しさん
08/07/14 19:10:53
NVIDIAの仕様書見てもcoalescedの意味がいまいちわからないのだけど、どういうこと?
どこかわかりやすく解説しているサイトない?
107:デフォルトの名無しさん
08/07/14 19:12:13
カーネル内で__shared__つけて配列を宣言するのと、つけないで配列を宣言するのでは何が違うの?
108:デフォルトの名無しさん
08/07/14 19:24:11
>>106
GPUはWarp単位で同じインストラクションが走るから、要は16人17脚みたいに考えればいい。
メモリアクセスを16人17脚によるパン喰い競争みたいに考えると、自分のパンが目の前にある状態がcoalesced。
そのとき、2,3人パンを喰う必要がなくてもスルーするだけだから問題ない。
処が、二人のパンが入れ違っていたらそこで入れ替える間、みんなが待たされることになるって感じ。
# 判っている人には判るけど、判っていない人には判らない説明だなw
>>107
共有メモリを使うかどうか違うだけだと思うけど。ptx出力させて較べてみたら?
109:デフォルトの名無しさん
08/07/15 00:26:16
>>108
その説明、いただいてもいいですか?
110:デフォルトの名無しさん
08/07/15 01:56:27
>>109
本にするならもっと書かせてくれw
Vipのwikiに載せるなら是非やってくれ
金取って講習するのに使うのなら分け前よこせw
111:デフォルトの名無しさん
08/07/15 07:52:19
>>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メモリの使い方はどこを調べればわかるんだ!ウンコ!
112:デフォルトの名無しさん
08/07/15 08:45:54
>>111
どんな数字を入れるとどんな結果になるか、Excelでも使って計算してみたら?
0x3f使うのは64の倍数にするためでしょ。
113:デフォルトの名無しさん
08/07/15 16:57:48
>>112
うーんそのなんで64にするのかがわからないのよ
関係ないかもしれんがシェアードメモリを使ってないSobelFIilterも
SobelTex<<ih, 384>>>でなぜ384かわからないお
こっちは1行1グリッドにして、1スレッド1ピクセルなんだと思うが
なんでスレッド数を画像の横幅iwにしないで384にするんだぁ
スレッド数が384だと計算が速い理由でもあるのかお!
114:デフォルトの名無しさん
08/07/16 00:40:50
単純に、warp数の適当な倍数になるからってだけじゃなかろか。
115:デフォルトの名無しさん
08/07/16 18:09:02
GPUコードではmemcpyは使えないの?
116:デフォルトの名無しさん
08/07/16 21:23:53
>>115
デバイス側で、デバイス間のコピーをしたいってことなら、自前で書くしかないんじゃないかな。
でも多分、そこがボトルネックになると思う。
117:デフォルトの名無しさん
08/07/16 21:43:54
global memoryからshared memoryへのコピーの時間と
shared memoryからglobal memoryへのコピーの時間が
異なるのは何でなんだぜ?
よく分からない…
118:デフォルトの名無しさん
08/07/16 21:47:14
>>117
グローバルメモリへの書き込みはcoalscedでも遅いと思う。
そうでないなら、ptxファイル見てみないとなんとも。
プロファイラを使うともう少し様子が掴めるかも知れない。
そだ、プロファイラの使い方って、日本語で書かれたものがWeb上で見つからないんだよね。
誰か、まとめてない?
119:デフォルトの名無しさん
08/07/16 23:29:54
>>113
指定できるスレッド数の上限は合計512までだからだろ
もしくは、スレッド数を多くすると使用するレジスタ数がパンクするから
120:デフォルトの名無しさん
08/07/17 00:02:48
>>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();
}
121:デフォルトの名無しさん
08/07/17 15:32:10
NVIDIAのサイトから、lameをCUDA化したサンプルコードをダウンロードして、コンパイルしてみたのだけど、
かえって遅くなるのだけど、速くなった人いる?
CPUはAthlon2.0GHzでGPUは8800GTXです。
コンパイルはサイトの指示通り、USE_GPU_HPFを有効にしてかつ、そのCPUパートはコメントアウトして実行しないようにしています。
122:デフォルトの名無しさん
08/07/17 15:37:52
>>121
nVidia「全然速くならないから誰か代わりにやってくれ。速くても賞金出すのは北米在住者のみな^^」
というコンテストだから当然
来週締切なのに今のところまともに投稿してるのが1チームという惨状
123:121
08/07/17 16:05:38
賞金はともかく、NVIDIAの書いたコードだから、勉強になると思ったのですが
124:デフォルトの名無しさん
08/07/17 16:14:53
>>123
甘いな。ハードの設計している連中とドライバを作っている連中と
CUDAを作っている連中とそれの応用を作っている連中が勝手にやっているのがNVIDIAだ。
125:デフォルトの名無しさん
08/07/22 17:35:55
URLリンク(3.14.by)
URLリンク(forum.insidepro.com)
URLリンク(www.insidepro.com)
MD5 Crack on CUDA で腕試しが流行っているみたいだよ。
日本も頑張らなきゃ。
126:デフォルトの名無しさん
08/07/23 00:13:58
だれかフォートラン仕様にしてくれ
127:デフォルトの名無しさん
08/07/23 13:37:55
>>122
2チーム目が来たね
2倍ちょっとって、これならquad core使った方がよくね?
128:デフォルトの名無しさん
08/07/23 16:41:58
GTX280でCUDAすると素人考えで単純にSP数増えてる分だけ
速くなりそうな気がしてしまうんだが実際はどーなの?
Warpとかいろいろ考えることもありそうなんだけど。。。
129:125
08/07/23 19:13:07
>>127
英語読めていないよ。こんな簡単な読みも出来ないとなると
本当に困る。
Teamの人数が2人だよ。
登録が200を超え、提出者が20組 トップは6回目のUploadだ。
提出者が20組だから結構な盛況だと考えるよ。
130:デフォルトの名無しさん
08/07/24 00:13:47
NVIDIA、GPGPUに関する説明会を開催 - 今年中のFortranのサポートを計画 (2) CUDA2.0正式版は近日中にリリース予定
URLリンク(journal.mycom.co.jp)
2.0とマルチCPUサポートはそろそろ、マルチGPU・多言語・デバッガ・プロファイラサポートは多分おそらくきっとそのうちだってさ
131:デフォルトの名無しさん
08/07/24 02:24:04
cudaを覚えはじめるのにぴったりな本ないかなぁ?
132:デフォルトの名無しさん
08/07/24 11:03:32
CUDAに特化した本は未だないだろ。
133:デフォルトの名無しさん
08/07/24 12:24:23
>>131
本家プログラミングガイド
変な日本語訳のプログラミングガイドもある。
134:デフォルトの名無しさん
08/07/24 16:45:20
>>132
出たら買うから 出してw
135:デフォルトの名無しさん
08/07/24 17:17:55
じゃ、書くからNVIDIAと出版社相手の交渉は任せた。
136:デフォルトの名無しさん
08/07/24 22:53:31
何か並列処理の本のおまけみたいな形じゃないと
正直売れないだろうなぁ
137:デフォルトの名無しさん
08/07/26 17:15:12
>>135
たぶん、猫でもできる・・・・とかと同じように、HPたててそこで解説したら出版社から来ると思う。
わざわざHPスペース借りてーとかでなくても、WIKIでもいいと思われ。
交渉とかは向こうが書いてくれっていうだけだと思う基本的に。
んで、ワードなり、一太郎なりで原稿書けばOKじゃね?
画面キャプチャなどの画像は、まとめて管理したほうがいいぽいよ。
なんか、ワードとかで埋め込んじゃったりすると、逆に手間かかるらしい。
138:デフォルトの名無しさん
08/07/28 09:38:41
NVIDIAの日本語サイトがようやく
更新されたw
139:デフォルトの名無しさん
08/07/28 11:32:06
gpuはcpuより劇早なんだから
使いたいよねえ
パスワード解析とかにも使えそう
140:,,・´∀`・,,)っ
08/07/28 11:37:29
速くねーよ。
CPUをセダンに例えるならGPUは貨物車だな。
荷物を大量に扱う場合だけは効率がいい。
141:,,・´∀`・,,)っ
08/07/28 11:49:02
それはそうとパスワードクラックなんて浮動小数演算能力が全く役に立たないものの代表格だな。
142:デフォルトの名無しさん
08/07/28 12:08:10
そうだね
並列処理すればいいじゃない?
143:デフォルトの名無しさん
08/07/28 17:46:54
そういや去年CUDA使ってそんなソフトが出てたな
144:デフォルトの名無しさん
08/07/28 19:41:23
>>141
MD5はやっていますが?
他に現在手に入るクラック専用マシンでCUDA以上のC/P
があればお話が成り立ちますがね。
145:,,・´∀`・,,)っ
08/07/28 20:22:47
>>144
それやってるのは整数論理算術演算じゃん
乗算すら浮動小数じゃないよ。
つーかMD5なんてすでにハッシュアルゴリズムとしては死んでるんだが。
ワンタイムパスワードとしてならまだまだ現役でやれるだろうけど
性質上瞬時にクラックできないとクラックの意味もなくなる。
ああFPGA焼いた方がいいんじゃない?
今各社が力入れてるのはAESなどの128ビットブロック暗号だな。
146:,,・´∀`・,,)っ
08/07/28 20:35:07
つーか個人でGPUで遊ぶ分にはアリかもしれんけど
でかいシステム組めるTeslaになると別に高CPでもなんでもないんだよね。
特に整数性能に関してはPS3でクラスタくんだ方がまだ良いくらい。
GTX280含め現行GPUの整数性能はまだまだ残念な印象。
147:デフォルトの名無しさん
08/07/28 20:44:09
パスワードクラックほどの極小のプログラムでも
Xeon>PS3だって自分のページで言ってなかったっけ?
148:,,・´∀`・,,)っ
08/07/28 21:02:00
暗号関連ならCellはAESみたいなブロック暗号はそれなりに速いよ。
例のvpermもどき命令で1コア16並列でエンコード/デコードできるからね。
ただAESアクセラレーションはWestmereでの強化点にもなるのでこの先はわからん。
まあいずれにしてもCUDAはいろいろ残念だな。
ATiとのGPUシェア争い無視してでも汎用性能をとるか来たるLarrabeeとガチるかの瀬戸際だし。
149:デフォルトの名無しさん
08/07/29 09:58:39
>性質上瞬時にクラックできないとクラックの意味もなくなる。
笑った。連れの言い分と全く一緒だからな。
流れているデータをリアルタイムで解読できなければ暗号解読ではない
というような考え。ある意味では正しいが、ネットのデータは記録可能
という点を考慮していないし大穴になる。
AESアクセラレーションハもう6年も前にある技術だし今さら。
しかし認証段階で破られたらAESは即死に体だよ。実際解読ソフト
使っているのでww
150:デフォルトの名無しさん
08/07/29 12:28:50
別にCUDAは整数性能ウリにしてないし
このスレってムチャを前提にケチつけてるだけで結局何も出来ない無能者の集いだな・・・
151:デフォルトの名無しさん
08/07/29 12:59:27
並列演算の応用として暗号解読の話題を語ってるだけだと思うんだが・・・
152:デフォルトの名無しさん
08/07/29 13:50:22
>>151
それがわかってないといわれてる原因なんでしょ…
153:デフォルトの名無しさん
08/07/30 12:07:16
そうなんだ
154:デフォルトの名無しさん
08/08/04 01:58:13
弱点だった倍精度もGTX280で改良された。しかし何かイマイチなんだよなぁ。
消費電力あたりの性能がたいした事ないからかなぁ
155:通りすがり
08/08/07 11:07:41
CUDAでレイトレーシングがいました。 CPUの15倍パフォーマンスが出ているようです。
URLリンク(noridon.seesaa.net)
156:デフォルトの名無しさん
08/08/07 13:30:10
SSEも使ってない2コアの速度と比べて15倍じゃ今一かも
157:デフォルトの名無しさん
08/08/07 16:48:46
URLリンク(www.nicovideo.jp)
URLリンク(www.nicovideo.jp)
レイトレーシング
158:デフォルトの名無しさん
08/08/07 20:25:59
>>154
このスレ的には正常進化じゃね?
159:デフォルトの名無しさん
08/08/08 01:46:07
だれか4850にCUDA載せてくれ
160:デフォルトの名無しさん
08/08/08 14:04:29
Core2Duo(3GHz)×2 = 最大 1.89FPS
CUDA(GTX 280) = 最大52.48FPS
27.7倍高速出ているようです。
URLリンク(noridon.seesaa.net)
161:デフォルトの名無しさん
08/08/08 14:36:26
それatiのgpuじゃ実行できないよね?
162:デフォルトの名無しさん
08/08/08 18:06:10
CUDAですから
OpenCLはCUDAを採用したと聞いたが…本当なん?
163:デフォルトの名無しさん
08/08/08 18:18:50
OpenCommonLisp
164:デフォルトの名無しさん
08/08/08 20:41:23
4 x 2.5 GHz の 3 fps と比べても 17倍
これは大したものだ
165:デフォルトの名無しさん
08/08/09 00:50:56
IntelC++ と比べてどうなんかな?
166:デフォルトの名無しさん
08/08/09 07:05:08
うごかねえw
167:デフォルトの名無しさん
08/08/09 16:57:34
>>162
んなわけない
既にMacOSXにもCUDAがあるのにわざわざ10.6でOpenCL採用を謳うかよw
168:デフォルトの名無しさん
08/08/09 22:58:33
インテルは0からスタートするわけだからどのぐらいの速度でIntel製GPU普及するのかは見ものだね。 NVDIAのGPUはすでにゲームやCAD、3Dグラフィックの用途にすでにかなり普及しているからね。今後はAMDやIntelのQuadcore CPUでCUDAが使えるようになるしね。
169:デフォルトの名無しさん
08/08/09 23:11:40
SPがPentium-Mベースだから割と簡単に追いつけるんじゃない?
170:デフォルトの名無しさん
08/08/11 09:30:11
>>168
そのとおりだけど、intelには事実上CPUと抱き合わせ、という裏技があるからなあ。
今回はかなり本気っぽいし、微妙かも。
171:デフォルトの名無しさん
08/08/11 13:44:37
質問です。
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
172:デフォルトの名無しさん
08/08/11 14:36:44
アハハ
173:デフォルトの名無しさん
08/08/11 17:33:27
夏だな
174:デフォルトの名無しさん
08/08/11 18:07:18
俺は夏かどうか確認するために書き込んだわけではありません。
馬鹿正直にくだらない質問に答えてくれる人を待っているだけですよ。
質問に答える気がない人間は不要。
スレ違いだと言わざるを得ませんね。
175:デフォルトの名無しさん
08/08/11 20:14:53
夏だね
176:デフォルトの名無しさん
08/08/11 21:15:36
そんな失礼な態度のヤシに親切に答えてくれる人がどれだけいるか
待ってるより自分の頭で英語の意味を理解した方がずっと早いと思うがw
>error: 'memcpy' was not declared in this scope
>error: 'memcmp' was not declared in this scope
普通の頭の持ち主なら必要なincludeファイルがincludeされてないと解釈するだろうな
177:デフォルトの名無しさん
08/08/11 21:22:37
>>176
一つ言わせてもらう。
お前らに 「聞いてやってる」 んだよ。
聞いてほしくてたまらないんだろ?ニートで暇だから。
お前らが答えないなら別のやつに聞くだけ。
聞く相手はいくらでもいるんだよ。
178:デフォルトの名無しさん
08/08/11 21:35:15
>>176
ああ、ちなみにもう解決したから。自力で。
お前らみたいに大した知識もないくせにもったいぶった無能に聞いたのが間違いだったな。
ちなみに、ソースに手を加えずに解決した。
解決方法なんて書かないよ。
同じ問題に出会ったやつは ググレカスw
179:デフォルトの名無しさん
08/08/11 23:04:59
>>178
で、どこがCUDAに関係してたの?
180:デフォルトの名無しさん
08/08/11 23:42:51
凄いのが来たなぁ
釣りであると信じたいが天然厨房っぽいな
181:デフォルトの名無しさん
08/08/12 02:26:14
いや、質問主はまだ困ったままだろw
質問横取り煽り厨だろ
182:デフォルトの名無しさん
08/08/12 05:57:35
>>160
俺もやってみた
4コアCPUが必死に動いても8800GTに25倍くらいの差付けられた
レンダリングの革命やでぇ~
183:デフォルトの名無しさん
08/08/12 06:15:06
GPGPUレンダはまだ無いから実用性のあるの早く作って売れば金になるんじゃね
今の外部レンダは綺麗だけど時間かかって静止画までしか実用性無いからな
お前らちゃんすだぞ
絶対金になるだろうから、多分作ってるところはもう取り組んでると思うけどね
ていうか俺が早く出してほしいんだよ
184:デフォルトの名無しさん
08/08/12 22:18:14
gemsもそうだが、なんでOpenGLやGPU関連の書籍というのべらぼうに高いんだ?
185:デフォルトの名無しさん
08/08/13 17:25:11
>>184
需要がない割りに書くのが大変だから。
186:デフォルトの名無しさん
08/08/13 17:32:40
OpenGLは需要あると思うが。
187:デフォルトの名無しさん
08/08/13 21:20:05
>>184
訳してる会社が(ry
「誰が買うんだこれ!?」って思う値段の本は大抵原著は安かったりする
酷いのは原著の2倍以上
188:デフォルトの名無しさん
08/08/14 00:20:49
OpenCLはIntel製GPUでも使えるのだろうか。
189:デフォルトの名無しさん
08/08/14 01:38:59
>>188
Larrabeeのこと?
190:デフォルトの名無しさん
08/08/14 02:57:25
>>188
使えるでしょ
自社が参加してるワーキンググループが発表するんだから
191:デフォルトの名無しさん
08/08/14 13:46:01
URLリンク(www.gpureview.com)
うーむ……
192:デフォルトの名無しさん
08/08/18 09:18:31
URLリンク(jp.youtube.com)
GeForceでスパコン終了
一時間が30秒に
193:デフォルトの名無しさん
08/08/18 10:23:41
ゲームをやってた助手が「あ、どうも」、わざとらしくてワラタw
194:デフォルトの名無しさん
08/08/19 14:09:07
こんちにわ。LinuxでCUDAのプログラムを作成している者です。
早速質問ですがMPIを用いて、複数のクラスタ(PC5台、そのうち1台が親機で4台が子機)で
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
195:194
08/08/19 14:09:55
続きです。
上記のようなMakefileをNVIDIAのCUDA関連のBBSで見つけたので参考に作成しました。
しかしコンパイル時に下記のようなエラーが出てきます、、
/usr/bin/ld: cannot find -lcutil
collect2: ldはステータス1で終了しました
make: ***[linux/release/mpicuda_matrixMul] エラー1
上記のエラー内容を調べたところ、libcutil.aというライブラリがリンクされていない、
という内容と思われるのですが、Makefileには正しい場所にリンクさせてあると考えています。
開発環境として、OSがはCentOS5、親機のGPUがGeForce7900、子機が8600です。
(親機でCUDAの演算は行わないので7900)
それでは、どなたかよろしくお願いします。
196:デフォルトの名無しさん
08/08/19 20:40:53
>>195
cutilはmakeした?
した積もりなら、ちゃんとライブラリファイルはできている?
OpenGLがインストールされていない環境だと、恐らく失敗している筈。
その場合、OpenGLをインストールするかOpenGL絡みのエラー修理辺りをコメントアウトすればmakeできるはず。
197:デフォルトの名無しさん
08/08/19 23:40:11
MPIでCUDAを試したことは無いけどGLUT,GREWがインストールしてないと
>>195のようなエラーが出たことがある
198:デフォルトの名無しさん
08/08/20 00:48:28
なあ、あのへんてこ日本語プログラミングガイドのリンクが消えているような気がするんだけど、
前からこんなんだっけ?
199:デフォルトの名無しさん
08/08/20 00:50:02
いいことだw
200:デフォルトの名無しさん
08/08/20 20:57:08
URLリンク(jp.youtube.com)
530万ドルのスーパーコンピュータ「CalcUA」で一時間
一台デスクトップPCで30秒
URLリンク(www.atmarkit.co.jp)
ベルギーのアントワープ大学では、それまで使っていたAMDベースの256ノードのクラスタサーバ「CalcUA」の性能を、8GPUを使ったデスクトップPCシステム「FASTRA」が上回った。「CalcUAは530万ドルのスーパーコンピュータ、FASTRAは7000ドルのデスクトップだ」(グプタ氏)。
NVIDIAによればGPUを使った同様のクラスタサーバは、NCSA、イリノイ大学、ノースカロライナ大学、マックスプランク研究所など、すでに十数の組織で使われているという。
201:194
08/08/21 15:43:36
>>196 さん 、>>197 さん
お返事ありがとうございます。
cutilのライブラリは確認済みです。CUDAインストール時に作成される、
~/NVIDIA_CUDA_SDK/lib の中にlibcutil.aがありました。
OpenGLは前からインストールしており、またそのプログラムは動くので、
GLUT等はインストールされているはずなのですが。。
202:デフォルトの名無しさん
08/08/22 00:33:34
/home/usr1/NFS/NVIDIA_CUDA_SDK/と~/NVIDIA_CUDA_SDK/libが違うって落ちじゃないよな。
203:デフォルトの名無しさん
08/08/22 01:02:50
自分からテスト専門です、って宣言してる派遣テスターって何なの?
将来プログラマとかSEになりたい、とかならわかるけど。
向上心ないよね、頑固だし。
そういう派遣テスターって、仕様書は読めない、
テスト仕様書も作れない、テストプログラムも作れない
やれることは「テキトーにプログラムを触る」ことだけ。
俺は派遣プだけどさ、こういう派遣テスターがいると
派遣全体がバカにされるんだよ。
テスト専門派遣なんて氏んで欲しいよ、まったく。
今日も正社員の人が派遣テスターに仕様書を読んで
テスト仕様書を作ってください、って説教してたよ。
その派遣は頑固に「何故、仕様書が必要なんですか?」って
反論してたから、きっとテスト専門派遣テスターだな。
仕様書も読まず、テスト仕様書も作らず、ただテキトーに
プログラム触るだけで給料もらおうなんて頭おかしいんじゃねーの?
あ~あ、あの派遣テスターが3ヵ月後に切られるまで、
仕様書も読まねーでテキトーにテストしたバグ票がまわってくんのかよ。
そんな糞なもん、読んで処理する派遣プの身にもなってくれよ。
うわ~、しかもそいつが切られる3ヵ月以内に中間納品あるじゃねーか!
テスト仕様書もなしにテキトーにテストして納品か。
中間納品後にソッコウクレームでデスマ必至だな。俺の休みも返上かよ。
派遣専門テスターさんよ、少しは向上心持てよ!
頑固な性格直して仕様書読めよ!テスト仕様書作れよ!
204:デフォルトの名無しさん
08/08/23 14:37:16
うちの客先はテスト専門部署があるけど
環境構築とかテストプログラム作ったり
海外のテスト専門会社へ行って研究したりしてるぞ
205:デフォルトの名無しさん
08/08/23 18:12:10
バカの作ったものをテストするなんて最悪な仕事の1つなんだぜ
206:デフォルトの名無しさん
08/08/23 18:24:28
以前入社してすぐやめた会社でテスターをやらされたが、
そこの会社、3Dの知識ないやつが、OpenGLで懸命に2Dの描画プログラム作ってた
まさにバカの作ったものでテストするのがバカらしくなって会社を辞めた
207:デフォルトの名無しさん
08/08/23 19:33:34
判ったからマ痛に行け。
208:デフォルトの名無しさん
08/08/23 19:54:30
>そこの会社、3Dの知識ないやつが、OpenGLで懸命に2Dの描画プログラム作ってた
こういうバカ、たまにいるよな
2Dだけ描画すればいい仕様なのにOpenGLしか知らないからとりあえずOpenGLで書いちゃうという
すぐに辞めさせないとプロジェクトが悲惨な結果になる
209:デフォルトの名無しさん
08/08/24 01:26:41
>>208
AutoCADは2DもOpenGLで描画していますが。
210:デフォルトの名無しさん
08/08/24 07:16:25
>>209
>207
211:デフォルトの名無しさん
08/08/24 10:49:13
3Dも描画するならいいだろ
2Dしか描画しないのにOpenGLはバカ
212:デフォルトの名無しさん
08/08/24 12:06:13
GDIで描画したらLinuxでWineが必要になっちゃうじゃん
213:デフォルトの名無しさん
08/08/24 12:27:07
いや、2Dだけ描くのにOpenGL使うようなバカはWindowsだけだろ
214:デフォルトの名無しさん
08/08/24 17:34:26
このスレでうだうだ3D話を続ける奴って頭弱いの?
それとも只の構ってちゃん?
しかも切っ掛けはコピペだって辺りが痛過ぎるよね。
215:デフォルトの名無しさん
08/08/24 17:41:37
きっかけはコピペでのバカの作ったものをバカにされるべき
GPU使っているはずなのにBITBLTしかしてないとかー
もーバカ杉
216:デフォルトの名無しさん
08/08/24 21:55:53
BITBLTの後ろ三つはベーコン・レタス・トマトの略だが、前のBITは何の略なんだ?
217:デフォルトの名無しさん
08/08/24 22:58:34
マルチプラットフォームで2Dを描くのに仕事で使えるレベルの表現力と速度を兼ね備えたライブラリって存在するの?
あればどんなものか教えてください。有償無償問わず。
218:デフォルトの名無しさん
08/08/24 23:02:35
>>217
OpenGL
219:デフォルトの名無しさん
08/08/24 23:23:16
マルチプラットフォームの仕事なんて普通ないだろ
プロならプラットフォームに最適化されたライブラリを選択しろよ
220:デフォルトの名無しさん
08/08/24 23:34:53
OpenGL専門学校の学生さんにマジレスするなよ
221:デフォルトの名無しさん
08/08/25 00:54:10
>>219
>マルチプラットフォームの仕事なんて普通ないだろ
?
222:デフォルトの名無しさん
08/08/25 00:56:55
>>217
URLリンク(d.hatena.ne.jp)
223:デフォルトの名無しさん
08/08/25 01:00:30
>>221
これから増えるだろう
お前は近視か?
224:デフォルトの名無しさん
08/08/25 01:01:51
間違った。>>219だ。
225:デフォルトの名無しさん
08/08/25 01:06:43
>>219
openofficeはいったい何なんだ?マルチプラットフォームのソフトウェアじゃないのか?
226:デフォルトの名無しさん
08/08/25 02:40:58
お前は仕事でopenofficeにかかわったのか?
違うだろ
派遣先のプラットフォームだろ
お前の会社がマルチプラットフォームの製品を開発してんのか?
違うだろ
派遣先のプラットフォームだろ
現実見ようぜ
227:デフォルトの名無しさん
08/08/25 02:51:05
>>226
まず、仮定からおかしい。
・マルチプラットフォームのソフトウェアはopenofficeだけではない。
・>>225の仕事の内容を断定する根拠が見当たらない。
・なぜ無数にある可能性から>>226の仮定を導き出したのか不明。
説明をお願いします。
228:デフォルトの名無しさん
08/08/25 02:56:05
>>227
>・マルチプラットフォームのソフトウェアはopenofficeだけではない。
まず仮定からおかしい。
実際に仕事でマルチプラットフォームを開発しているのかを問題にしている。
>>225がopenofficeを開発したというのなら謝る。おみそれした。
>・>>225の仕事の内容を断定する根拠が見当たらない。
まず自分で開発してそうにないopenofficeを出してくるあたり、
ただのバカ=派遣ってこと
>・なぜ無数にある可能性から>>226の仮定を導き出したのか不明。
派遣に可能性なんてない
229:デフォルトの名無しさん
08/08/25 03:04:17
荒らしは放置の方向で。
230:あぼーん
08/08/25 08:29:00
219のほうがバカだと思います。
マルチぷラットフォームの仕事はあると思います。
231:デフォルトの名無しさん
08/08/25 09:16:36
だからお前はマルチプラットフォームの仕事をしたことあるのかと
232:デフォルトの名無しさん
08/08/25 11:23:15
サターンとプレステで同じゲームを同時に作ったぞ
同時に作っただけだけど
233:デフォルトの名無しさん
08/08/25 12:41:48
>>231
Firefoxはマルチプラットフォームじゃないのかね。
234:デフォルトの名無しさん
08/08/25 12:43:22
マルチプラットフォームの要件はオープンソースではかなり多い。
オープンソース系の企業ならそういう仕事していてもおかしくないね。
235:デフォルトの名無しさん
08/08/25 12:48:10
日本語読めないの?
聞かれてるのは「お前は仕事でマルチプラットフォームをやったことがあるのか?」だぞ
236:デフォルトの名無しさん
08/08/25 12:51:51
>>235
話のすりかえ。
俺自身は学生時代にFirefoxの開発に携わったが、
今の話題は「マルチプラットフォームの仕事があるかないか」。
お前が間違っていることに気づけ。
237:デフォルトの名無しさん
08/08/25 12:53:16
そもそも、ここの住人が開発に携わったことがあるかどうかなんて誰も興味を持っていないし、
信憑性もない。
もっと有益は議論をすべきだと思うけどね。
238:デフォルトの名無しさん
08/08/25 12:54:24
スレ違い お前らネットイナゴは佃煮にされて死んでしまえ
239:デフォルトの名無しさん
08/08/25 12:58:38
派遣先企業で派遣先企業のプラットフォームに合わせたものを実装するのがおまえら派遣プログラマの仕事だろ
派遣先正社員からマルチプラットフォーム対応なんて指示されてないだろ
くだらないこと考えて余計な工数使うなよ
240:デフォルトの名無しさん
08/08/25 13:13:32
>派遣先企業で派遣先企業のプラットフォームに合わせたものを実装するのがおまえら派遣プログラマの仕事だろ
派遣先がパッケージベンダってこともあるだろう
なにムキになってるの?
241:デフォルトの名無しさん
08/08/25 13:23:54
>>239
何故派遣にこだわる?
242:デフォルトの名無しさん
08/08/25 13:49:24
ID出すとしょうもない素人以下の奴がいろんなスレに粘着してた
243:デフォルトの名無しさん
08/08/25 14:17:22
>>242
この板はID出ないじゃん
244:デフォルトの名無しさん
08/08/25 14:18:21
なんだみんな派遣だったのか
245:デフォルトの名無しさん
08/08/25 15:59:39
派遣に異常なコンプレックスを持っている人がいるようですが、
池沼の方なので基本放置でおねがいします。
246:デフォルトの名無しさん
08/08/25 16:06:10
夏の間だけ2chもアカウント認証式にすればいいと思う。
247:デフォルトの名無しさん
08/08/25 21:59:49
正社員にコンプレックスでしょ?
派遣のどこにコンプレックスをもてと
248:デフォルトの名無しさん
08/08/25 23:45:38
ひどい脱線ぶりだな。
249:デフォルトの名無しさん
08/08/25 23:49:36
派遣野郎はスレ違い
250:デフォルトの名無しさん
08/08/26 00:13:24
>>232
それありそう。PS3 と Xbox360 とか。どう似せるかに力が注がれる・・・
251:デフォルトの名無しさん
08/08/26 00:59:55
ここはなんのスレなんだ
252:デフォルトの名無しさん
08/08/26 04:09:43
くだすれ
253:デフォルトの名無しさん
08/08/26 06:23:56
CUDAはマルチプラットフォーム。
254:デフォルトの名無しさん
08/08/26 07:35:34
アホかい。
255:デフォルトの名無しさん
08/09/25 14:29:37
━━━━━━━━━━━━━━━━━━━
■ 今週の名言
───────────────────
「HPCは“ハイ・プロダクティビティ・コンピューティング”」
日本SGIのHPC・サービス事業本部で本部長を務める田坂隆明執行役員)
256:デフォルトの名無しさん
08/09/28 22:50:51
CUDAってSIMD並列といいうことは、過去のスパコン用のコード
移植したら激速ってことでOK?
257:デフォルトの名無しさん
08/09/29 11:43:54
>>256
SIMDじゃないよ。仮にそうだとしても、そんなに単純な話じゃない。
258:デフォルトの名無しさん
08/09/30 05:16:33
SIMDだよw
259:デフォルトの名無しさん
08/09/30 16:35:15
SIMDじゃないなら、何なのか問いたい。
260:デフォルトの名無しさん
08/09/30 17:38:27
SIMDだったりMIMDだったり
261:デフォルトの名無しさん
08/09/30 21:12:21
実行はSIMDだけど
コードはスカラだから面倒臭いよ
262:デフォルトの名無しさん
08/10/01 10:37:52
NVIDIAはSIMTと言っているね。SIMD+αくらい。
オンボメモリはSX8i並だからやっぱ速いんじゃねーの。
263:デフォルトの名無しさん
08/10/01 13:58:23
理屈はいいからなんか作れよ
>>過去のスパコン用のコード移植したら激速ってことでOK?
OKの一言で終わることをなにぐだぐだやってんだか
264:デフォルトの名無しさん
08/10/01 18:46:54
要するに、「移植できるもんならしてみろ」ってことだろw
265:デフォルトの名無しさん
08/10/01 23:50:07
だね。単純に移植できるとは思わない方がいい。
266:デフォルトの名無しさん
08/10/02 00:27:50
たしかにブロック切り分けの所もう少しスマートにならないのかと思うのだが
あんなもんデバイスの種類で最適化した組み合わせを自動算出すりゃいいだろ
267:デフォルトの名無しさん
08/10/02 00:29:52
>>266
CUDA2でストリームプロセッサ数なんかも取得できるようになったから、
是非とも自動算出関数を作ってみてくれ。
268:デフォルトの名無しさん
08/10/02 01:08:48
いや作るのは簡単なんだけどCUDAってクラスも使えないしラップ出来ないからだめじゃん
269:デフォルトの名無しさん
08/10/02 01:11:17
あいや出来るのか
使うのなら作るぞ
270:デフォルトの名無しさん
08/10/06 20:37:38
CUDAに興味があっていろいろ調べてるんだけど
URLリンク(slashdot.jp)
こんなにメモリアクセスがきついのか…
へたれの俺はCPUで頑張ろうって気になってきた
このあたりの条件が和らぐのはどれくらい後だろう?
271:デフォルトの名無しさん
08/10/06 21:10:43
それでもビデオカードは一世代速いメモリを積んでいるから、レイテンシを誤魔化せれば、CPUより多少はいいんじゃないかな。
CPUもGPUも、演算能力よりメモリのスピードがネックなんだよねえ。
272:デフォルトの名無しさん
08/10/06 22:18:57
>>270
そうそう、グローバルメモリはコヒーレントな読み出しなら4クロックなのに
ランダムアクセスすると100倍のクロックを喰ってしまう。
スラッシュドットの記事の書き方はちょっと言葉が足りてなくて、
実際にはいかにメモリをコヒーレントな読み書きで済ませるかが鍵ということになる。
例えば、CUFFTの2Dを1000回回すのと1Dバッチと自前の転置を2回ずつ動かすのとでは、
64x64辺りだと余り変わらないのに128x128になると途端に2Dの方が速くなる。
プロファイラで調べると、CUFFTの内部で用意されている転置関数はコヒーレントな読み書きしかしていない。
一方、自前の転置はコヒーレントな読み出しとインコヒーレントな書き出しになっている。
つまり、1000回呼び出すコストをインコヒーレントな書き出しコストが上回ってしまうということだ。
273:デフォルトの名無しさん
08/10/07 01:53:43
CPUより遅くなるということはない
電気代の割りに効果が少ないというのはあるw
274:デフォルトの名無しさん
08/10/11 06:33:45
TMPGEncで軽い処理はCPUより遅い。
VRAMとのメモリ転送あるし。
275:デフォルトの名無しさん
08/10/15 22:02:44
メインメモリと共有するタイプのGPUで処理した方が速い事も状況によってはあるのかな
276:デフォルトの名無しさん
08/10/15 22:26:03
CUDAでは有り得ないからスレ違い。GPGPUスレにでもどうぞ。
277:デフォルトの名無しさん
08/10/26 18:43:49
すいません、超低級な質問です。
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
278:デフォルトの名無しさん
08/10/26 20:58:06
>>277
>ホスト側のスレッド数が128個生成されるみたいです。
そう判断した根拠は?
それとは別になるが、128x128が妥当かどうかは検証が必要だと思う。
279:デフォルトの名無しさん
08/10/27 17:11:09
質問です。
__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」
となるんですが、これはどういうことなんでしょうか。
280:デフォルトの名無しさん
08/10/27 17:47:00
__global__な関数から__device__の関数を<<<>>>で呼ぶことはできない。
281:デフォルトの名無しさん
08/11/01 09:40:52
ちょっと質問させてくらさい。。
8ストリーミングプロセッサ(SE)単位で構成されている
ストリーミングマルチプロセッサ(SM)でのことなのですが、
グローバル(ローカル)メモリ上へのデータのロード、
ストア命令は100サイクル以上かかってしまいますよね。
その間、8つのストリーミングプロセッサはおお休みしている
のでしょうか?
それとも、他のWarpのインストラクションが割り当てられたり
するのでしょうか。そうであるといいのですが、
そうするとSharedメモリがあとからのWARPにのっとられてしまって
まずそうですよね???
レジスタ的には32bit×2048本あるようなので(280の場合)
OKそうですが、Sharedメモリは16KBytesしかないし。
実際LD、STに160サイクルぐらいくったら、
160命令ぐらい無駄にしちゃって効率をかなり落としてる気がするので、
なんかやです。
24WARPぐらい、バッファリングしていて、割り当てられるWARPを
out of oderで実行するという記事も、どっかで見たような
気がするのですがSharedメモリがじゃまして無理があるような・・・
なんか良い方法があるのかな???
っていうか、なにか根本的な部分で俺の勘違い?
282:デフォルトの名無しさん
08/11/01 10:12:52
・グローバルメモリアクセスは、最大400(?)クロック掛かるが、最短では4クロックで済む。
# そのためには、coalescedにアクセスできるように工夫する必要がある。
・各ストリーミングプロセッサは、独立して動作する。Sharedメモリも同様。
例えば、行列の転置のような処理の場合、普通に書くとcoalescedに読んでincoherentに書かざるを得ない。
# 或いはその逆か。
そこで、CUFFT内で行なっている転置処理では、(プロファイルで見る限り)一旦共有メモリにおいて同期を取ることで、
読み書き共にcoalescedアクセスを維持しているようだ。
283:デフォルトの名無しさん
08/11/01 18:40:52
>>282
早速のレス、さんくす。
なるほど、最短4クロックですか。coalescedにしてもレイテンシー
だからだめかなって思ったけど、よくよく考えると、DRAMへの直接
アクセスに数百クロックってのはおかしいことに気づいた。
各ストリーミングプロセッサは、独立だけど1インストラクション単位で
同期してるんじゃないんかな(んなことぁない?)
・・・っておもったけどブランチがあるから、TPXレベルから見ると独立か???
最初見たときSharedメモリも他のSPのregも0レイテンシで使える
ようだったんで、演算と独立にグローバルメモリとのロード、ストアができん
16K程度のSharedメモリ、あまり意味ないじゃんとかおもったけど、
そんなことなさそだね。
なかなか面白そうなチップだ。
分岐粒度が32って実際、汎用でどうなんだろ。
ベクトルつかったことないんでわからんが、
適当にセルオートマタでも乗っけて遊んでみまつ。
284:デフォルトの名無しさん
08/11/12 22:42:34
CUDA-Zなんて便利なものがありました。
forum.nvidia.co.jp
285:デフォルトの名無しさん
08/11/12 22:55:32
>>284
kwsk
286:デフォルトの名無しさん
08/11/12 23:19:24
URLリンク(cuda-z.sourceforge.net)
287:デフォルトの名無しさん
08/11/14 19:39:30
あげてすまん。
CUDAのためにGTX280つきのPCかったのだが
ほかにもう一枚グラボ買わないといけんのかな?
あとPCIバスしか残ってないのだが・・・
LINUXで使う予定。
288:デフォルトの名無しさん
08/11/14 22:57:10
>>287
別に一枚でも大丈夫。二枚挿しても(開発には)却って面倒だったり。
まぁ、二枚あれば表示に足引っ張られる心配なくなるから安定はするけどね。
289:デフォルトの名無しさん
08/11/15 00:41:03
>>288
そうなのか、安心した。
ありが㌧
290:デフォルトの名無しさん
08/11/15 17:39:36
質問です。
float にてゴリゴリ計算して、結果を返すプログラムを書いてみたのですが、
普通のアルゴリズム、CUDA+GPU、CUDA+CPU(deviceemu) の2つで比較して
みたところ、思ったより差が大きいのですが、こんなもんですか?
もしくは、機種依存があったりするんでしょうか?
291:290
08/11/15 17:40:23
>>290
すいません、間違えました。
「2つで比較して」->「3つで比較して」
292:デフォルトの名無しさん
08/11/15 19:43:01
何の差?
演算精度なら、そりゃぁあるさ、floatなんだもの。
293:デフォルトの名無しさん
08/11/15 20:51:43
ホストは、デフォルトではfloatをfloatで計算してくれないからな。
SSEを使うなりで、ホストがちゃんとfloatで計算すれば、結果は一致するんジャマイカ?
294:デフォルトの名無しさん
08/11/15 20:57:47
>>293
一致する保証はないよ。CUDAのドキュメントにもあるけれど、超越関数はGPU内部の組み込み版を使うと若干誤差が残る。
いずれにしてもfloatの想定の範囲内だから、実用上は問題にならないけどね。
295:,,・´∀`・,,)っ-●◎○
08/11/15 22:26:10
x87 SSE CellB.E. CUDA の浮動小数サポートの対比表みたいなのがCUDAのマニュアルにあったな。
確かに完全じゃない。
糞まじめに準拠してるのはx87くらいだ
296:290
08/11/16 02:56:28
>>292-295
なるほど。出てきた結果が、CPU上の double で計算した結果と、
GPUの float で計算した結果が、最大1%程度違ったから、正直
驚きました。普段doubleしか使って無くて、誤差なんかほとんど
気にしなくてよかったので。
誤差が蓄積してくようなタイプのアルゴリズムではないと思っていた
だけに、少し驚きました。
297:デフォルトの名無しさん
08/11/20 18:53:51
Cで使っていた自作ライブラリは、
nvccでコンパイルし直さなきゃダメなの?
298:,,・´∀`・,,)っ-○◎-
08/11/20 18:58:40
食べてみたけどおいしかったよ。
一つどうぞ、あーんして
299:,,・´∀`・,,)っ-●◎○
08/11/20 19:44:47
飾りだ、食うなボケ。
300:,,・´∀`・,,)っ----
08/11/20 21:36:00
僕のもうないです
>>299から貰ってね
301:デフォルトの名無しさん
08/11/20 23:10:32
>>297
ホスト(CPU)上で実行させるものならVC++と一緒にリンクできる。
(cppIntegrationサンプル参照)
GPU上で実行させたいならそれなりにいじくらないと無理。
302:デフォルトの名無しさん
08/12/01 15:16:29
-deviceemuでは動くプログラムが、GPUを使うと誤動作します。
__device__関数に引数が一部しか渡らなくなります(float4のz成分だけしか渡せない)
ループ回数を極端に減らすと改善されるのですが、これはregisterメモリがパンクしてしまったということでしょうか?
303:デフォルトの名無しさん
08/12/01 17:15:56
ソースを見ないとなんとも言えませんが、ループ回数とregisterは依存関係にはありません。
nvcc -ptxでptxファイルを出力して眺めてみては如何でしょう。
304:デフォルトの名無しさん
08/12/01 23:02:30
ptxファイルですか。。。
正直どこをどう見たらいいのかわからないので敬遠していましたが、
遂に避けられないとこまで来てしまいましたかねえ。
見るべきポイントとかありましたら、よろしければ教えてください。
305:,,・´∀`・,,)っ-●◎○
08/12/01 23:09:56
ptxもネイティブコードじゃなしに中間コードの断片だからな。
レジスタが何本割り当てられるかとかそういった情報すら持ってないから困る。
その点Larrabeeは16本+16本+αとかレジスタ本数が決まってて
コード生成時点で割り当てが決まってしまう。
結果的に静的な最適化がしやすい。
良し悪し。
306:デフォルトの名無しさん
08/12/02 13:02:50
リョウテニ●ヽ( ・∀・)ノ● CUDA!
307:デフォルトの名無しさん
08/12/03 17:58:11
.oとかにしてgccでリンクすることってできますか?
308:デフォルトの名無しさん
08/12/04 00:09:51
>>307
何を? どこで? なんで質問もろくにできないの?
ちなみに、エスパー募集ならお門違い。
309:デフォルトの名無しさん
08/12/04 00:52:30
>>307
普通にできますよ。
-l/usr/local/cuda/lib -lcudart
等のオプションは当然自分で追加する必要があります。
310:デフォルトの名無しさん
08/12/04 06:13:34
Windowsで、とかいう話だったりしてね。
311:デフォルトの名無しさん
08/12/04 10:05:39
.objじゃなくて.oって書いてあるんだからそれはないんじゃない?
312:デフォルトの名無しさん
08/12/04 22:31:14
それだったらわざわざ聞くかねぇ。まぁいいか。
313:デフォルトの名無しさん
08/12/12 18:47:36
【GPGPU】CUDA/ATI STREAM 速度・画質検証スレ
スレリンク(jisaku板)
314:デフォルトの名無しさん
08/12/14 19:30:29
これからCUDAを始めようといろいろサイトを巡回しているのですが、
グリッドサイズとブロックサイズはどのようにして決めたらいいのでしょうか?
同期をとるかとらないかで変わると思いますが、とりあえずとらない場合でお願いします。
315:デフォルトの名無しさん
08/12/14 20:41:30
>>314
BlockSize(スレッド数)は32の倍数で余り大きくない辺り。GridSize(ブロック数)は充分大きく。
316:デフォルトの名無しさん
08/12/14 21:31:49
>>315
ありがとうございます。
もう一つ質問させていただきたいのですが、たとえば、512x512のようにブロックサイズ、グリッドサイズ
ともに最大値未満の場合、<16, 32, 1>のように複数の次元に分けるのと、<512, 1, 1>のように1つの次元
にまとめるのとどちらがいいのでしょうか?
多分通常はブロックサイズは各次元の最大値<512, 512, 64>をオーバーすることが多いと思うので、こちらは
分けることになると思いますが、グリッドサイズの最大値は<65536, 65536, 1>もあるので、
次元を分けなくてもいいときは分けない方がいいのでしょうか?
317:デフォルトの名無しさん
08/12/14 22:01:51
threadサイズが実際の分散処理の単位になる
つまりthread数だけ並列演算が行われる
例えばthread(1)にしてblock(10)とかだと順次処理と変わらない
ただしthreadは256までしか出来ないのでblockという単位が容易されてる
blockは単にプログラムがやりやすいように用意されただけのもので
実際に分散処理には影響しない
318:デフォルトの名無しさん
08/12/14 22:15:42
>>316
言いたいことがよく判らんが、演算量が512x512ならブロックを8x8にしてグリッドを64x64でいいんでない?
或いはxyの区別が重要でないならブロックを64、グリッドを4096とか。
私は後者の方針でptxファイルの出力とプロファイルの結果を睨みながら決定することが多いけど。
# 横着するときは前者だなぁw
319:デフォルトの名無しさん
08/12/14 22:17:30
>>316
単に内部ループが
for(x=0;x<?;x++){}
_global_
か
for(x=0;x<?;x++){
for(y=0;y<?;y++){
_global_
}
}
になるかの違い
320:デフォルトの名無しさん
08/12/14 22:24:34
元の位置を特定する為にglobal関数でインデックスを計算することになるから
もしxやyのどちらかで収まるなら1つだけ使うのが良い
global関数内部の計算コストが一番ネックになる部分だからね
インデックスの計算部分が一番邪魔で削れるだけ削った方が良い
321:デフォルトの名無しさん
08/12/14 22:35:25
>>317,319
ありがとうございます。ようやく喉のつっかえが取れました。
>>318
試しに昔書いた画像処理のコードを移植しようとしていたのですが、入力によって
いろいろサイズが変わるので、512x512で固定するのはちょっとやりづらいかなと…
ただ、ここまでの話からするとchannel(RGB)*height*widthの三次元配列に画像をつっこんで、
grid = <height, channel, 1>
block = <32, width / 32, 1>
で処理を回せそうですね。画像の幅をを32の倍数に合わせる必要はありそうですが。
>>320
ありがとうございます。実際に使うときは上のような感じにブロックとグリッドの数を決めようかなと
思いますが、どうでしょうか?
322:デフォルトの名無しさん
08/12/14 22:59:47
heightは超えてもいいの?
範囲は内部でインデックスが範囲かどうか判定するよりは
端数を含めたメモリを確保して必要無い部分も計算させてしまった方が速いよ
323:デフォルトの名無しさん
08/12/14 23:22:38
>>322
そこまで大きな画像を入れるかどうかは分かりませんが、確かに高さが65536を
超える場合は分割してやる必要がありますね。
>端数を含めたメモリを確保して必要無い部分も計算させてしまった方が速いよ
ということは、配列にコピーするときに幅を合わせてコピーするのではなく、後ろに空白部分を
まとめた方がいいということでしょうか?
#これって多次元配列を引数に渡すのってどうするんだろう…
324:デフォルトの名無しさん
08/12/15 00:01:31
いや、最終的には全て一次元の配列でそれが適切な間隔でアクセスされるように並んでいるのが一番いいことになる。
だから、画像なんかの場合だと座標値に意味がないなら詰め込んで構わない。
実際には、コンボリューションフィルターなどで隣の画素を参照しないといけないことが多いだろうけど、
その場合はどうせ共有メモリを使ってアクセスを減らすなどの工夫が必要だからどの途32の倍数に拘る必要はない気がする。
325:デフォルトの名無しさん
08/12/15 00:50:40
論文用のならともかく普通の画像処理で画像サイズが固定されてることはまずないし
サイズも共有メモリじゃとても足りない
共有メモリを使うならCPUサイドであらかじめ画像を分割して渡すという方法しかないけど効率悪いよ
毎回隣接した画素情報も含めないといけないから
326:デフォルトの名無しさん
08/12/15 01:06:12
>>324
カーネルサイズにもよりますが、sharedメモリはちょっとサイズ的にきつそうですね…
constantメモリならブロックサイズを調整すれば何とかなるかも?
>>325
単なる趣味グラムなので、ぶっちゃけ全部globalメモリからとってきてもいいような気がしたのですが、
逆にそれだとCPUで処理するよりも遅くなりそうですね
327:デフォルトの名無しさん
08/12/16 01:54:49
遅くなりそうというかどんだけがんばっても最終的にはglobalメモリを使うのが一番早いという結論になるよ
まあそれがGPGPUの一番のネックなんだけどね
CPUサイドのメモリを共用するようにいずれはなるんだろうけど今の限界はそこ
328:デフォルトの名無しさん
08/12/16 06:11:12
一番(手っ取り)早いのは同意するが、一番速いとは限らない。
それこそ、思考放棄じゃないかな?
329:デフォルトの名無しさん
08/12/19 00:31:55
詰まってます、助けて・・・
私はVS2005で開発しようとカスタムウィザードを入れ、ツールキットとSDKをインスコし
パスも通したのにLNK1181エラーでるんです・・・存在しないオブジェクトファイルを指定するのはなぜでしょう?
ライブラリパスがみつからないってどういうことなんでしょうか?
330:デフォルトの名無しさん
08/12/20 01:37:50
"Program" "Files"と認識してるとか
331:デフォルトの名無しさん
08/12/20 02:38:31
Windowsってなんで重要なディレクトリ名にスペースなんか入れるんだろう。
MSって頭悪いの?
332:デフォルトの名無しさん
08/12/20 04:16:06
低脳を排除するためかもねw
333:デフォルトの名無しさん
08/12/20 05:06:25
プログラマお断り。ずっと消費者で居てくだちい。
334:デフォルトの名無しさん
08/12/20 06:27:59
デスクトップやマイドキュメントのスペースや半角カナフォルダはなくなったが、Program Filesはそのままだな。
335:デフォルトの名無しさん
08/12/20 09:43:22
>>331
その昔16ビットOSが主流だったころM$は8文字ファイル名しか扱えなかった
そしてUnix系OSは既にロングファイルネームに対応していた
そこで登場したのが32bit Window 95だった
長いファイル名が使えるんだぞ、凄いんだぞと新発明でもしたかのように宣伝していた
そして自慢げにフォルダ名を無駄に長くした
336:,,・´∀`・,,)っ-●◎○
08/12/20 12:23:53
VistaはUNIXerには優しくなったな。
上位エディション限定だけどSUAも使えるし
337:デフォルトの名無しさん
08/12/20 20:39:06
CUDAをさ、VS2005で使えてる人っているの?
つか、なんで英語のマニュアルしかないんだよ。わかんねーよ
338:,,・´∀`・,,)っ-○◎●
08/12/21 00:44:02
むしろ2005以外でどうやって使うんだ?
2.1βで2008対応したけど
339:デフォルトの名無しさん
08/12/21 00:58:06
>>338
おお、それはいいこと聞いた
やっと2005アンインストールできる
340:デフォルトの名無しさん
08/12/21 01:37:02
>>338
どこか日本語で環境設定を教えてくれるサイト知ってたら教えてくれ・・・
341:デフォルトの名無しさん
08/12/21 01:55:48
環境設定なんか何も要らんがな。スクリプト走らせるだけ。
bashのバージョンによってはウォーニング出るけど、ちゃんと動く。
スクリプト猿にも成れないとは、どんだけゆとり脳なんだよ?スポンジなの?
342:,,・´∀`・,,)っ-●◎○
08/12/21 10:35:26
どのみちあの程度の英語が読めない人に優れたプログラミング書くのは無理だと思うんだ。
俺の私的な業務の手伝いしてくれるなら日本法人さんに代わりにチュートリアルとかの要求
してあげてもいいんだけど、職権乱用か。
>>340 ここで解らなかったらもう諦めたら?
URLリンク(chihara.naist.jp)
343:デフォルトの名無しさん
08/12/21 11:55:12
どうでもいいけど、↑のURLは「disp_content」を削らないとエラーになった。
344:デフォルトの名無しさん
08/12/21 12:58:13
>>340
サンプルプロジェクトを開いて設定を見て勉強するのがいいよ
345:デフォルトの名無しさん
08/12/24 12:51:08
英語読めないのにプログラミング言語してる男の人って・・・
346:デフォルトの名無しさん
08/12/24 18:35:50
英語に優越感もってるやつって、さもしいな
347:デフォルトの名無しさん
08/12/24 19:03:01
>>345
可愛いもんだよねv
348:デフォルトの名無しさん
08/12/24 23:58:39
英語なんて誰でも1ヶ月もあればある程度は読めるようになるよ
出来ないと思い込んでやろうとしないだけで
349:デフォルトの名無しさん
08/12/26 12:46:32
vs2005の環境で拡張子cuの場合、定義参照はできるのですが、
インテリセンスが機能しません。freeとかvcのランタイムも同様に
機能しないのですが、やり方分かる方いますか?
ツール->オプション->テキストエディタ->ファイル拡張子にcu c++として
登録するとできるのは上記だけです
350:デフォルトの名無しさん
08/12/26 22:27:18
いっそ、nvccの拡張機能以外は全部cppでやったら?
351:デフォルトの名無しさん
08/12/27 05:03:22
それがスマートだね
352:デフォルトの名無しさん
08/12/27 17:53:02
質問です。
VC++2005で開発する際に、コマンドプロンプトにデバイス名も含め一切文字を表示させたくないのですが、どのようにすればよいでしょうか?
353:デフォルトの名無しさん
08/12/27 18:38:27
>>352
長いパス名が邪魔なだけなら set PROMPT=$G とか?
354:デフォルトの名無しさん
08/12/27 18:47:25
>>352
CUTILを使うと標準出力があるって話なら、リダイレクトするかCUTILを使わなければいいと思うのだが。
355:デフォルトの名無しさん
08/12/28 01:33:08
プロンプト出したくないって話なら_tmainじゃね
356:352
08/12/28 02:53:51
>>353-355
回答ありがとうございます。
>>353
質問があいまいですみません。
パス名が邪魔、ってわけではないです。
>>354
おそらくこの指摘にあてはまるのだと思います。
>リダイレクト
リダイレクトしても、コマンドプロンプトに表示が出てしまいます。
>CUTIL
cutil.hを使わずに、cutil.hで提供されている関数を使わない方法、もしくは、代替可能なヘッダーファイルはありますか?
>>355
_tmainを使ってしまうとWindowsでしか使えないソースコードになってしまうので、別の方法があれば教えてください。
357:352
08/12/28 02:55:28
自己レスですw
誤:cutil.hで提供されている関数を使わない方法
正:cutil.hで提供されている関数を使う方法
の間違いです。
358:デフォルトの名無しさん
08/12/28 03:10:48
system("cls");
とか
359:デフォルトの名無しさん
08/12/28 10:49:28
>>356
> _tmainを使ってしまうとWindowsでしか使えないソースコードになってしまうので、別の方法があれば教えてください。
Windows かどうかで #ifdef するくらいは許されるんじゃないの?
360:デフォルトの名無しさん
08/12/28 14:29:28
プロジェクトの設定でWindowsアプリにすればいいだけでしょ
361:デフォルトの名無しさん
08/12/28 14:32:32
>>_tmainを使ってしまうとWindowsでしか使えないソースコードになってしまうので、別の方法があれば教えてください。
mainだろうがWindows依存部分のコードが発生するのは当たり前でしょ
CUDAの部分とC++の部分とWindows依存部分とファイルを分離するのが鉄則だよ
複数の環境で使う場合は#ifdefプラグマを使って分離したWindows依存部分の
#includeを切り替えるようにするだけ
362:デフォルトの名無しさん
08/12/28 17:29:43
コンパイラのオプションで指定できるでしょ
環境ごとにmakefile書くだけじゃね
363:デフォルトの名無しさん
08/12/28 17:48:49
>>361
#ifdefはいつからpragmaになりましたか?
364:デフォルトの名無しさん
08/12/28 17:53:47
>>363
#ifdefディレクティブと言いたかったんじゃないの?
で、使い慣れない言葉なんでつい間違えたと。
句読点もちゃんと使えないようだし、日本語に慣れていない三国人なんでしょ。
365:デフォルトの名無しさん
08/12/28 22:36:27
cudaのMDコードどこかに追ってませんか?
366:デフォルトの名無しさん
08/12/29 12:11:23
たった1行でここまで意味不明なのも凄いな。
367:デフォルトの名無しさん
08/12/29 12:24:23
×追ってませんか?
○オッドアイいませんか?
368:デフォルトの名無しさん
08/12/29 13:39:38
>>365氏はcudaで書かれたMD(分子動力学)コードどこかに落ちてませんか
と聞いてるのかと。当方も知らないので答えられませんが。
Fortranで書いたものを全部はアレなのでGPU上で実行したいサブルーチンだけ
Cに変えてCUDAで動かしたいのですが、そんな例とかは落ちてないですかね。
mainルーチンその他関係ないところまで全部Cに移植するのが嫌ってだけな
んですが。あ、当方はIntelFortran使用。
当方まだCUDA触りたて、試しにSTREAM BMTのtriadだけ手で適当に書いて
GeForce9600GTで40GB/s弱(効率7割弱)のメモリバンド幅。あ、じゃもうでき
るじゃんとか勝手に思い、Fortranのコードに挑もうとしてあえなく止まってますorz
369:デフォルトの名無しさん
08/12/29 14:08:06
つURLリンク(www.easize.jp)
370:368
08/12/29 14:20:59
>>369
ありがとうございます。m(_ _)m
多次元配列はどのみちGPU上では一次元化するつもりでしたがなるほど。
参考にさせていただきます。
371:デフォルトの名無しさん
08/12/29 14:45:46
>>369
そのページ、ポインタの扱いがメタメタだな
わかってるなら別にいいが
372:デフォルトの名無しさん
08/12/29 14:49:43
うん、ぶっちゃけ「"cuda fortran"でぐぐって一番上に来たリンク貼り付けただけ」なんだ、すまない
373:デフォルトの名無しさん
08/12/29 14:50:01
何か普段使うものをCUDAに移植しようと思いつつ適当な物が見当たらない
374:デフォルトの名無しさん
08/12/30 02:15:09
ぶっちゃけ、俺もそうなんだよね。おまけで付いて来たサンプル書き換えて
動かしてFLOPSベンチしてるだけ。
CUDA動かす環境をコスパ良く構築するには?と考えて色々やってみたが、
構築した環境で動かすモノって、結局サンプルの改造ばっか。まんまと
nVidiaの販促戦略に乗せられたぜw
375:デフォルトの名無しさん
08/12/30 08:50:29
zip圧縮とかjpg圧縮とかを移植したらライセンスの関係はどうなるの?
376:デフォルトの名無しさん
08/12/30 12:52:04
圧縮アルゴリズムに関しては問題ないんじゃね?
377:デフォルトの名無しさん
09/01/19 15:08:44
CUDAプログラミングモデルの概要
URLリンク(http.download.nvidia.com)
CUDAプログラミングの基本(パート I - ソフトウェアスタックとメモリ管理)
URLリンク(http.download.nvidia.com)
CUDAプログラミングの基本(パート II - カーネル)
URLリンク(http.download.nvidia.com)
378:デフォルトの名無しさん
09/01/21 03:21:23
デバイス側に確保したメモリにホスト側から
cudaMemcpyのように一括でコピーするのではなく、
一部分だけコピーする、又は書き換える良い方法がありましたら教えてください。
379:デフォルトの名無しさん
09/01/21 03:28:19
・cudaMemcpy()で一部分だけコピーする。
・cudaMemset()で(ry
・そういうカーネル関数を用意して呼び出す。
380:デフォルトの名無しさん
09/01/21 12:27:16
一部分だけってのはたいてい順次処理になるからCPUにやらせたほうが有利だよ
381:デフォルトの名無しさん
09/01/22 01:33:54
CUDAのプログラミングでZIPアーカイブのパスワード解析とか、早くなりませんかねぇ。
エンコード/デコードに使えるんだから、どうなんでしょう?
382:デフォルトの名無しさん
09/01/22 06:38:19
CUDAのfortranサポート予定ってGPUカーネルの部分をfortranライクに
書けるようになるって理解でいいのかな?2.0からって見かけたけど、それっぽい記述が全く無い…
383:デフォルトの名無しさん
09/01/22 08:47:26
>>379-380
ありがとうございました。
参考にさせていただきます。
384:デフォルトの名無しさん
09/01/22 19:18:17
>>381
総当りか辞書型か知らんが演算能力よりI/Oの問題だろ普通。
385:デフォルトの名無しさん
09/01/22 20:12:54
>>381
早くなるよ
386:デフォルトの名無しさん
09/01/23 00:20:22
>>384
ファイル数が多い場合はI/Oも問題になるかもしれませんが、
ある程度のサイズであればWindowsでもメモリに読み込むので、
それほど問題にならないのでは?
毎回物理的に読みに行くわけではないし。
387:デフォルトの名無しさん
09/01/26 09:28:38
ファイルI/OじゃなくメモリI/Oなのでは?
388:デフォルトの名無しさん
09/01/28 11:49:27
>>387
お前は何を言っているんだ
389:デフォルトの名無しさん
09/01/28 12:47:12
>>388
pciバスつかうんだろ?
立派なioとおもうが
つかクーダってほんとマイナーなんだね。
390:デフォルトの名無しさん
09/01/28 13:18:14
そりゃNVIDIA限定だからだろ?グラボを選ぶのなら汎用的なアプリは作れない。
CUDAよりもっといろんなグラボに対応しているOpenCLへのつなぎに過ぎないよ。
391:デフォルトの名無しさん
09/01/29 00:08:14
ZIPのパスワード解析にZIPファイル全体へのアクセスが必要だと勘違いしてる馬鹿w
392:384
09/01/29 08:00:22
>>381-391
ああ、俺がぼけてた。なんか知らんけど自分でもなんであんなレスつけたんだろ orz
わけわからない流れにしちまって、すまん。
正しくはこうだな。↓
パス解析では総当りにしても辞書にしてもforなどのループ速度が求められるだけで、
CUDAによる計算は意味があるのかい?と思った。
MD5とかデータから導き出す数値ならば本当の意味での解析だから意味ありそうだけど。
393:デフォルトの名無しさん
09/01/29 23:02:34
こんな記事もあるしな
GPUを利用した無線LANのパスワードクラッキングが主流に
URLリンク(itpro.nikkeibp.co.jp)
394:デフォルトの名無しさん
09/01/30 03:18:20
>>393の言ってるのはまさに本当の意味での解析ってやつだしな。
WPA/WPA2-PSKからの復元は総当り/辞書比較と違ってGPGPUにする意味がある。
395:デフォルトの名無しさん
09/02/01 15:41:49
Catalystだけど、トリップ検索は随分速くなるみたい
URLリンク(sourceforge.jp)
396:デフォルトの名無しさん
09/02/02 21:51:40
ホスト側でconstantメモリを動的に確保できないのでしょうか?
方法がありましたら教えてください。
397:デフォルトの名無しさん
09/02/02 22:44:30
>>396
目的が良く分からないのでなんとも言えませんが、
どうせ64KBしかないのだから静的に64KB確保しておくのではだめなのでしょうか。
398:デフォルトの名無しさん
09/02/02 22:54:45
>>396
あるけど、なんでprograming guide読まないの??
399:デフォルトの名無しさん
09/02/02 23:59:28
>>397
静的に確保しておいてもよかったんですが,綺麗にかけるならとw
>>398
一度目を通したつもりでしたが,素通りしていたようです。
もう一度読み直してきますm(_ _)m
400:デフォルトの名無しさん
09/02/06 00:08:13
誰かZIPの暗号解析ツール作って><
それか解析のループ部分だけうpしてくれたら、カーネルは俺が書きます。
401:デフォルトの名無しさん
09/02/06 06:54:02
総当りするだけならどっかのオープンソース実装の
パスワードチェックと復号ルーチンとってくればいいだろ
402:デフォルトの名無しさん
09/02/06 10:20:56
ただでさえ消費電力がデカイのにパスワード解析に何週間も動かしてられっかw
403:デフォルトの名無しさん
09/02/07 16:57:19
超解像のような考え方は昔からあるようだな
フリーソフトImageD2
URLリンク(www.tiu.ac.jp)
これは時間軸方向にも参照するものだ。
NVIDIAのMOTIONDSPと同じ考え方だな。
西川善司の大画面☆マニア 第113回 超解像
URLリンク(av.watch.impress.co.jp)
この方法は数フレームを参照することで低解像な映像のブレから情報量をアップさせるようだが、
この方法だとシーンチェンジやめまぐるしく動く映像では逆効果でめちゃくちゃな映像になるのではないか?
(これは上記のフリーソフトの別ページでも注意点として載っていた)
でも東芝などの日本の各社がやってる1フレームだけで行う超解像はそもそも無理がある。だから不自然な画質になったり、情報量が逆に消失したりする。
Lanczosなどでそのままアプコンしたほうがずっと情報量あるし自然な画質だ。比較してみれば一発で分かる。
URLリンク(plusd.itmedia.co.jp)
ここの元画像を720×480にし、それをAVIUTLなどでLanczosでフルHDにしたもののほうがずっと綺麗。
超解像は単純な処理だから柵とか崩れてるし、文字も駄目だし、元からあった情報を処理によって消しちゃう副作用のほうが強い。
超解像、超解像と目新しくいって盛り上げようとしたいのは分かるが、こんなのはまやかしだよ。
URLリンク(www1.axfc.net)
比較用画像もうpしておいた
一方、数フレームでやる方法は計算が大変だが、シーンチェンジや盛大な動きの問題さえクリアすればかなり使えそうではある。
MOTIONDSPや↓はそのあたりちゃんとクリアしているのだろうか?
URLリンク(www.flashbackj.com)
404:デフォルトの名無しさん
09/02/07 17:01:19
10年ぐらい前にカーネギメロン大学で勉強してた頃に超解像のプログラム作れっていう課題を出されたことがある
405:デフォルトの名無しさん
09/02/07 17:34:44
ちょー解像は判ったからマルチすんなや。
406:デフォルトの名無しさん
09/02/07 20:13:09
>>403
スレ違いじゃない?
GPGPU 向きな処理なのは確かだし、そもそも Cell なり SpursEngine は GPU じゃないし。
画像処理スレとかあったと思うよ?
407:デフォルトの名無しさん
09/02/07 23:45:41
アメリカのメロン大学は一応名門だけど日本校はどうなんだ?
408:デフォルトの名無しさん
09/02/08 01:05:45
>>404
なにその高レベルな課題
俺の大学時代と雲泥の差
409:デフォルトの名無しさん
09/02/08 02:58:06
超解像の原理が未だに理解出来ません
複数の映像フレームの同じポイントの色の出現頻度に一番高い色を適用するってことですか?
410:デフォルトの名無しさん
09/02/08 13:45:42
は?今の超解像って、時間軸の補完までやってるのか?
411:,,・´∀`・,,)っ-○◎●
09/02/08 13:52:05
MPEGの原理は主にフレーム間差分をとってJPEG圧縮なんで、それを改めてチェックしたところで・・・
412:,,・´∀`・,,)っ-○◎●
09/02/08 21:06:32
>>407
神戸のハーバーランドの近くにあるあれか?
兵庫県の財政圧迫してるらしいよ。
学生少なくて撤退の噂もあります
同じビルの同じ階に神戸電子専門学校の学校法人が経営する大学院がある。
何を勘違いしたのか兵庫県立大学も情報科学の大学院を神戸市内に置いてる。
神戸という都市に何を求めてるのか、理解不能である。
たしかに古くからの工業都市で組み込みソフト屋の数がそれなりにいるのはわかる。
最大の勘違いは、日本のITドカタは職場を離れて大学院に通えるほど裕福ではないことだ。
413:デフォルトの名無しさん
09/02/08 21:49:03
俺はてっきり、ラーメン大学とか、洗車大学と同じノリで、
神戸電子専門学校が作ったヒト寄せパンダだと思ってた。
ってか別法人なのか?
学生少なくて、って言う前に、派遣のワーキングプア、
デジタル土方に成る為にわざわざ苦労して大学通う馬鹿
が神戸のどこに居るんだよっ?
俺が居るよ…orz
414:,,・´∀`・,,)っ-○◎●
09/02/08 22:51:22
兵庫県もネギメロンなんて誘致するくらいなら神戸大に寄附講座でも作ったほうがよかったんじゃね?
現状兵庫県内で優秀な学生がいるんだし。
おっと、「学生」の枠で括ると近くの中高一貫校のほうがよっぽど・・・