【GPGPU】くだすれCUDAスレ part5【NVIDIA】
1 :
デフォルトの名無しさん :
2011/08/23(火) 22:08:06.09
おい、pertだろ まちがえんな
>>2 個人blogのは情報古すぎなような…
「cudaさわってみた」はリンク切れだし。
>>4 やっぱりCUDAやってる人は荒んでるんだな
ほんとCUDAらないことで騒ぐよな
ケンカするのはやめてCUDAさい><
そうだよ。喧嘩はやめ…GeForんゲフォン
まだスレが始まっTeslaいないのに
ときに口論になるのもみんな自分の技術にそれだKeplerイドを持ってるってことだ。
amazon.comで検索して出てくるCUDA本ってことごとく出版がずいぶん先か延期になったりするな なんでだろ
無効とは情報の登録ポリシーが違う…ような気がする。 それがAmazonのポリシーなのか出版社のポリシーなのかは知らないけれど。 多分出版社のサイト直接見た方が正確。
cudaの本が必要、って訳でも無いが出版されるなら買ってしまうと思う
15 :
デフォルトの名無しさん :2011/08/31(水) 20:00:08.40
4.0になってdeviceemuが使えんのだが、CUDA載せてないPCで動かす方法ない?
3.2を使う。
>>17 現実的じゃないか。どうせemuで正確なエミュレーションができるわけじゃないし。
>>17 4.0でdeviceemu使う方法なんてねーよ、ってことだよ。
言わせんな恥ずかしい
21 :
デフォルトの名無しさん :2011/09/01(木) 15:33:21.69
理屈の上ではKnights Ferryでも動かせるのかな?
ubuntu 11.04 tesla c2050 cuda sdk 4.0の環境でdoubleのアトミック演算が正しく動作しません オプションは -arch=compute_20を指定してます。 本家のプログラミングガイドに書いてある、以下のコードを実行してます。 __device__ double atomicAdd(double* address, double val) { unsigned long long int* address_as_ull = (unsigned long long int*)address; unsigned long long int old = *address_as_ull, assumed; do { assumed = old; old = atomicCAS(address_as_ull, assumed, __double_as_longlong(val + __longlong_as_double(assumed))); } while (assumed != old); return __longlong_as_double(old); } 引数のadrees*には共有メモリのアドレスを渡してます。 一つのスレッドのみが共有メモリにアクセスする場合は正しく加算は行われますが、二つ以上のスレッドが衝突する場合は 正しく動作しません。printfで見てみたら、二つ目のスレッドがループを抜けれずに何度も共有メモリの値に 加算しているようでした。上の文をdouble→float,long long → int に書き換えた場合は正しく動作しました。 また、atomicCASのところを二つのprintf、 printf("%d %f %f\n",threadIdx.x,__longlong_as_double(assumed) ,__longlong_as_double(old)); ではさむとなぜか正しくループを抜け、正しくsharedメモリに値を書き込みます。 原因がわかる方、教えていただけないでしょうか。
>>22 の続きです
assumed oldの値をcpu側に持ってきて表示してみたら、
oldの値(atomicCASの返す値)とassumedの値が末尾の数ビットが
異なってました。50回ほどループしてやっとビットレベルで値が一致して
ループを抜けているようでした。その間に過剰に加算が行われていたようです。
カードの計算が正しいかチェックするベンチマークってありませんか? 四則演算と√と三角関数とかを全ビットを演算器全部チェックしろと云う無茶ぶりがきた。 計算機が計算を間違うなんてありえないのに…
Teslaの存在意義全否定デスカ
>>24 過去スレを読むとわかってもらえると思うが、みんなそれを求めているんだ。
そのチェックにおける最先端はおそらく長崎大学。
1スレッドの実行がどのコアかがわかる仕組みがどのような方法か分かればいいわけだよな 誰か調べてみろよ
>>22 ,23
もしかしてブロック辺りのスレッド数が32以下なら動いたりする?
>>29 コアと計算結果の表が出来上がればコア間の相違が分かるじゃん
やり方知ってるなら教えてください
煽っても教えてやらん 29じゃないけどw
32 :
デフォルトの名無しさん :2011/09/02(金) 02:12:08.96
NvidiaのCUDAをやろうと思い、Nvidiaのカードの刺さったマシンにFedoraを 入れてCUDAのためのNvidiaのドライバを入れて使ってみたが、OSのカーネル のアップデートがあるたびに、コンソールが再起動時にブラックアウトした。 そのたびごとに、別のマシンからリモートログインをしてCUDAのための Nvidiaのドライバを再インストールする必要があった。 CUDA3.xまではそれでなんとかなったが、現在のCUDA4.0になると、 幾らやっても駄目になってしまった。どうやらFedora15ではまともに サポートされていないようなのだ。Nvidiaの提供しているCUDAのための ドライバがインストール時におかしなメッセージが出てインストールが 出来ない。
ここはおまえの日記帳だから好きに使いなさい
>>28 ブロックあたりのスレッド数を32以下にして試してみましたが、
結果は変わらなかったです。
>>23 とりあえず、プログラミングガイドのコードそのままだし、
実際アルゴリズム自体には問題ないから、
意図どおり動かないとしたら、ドライバかSDKのバグか
ハードの故障じゃね。
まあ、ここで詳細なバージョンとか晒してみたり、
PTXやアセンブリを貼ってみるのもいいけど、
本家のフォーラムで聞いてみたほうが早いんじゃね。
長崎大の言う演算間違えって、具体的にどういう演算を間違えたのかの情報ってないの? 間違えますって情報だけ垂れ流して具体例出さんって、どういうことなんだべ。
×演算間違え ○演算間違い
CUDAを使用するプログラムで巨大なデータを扱いたいのですが、nvcc でのコンパイル時に relocation truncated to fit: R_X86_64_32S against `.bss' とエラーが出て完了しません。gcc ではデータ量制限を解除するのに-mcmodel=medium というオプションがあるようですが、nvcc でこれにあたるものはないでしょうか。
巨大データってどれくらい? CUDAは基本グローバルメモリ使用の制約があるから”巨大”は決して得意ではないのだけど
40 :
38 :2011/09/07(水) 22:05:04.34
>>39 今確認したところ、要素が数十万個の配列を使用していて、それらの合計は200MBほどです。
そのうち、実際にGPUに計算させるためGPUに渡すデータは50MB程度です。
上記のエラーメッセージでググったところ、データ量制限が問題だろうということだったのですが…それほど大きなデータでもないですよね。
なんなんでしょうこのエラー。
プログラミングに関してはまだ知識が浅いためチンプンカンプンなこと言ってたらすみません。
どうせデータをソース中に書き込んでるんだろw
>37 日本語の間違えより、長崎大の言ってる間違いの詳しい情報くれよぅ…。
要素数が大きすぎるなら、mallocを使ってみたら?
>40 配列は何次元? CやC++だと配列は連続したメモリ領域と規定されてるので、 確保可能なメモリ量は物理メモリの絶対量ではなく、 アドレスを連続で確保できるメモリ量に規定されるよ。 例えば大きめの2次元配列を確保する場合は double* pBuf = new double[row * column]; とするより double** ppBuf = new (double*)[row]; for(int i = 0; i< row; ++i){ ppBuf[i] = new double[column]; } とかにしたほうがいいよ。
ファイル読み込みで、cudaMallocHost(だったかな??)でメモリ確保すればいいじゃん
46 :
38 :2011/09/07(水) 23:11:33.99
>>43 >>44 gccと同様にオプションでエラー回避できないかな…と思ったのですが、配列の取り方を変えるほうが良さそうですね。
特に
>>44 は知りませんでした。
ありがとうございました!
良く分からんが、-Xcompiler -mcmodel=mediumじゃ駄目なん? -Xlinkerも要るかもしれないが。
>>44 普通はありえない。そんなことしたら、ポインタアクセスが一回余計に要るから遅くなるし、
キャッシュアウトしやすくなるから更に遅くなる。
まぁ、GPUに転送するのは1スライス分だけだろうからGPUに任せてCPUは暇こいていられるならいいかも知れんが。
>>44 それは無い。
・メモリも無駄に消費
・アクセス性能がた落ち
・バグの温床になる
つーかC++ならvectorで初期化時に
サイズ指定したほうがいろいろ便利
配列でやる場合、 配列の添え字がINT_MAXを越える大きさならそうするしかないと思うが・・
キャッシャがアクセスがとドヤ顔されても・・・。 要素のデカイ配列が確保できんって話なわけだし。
コンパイル時の話なのに
ヒープのフラグメンテーションの訳無いだろう。
nvccからgccに自由にオプション渡すやり方は
知らんが、最悪c(++)部分のコードを分割して
gccでコンパイルすりゃ良い。
仮にヒープのフラグメンテーションが酷い場合でも
>>44 みたいなやり方は下策だよ。
素人におかしなやり方を覚えさせるだけ。
大きな配列確保する前に小さな物についてShrink-to-fitやればデフラグされるから
その後巨大配列を確保すれば良い。
特にGPGPUで巨大な配列扱うような用途なら
配列の配列なんて転送も困るだろうに。
>大きな配列確保する前に小さな物についてShrink-to-fitやればデフラグされるから >その後巨大配列を確保すれば良い。 コレ具体的にどんなコード書けばいいのん?
メインの計算に必要な必要な小容量配列 保持用のポインタなりvectorの変数だけ始めに ヒープに確保した後、メインの計算前の前処理をする。 この処理で、一時的なデータと最終的に使うデータの ヒープ確保が入り乱れる所為でフラグメンテーションが発生する。 最終的に使うデータについて、大体処理した順番を守って ひたすら同容量確保してはコピーを繰り返せば、 こういうデータは全て前方に圧縮されて配置されるようになる。
? まあスレ違いなんで他でやってよ。
kernelへの引数に構造体は使えんの? 今やってみたらコンパイルでエラーになった
あ、間違えてただ 構造体引数使えるべな
cuda4.0、gtx580で Maximum dimension of a grid( 65535 , 65535 , 65535 ) ってなってるんですがz軸のブロックの取れる数は1じゃなくなったんですか?
デバイス上の処理が長時間になる場合、途中経過を取得する方法ってないかな? 例えば100万回のモンテカルロを実行して、途中の10万回、20万回が済んだらその情報を取りたいんだけど…
streamingってなに? SMのSじゃないよね?
カーネル呼び出し時の<<<>>>の第3引数がstreamingの何かじゃなかったっけ?
俺は
>>61 を見てそういうのがあったなと思い出したレベルなので使い方は知らないけど。
第3パラメータはsharedMemoryのサイズだよ。そこで指定しなくても使えるけど。
ごめん第4引数だったみたい。それとAsync系関数がうんたらかんたら。
あー、その辺だね。
個人でCUDAプログラミングやってうれしいのはその程度なのか? と逆に不安になるんだが…
世の中なんでも膨大なデータをいかに早く処理するかが勝負の世界に なってきているから、CUDAプログラミングの旨味はいくらでもあるだろ。
いつかやってみようと思って半年くらいこのスレROMってるが いまだに開発環境のセッティングの仕方すら調べていない・・・
>>71 おれがそれくらいのコード書けるようになるころには、自ずとましな使い道が浮かんでくることだろう、
っていうこと?
73 :
デフォルトの名無しさん :2011/10/02(日) 08:23:54.81
>>60 試してないけど、長い処理の途中でグローバルメモリに書き込んで、
別streamでその値を読むカーネルを起動すれば取れると思う。
キャッシュからグローバルメモリに書き込むのは特別な関数を
使うと思ったけど名前忘れた
74 :
デフォルトの名無しさん :2011/10/02(日) 09:01:54.77
>>73 思い出した。atomicを使うか、書き込む側で __threadfence()
75 :
デフォルトの名無しさん :2011/10/04(火) 20:59:53.36
失礼します、先輩方お助けを。 atomicを使ったソースを"-arch compute_20"オプションでコンパイルすると次のようなエラーが出てしまいます。 「'-arch' は、内部コマンドまたは外部コマンド、操作可能なプログラムまたはバッチ ファイルとして認識されていません。 C:\Program Files (x86)\MSBuild\Microsoft.Cpp\v4.0\Microsoft.CppCommon.targets(103,5): error MSB3073: コマンド "-arch compute_20 error MSB3073: :VCEnd" はコード 9009 で終了しました。」 -archコマンドはプロパティからビルド前イベントのコマンドラインに追加してあります。 -archコマンドがあるディレクトリへパスが通ってないのかなと思うのですが、このようなコマンドがあるディレクトリの追加の方法がよくわかりません。 環境はcuda4.0 + VC2010 + teslac2070(cc2.0) + windows7 proです。 どなたかアドバイスください。
-archコマンドってなんだそら? compute capabilityはnvccなりcl.exeのオプションに渡すものであって、 -archコマンドなんてものはないよ。
自分でオプションでって書いてるけど、メッセージはそれで実行した系だよな
78 :
デフォルトの名無しさん :2011/10/04(火) 23:16:18.47
>>76 >>77 ありがとうございます。「cuda by example」にもatomic関数を使ったソースには-archコマンドラインオプションを指定する必要があります。
とあるのですが、みなさんはatomicを使ったソースをコンパイルするときなにかコマンドつけませんか><?
79 :
デフォルトの名無しさん :2011/10/04(火) 23:19:48.93
CUDA以前の話だな。コンパイラオプションの意味がわかってない。 というか、VCなら、カスタムビルド規則でCUDA Runtime API Build Ruleを追加したほうが早い
-gencode=arch=compute_20,code=\"sm_21,compute_20\" になってるなぁ俺の環境だと。VC2010ならプロジェクトのプロパティでCUDA関係の設定するところがあるんじゃない?
82 :
デフォルトの名無しさん :2011/10/05(水) 11:51:04.38
>>80 ありがとうございます。
えーと、それはおそらくVC2008の話ではないかと思います。2010にカスタムビルド規則という名前はなくなり、その代りCUDA4.0というビルドカスタマイズに統一されています。
83 :
デフォルトの名無しさん :2011/10/05(水) 11:58:44.51
>>81 ご返答ありがとうございます。
そのコードはどこに書いていますか?
はい、cuda c/c++という項目にいろいろいじるところはあるのですがコンパイラオプションを設定する項目はないように思えます。
私がオプションを追加しているのはビルドイベントという項目ですね。
以前にもatomic関連の話が出ていたようです。
>>22
-arch=compute_20 にすればいいんじゃねえの?
85 :
デフォルトの名無しさん :2011/10/05(水) 12:17:32.55
>>84 ありがとうございます。
おっしゃるように
>>75 のように追加するとエラーを吐きます。
追加の仕方が違うのでしょうか。
86 :
80 :2011/10/05(水) 13:36:15.15
VS2010は色々変わってるんだな。 知らんかったすまん。 とりあえず「CUDA VS2010」でググって出てくる一番上のサイト参考にして、CUDA4.0のビルドカスタマイズファイルを導入した後、プロパティからCode Generationをcompute_20,sm_20にしたらいけた
GPUのリソースを全部CUDAに割り当てるため,グラボを2枚さしてディスプレイ出力用とCUDA用に分けるというのはありますが 同じようなことをsandy bridgeとグラボ1枚でも出来るのでしょうか?
>>87 Z68マザーなら、マザー側DVIにディスプレイつないでおけばいい話じゃ?
CUDAで計算したデータをintelで表示したいなら、
1.DirectXかOpenGLのリソースをCUDAにバインドしてデータ書き込み
2.メインメモリに転送
3.intel側に転送
って形になるかなぁ。多分。
89 :
81 :2011/10/05(水) 21:07:18.54
>>83 ゴメン、俺もVC2010じゃないや。VC2008。
ちなみに
>>81 は構成プロパティの中に「CUDA Runtime API」があってそこの「コマンドライン」にある。
CUDA Runtime APIの中にGUIでGPU Architecture(sm_21)を指定するところがあるから、
コマンドを打ってるわけじゃないんだけど、結果としてnvccには
>>81 のようなオプションが渡ってる。
90 :
89 :2011/10/05(水) 21:17:54.24
そもそもなぜCUDA Runtime APIという項目があるのかというとサンプルのプロジェクトをコピったから。
91 :
デフォルトの名無しさん :2011/10/06(木) 16:35:52.62
>>80 >>81 ありがとうございます!code generationを
>>80 のように書きなおしたらオプション指定することなくコンパイル通りました。
どうやらVS2010からはcommand line指定しなくてもcode generationでオプションを渡せるようです。
本当に助かりました。ありがとうございます!
最近巷で流行ってる error MSB3721: コマンド ""C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.0\bin\nvcc.exe" -gencode=arch=compute_20,code=\"sm_20,compute_20\" --use-. ... template.cu"" はコード 2 で終了しました。 のエラーなんですけど解決法どなたか知りませんか?? CUDA FORUMにもかなり話題が上がってるっぽいんですが...私のところにも発病してしまいました。
そこだけ書かれても分からないと思う。 nvccのエラーメッセージ載せないと。
>>92 敵性語なんでよくわかんなかったけど、
VisualStudio を管理者として実行したら動く?
Visual C studioは下手にGUIに頼るよりも、コマンドラインでやった方がはやいこともある。コマンドラインでnvccやってみよう。
単精度のmadってmulより遅いやないかー プログラミングガイド嘘つきやないかー ちくしょー
>>96 そりゃ命令実行時間は同じだけど、
必要な帯域は違う。
>>92 コマンドラインのオプションが何か1つ間違ってた気がする。
コマンドプロンプトで同じコマンド実行するとエラーが出たので、
そのエラーのひとつ前の引数が問題だった、はず。
CUDAとRTミドルウェアを組み合わせてみた。 行った方法は、CUDAのサンプルをRTコンポーネントに パッケージングする方法。 動作周期にあわせて、行列計算が実行されたのを確認。
カーネル側のソース分割ってできないんだっけ? 環境はLinux cuda4.0
できるけどコンスタントメモリが変な事になったような気がする
このままだと5万行くらいになっちゃいそう・・・
>>102 正直処理対象がCUDAに向いてないような。。。
とある文献でレイテンシを表す際, 種類によってclock, cycleが使い分けられていました。 clock : texture cache, constant cache cycle : global memory, shared memory なんとなくメモリ, キャッシュで使い分けられているのはわかりますが, このままでは速度の比較ができません。 二つの違いは何なのか, またcycle-clock間で単位の変換が可能なのか, もしご存知の方がいたらご教授ください。
カーネル関数内で立てたフラグをホスト側で判断する事って出来る?
グローバルメモリ介してできるだろ
はじめてのCUDAプログラミングの9章に載ってるcublas使うプログラム実行したんだけど、 cublasInit()あたりで「〜0x7c812afb で初回の例外が発生しました: Microsoft C++ の 例外: cudaError_enum〜」って出てうまく動かないんだけどなして?
かなり初歩的な質問で申し訳ないのですが 構造体hoge_tにバッファを持たせPtにバッファのアドレスが入るようにしたいのですが 以下のようにするとsegmentation faultになってしまいます。 このようなことは可能なのでしょうか。 struct hoge_t{ int* Pt; int a,b,c,d; } int main(){ hoge_t *hoge_d; cudaMalloc((void**)&hoge_d,sizeof(hoge_t)); cudaMalloc((void**)&hoge_d->Pt,sizeof(int)*100); //終了処理は省略 }
適当なポインタで確保した後に, そのアドレス転送してやればいいんじゃない?
>>108 >cudaMalloc((void**)&hoge_d,sizeof(hoge_t));
意味のおさらい。
& hoge_dをcudaMalloc()に渡すことで、ローカル変数hoge_dにデバイス上のアドレスが格納される。
>cudaMalloc((void**)&hoge_d->Pt,sizeof(int)*100);
& hoge_d->Ptは、意味としては & (* hoge_d).Pt に等しい。
つまり、hoge_dがデバイス上のアドレスを保持しているにも拘らず
ホスト上で* hoge_d するからsegmentation faultを引き起こす。
データ効率を考えると固定長配列にできないか検討するなどして見直すべきだが、
どうしてもこのままの構造でやりたいなら sizeof(int) * 100 確保したデバイス上の
アドレスをcudaMemcpy()でhoge_d->ptにコピーする必要がある。
って、この下りは>109が書いていることと同じだね。
111 :
108 :2011/10/19(水) 17:14:09.68
>>109 ,
>>110 お忙しい中、レス有難うございます。
>>110 データ構造については検討しようと思いますが、
いただいたアドバイスを参考に以下のようにしてみたところsegmentation faultがなくなりました。
109,110さんの意図していたものとなっておりますでしょうか?
int main(){
hoge_t *hoge:
hoge=(hoge_t*)malloc(sizeof(hoge_t));
cudaMalloc((void**)&(hoge->Pt),sizeof(int)*100);
hoge_t *hoge_d;
cudaMalloc((void**)&(hoge_d),sizeof(hoge_t));
cudaMemcpy(hoge_d,hoge,sizeof(hoge_t),cudaMemcpyHostToDevice);
//終了処理は省略
}
あるいはこういう感じでしょうか・・・
int main(){
int *tmp_pt;
cudaMalloc((void**)&(tmp_pt),sizeof(int)*100);
hoge_t *hoge_d;
cudaMalloc((void**)&(hoge_d),sizeof(hoge_t));
cudaMemcpy(hoge_d->Pt,tmp_pt,sizeof(int),cudaMemcpyHostToDevice);
//終了処理は省略
}
概ねいいんでない? 前者なら、私はhoge_t hoge;で定義しちゃうけど。
113 :
109 :2011/10/19(水) 22:04:08.49
俺が考えてたのは後者だな。 合ってると思うよ。 でもmemcpyのsizeof(int)はsizeof(int *)にした方が良いと思う。
114 :
108 :2011/10/19(水) 23:57:45.47
未だにポインタが使いこなせてないと痛感しました。 このたびはアドバイス有難うございました。
ポインタ使える男の人ってモテそうだよね
ポインタでいろんな所を指して、 覗いたりいじくり回したりするんですね。
トラックポイントではない
デバイスから出る「the launch timed out and was terminated.」とかのメッセージを CPU側で受信する方法ってどうやるのかな?コンソール出力されるんだからできると思うけど
>>118 デバイス側といっても、cudart辺りがドライバからのステータスを見て出していると思われ。
従って、当然リダイレクトできる。
120 :
デフォルトの名無しさん :2011/11/07(月) 04:01:37.49
CUBLASとCUSPARSEを使って倍精度で計算を行うCG法のプログラムを作成したのですが、 CUSPARSEのcsrmv関数だけが異様に遅く速度が出ません。 この原因としては何が考えられますでしょうか? 自分としてはCUBLASとCUSPARSEでそれぞれ別のhandleを使っているせいかと思っているのですが それを確かめる方法や、CUBLASとCUSPARSEでhandleを共有する方法がわかりません...
CG法と書いてあるので問題行列は対称疎行列でCUSPARSEからはcsrmv関数だけ、CUBLASからは Level-1 BLASルーチンしか使っていないのだと思いますが、 実際にどの程度遅いのでしょうか? ランダムアクセスを含むcsrmvはLevel-1 BLASルーチンとくらべてメモリアクセス効率が極端に下がるのは 仕様というか当然の結果です。 また古いバージョンのライブラリを使っていたりしませんか? handle云々は基本的に関係ありません。関係するとすればよほど小さな行列を扱っている場合に限られます。
仰る通り対称疎行列でCUSPARSEからはcsrmv関数のみです。 行列のサイズはn×n n=122187で、ライブラリのバージョンは4.0です。 BLASでの総計算時間が12.66sに対し、 CUSPARSEでの総計算時間が80.36sという結果です。 このCUSPARSEでの計算時間はCPUで計算を行った場合よりも遅く、 nVidiaのCUDA toolkit 4.0 Performance Reportの結果と比較しても 非常に遅いと思います。
構造体の中にある配列はどのように確保してmemコピーすればいいんでしょうか? #define N (32) struct var_box{ int boxi; float boxf; }; struct var_str{ int *vari; // ~N*N float var_g; struct var_box *vb; // ~N }; のような構造体がある時 struct var_str *vvv,*vvv_d; vvv = (struct var_str*)malloc(sizeof(struct var_str)); vvv->vari = (int*)malloc((N*N)*sizeof(int)); vvv->vb = (struct var_box*)malloc(sizeof(struct var_box)*(N)); 値代入 cudaMallocHost((void**)&vvv_d,sizeof(struct var_str)); cudaMallocHost((void**)&vvv_d->vari,sizeof(int)*(N*N)); cudaMallocHost((void**)&vvv_d->vb,sizeof(struct var_box)*(N)); cudaMemcpy(vvv_d, vvv, sizeof(struct var_str), cudaMemcpyHostToDevice); cudaMemcpy(vvv_d->vari, vvv->vari, sizeof(int)*(N*N), cudaMemcpyHostToDevice); cudaMemcpy(vvv_d->vb, vvv->vb, sizeof(struct var_box)*(N), cudaMemcpyHostToDevice); GPUに送ってCPUに戻しただけですが float var_g; に関しては問題なくできていますが配列にした部分が送れていないみたいです。 そもそも確保の部分だけで0.5sかかっているのでちゃんとできてるかどうか怪しいです。
構造体にポインタがあるのはダメ 理屈で考えればわかるよな
馬鹿でプログラミング初心者には難しいれす(^ρ^)
無駄にでっかくメモリ確保された ただのポインタがコピーされてるだけに見えるのは気のせい? 実体をコピーしようとしている意図は読み取れるけれど、 コードはその意図を反映してないような。。。
あ、考え違いか、すまん
>>124 わからないのでその理屈を教えていただけないでしょうか
cudaMallocHostはあるけどcudaMallocはしてないの?書き洩れ?
デバイス側のcudamallocが無いよね? cudaMallocHost((void**)&vvv_d->vari,sizeof(int)*(N*N)); ってvvv_d->variをホスト側にメモリ確保しちゃってんるんだけど 必要なことは例えば cudaMallo(reinterpret_cast<void**>(&vvv_d->vari,sizeof(int)*(N*N)); じゃないの?
取り敢えず、ホスト側とデバイス側の二種類のメモリ空間を扱っていることを肝に銘じよう。 「構造体にポインタがある」ある場合、そのポインタがホスト側のメモリを指しているのか デバイス側のメモリを指しているのか常に意識しないといけない。 それがいやだったら、ポインタではなく配列にしてしまえ。 cf. struct var_str { int var[N * N]; float var_g; struct var_box vb[N]; };
それ以前に、そんなデータ構造を渡しても素直に並列化できなくて速度が出ない悪寒。
135 :
デフォルトの名無しさん :2011/11/08(火) 13:04:40.64
C# (.NET Framework 4.0)で動作するCUDAのライブラリってありますか? CUDA.NETがそうだとは思うのですが、.NET Framework 2.0 or newerとしか書いていないので、 結局のところ.NET 4.0で動作するのかどうかわかりません。
おまえ、そのレベルでよくこのスレに辿り着けたな。。
とりあえず動かしてみろw
4.1RC出てるんだな デベロッパプログラムに登録したいけど審査あるから、 他人のソースコンパイルして遊ぶだけの俺には無理
試しにデタラメ並べ立てて申請したら通った んでSDKのサンプルを上から順番に走らせて遊んでたらブルスク食らった 危ねえ・・・
>>138 審査あんの?俺普通に登録してOKでたけど
審査ってなにを調べるの?
ウンコの重さとか
Parallel Nsightのアカウントとは別なのか。 面倒くさすぎだろ。
CUDA.NETってCUDA 3.0までで止まってるよね?
>>144 くだんねーと(CUDANET)思ったんだろ作ってる方も
【審議中】 ∧,,∧ ∧,,∧ ∧ (´・ω・) (・ω・`) ∧∧ ( ´・ω) U) ( つと ノ(ω・` ) | U ( ´・) (・` ) と ノ u-u (l ) ( ノu-u `u-u'. `u-u'
kernelの中でどの部分でレジスタがMAXになっているか、どの値を保持しているかっていうのは簡単にわかるものですか?
細かいことは兎も角、取り敢えずptx出力を見てみたら医院で内科医。
shrQATest.hって4.0のtoolkitとSDK入れたらどこにあるはずですか?
%ProgramData%\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.0\shared
\inc
見つかりました ありがとうございます
linux kernelからGPGPU使いたいんだけどKGPU以外に何かいい方法ある?
CUDA developer用としてCUDAサイトで配布されてるドライバと普通のドライバってどんな違いがあるの? CUDAに必須と思ったら、普通のドライバインストールしてもCUDA開発できると書いている所があった。
一時期、developer用と一般用のバージョンが同じ時があったけど md5sumは一致した。 おそらく、一般用ドライバのバージョンで対応が間に合っていないときに beta版を出してるだけだと思う。 だから、両者でバージョンの新しい方を使えば問題ない。
158 :
デフォルトの名無しさん :2011/11/22(火) 17:29:12.79
すいません、経験者の方教えてください。 ホストメモリをページロックメモリで確保して、 ストリームを300個くらい起動して MemcpyAnsyncでh->d ,kenrel, h->dで非同期実行させてるんですけど、 どうやらkernelでセグって停止してしまいます。 しかし、cuda-gdbで見てみるとどうもセグるようなポジションじゃないんです。 そこで質問なんですが、MemcpyAsyncでコピーリクエストをだしてコピーが始まり、 制御がカーネルに入り、コピーが完了したストリームを実行してる途中に、 まだコピー中の領域でセグった場合、どこでセグったように見えるのでしょうか? やはりその時実行しているカーネルの中でセグっているように見えるのでしょうか?
全スレッドの計算結果を1コアに集約して1スレッドで処理したいんだけどいい方法あるかな? リダクションのサンプルだと最後はCPUで合計だしこれでは都合悪いんだよね 全てGPUでやりたいのよね
しかたないじゃん、CPUで合計出した方が速いんだから。
>>159 バイナリツリーで足していくしか。
Ex. thread数が64のとき
if (threadIdx < 32) data[threadIdx] += data[threadIdx + 32];
__syncthreads();
if (threadIdx < 16) data[threadIdx] += data[threadIdx + 16];
__syncthreads();
if (threadIdx < 8) data[threadIdx] += data[threadIdx + 8];
__syncthreads();
if (threadIdx < 4) data[threadIdx] += data[threadIdx + 4];
__syncthreads();
if (threadIdx < 2) data[threadIdx] += data[threadIdx + 2];
__syncthreads();
if (threadIdx < 1) result = data[threadIdx] + data[threadIdx + 1];
>>160 データ量によってはそうでもないよ。
>>161 >データ量によってはそうでもないよ。
>>159 は「1スレッドで処理したい」と言っている。
それに、「計算結果を集約」と言っているだけで足し算とは言っていない。
この条件の下では、データ量が多くなれば多くなるほどCPUで処理した方が速い。
データ数が極端に少なければ、転送時間の関係でGPUでしたほうが速いかもしれないが。
あーそうか、まさか1スレッドだけ動かす方法が判らんと言う話とは思わなかった。 単純に、「CPUで合計」の代わりをやりたいのかと思ったよ。
164 :
159 :2011/11/27(日) 23:55:11.26
>>160-163 レスサンクス
一度集約して最少二乗法を適用してから再度マルチスレッドで処理しようとしてるんだけど
一度CPUに戻すしかないかな・・・
最小二乗法ならマルチスレッドでできるじゃん。
リダクションしろよ 1ブロック分しかできないってこと?
サイズが決まっている3次元配列どうしの演算で演算回数も決まっているんだけど、 ブロック数とスレッド数を調整して演算回数ぴったりにするとコアレッシングしなくて効率悪いんだ。 こーいう時ってどゆアプローチ?
CPUで計算させてみる
もちろんCPUでは計算させてみてるよぉぉん(><) 1つ目の演算対象配列はギリギリL2に乗らないサイズ、2つ目の配列はVRAMにも乗らないサイズなので悩んでる。
170 :
159 :2011/11/29(火) 16:52:31.44
もっと良いCPUかGPU買え
>>159 同じブロック内の複数スレッドは共有メモリで集約できるけど、
ブロックをまたがっての集約ができないってこと?
手元のコードではブロック数の一時保存のための領域を確保し
そこにブロックごとの結果を放り込むカーネルを起動し、
次にブロック数が1、スレッド数が前のカーネルのブロック数の
カーネルを起動。その中でさらに集約して結果を1つにまとめている。
カーネル2回起動よりもCPUでやらせた方が速いかもしれないが。
カーネルの中で関数作って足せばいいんじゃない threadが終わったら自作関数に投げて足すとか バリアがいるけど
微妙に話が噛み合っていない悪寒。 つーか、>159が混乱に輪を掛けているし。
175 :
159 :2011/11/30(水) 10:04:46.33
各スレッドでサンプリングされたXiとYiを使って Σ(Ai*Xi-Yi)^2を計算して最小化するパラメータのAiを求めたいんだよ 求まったAiを使ってまたGPUでマルチスレッドで処理をするから CPUに戻したくなかっただけ 戻せばできるのは分かるけど、戻さないでできるならその方がいいと思ったわけね
なんか式がおかしい、Ai=Yi/Xi なんて答えを求めたい訳じゃなかろう?
177 :
159 :2011/11/30(水) 11:22:19.91
>>176 うん
正確には多項式近似の最小二乗法
めんどいので略した
178 :
171 :2011/11/30(水) 14:43:43.94
>>177 非線形最小二乗法ってこと?非線形最適化法しか知らないけど、目的関数が
特殊な場合にそれを上手に利用したのが非線形最小二乗って理解。
手元では対数尤度(二乗和みたいなもの)を171の手順で計算する関数を
CUDAで書いて、MCMCの中で使って推定値を求めているけど、同じことは
最小二乗法でもできると思うんだが…。
179 :
159 :2011/11/30(水) 14:50:41.39
>>178 カーネルの二度起動だよね?
その方法でもできると思うけど、一度で済ませたいというのがニーズ
それよかMCMCをマルチスレッドでやるってのが分からんのだが…
あれはマルコフ連鎖だからほとんどマルチスレッドで効果上がらんだろw
180 :
178 :2011/11/30(水) 15:21:59.61
>>179 こちらも一度の起動で済ませたかったけど方法が分からなかったので、
分かったら是非報告よろしく。
尤度の計算対象とするデータ数があまりにばかでかくて…。マルコフ連鎖は
順次計算だけど、その中の並列化は相当のメリットがあってね。
CUDA Visual Plofiler への入力ファイル(cpjファイルとcsvファイル)をコンソールを使って作成したいのですが、可能でしょうか。 GUI環境ではメモリが足りず実行できないプログラムのビジュアルプロファイルを取りたいのですが…
linuxで4.0のtoolkitを導入した時に古いものをリムーブするかと出てきたのでyesを選んだのですが、もしかして同じ階層のものを全部削除とかしています? X windowすら上手く開かなくなったのですが…
unixは知らんけど、windowsだとインストールしたファイルしか消されないよ 新しいの入れる度に毎回cl.hppを手動で追加してるけど、 アンインストールしてもそれだけ残ってる
Toolkitを入れたからじゃなくて、Driverを入れたからじゃね? Driverを入れるときにroot権限がないとXのconfigを更新できないから。
186 :
デフォルトの名無しさん :2011/12/06(火) 06:03:27.50
蔵人にメル栓抜きツイスターきたのか
コンパイルどれくらい早くなってんのかな VCでコンパイルしてて小さなcuファイルでも数秒待たされるのがガチでイライラするんだけど
vcでコンパイルしているなら、変わるわけないだろ。 vsからnvccを起動しているなら兎も角も。
cuファイルとcppファイルに分けるのはどういう意図があるんでしょうか?thrustって何に使えるんですか?
>>190 ・nvccに任せておきたくない部分は、.cppに書く。
・抽象的に取り敢えず書くのに便利。
cuにkernel_foo()を書いて cppに書いたfoo()からkernel_foo()を呼ぶのが定石かと思ってた そうすればopenclに移行したりしてもocl_kernel_foo()を呼び出すようにすれば変更すくないし。
>>192 その定石に則れば、>191の前半を満たせるから問題ないよ。
ただ、それだとCPU側とGPU側で共通のロジックを更に別のソースに書かないといけなくなるから適材適所だと思う。
194 :
デフォルトの名無しさん :2011/12/07(水) 00:20:58.16
青木氏の本読んでも、ガイド読んでも、くすだれcudaを見ても全くわからなかったので質問します。 おそらくwarpの話で、cudaをまったく分かっていないゆえの質問だと思います。 cudaのデータの取扱いを調べるために、以下のような構造のデータをGPU側に送り、 typedef struct data{ int i,j; //初期化でi=1,j=0 }DATA; Grid1コ,ThreadBlock1コ,総Thread数512コと指定して、以下のようなコードを実行させました。 __global__ void test(DATA *A){ int i = blockDim.x*blockIdx.x + threadIdx.x; if(i%2==0){//threadIdの奇遇でiの値及び加算値を変更 A->i=2; A->j+=1; }else{ A->i=3; A->j+=2; }} 5,6回実行して、iの値はthreadの総数を奇遇どちらにしても3で不変でした。 jの値は実行するたび値が異なり、j=3,5,7,9のいずれかになりました。 iの結果は各warpの32コのThreadが順次if文を実行してて、 idが奇数のときの場合が後に実行されるから、結果がi=3となるのか?という理解でよろしいのでしょうか。 また、jの結果は青木氏の言う「加算命令を実行した結果は有効なスレッドに対してのみ反映される」 の理解がいると思うのですが、そもそも有効なスレッドがどう判定されているのかわかりません。 また512コのthreadがあるのに、jの値の結果が10以下になるのはどうも腑に落ちません。 i,jの値を決めているものは何か、ご教示願います。
+=で同じアドレスに同時書き込みしてるから
196 :
デフォルトの名無しさん :2011/12/07(水) 01:09:59.92
>>195 +=で同じアドレスに同時書き込みすると、内部で何が起こるんですか?
競合状態が発生してんじゃないの? atomic演算とか同期が必要だと思うよ。 512スレッドで同一アドレスの変数の読み書きしてんだから。 まず512個の要素の配列作って、添え字にスレッド番号(上のi)を指定して確認してみたら? 書籍ならcuda exampleも買って読むといいかもね
198 :
デフォルトの名無しさん :2011/12/07(水) 02:05:18.21
>>197 >競合状態が発生してんじゃないの?
>>195のコメと合わせて考えるに、なんとなく予想はしてましたけど、取り合いになってるんですね...
>atomic演算とか同期が必要だと思うよ。
まだザックリとしか勉強してないので、atomic演算は知らなかったです。あとで試してみます。
同期を行う場合だったら、どうすればいいのだろう。
>まず512個の要素の配列作って、添え字にスレッド番号(上のi)を指定して確認してみたら?
それは分岐条件がきちんと実行しているのか見るためのテストをしたとき確認しました。
根本的にプログラミングの基礎が抜けている悪寒。
>>198 atomicという排他的処理を知らないのなら並列化プログラムをやるのは早すぎる。
増してCUDAという特殊な並列化をやるのはもっと早すぎる。
201 :
デフォルトの名無しさん :2011/12/07(水) 12:02:55.64
>>199 どこまで戻ればいいんでしょうか...
>>200 ...これ使って課題提出しなきゃならんので、早すぎると言われても後に引けないです...
202 :
デフォルトの名無しさん :2011/12/07(水) 13:13:17.16
>>198 >>201 ですけど、
>>200 の"atomicという排他的処理を知らないのなら並列化プログラムをやるのは早すぎる。"
では並列化プログラムをやるにあたり、
どういったことを勉強して、どういった手順でやればいいんでしょうか?
そこがよくわからず、私はいきなりcudaに突っ込んだので四苦八苦してるわけですが...
青木本読んだんじゃないのん? x+=1ってのは 1.xを読み出す 2.読み出した値に1を加算する 3.結果をxに格納する って手順なんだけど以下略
205 :
デフォルトの名無しさん :2011/12/07(水) 16:45:28.63
>>203 同期とロックですか。勉強します。
>>204 青木本は読みましたけど、ものにしたっていう状態じゃないです...
>x+=1ってのは...
普段、なんとなく使ってるので... 勉強不足ですみません。
206 :
デフォルトの名無しさん :2011/12/07(水) 17:29:38.47
207 :
デフォルトの名無しさん :2011/12/07(水) 17:58:07.38
209 :
デフォルトの名無しさん :2011/12/08(木) 00:42:53.72
>>208 ありがとうございます。明日早速本屋行ってきます。
210 :
デフォルトの名無しさん :2011/12/08(木) 17:37:10.02
CUDAというかVisualStudioが分からない。鬱だ死のう。
VSがわからないんじゃなく、C++がわからないんだろ
fortranもここでいいかな? pgiのコンパイラ買って、三重ループの前後に指示行追加してやったけどまったく速くならない。憂鬱。。
fortran使う人って量子論とかやってる人? どの世界の人が使うのかしら?
fortranを使うのは過去の遺産のためだろう。 fortranを使えば速くなるわけでもないしなあ。
fortran懐かし過ぎる 先日、二度と使うことは無いと思って本を捨てたばっかりだわ
高レベル言語使うのは時代の趨勢じゃないかな。 LAPACKはfortranで書かれている。 fortranはnvccよりもアプリケーションに近い。 CUDAを使うのにnvccのレベルでなきゃダメということは無いと思う。 逆に、nvcc使わない人をバカにするヤツは、本物のバカだと思う。
219 :
デフォルトの名無しさん :2011/12/12(月) 20:20:14.55
道具にこだわることが目的となって、肝心の結果を出せない人間を馬鹿というでは?
TSUBAME2.0でnvccを使うのと京でfortran使うのと比べるとどうなんだろう。 京はSPARCだけで計算して、TSUBAME2.0の10倍くらいの速度みたいだ。 東工大は教育機関なのに対して、理化学研究所は結果を求められる独立行政法人。
既存のコードを使う限りでは京のほうが速いだろう。 ただ、富士通のコンパイラがダメダメだと聞いたけど、 それは解消されたのかな?
222 :
デフォルトの名無しさん :2011/12/12(月) 22:29:03.43
富士通のコンパイラは、8コアCPUへの割り当ては自動らしい。 これ、すごいことだと思うんだけど、 このコンパイラは完成しているのだろうか。
>>222 openMPに対応してたら各CPUのコアへの割り当ては普通にコンパイラ側でやると思うけど?
>>223 OpenMPがそんな事やるかよ。やるのはOSだぞ。
>>222 が言っているのはNUMAのことだ。
現時点でも普通のLinuxならnumactlを使えばかなりのケースでノードの割り当てができるようになる。
CUDAを使う場合でも、複数GPUを使う場合に、MPIを使う場合に有効だ。
なんだ?OpenMPやNUMAを何のことだと思ってるんだ? というか釣られたか?
>>222 > これ、すごいことだと思うんだけど、
スパコンの世界では、ずうううっと前からやってるんじゃない?
じゃあ、そろそろ俺らも本気出すか
229 :
デフォルトの名無しさん :2011/12/14(水) 20:16:13.92
4Gamer.net ― NVIDIA,CUDA 4.1をリリース。CUDAコンパイラのソースコード公開も
http://www.4gamer.net/games/076/G007660/20111214033/ 中国時間2011年12月14日,NVIDIAは「CUDA 4.1」をリリースした。
最大の注目点は, LLVMベースとなるCUDAコンパイラが搭載されそのソースコードが公開された点だ。
LLVM(Low Level Virtual Machine)はAppleなどが参加する“言語非依存”のコンパイラ環境だ。
NVIDIAは同時に「CUDAプラットフォームのオープン化」を明言し,
LLVMベースとなるCUDAコンパイラのソースコードを研究者やツールベンダーに公開している。
要するに,NVIDIAのGPUでしか利用できなかったCUDA環境がほかのCPUやGPUに広がっていく可能性が出てきたのだ。
たとえばCPUや,それこそRadeonなどの他社製GPUでもCUDAを利用できる可能性が出てきたわけで,
CUDAの標準化をさらに推し進める起爆剤となり得るのである。
これって、個人でもソースもらえるのかな
へえ、ソースを公開したんだ。 でも思いっきりOpenCLとかぶるな。 OpenCLが死亡か。
なんで、OpneCLが死亡するんだ?
まだ来てなくない?
Radeon用CUDA作っても、特定の機種用になりそうだな
そもそもRADEON上でCUDAってマトモに動くのか? 一応HD3000世代のRV670からはFP64対応らしいが
動かないよ 動かそうとする人も現時点ではいないだろ まだまだGPGPU向けには問題ありそうだし
cudaスレだが、 研究室や自分の周りでけで高速化すればいいんだったらcudaでいいんだろうけど 製品に組みこむならintelの普及度やなるべく多くの人に使ってもらえることを 考えるとOpenCLのほうがいいんじゃないかと思ってる。 photoshopのGPU対応もOpenCLでやってるみたいだし。
Photoshopは現状CUDAとATi Streamです 高速化しようとすると結局はデバイス毎に最適化したプログラムを書くことになるので、 デバイスに特化したライブラリの方が融通が利きます Parallel NsightでのデバッグはOpenCLではできません
239 :
デフォルトの名無しさん :2011/12/20(火) 03:38:50.36
CUDAのハードウェア種別を問わない展開が可能になったという事は、 CUDAの上に乗っかる形で実装されているPhysXやOptiX等もまた ハードウェア種別を問わず動作可能になる可能性があるという事だよな 期待したい
別にソースコードなくたって、ある程度の知識を持った人なら CUDAを移植できたと思うけど
RADEONでPhysXはマーケティングの理由で使えないんだろ?
AMDがCUDAを実装されるの嫌がってPhysX蹴ったんだっけ
244 :
デフォルトの名無しさん :2011/12/21(水) 19:57:27.52
質問: GPU側に2次元配列を送るには、 1.2次元を1次元に直してからcudaMallocでメモリ確保→cudaMemcpyで送信 2.cudaMallocPitchでメモリ確保→cudaMemcpy2Dで送信 という2つの方法がありますけど、 構造体のメンバーに2次元配列があるものをGPU側に送る場合はどのようにしたらいいんでしょうか?
ポインタじゃなく配列ならそのままでうまくいかんか?
246 :
デフォルトの名無しさん :2011/12/21(水) 21:23:30.55
>>244 再度検証してたけど、サジ投げた。
最終的には以下の構造体をGPU側に送信して処理した後、CPU側に戻したい。
typedef struct test{
int a; int b[5];
}TEST;
typedef struct data{
int c; TEST d[1][10];
}DATA;
上をcudaMallocでメモリ確保&cudaMemcpy送信しようとしてるけど、違うのかな。
教示お願いします。
そのまま転送すればいいじゃん
248 :
デフォルトの名無しさん :2011/12/21(水) 21:49:19.70
>>246 です。
すいません。できました。お騒がせしました。
>>229 このコンパイラって、謎のCをPTXに変換する部分だけじゃないの?
今のフロントエンド(cudafe)は、他社製品のEDGベースだから、出てこないと思うが。
cudafe -> PTXへのコンパイラ(今だとnvopencc) -> ptxas
のまんなかが出来るだけなので、この部分のソースがあったところで、CUDAコンパイラを自分で作れるとか、
CUDAを他のプロセッサで動かせるようになるとかいうものではない。
.cファイルがコンパイルできません たすけて
んvccじゃできんだよ
252 :
デフォルトの名無しさん :2011/12/26(月) 15:17:58.63
あの、OpenCLってオワコンなの?IntelとAMDが支持してるのに全然盛り上がってないんだが。
科学技術計算向けがtesla一択だから他を使う理由もないってっところなのかね。 わざわざcudaで組んだものを他に移植する理由もないだろうし。
254 :
デフォルトの名無しさん :2011/12/26(月) 21:42:01.58
OpenCLで書くと長くなってしまう。 それだけのこと
C++ Bindings使えよ。
ちょっとしたラッパーかくのも嫌なのか
257 :
デフォルトの名無しさん :2011/12/27(火) 11:26:29.65
超素人の質問です。 コマンドラインでnvccを使いコンパイルを行っているのですが、OpenGLのサンプルコードをコンパイルできません。 freeglutを入れてます。 gcc -lglut program.cではコンパイル可能でした。 しかしnvcc -lglut program.cuとするとコンパイル不可となります。 -Lや-lによってライブラリやヘッダファイル先を指定したところで、コンパイルできませんでした。 阿呆な質問かもしれませんが、自己解決できそうにないのでよろしくおねがいします。
凄いアフォだな エラーメッセージすべてコピペし
259 :
帝徒=繪璃奈=啓北商業の野島えり :2011/12/27(火) 14:53:52.03
主犯 少頭劣一族=蔗冽一族とは 中国 華喃の山の梺の村八分の家。 二間位の横長で玄関の右側がお勝手。 大正に猿のままで生まれたのが 鈴木あゆみ(网(アミ) 范襤の子) フィリピン人の范襤と 同じくフィリピン人のモンゴルに逃げた『シバ』との間の男児。 日本名 鈴木ひろしと聞いた。 その後、親戚の鈴木大樹を殺し 戸籍を使用。 今は一文字 雉。
母体になるプログラムなしで、 CUDAのコード単体で動かせるツールって無いかねぇ。 外出先でCUDAのサンプル見せるのがメンドイ。 例えば ./cuda-run < source.cuda みたいな感じで実験できるといい感じ。 codepadみたいな感じで動かせればなおよしだけど。
261 :
デフォルトの名無しさん :2011/12/29(木) 22:01:29.67
そういうスクリプトかけばよい
>>260 >母体になるプログラムなしで、
それってなんのこと? 外出先のPCで動かしたいってこと?
コンパイラのことなら、事前にコンパイルしておくかリモートでアクセスすればいい。
ドライバのことなら、外出先の端末にインストールさせて貰うしかない。
外出先のPCなんか頼らず、自前のPCを持っていけばいいじゃん。
263 :
デフォルトの名無しさん :2012/01/05(木) 13:36:51.62
>>257 手元に開発環境がなく、エラーメッセージをコピー忘れしてました。すみません。
# cc Drawtest.c -lglut
# ./a.out
(空のウィンドウが製作される)
(Drawtest.cをDrawtest.cuに中身は同じで拡張子だけ変えて)
# nvcc Drawtest.cu -lglut
/usr/bin/ld: skipping incompatible /usr/local/cuda/bin/../lib64/libglut.a when searching for -lglut
(-lglutを見付かられなかった?)
# nvcc -L /usr/local/cuda/lib64/ -l /usr/local/cuda/include/ Drawtest.cu -lglut
/usr/bin/ld: cannot find -l/usr/local/cuda/include/
collect2: ld はステータス 1 で終了しました
となりました。
インクルードパスの指定は -l じゃなくて -I な
/usr/bin/ld: skipping incompatible〜って64/32ビット混在時のエラーメッセージじゃなかったかな
266 :
デフォルトの名無しさん :2012/01/05(木) 15:09:35.86
>>264 なんというミス。
目で見て違いが分からない。
L(小文字)→i(大文字)に変えて同じようにしてみました。
# nvcc -L /usr/local/cuda/lib64/ -I /usr/local/cuda/include/ Drawtest.cu -lglut
/usr/bin/ld: skipping incompatible /usr/local/cuda/lib64//libglut.a when searching for -lglut
/usr/bin/ld: skipping incompatible /usr/local/cuda/bin/../lib64/libglut.a when searching for -lglut
結局パス指定なしと同じ見たいですね。
おれも64/32の違いかと思ったが、パス指定から指摘してみた 実行している環境は? -m64 オプション付けてみ
268 :
デフォルトの名無しさん :2012/01/06(金) 13:53:42.06
# nvcc -L /usr/local/cuda/lib64/ -I /usr/local/cuda/include/ Drawtest.cu -lglut -m64 /usr/bin/ld: skipping incompatible /usr/local/cuda/lib64//libglut.a when searching for -lglut /usr/bin/ld: skipping incompatible /usr/local/cuda/bin/../lib64/libglut.a when searching for -lglut ううむ・・・・。 centOS 64bit Device 0: "Quadro 4000" Device 1: "GeForce 9500 GT" パスは LD_LIBRARY_PATH = /usr/local/cuda/lib64 です。 openGLのないコードはコンパイル及び実行まで問題ありません。 openGLを使うにあたって /usr/local/cuda/include/にglut.h freeglut.h freeglut_ext.h freeglut_std.h /usr/local/cuda/lib64/にlibglut.a をコピペしましたが、まだ足りないのでしょうか?
libglutは64ビット? 32ビットなんじゃねか?
270 :
デフォルトの名無しさん :2012/01/06(金) 15:03:44.65
glut.h freeglut.h freeglut_ext.h freeglut_std.h libglut.aをccが見ているであろう先のものと置き換えてしてみました。 # nvcc -L /usr/local/cuda/include/ -I /usr/local/cuda/include/ Drawtest.cu -lglut -m64 /tmp/tmpxft_000030a4_00000000-12_Drawtest.o: In function `disp()': tmpxft_000030a4_00000000-1_Drawtest.cudafe1.cpp:(.text+0x14e6): undefined reference to `glClearColor' tmpxft_000030a4_00000000-1_Drawtest.cudafe1.cpp:(.text+0x14f0): undefined reference to `glClear' /usr/local/cuda/bin/../lib64/libglut.a(libglut_la-freeglut_init.o): In function `glutInit': … (中略) … (.text+0x898): undefined reference to `glPopClientAttrib' collect2: ld はステータス 1 で終了しました openGL系の関数が読まれていないみたいです。
いや、だからfreeglutは64ビットを使用してるのかと聞いてるんだが 質問に対しきちんと回答できないようでは教えようがない罠
272 :
デフォルトの名無しさん :2012/01/06(金) 18:39:39.36
>>271 freeglutは64ビットのものを入れ直してやってみましたが、同じ結果になりました。
>>257 そこでCUDAのヘッダフォルダとライブラリフォルダに入れていたものを消し、-Lと-Iで直接freeglutのフォルダを指定したら
解決しました。
ライブラリが足りなかったのか、それともCUDAと一緒くたにしたのがダメだったのか解かりませんが、ちゃんとPATH通せってことみたいです。
お騒がせしました。
274 :
デフォルトの名無しさん :2012/01/06(金) 20:14:34.67
環境
Windows7 Professional 64bit
Microsoft Visual C++ 2010 Express Version 10.0.40219.1
Microsoft .NET Framwork Version 4.0.30319
GeForce GTX 580
CUDA Toolkit 4.0.17
SDK 4.0.19
devdriver_4.0_winvista-win7_64_270.81_general
このサイトを参考に環境を構築しました。
http://feather.cocolog-nifty.com/weblog/2011/07/visual-studio-2.html そして以下のサイトのサンプルプログラムを実行してみました。
http://www.gdep.jp/page/view/218 Hello,Worldと99 bottles of beerはcpu、gpu共に実行できました。
しかし、Matrixのプログラムはcpuの方は実行できるのですがgpuの方が実行できません。
以下のエラーを吐きます。
matrix_gpu.cu(5) : fatal error C1083: Cannot open include file: 'cutil_inline.h'
: No such file or directory
どうやらVisualStudioのパスがうまく通ってないということまでわかり、以下のサイトなどを参考にCUDA_INC_PATHなどを変えてみましたが、一向に変わりません。
http://d.hatena.ne.jp/Levi/20090921/1253535802#c SDK内のcutil_inline.h自体をtoolkitのincフォルダにコピペすると、他の.hファイルもいくつか同じエラーが出たのでエラーになったものをすべてコピペしたところ、LNK2019"link.exe"というエラーで先に進めませんでした。
一度VisualStudioを再インストールしてみましたが、状況は変わりません。
Nvidia GPU computing SDK Browserではサンプルプログラムを実行できているので、CUDAの環境は整っていると思われます。
どうすれば解決できますでしょうか・・・。かれこれ1週間以上格闘しています。
VisualStudioは2010よりも2008にした方がいいでしょうか?
C初心者にはきついと思うんだが… とりあえず 'cutil_inline.h'のある場所を見つけて そこを-I /'cutil_inline.hのある場所'と指定する 意味わからなければCのコンパイルを勉強すること
274ではありませんが、「Cのコンパイル」の勉強にオススメの書籍とかありましたら紹介していただけるとうれしいです
275だけど、お薦めは他の人に任せるけど 一つ言えるのはCUDAはC/C++より数段難しい C/C++の質問をしてるようでは先が思いやられる、と正直思うんだが… 超低レベルな質問もOKなくだすれだけど、その程度は自力で調べる能力が必要ではないだろうか?
>>275 >>274 です。ありがとうございます。
今実行できる環境にないので後で試してみます。
また結果報告しにきます。
279 :
デフォルトの名無しさん :2012/01/07(土) 12:14:43.10
一時期は同価格帯のCPUとGPUを比較してGPUは50倍の性能があるって言われてたけど 今10〜15倍程度しか無くない? 3930k 5万円 158GFlops GTX590 6.3万円 2480GFlops GTX 580 3.8万円 1500GFlops 次に出ると言われてるGTX700系は500系の2倍の性能らしいけどそれでも50倍には遠く及ばない。 この「同価格帯でGPUはCPUの10〜15倍程度の性能しかない」という評価は正しいだろうか?
>>279 それは単にCPUは1コアとして比較しているからGPUを使えばうん十倍と言われるんだよ。
単純にメモリ帯域で3〜4倍、演算器としては、単純にAVXの8並列x6コアxクロック2倍で比較すると5倍にしかならない。
CPUは大きなキャッシュがあるから局所性の高い処理の場合はメモリ帯域の差が縮まるので、10倍差がつくならいい方。
もちろんアプリによるけど、CPU向けのコードがカリカリに最適化されていた場合は思ったほど差がつかない印象。
それでも3〜5倍位は早くなるからいいんだけどね。
問題はメンテナンスだな。。。CUDAのみと割り切れればいいんだけど、なかなかそうは行かないからなあ。
>>279 GTX580で500GFLOPS程度。そのデータはおかしい
CPUなんてE8500で25GFLOPSくらいだから20倍程度だな。
もちろん最近のコア多いやつならもっと差は縮まるだろうな。
それに理論値の話だから多数のスレッドでまわしたらコンテキストスイッチの
オーバーヘッドとかGPUプログラミングは難しくて効率上げるの難しいとか
そういうので実効は変わってくるんじゃね?
あ、同一価格帯か、E8500が2万しなかったことを考えると 580にそろえて2倍したとして10倍程度だから、まあおおざっぱには 「同価格帯でGPUはCPUの10〜15倍程度の性能しかない」は正しいんじゃないかと思う。 580と同時期の4万弱のCPUならもっと差は縮まるだろうし。
32nmと40nmで比較とか
void *p; CUDA_SAFE_CALL(cudaSetDevice(0)); CUDA_SAFE_CALL(cudaMalloc(&p, 16)); でcudaMallocからcudaError_enum例外が飛ぶのですが 原因はなにが考えられますか? ちなみにSDKのサンプルは動きます。
285 :
284 :2012/01/07(土) 22:23:52.14
異なるディレクトリに複数のCUDAのライブラリが入ってて 古いバージョンが参照されてるのが原因だった・・・
CUDA4.1でcomputeProfはnvvcになったんですか? だいぶ代わってしまって驚きますた
nvvp
>>282 GPUの論理性能ってFMAとかMADとかの積和で倍になっているから、
GTX 580 1500GFlopsとかって、それらを使わなければ750GFlopsとかになるんだよな。
だから実効れべるではもっと縮まるんだろうな。それでも5倍以上差がでるとかなり有効なんだけど。
いる。というか俺が買う。
こういう需要が限られていてかつ堅くて高い製品はamazonに出した方が高く売れるんじゃないかねえ。
>>275 274です。
-Iオプションを付けてヘッダファイルのある場所を指定してみましたが、相変わらず同じエラーです。
C:\cuda_practice\test>nvcc -o matrix_gpu.exe matrix_gpu.cu -I C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.0\C\common\inc
matrix_gpu.cu
matrix_gpu.cu(5) : fatal error C1083: Cannot open include file: 'cutil_inline.h': No such file or directory
自分でも試行錯誤してみますが周りにcudaがわかる人がいなくて行き詰ってます・・・。
もっとコンパイル早くなんねーかなぁ… thrust::sortとか使うようにしただけでコンパイルに十数秒かかってイライラする
速いマシンでやる。
nvccコンパイラをGPGPUで実装すればいい! CUDAがオープンになるどさくさにまぎれてそういうのを作る人が出るかもしれないこともないかも
コンパイラの動作は現状GPU向きとは言えないんじゃね? GPU自体が単純な性能向上だけではない、思いもよぬ進化でもしない限り
コンパイルなんて依存関係ありまくりで、並列処理が苦手な最たるものだからまずやる意味がないだろう。 CUDAにすれば何でも速くなると思ってんじゃね?
sharedメモリって同一WARP内でのやり取りだったら同期なしで大丈夫なのでしょうか? たとえば全部で32スレッドで次のkernelを実行した場合、 WARP内で一斉にsharedメモリに書きに行くので同期しないで大丈夫かと 思ったのですが、実際にはうまくいきません。 globalメモリに読みに行く段階でコアレッシングが発生していないので それが原因なのでしょうか? どなたか教えてください。 __global__ void kernel(float *g_v, float *g_x){ float x = 0.0f; int i = blockDim.x * blockIdx.x + threadIdx.x; __shared__ float s_v[32]; s_v[i] = g_v[i+i%3]; __syncthreads(); // これが必要かどうか? x = s_v[(i+3)%32]; g_x[i] = x; }
>>298 iが32以上になることはないかい??
ブロックの数は??
301 :
274 :2012/01/10(火) 17:52:39.67
>>299 書き込みしてからそのミスに気付きました・・・。
それを直してもうまくPathが通らなかったので、CUDA4.0、VisualStudio2008でやることにしました。
303 :
298 :2012/01/10(火) 17:59:16.44
>>300 ごめんなさい、書いたコードが一部誤ってました。
iが32以上になることがあります。
ブロック数は数100程度になります。
このときは、下のソースのようになると思うのですが、
やはりうまくいきません。
__global__ void kernel(float *g_v, float *g_x){
float x = 0.0f;
int i = blockDim.x * blockIdx.x + threadIdx.x;
__shared__ float s_v[blockDim.x];
s_v[threadIdx.x] = g_v[i+i%3];
__syncthreads(); // これが必要かどうか?
x = s_v[(threadIdx.x+3)%32];
g_x[i] = x;
}
>>303 コンパイル後のバイナリを見たわけじゃないから予測だが、
s_v[threadIdx.x] = g_v[i+i%3];
x = s_v[(threadIdx.x+3)%32];
g_x[i] = x;
は
1.1:tmp=g_v[i+i%3]; //グローバルアクセスなので遅い
1.2:s_v[threadIdx.x]=tmp; //
2.1:x = s_v[(threadIdx.x+3)%32];
3.1:g_x[i] = x;
みたいに解釈される。
で、単純な式だからコンパイラは
s_v[threadIdx.x]とs_v[(threadIdx.x+3)%32];
が必ず別のアドレスになる事を検知して
1.1 2.1 3.1 1.2のように命令を入れ替えてしまう。
だから__threadfence_block()がいるはず。
305 :
298 :2012/01/10(火) 18:17:13.87
>> 305 だいたい x = s_v[(threadIdx.x+3)%32]; の前にすべてのs_v[??}が定義されていないといけないだろ?? スレッドが32に設定されているのならば (threadIdx.x+3)%32 は0か1にしかならないが...
>>301 -lcutil32 とその参照パスを追加すれば Visual Studio 2010 でもコンパイル,実行できる♪
ちなみに私の場合は
nvcc matrix.cu -I "C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.0\CUDALibraries\common\inc" -L "C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.0\CUDALibraries\common\lib\Win32" -lcutil32
tmpxft_000024b8_00000000-3_matrix.cudafe1.gpu
tmpxft_000024b8_00000000-8_matrix.cudafe2.gpu
matrix.cu
tmpxft_000024b8_00000000-3_matrix.cudafe1.cpp
tmpxft_000024b8_00000000-14_matrix.ii
C:\cuda> a
Processing time: 1936.354858 (msec)
Press ENTER to exit...
ちなみに家のはGeForce GT 220なので非力....
309 :
306 :2012/01/10(火) 22:21:29.79
間違えた ↓ (threadIdx.x+3)%32 は0か1にしかならないが... 32で割った余りであった...
310 :
298 :2012/01/11(水) 09:16:02.50
>>306 >>309 いろいろ試してみたところ、とりあえず上記の問題は解決しました。
回答ありがとうございました。
すごい。 見た感じ、ちょっと不安だったから 参考になったみたいでよかった。
312 :
274 :2012/01/11(水) 19:30:44.93
>>307 VS2008でもうまくできませんでした。
再度VS2010を入れてみるも手順通りいかず、当然実行もできませんでした。
リカバリしたのでこれからもう一度2010でやってみます。もう心が折れそうです・・・。
ちなみに2010expressとprofessionalの違いで実行できないこととかあり得ますか?
別のPCでprofessionalで試してみてくれた方がいて、うまくできたようなのですが・・・。
>>312 どんなエラーがでますか?
そもそもnvccでコンパイルするのだから,Visual Studioは関係ないでしょう?
VisualStudioのないLinux上でもコンパイル,実行できるはず.
もう一度やるのならばライブラリ lcutil32を付ければ良いかと(32ビットならば)
314 :
274 :2012/01/12(木) 01:51:21.33
>>314 はじめはVSのプロジェクトを構築するのではなく、VSの(32bit用の)コマンドプロンプトを立ち上げて、そこで
コマンドを打ち込み、コンパイル、実行した方がよいのでは。
これがうまく行ってから、プロジェクトの構築をした方がよいでしょう。
cutil64.libがないと言うことは64bit用のSDKがインストールされていないのでは?
ここにあるはず。
C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.0\CUDALibraries\common\lib\x64
プロジェクトを構築するときに32bit, 64bit、あるいは両方の実行ファイルを作るかどうか設定できますよ。
>>314 64bit版のcl.exeがある。
見つからないなら-ccbinオプションでcl.exeの場所を指定。
cutil32.libとcutil64.libを両方指定してはいけない。どちらか必要なものを。
で、cutil64は確か自分でビルドする必要がある。
SDKのディレクトリの中にプロジェクトファイルがあるからそれをビルド。
http://codepad.org/Kyns70af 上記のプログラムが期待と違う結果を出力するのだが原因がわかる方がいれば教えていただけないだろうか.
(CUDA 3.2 + GeForce 320M + Fedora 14 環境と CUDA 3.2 + Tesla C1060 + CentOS 5.5 環境でテスト)
-- 期待した出力 --
1.00
0.00
1.00
0.00
1.00
0.00
1.00
0.00
-- 実際の出力 --
0.00
0.00
0.00
0.00
0.00
0.00
0.00
0.00
ちなみに,11行目の右辺を 2.0f に変更すると期待通りの結果を出力してくれる.
-- 出力 --
1.00
2.00
1.00
2.00
1.00
2.00
1.00
2.00
318 :
317 :2012/01/12(木) 08:37:35.44
やってみたけど、問題は再現されなかった cuda_funcの中に printf(" %d %d %f\n",threadIdx.x,threadIdx.x & 1,a[threadIdx.x]); を入れて、チェックしてみれば?? ちなみにこんな結果となる 0 0 1.000000 <- cuda_funcの中では問題はなさそう 1 1 0.000000 2 0 1.000000 3 1 0.000000 4 0 1.000000 5 1 0.000000 6 0 1.000000 7 1 0.000000 1.00 <-- 外でも問題はなかった 0.00 1.00 0.00 1.00 0.00 1.00 0.00
>>317 同じく問題再現されず(CUDA4.1 + GF560Ti)
そのデバイス両方とも compute capability 1.x だし
2.x 環境で試してみたらどうだ
321 :
317 :2012/01/12(木) 10:56:51.02
>>319 >>320 コメントありがとうございます
現在の私の環境ではカーネル関数内で printf は使えないですが
変なことになるのは私の環境だけのようなので Compute capability 2.x 環境で後日試してみます
ちなみにcuPrintf.cuh とcuPrintf.cu をネット上からひらってくれば、printfと同等の関数(cuPrintf)が使えますよ♪
CUDA 3.2のコード出力に問題がある気がします。 カーネル部分をどうも a[threadIdx.x] = (float)(FUNC(-(threadIdx.x & 1)); に変換しようとしている気配が。 FUNCはISET.S32命令による処理だけどそれがミスっているとしか思えない。 自分で以下のように書くか、あるいはFermi以降またはCUDA4.0以降を使うべし。 a[threadIdx.x] = (float)(1 - (1 & threadIdx.x)); ちなみに実行はしてないので同じ問題が起こるかどうかは知りません。
324 :
274 :2012/01/12(木) 13:42:05.81
>>315 >>316 ありがとうございます。
とりあえず315さんの通りコマンドライン上で実行できるところを目指そうと思います。
cl.exeを指定するオプションがあるんですね。
自分はコンパイラオプションの勉強不足のようです・・・。
cutil64は自分でビルドするんですか、それは今までやってなかったと思います。
恐縮ですが、これから環境を整えるので何か参考になるサイトがあれば教えていただけますか?
いくつかサイトを見て設定したのですが、どれもうまくいかなかったので・・・。
何度もすみません。よろしくお願いします。
質問です。
同じkernel関数をfor文で何度も実行するのですが、
回数によって1回あたりのkernel関数の実行時間が異なるみたいです。
具体的には
for(i=0; i<10000; i++){
kernel_foo <<<grid_size, block_size >>>(...);
}
とすると300,000msかかるのが
for(i=0; i<100; i++){
kernel_foo <<<grid_size, block_size >>>(...);
}
では1,000msくらいになります。同じ関数を実行しているので
単純に考えれば10000回ループを回したときは100回の100倍時間がかかるはずですが、
実際には300倍程度になりました。
もう少し細かく見てみると、最初60回程度はkernel関数を1回実行するのに0.04ms程で済んでいたのですが
それ以降は1回あたり25ms程になっていました。
似た事例として
ttp://d.hatena.ne.jp/ymizushi/20091214/p1 という記事を見つけましたが、この記事内でも原因はよく分かっていません。
どなたか詳細をご存じないでしょうか?
327 :
274 :2012/01/12(木) 17:37:04.49
>>326 ありがとうございます。
全て拝見しましたが既読のものでした。やはり普通に設定すればできるようですね・・・。
今格闘しているので、また報告します。
328 :
325 :2012/01/12(木) 17:50:39.27
>>326 ありがとうございます。
全項目について確認してみましたが、
1. 演算数はどれも一定
2. ループ回数によらず正しい値が得られている
3. 測定時間中にcudaMemcpy等の関数は呼ばれていない
という状態です。
>>328 1プロセスで2つのカーネル実行してるなら順番入れ替えて試してみた?
最初の実行は起動コストだかで時間かかるのが普通なんだが
時間の計測が間違っている。 cudaThreadSynchronize()を読んだあとに終了時刻を取得する必要がある。 1024個までのカーネル呼び出しはキューに積むだけですぐに次に進んでしまう。 30000回ループの方はキューが詰まるので(30000-1024)回くらいの実行完了を待つことになり、比較的妥当な計算時間に見えるはず。 もちろん正確ではないのだが。
331 :
325 :2012/01/12(木) 19:35:41.01
>>329-330 回答ありがとうございました。
cudaThreadSynchronize()を入れるのを失念しておりました。
cudaDeviceSynchronize() がナウいらしい
333 :
sage :2012/01/13(金) 08:22:49.98
>>330 便乗して質問ですが、
キューに一度に入るカーネル数が1024個で、
それ以上はカーネル呼び出しが待ちになるというのは、
CUDAのマニュアルかどこかに書いてありますか?
>>333 実験したらそうだったけど、文書化されているかどうかは知らない。
これだけ多ければ普通は問題ないし、バージョンアップで変わるかも。
>>334 了解しました。
どうも有難うございました。
それにしても きゅー にカーネルが溜まっているからと言って,GPUの処理能力が落ちるとは思えないぞ
きゅー にカーネル田丸? それってstream使ってる時のこと? 文脈呼んでないからわからんけど
キューに空きがあれば積んで即制御が戻る。 空きが無ければ空くまで待ってる。 詰まってるから遅くなるわけではなく、遅い方が正しい実行時間なだけ。
CUDA4.0からデバイスエミュレーションがなくなったらしいけど 実機で実行するとOSごと固まることがある みんなどうしているのでしょうか
デバイスを物理的に叩き割る
2.3くらいからなかったような気が
>>342 メモリアクセスのバグがあるとWindowsが固まって動かなくなったり
ディスプレイにノイズが出てくる
OSが画面表示に使っている領域を壊しているようだけど
普通は起こらないことなんでしょうか?
私の経験では、CUDAのプログラムをチェックするために printf を使って変数を書き出したところ、 その量が多すぎて、画面が真っ暗になったり、PCがシャットダウンしたことがあった。 書き出し量を減らしたところ、問題はなくなった。 こんなことかな??
345 :
344 :2012/01/21(土) 11:00:22.89
追加です。 グラフィックカードを2枚刺して、一枚をCUDA用、もう一枚をディスプレイ用にすればよいとのことを聞いたことがあるが、 私のはGTX590でカードが2枚、入っているとのこと。 CUDAの計算に Device(0)にしても(1)にしても、前述の問題が起きたので、枚数には関係ないこともあるらしい??
printf使えるとか便利になったわ
>>344 fermiではないのでprintfは使えないです。
GTX260なんですけど、それが悪いのかも。
計算用GPUを別に刺すってのは考えていました。
ただ、そうするべきという常識みたいなものがあるのかなっていう疑問があったので。
あとこういう状態でデバイスエミュレーション外したならひどい話だなと思ったのですが
みなさんあまり起きていないようで。
グローバルメモリの添え字計算計算が間違っていたりで確保した範囲外にアクセスすると Windowsが固まったりディスプレイにノイズが出てくることがあるって状態です。
グラフィックカード二枚差しにすると、 nSghtでハードウェアデバッグができるとかじゃなかったっけ?
350 :
デフォルトの名無しさん :2012/01/22(日) 17:02:34.79
YES もしくはPC2台で片方をデバッグ用にしても良い。 実際に動かすGPUがCUDA対応なら、開発側は非nVIDIAでもOK
streamについて分からない点があるので どなたか教えてください。 以下で「カーネル」はカーネル関数の呼び出し、 「転送」はcudaMemcpyAsync、()内はstreamだとします。 (質問1)以下の状態のとき、転送(1)は、カーネル(1)と同じstreamなので 実行できませんが、キュー内でそれより後にある転送(2)は実行できますか? カーネル(1) 実行中 転送(1) キュー内の最初のタスク 転送(2) キュー内の2番目のタスク (質問2)以下の状態のとき、転送(2)は実行できますか? (キュー内で転送(2)の前にstream(0)のタスクがあります) カーネル(1) 実行中 カーネル(0) キュー内の最初のタスク 転送(2) キュー内の2番目のタスク (質問3)以下の状態のとき、転送(1)は実行できますか? (実行中のタスクはstream(0)です) カーネル(0) 実行中 転送(1) キュー内の最初のタスク 「CUDA C Programming Guide」を読んでもよく分かりませんでした。 宜しくお願いします。
1はFermiならOK。古い世代ではダメ。 2と3は絶対ダメ。stream0は前後をぶった切る役目がある。
>>352 1の動作がFermiとそれ以前で異なるのは知りませんでした。
回答どうも有難うございました。
Fermiでプロファイラのtimestampを使ってテストしたところ カーネル(1) 実行中 転送(HostToDevice)(1) キュー内の最初のタスク 転送(HostToDevice)(2) キュー内の2番目のタスク だと転送(2)はカーネル(1)と同時に動かず、 カーネル(1) 実行中 転送(HostToDevice)(1) キュー内の最初のタスク 転送(DeviceToHost)(2) キュー内の2番目のタスク だと転送(2)はカーネル(1)と同時に動きました。
355 :
デフォルトの名無しさん :2012/01/24(火) 19:53:55.48
コンパイル時に以下のエラーメッセージが出ているのですが、 ptxas error : Entry function '...' uses too much local data (0x16054 bytes, 0x4000 max) これはローカルメモリーと何か関係あるような気がするのですけど、 ちょっとわからないので教えて頂けないでしょうか。 宜しくお願いします。
356 :
デフォルトの名無しさん :2012/01/24(火) 21:03:52.65
それともう1つ。 externを使って、グローバル変数を共有させたいのですが、 ホスト・デバイスともに共有させたい場合はどうしたらいいのでしょうか?
> ptxas error : Entry function '...' uses too much local data (0x16054 bytes, 0x4000 max) 単に配列の次元が大きすぎるのでは? 最大 0x4000バイトのところを0x16054 bytes使おうとしている? > ホスト・デバイスともに共有させたい場合はどうしたらいいのでしょうか? 物理的に無理でないの? ホスト←→デバイス間でコピーしないといけないのでは??
CUDAの思想はいいと思うけど、ターゲットデバイスを完全には 隠蔽しきれないのが残念だな。
スピード勝負の世界で、しかもかなり複雑なアーキテクチャなのに、 ハードウェアを完全に隠蔽する訳にいかないだろ。
361 :
357 :2012/01/25(水) 12:34:12.36
> ホスト・デバイスともに共有させたい場合はどうしたらいいのでしょうか? ヘッダーファイルに #DEFINE などで定義して、ホストとデバイスのプログラムに そのヘッダーファイルをinclude する手があった♪
362 :
355,356 :2012/01/25(水) 15:21:00.98
>>358 青木氏の「はじめてのCUDA」に似たようなエラーメッセージが載ってるのですが、
local dataがどの部分の話なのか、ちょっとわからない状態でして...
>>357 >単に配列の次元が大きすぎるのでは?
予定だと1×256スレッドしか使うので、足りないはずないような... 意味が違うか。
ちょっとわからないので詳細お願いします。
>ヘッダーファイルに
>#DEFINE
>などで定義して、ホストとデバイスのプログラムに そのヘッダーファイルをinclude する手があった♪
ちょっと試してみます。それができなかったら、コピーするしかないかもしれません。
__device__ extern TEST test;
のような宣言してて、実行ファイルができたときがあったのですが、何か違うよな、と引っかかってたので、
ここで質問しました。ありがとうございます。
差し支えなければプログラムをアップしてくださいませ♪
365 :
355,356 :2012/01/25(水) 18:04:09.31
>>362 です。
>>363 申し訳ないのですが、プログラムのアップは出来ないです。すみません。
なんかもう分からなくなったので、ホストとデバイスの住み分けを行ったところ、
「ptxas error...」云々のメッセージが消え、コンパイルできました。
(さっきまでhost,device,global宣言関数がごちゃ混ぜな状態だった。)
何が原因だったのか分からず仕舞いです。
グローバル変数の共有はとりあえず、コンスタントメモリにデータコピーで様子見することにしました。
元コード(.c)をもう一度読み直したところ、デバイス側の方はReadOnlyで十分なようでしたから。
皆様回答ありがとうございました。またよろしくお願いします。
366 :
365 :2012/01/25(水) 21:05:50.75
>>364 資料ありがとうございます。
うまくいったと思った途端、また同じエラーが出てきてしまったので、確認します。
.cファイルに.cuファイルを組み込ませるようにしたら、
__host__ __device__で修飾した関数が定義されてないと.cファイル側に言われ、
.cファイルと.cuファイルそれぞれ単独で動かせるようにしたら、
(同じ内容の関数を.cと.cuそれぞれ別名で実体作った。.c:test .cu:ktestみたいな)
今度はさっきと同じエラー。
実行ファイル作るだけなのに難しい...
Host側とDevice側で型を別にすれば開発もっと楽になると思うんだけど
取り敢えずCudaのサンプルは捨てて、インクルード関係とオブジェクトの生成手順を確認するんだ。
見せられないところは削除して、問題の部分だけ残して、アップすれば??
370 :
366 :2012/01/26(木) 16:19:08.01
>>368 >取り敢えずCudaのサンプルは捨てて、インクルード関係とオブジェクトの生成手順を確認するんだ。
のように手順を踏んでコンパイルしたところ、実行ファイルができました。
皆様ありがとうございました。
371 :
370 :2012/01/26(木) 20:33:24.55
お恥ずかしながら、また戻ってきました。 あれからセグメンテーション違反が出てきてしまったので、あれこれ探していた結果、 どうやらデバイスのメモリ確保&送信に失敗していたようです。 しかし解せないことがあって、 構造体A a,構造体B b[100],構造体C c[100],...(以下略 をデバイス側に送るのですが、 (1つ1つのサイズは結構大きめ。グローバルは1GBあって余裕で確保出きるハズ...) void main(){ ・・・ test(&a,b,c....); } void test(A *a,B *b,C *c...){ A *d_a; B *d_b; C *d_c; CUT_SAFE_CALL(cudaMalloc((void**)&d_a,sizeof(A))); CUT_SAFE_CALL(cudaMemcpy(d_a,a,sizeof(A),cudaMemcpyHtoD)); CUT_SAFE_CALL(cudaMalloc((void**)&d_b,sizeof(B)*100)); CUT_SAFE_CALL(cudaMemcpy(d_a,b,sizeof(B)*100,cudaMemcpyHtoD)); ・・・ } はじめてcuda,cuda_by_exampleで確認したところ、 文法的ミスはないはずなのに確保ミスしてるらしく、中断してます。 この原因は一体全体なんなんでしょうか。
372 :
デフォルトの名無しさん :2012/01/26(木) 21:28:56.32
構造体の中身がよく分からんが,allocateのところで *100は必要なのかな? 構造体Bのメモリーを100個分用意しようとしている?? 構造体Bの中にすでに [100]個の配列を取っているのに??? 数は合っておるのか????
rhel6.0のがリンクミスかでダウンロドできん
>>372 遂にRC取れたか!4.0からどう変わってるか知らんけど。
Kepler向けのプログラムを作れるのは5.0とかになるんだろうか。
そしてそれはいつ出るんだろう。
>>371 1GBのVRAM積んでるカードでも400MBぐらいの確保で失敗したことがあったような覚えがある。
確保するサイズを小さくしてもエラーが出るってんならこの話は忘れてくれ。
100x100の行列の積を1万回くらい実行したいのですが、 CUBLASだと小さすぎてCPUより遅くなってしまいます。 この1万回は相互に依存関係は無いので、並列化が可能なのですが、 このプログラムは自分で書くしかないのでしょうか。 TESLA C1070があるので、丸ごとVRAMに載ります。
4.1にcublas{S,D,C,Z}gemmBatchedってのが追加されてました。 especially for smaller matricesに効くってあったので、使ってみます。
一つのブロック内で100x100の行列の積を1回行う?
間違った
あれ?
一つのブロック内で100x100の行列の積を1回行う?
たびたび,失礼 一つのブロック内で100x100の行列の積を1回行う? そのブロックを1万個用意する? それらを並列に計算する? と言うコンセプトかな??? 自分でプログラムを作った方が,よほど勉強になると思う
383 :
371 :2012/01/27(金) 18:15:48.81
>>373 CUT_SAFE_CALL(cudaMalloc((void**)&d_b,sizeof(B)*100));
CUT_SAFE_CALL(cudaMemcpy(d_a,b,sizeof(B)*100,cudaMemcpyHtoD));
これはミスですね....
正しくは
CUT_SAFE_CALL(cudaMalloc((void**)&d_b,sizeof(B)));
CUT_SAFE_CALL(cudaMemcpy(d_a,b,sizeof(B),cudaMemcpyHtoD));
です。すみません。
>>375 4000MBの領域確保でも失敗することがあったんですか...
ちょっと試してみます。
d_aに送信してるあたりが
d_b のサイズは構造体Bと同じ?? d_aのサイズはbと同じ?すなわち構造体Bと同じ?? 違っていたらエラーが出て当たり前では???
それから,ホスト側で も メモリーは確保しているのか????
387 :
383 :2012/01/28(土) 17:54:35.37
>>371 なんかサンプルがめちゃくちゃなんで書き直します。
388 :
387 :2012/01/28(土) 20:55:21.49
確認なのですけど、カーネル関数の引数はポインタ限定ですか?
配列はポインタで,1変数はそのままではなかったかい? なぜか知らんけど
夜の埼玉は
>388 ホスト側のポインタ以外なら何でもOKのはず。少なくともintは渡せる。 クラスを渡せるかどうかは知らないけど。
ホストのポインタ送って、GPU側のload/store時に メインメモリから自動転送ってこともできるから、 その説明は正しくない。
393 :
388 :2012/01/30(月) 15:13:56.68
>>391 セグメンテーション違反してた部分の引数に(デバイスで使用する)double(実体)を指定してたのですけど、
(デバイスで使用する)ポインタに変えたら、問題なく実行できました。
dtempはcudaMalloc/cudaMemcpy使ってます。
そもそもな話cudaはポインタ宣言したdtempの実体は*dtempではないのかな。
>>393 何を言っているのか判らない。問題が再現する最低限のコードでも貼ってみていただきたい。
395 :
デフォルトの名無しさん :2012/01/30(月) 16:19:41.78
>>393 __global__ void kernel(double dtemp){
}
int main(){
double temp=0.0;
double *dtemp;
cudaMalloc((void**)&dtemp,sizeof(double));
cudaMemcpy(dtemp,&temp,sizeof(double),cudaMemcpyHostToDevice);
kernel<<<1,1>>>(*dtemp);//(←※1)
cudaMemcpy(&temp,dtemp,sizeof(double),cudaMemcpyDeviceToHost);
return 0;
}
※1の部分でsegmentation fault。
__global__ kernel(double *dtemp){
}として
kernel<<<1,1>>>(dtemp);
とすると問題無し。
dtempのアドレスを引数にするのはいいけど、dtempの中身(0.0)を引数にしようとするのはダメ?
そもそも※1の宣言自体間違ってる?どうなんだろう。
393じゃないけど試してみた __global__ void kernel(double dtemp){ } int main(){ double temp=0.0; double dtemp; // cudaMalloc((void**)&dtemp,sizeof(double)); // cudaMemcpy(dtemp,&temp,sizeof(double),cudaMemcpyHostToDevice); kernel<<<1,1>>>(dtemp);//(1) // cudaMemcpy(&temp,dtemp,sizeof(double),cudaMemcpyDeviceToHost); return 0; } これでコンパイルも実行もおkだったけど
>>395 dtempはdevicePointerなんだから、Hostコード中でデリファレンスしちゃダメ。
double値を渡したいだけなら引き数はdoubleで充分。
それと、※1は宣言じゃなくて呼び出しだから構文としてはあっている。
そもそも何がしたいのだろう。値を渡したいなら値を渡せばいいし、
ポインタを渡したいならポインタをそのまま渡せばいい。
devicePointerをHostでデリファレンスしちゃダメだし、
(その後どうせdevice側でデリファレンスできないから)HostPointerをdeviceにコピーしてもダメ。
398 :
デフォルトの名無しさん :2012/01/30(月) 17:04:05.78
これからCUDAの勉強を始めようと思っています。 GPGPUプログラミングはシェアードメモリを有効活用できるかがポイントであるように思います。 そこで、あえて最新のグラボに買い換えず、キャッシュのないTesla世代のGTX285を所持し続けています。 キャッシュがあると、シェアードメモリをうまく使えてなくても、それをカバーしてしまって、 うまくプログラムできているかどうかを把握し辛いと考えたからです。 ただ、心配なのが、Fermi世代以降に発売された書籍だと、内容がTesla世代のGPUにそぐわない といった状況が起きてしまうことです。 GTX285のまま勉強するに際して気をつけること等あれば教えてください。 問題がなければ関連書籍(日本語)を全部購入して読もうと思います。 よろしくお願いします。
>>GPGPUプログラミングはシェアードメモリを有効活用できるかがポイントであるように思います。 そんなの計算する内容によるわ 活用しようがないのもあるからな 計算(使用)目的は明確なの?
401 :
399 :2012/01/30(月) 21:39:43.45
>>400 信号処理に使う予定です。
FIRフィルタリングやFFTなどを大きな1次元、あるいは2次元配列に適用したいと思っています。
俺は400じゃないが、気をつけることは、勉強がひと段落していざ使うときになって知識が陳腐化してても泣かないことかな。 ハードウェアはどんどん変わるし、TeslaやFermiで良しとされたことが数年後に通用するかはわからないからね。 せっかく今から始めるならFermi・Keplerを使ったほうがいいと俺は思うけど、 10年後にはどれも等しく陳腐化してるだろうから長期的に考えると確かにTeslaでもいいのかもしれない。 ただ長期的に考えること自体がGPGPUでは…と1行目にループする。
古い方がいろいろ制限があって,勉強になるとは思う.
fftはcpu側からはインターフェースあるけどディバイス内では自前で作るしかない
> キャッシュがあると、シェアードメモリをうまく使えてなくても、それをカバーしてしまって、 > うまくプログラムできているかどうかを把握し辛いと考えたからです。 シェアードメモリーを使うようにプログラミングすればよいかと どのメモリーを使っているのか分からない,と言うことはないと思う あれば,そこをしっかり勉強すべきかと
そう言えば古いのは倍精度の計算ができないのでは??
まあ,デバイスメモリーだけで計算してみるとか, シェアードメモリーも使って計算してみるとか, いろんなプログラムを作って比較するのが良いかと
最初はシェアドするスレッドはどれだってだけでも頭が痛くなった
409 :
402 :2012/01/30(月) 23:28:19.90
あとTeslaでやるなら、FermiとTeslaでバンクコンフリクトの発生条件が違ってるから Teslaでのこれの回避には深入りしないほうがいいかと。 他にシェアードメモリの大きさなどの量的な違いがあるけど、そういった量的な 制約のために質的なというかアルゴリズムの種類を変えるようなことまで深入りするのも 避けたほうが。だってFermiやそれ以降で量が変わると一から見直しになるから。 これはFermi使ってる人もいずれ同じことなんだけど、勉強じゃなくて性能が欲しいんだからしょうがない。
411 :
399 :2012/01/31(火) 00:31:29.61
>>403 やはり、そうですか!
>>404 参考になります。
FFTは重要になるのでしっかりと取り組みたいです。
>>405 >>407 >>408 なるほど、どのメモリかをよく意識してプログラミングします。
CPUの最適化に取り組んでいて思ったのですが、
キャッシュは便利な反面、こちらから意図的に挙動を制御できないことが
パフォーマンスの考察を難しくしてしまう側面もあると感じました。
キャッシュなしのTeslaで鍛えたいです。
>>406 単精度しか使わないので問題なしです。
もし半精度でパフォーマンスが上がるなら、そのほうがイイくらいです。
>>402 >>409 世代を跨ぐにあたって非常に参考になるご助言をありがとうございます。
各種メモリの特性や帯域を意識して取り組むことで、
固有のデバイスに限定されない定性的なところを理解したいと思います。
Keplerではこんな組み方がイイかもしれないな、と思索できたら楽しそうです。
>>410 興味深い資料をありがとうございます。
各部のレイテンシ、スループットを頭に叩き込みます。
>>411 最終的に動かす機器で最もよい性能が出るように最適化するべきだし
自動でやってくれることを手でやるってのは勉強にはなるかもしれないけど意味ないのではって思うけど。
コンパイラの最適化を切ってアセンブリ書くようなものでしょ。
時間がいくらでもあるorそういう研究ならともかく 短期間で組めてそこそこのパフォーマンスが出せればいいなら キャッシュのあるFermiで適当に書いてもいいと思うんだけどね
414 :
399 :2012/01/31(火) 21:32:30.15
>>412 おっしゃる通りです
そのあたりは目的をよく考えて判断したいです。
今回はどちらかというと、GPUのポテンシャルを見極めたいというのが主です。
コードの可読性を優先して、性能はコンパイラやキャッシュ等にできるだけ委ねるというスタイルも
今後検討していきたいと思っています。
>>413 今回は研究用途ですが、
将来的にはFermi以降のGPUで、ちゃっちゃと書いたコードが快適に動くというところを目指したいです。
ただ、あまりにも下手なプログラムにならないよう、最初に少し深くやっておきたいです。
レジスタといえば、GK104では強化されるのか、それともさらにピーキーになるのか気になる。 GF104はグラフィック特化で、レジスタ数やワープ数の上限に悩まされたし。 とりあえずアーキテクチャが一般公開されるまで待つしかないかな。
誰かいるかな…
http://includehoge.blog.fc2.com/blog-entry-71.html このブログのソースをコンパイルして動かしてみようと思ったのですが
1>cudamatrix.obj : error LNK2019: 未解決の外部シンボル _cufftDestroy@4 が関数 _main で参照されました。
1>cudamatrix.obj : error LNK2019: 未解決の外部シンボル _cufftExecC2C@16 が関数 _main で参照されました。
1>cudamatrix.obj : error LNK2019: 未解決の外部シンボル _cufftPlan1d@16 が関数 _main で参照されました。
1>../../bin/win32/Debug/template.exe : fatal error LNK1120: 外部参照 3 が未解決です。
というエラーが出ます。
cufft.h自体は読み込んでいて、引数の過少があるときはそれに対してエラーを返すので
定義されていないという扱いになっているわけではないみたいなんですけど
解決法があったら教えて下さい
誰かいるかな…
http://includehoge.blog.fc2.com/blog-entry-71.html このブログのソースをコンパイルして動かしてみようと思ったのですが
1>cudamatrix.obj : error LNK2019: 未解決の外部シンボル _cufftDestroy@4 が関数 _main で参照されました。
1>cudamatrix.obj : error LNK2019: 未解決の外部シンボル _cufftExecC2C@16 が関数 _main で参照されました。
1>cudamatrix.obj : error LNK2019: 未解決の外部シンボル _cufftPlan1d@16 が関数 _main で参照されました。
1>../../bin/win32/Debug/template.exe : fatal error LNK1120: 外部参照 3 が未解決です。
というエラーが出ます。
cufft.h自体は読み込んでいて、引数の過少があるときはそれに対してエラーを返すので
定義されていないという扱いになっているわけではないみたいなんですけど
解決法があったら教えて下さい
誰か居るかな・・・いたら助けてください
http://includehoge.blog.fc2.com/blog-entry-71.html このブログのソースをコンパイルして動かしてみようと思ったのですが
1>cudamatrix.obj : error LNK2019: 未解決の外部シンボル _cufftDestroy@4 が関数 _main で参照されました。
1>cudamatrix.obj : error LNK2019: 未解決の外部シンボル _cufftExecC2C@16 が関数 _main で参照されました。
1>cudamatrix.obj : error LNK2019: 未解決の外部シンボル _cufftPlan1d@16 が関数 _main で参照されました。
1>../../bin/win32/Debug/template.exe : fatal error LNK1120: 外部参照 3 が未解決です。
以上のようなエラーが出てコンパイルが通りません
構造体の生成や関数の過少には反応するので、cufft.hが読み込めてないってことはないと思うんですが…
解決法があったら教えて下さい
誰か居るかな・・・いたら助けてください
http://includehoge.blog.fc2.com/blog-entry-71.html このブログのソースをコンパイルして動かしてみようと思ったのですが
1>cudamatrix.obj : error LNK2019: 未解決の外部シンボル _cufftDestroy@4 が関数 _main で参照されました。
1>cudamatrix.obj : error LNK2019: 未解決の外部シンボル _cufftExecC2C@16 が関数 _main で参照されました。
1>cudamatrix.obj : error LNK2019: 未解決の外部シンボル _cufftPlan1d@16 が関数 _main で参照されました。
1>../../bin/win32/Debug/template.exe : fatal error LNK1120: 外部参照 3 が未解決です。
以上のようなエラーが出てコンパイルが通りません
構造体の生成や関数の過少には反応するので、cufft.hが読み込めてないってことはないと思うんですが…
解決法があったら教えて下さい
うわなんか四回も書き込んじゃったごめん
コンパイル詳しくないけどlibcufft.so(linuxの場合)のリンクしてる? windowsならlibcufft.dllだとおもう
コンパイルじゃなくてリンクの問題でしょ
nvcc -lcufft sample.cu
このスレでときどき思うのはC/C++をほとんど知らないのにもかかわらず CUDAに挑戦しようとする勇壮な挑戦者が多いな、ということ 彼らはその後うまくいってるのだろうか?
CUDA.NETも3.0で止まってるしなぁ・・・
コンピュート・ケイパビリティによって何が変わったかを俯瞰できるような一覧を掲載したサイトってご存知ありませんか??
>>417 追加のライブラリファイルにナントカ.libを追加すればいけそう。
ナントカは
>>422 さんのレスから察するにlibcufft.libかなぁ。
このへんは
>>425 さんの言うとおりcuda云々じゃなくてVisual Studio(C++)を使う上で頻繁に設定するところだねー。
>>427 俺も知りたいなぁそれ。その手の情報はSDKpdfのProgramming GuideとBest Practiceくらいしか知らない(片方だけかも)
cufft.libなのだが。
てすてす
コマンドなら
あれ?
以下のよう
nvcc fft.cu --compiler-bindir="C:\Program Files\Microsoft Visual Studio 10.0\VC\bin" -arch=sm_10 -L"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.1\lib\Win32" -lcufft
ファイル名をfft.cuとしてみた
結果の一部 1014 6.347414 0.587785 1015 6.318954 -0.000000 1016 6.293659 -0.587785 1017 6.271449 -0.951057 1018 6.252275 -0.951057 1019 6.236131 -0.587785 1020 6.222946 0.000000 1021 6.212736 0.587785 1022 6.205431 0.951057 1023 6.201046 0.951057
-arch=sm_10 は適宜,変更のこと
これでも行けた♪ nvcc fft.cu -lcufft
既出でした.失礼 >> 424
GPUの勉強を始めたばかりなのですが、スパース行列の格納形式種類で ELLの格納方法とメリットが分かりません 調べたらlisのマニュアルに配列方法はあるのですが、説明と図とが噛み合ない点があり 理解できずにいます。 だれか教えてください
>>426 CUDA#とかCUDAfy.NETとか代替のライブラリはあるっぽいよ
使ったことないから詳細は知らないけど
sparse matrixの格納形式はCUDAの問題じゃね〜だろ
>>440 コアレッシング状態で転送できるようになる。
CUDA.NETは使えなくはないぞ♪
>>440 スパース行列の格納方法がのっているのがLisのオンラインマニュアルくらいだったので
>>443 ELLはGPUに適した格納方法と、ネットのどこかにあったので関連があるかと思い書き込みました
たしかにCUDA以前の問題ですよね
>>444 なるほど、見えてきました!ありがとうございます
>>422-
>>439 こんなにたくさん答えてくれるとは・・・
おかげで動きましたーありがとう
おお、よかった^^
>>441 CUDAfy.NETは知らんかった
最近出てきたのね、ありがとう
2次元配列にアクセスするときのインデックス計算で gy = blockDim.y * blockIdx.y + threadIdx.y; gx = blockDim.x * blockIdx.x + threadIdx.x; というのを頻繁にやると思うのですが、 この積和演算ってCUDAコアが使われるのでしょうか? それとも、CUDAコアとは別にアドレス演算器(?)的なハードウェアがあって、 それで計算してくれるのでしょうか? 後者だとうれしいです。
452 :
デフォルトの名無しさん :2012/02/10(金) 21:37:17.01
アドレスもプログラムもデータ
>>451 CUDAコア上に展開された(定義された)ブロックの中で演算をしていると思う
454 :
451 :2012/02/11(土) 01:18:08.05
>>453 CUDAコアのサイクルを消費して計算しているということですね?
ありがとうございました。
インデックス計算かどうかをGPUがどう見分けるというんだ
456 :
451 :2012/02/11(土) 10:05:14.94
>>455 単にCUDAの予約語であるblockDimやthreadIdだけで構成されていて、
○○ * ○○ + ○○ の積和の形になっていれば、
専用の回路に計算させる、とかできそうなものですが。
CPUではメモリアドレス式のカタチを解析してアドレス演算器に割り当てています。
そもそもグリッド,ブロック,スレッドの概念を理解しているのかな?
○○ * ○○ + ○○ だけではプログラムは作れないと思うのだが
459 :
デフォルトの名無しさん :2012/02/11(土) 11:55:45.33
460 :
デフォルトの名無しさん :2012/02/11(土) 19:16:02.78
環境: CUDA 4.1 + Visual Studio 2010 Pro. + Windows 7 Pro. 64bit C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.1\C\common\cutil_vs2010.sln をバッチビルドして、 cutil64.lib, cutil64.dll, cutil64D.lib, cutil64D.dll を作ろうとしたところエラーがでてしまいます。 エラーは出るのですが、cutil64.lib 以外は生成でき、 そのライブラリを使った行列積のサンプルプログラムもコンパイル&実行できました。 #ちょっと気持ち悪いですが。 しかし、cutil64.lib はどうしても生成できません。 何が悪いのでしょうか。 cutil64.lib, cutil64.dll, cutil64D.lib, cutil64D.dll をビルドしたときのエラーメッセージは以下の通りです。
461 :
デフォルトの名無しさん :2012/02/11(土) 19:16:34.66
------ すべてのリビルド開始: プロジェクト: cutil, 構成: Release x64 ------ C:\Program Files (x86)\MSBuild\Microsoft.Cpp\v4.0\Microsoft.CppClean.targets(74,5): warning : パス 'C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.1\C\common\lib\x64\cutil64.dll' へのアクセスが拒否されました。 bank_checker.cpp cmd_arg_reader.cpp cutil.cpp stopwatch.cpp stopwatch_win.cpp Generating Code... LINK : warning LNK4075: /INCREMENTAL は /OPT:ICF の指定によって無視されます。 LINK : fatal error LNK1104: ファイル 'lib\x64\\cutil64.dll' を開くことができません。 ------ すべてのリビルド開始: プロジェクト: cutil, 構成: Debug x64 ------ bank_checker.cpp cmd_arg_reader.cpp cutil.cpp stopwatch.cpp stopwatch_win.cpp Generating Code... ライブラリ C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.1\C\common\lib\x64\\cutil64D.lib とオブジェクト C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.1\C\common\lib\x64\\cutil64D.exp を作成中 LINK : fatal error LNK1158: 'cvtres.exe' を実行できません。 ========== すべてリビルド: 0 正常終了、2 失敗、0 スキップ ==========
462 :
デフォルトの名無しさん :2012/02/11(土) 19:21:14.69
追記です。 cutil64.lib がないので 行列積サンプルプログラムは、Releace モードではビルドできていません。 ビルド&実行でてきているのは、debug モードです。
Releace + x64モードでcutil64.dllはできるが,cutil64.libはできないと言うことですか?
エラーが出てもライブラリはできている,と言うことはないですか? この中に↓ C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.1\C\common\lib\x64
失礼 Releace -> Release
まあ、CUDAのインデックス計算ほど アホ臭いものは無いと思うがな。 テクスチャアドレスではもっと複雑な計算しているのに。
CUDAを数値計算に使おうという話だよ
468 :
デフォルトの名無しさん :2012/02/11(土) 22:18:37.30
>>463 >>464 お返事ありがとうございます。
C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.1\C\common\cutil_vs2010.sln
のバッチビルドの前後で以下のようにライブラリが増えます。
ビルド前
common\lib\x64\ には cutil64.dll のみ存在
ビルド後
common\lib\x64\ には cutil64.dll cutil64D.dll cutil64D.lib の3つになる
cutil64.libは生成されません。ビルド時に
>>461 のエラーのためと思われます。
変だね.こちらも同じようなエラーが出るが(xx.dllを開始できません),xx.libはできている. 試しに,ソルーションエクスプローラーのcutilを右クリックして,「追加」,「既存の項目」で,cutil64.dll を追加してみる....とか
で,何でビルド前に「cutil64.dll のみ存在 」するのかな? 一度,全部クリーンしてみる....とか
471 :
デフォルトの名無しさん :2012/02/12(日) 01:08:37.66
>>469 お返事ありがとうございます。
エラーは出るんですね。貴重な情報です。
エラーを気にしないですみます。
>cutil64.dll を追加してみる....とか
無理を承知でやってみました。当然ですがダメでした。
>>470 そういうものだと思っていましたが、確かにそうですね。
クリーンインストールしてみます。
472 :
デフォルトの名無しさん :2012/02/12(日) 02:45:15.54
クリーンインストールしてみました。 C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.1\C\common\lib\x64 に最初から cutil64.dll が作られるみたいです。 なお、インストールたCUDAファイルは、 devdriver_4.1_winvista-win7_64_286.16_notebook.exe cudatoolkit_4.1.28_win_64.msi gpucomputingsdk_4.1.28_win_64.exe の3つです。
cutilのクリーンを実行すれば,すべてのdll,libが消えるはずなのに,もし, cutil64.dllが残っているとすれば, 書き込み,読み込みの権限がないと言うことでは? cutil64.dll はマニュアルで消せますか?
474 :
デフォルトの名無しさん :2012/02/12(日) 15:39:25.13
>>437 お返事ありがとうございます。
cutil64.dll をマニュアルで削除して、再度
C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.1\C\common\cutil_vs2010.sln
でビルドしたら、すべてのライブラリが生成されました。
ご推察の通り、cutil64.dll にアクセスできなかったのが原因だったようです。
本当にありがとうございました。
<追記>
adminユーザーで環境構築作業をしていたので、
cutil64.dll にアクセスできなかった原因は謎のままです。
自分で構築してしまえば、その後は、上書きビルドは問題なくできます。
同じく、CUDA-SDKの中に、simpleTemplate というソースコードがあるのですが、
このソースコードを書き換えて、Releaseモードでビルドすると、
インストール初期からある、simpleTemplate.exe にコンフリクトするらしく、
ビルドエラーが出ました。この場合はも、初期からある、simpleTemplate.exe を削除すれば
問題なく繰り返しビルドができるようになりました。
同じ原因だと思います。
475 :
デフォルトの名無しさん :2012/02/12(日) 15:40:34.77
↑ 番号ミスです。ごめんなさい。
>>473 が正しいです。
よかった,よかった♪
477 :
デフォルトの名無しさん :2012/02/12(日) 22:34:46.87
OS:32bit CUDA:32bit VisualStudio:32bit でCUDAプログラムをしている人が、速度を上げることを目的に OS:64bit CUDA:64bit VisualStudio:64bit に変更することって意味あります?
ないと思う。ホスト側で2GB以上のメモリを使いたいわけでもないのだろうし。
64bitビルドして動かすとポインタがレジスタを2つ使ってしまうという話があるので、逆に遅くなる可能性も。
本当に64bitビルド時にデバイス側でもポインタが64bitで扱われるかどうかを
自分の目で確かめたわけではないけど、少なくともここではそういうことを言っている。
高速演算記 第3回 「チューニング技法その1 CUDAプログラミングガイドからピックアップ」 | G-DEP
http://www.gdep.jp/column/view/3 >ここでレジスタ数は32bitレジスタの数で、long long, double値や64ビット環境を対象としたCUDAプログラムでは
>ポインタなどが64bitとして扱われますのでレジスタを2つ使用することになります。
479 :
デフォルトの名無しさん :2012/02/13(月) 10:45:03.24
>>478 なるほど。ありがとうございます。
これと少し環境が違う場合ですが、
OS:64bit CUDA:32bit VisualStudio:32bit
と、
OS:64bit CUDA:64bit VisualStudio:64bit
では、速度差はどうなると思われますか?
ネット上でよく目にするmatrix積のサンプルコードで実験したところ、
同程度か、どうかすると前者の方が速い場合がありました。
前者は、WOW上で動いていると思われるので、
遅いだろうと予測したのですが、それに反する結果でした。
これをどう解釈すればよいか謎なんです。
480 :
デフォルトの名無しさん :2012/02/13(月) 10:46:42.16
479です。OSは、Windows7 64bit Pro です。
プログラムのどの部分を計測したのであろうか? CUDAの部分のみ??
482 :
デフォルトの名無しさん :2012/02/13(月) 12:26:28.94
CUDAの部分のみです。
483 :
デフォルトの名無しさん :2012/02/13(月) 12:35:43.06
CUDAはOSに依存せずに動くとかで、OSのbit数には無関係ということですか?
そのため、
>>478 で説明して頂いたように、同程度もしくは32bitが若干速くなる場合もあるということでしょうか。
ちなみに測定時間はいくらでしょうか? 演算量を増やすとどうなるでしょうか?
485 :
デフォルトの名無しさん :2012/02/13(月) 16:09:51.62
ソースは、以下のHPの最後の「Matrix_GPU」です。
http://www.gdep.jp/page/view/218 計算は、スレッドサイズ一杯で組まれているようですので、
#define MATRIX_SIZE 1024/*行列1辺の数*/
をこれ以上上げることができませんでした。
ちなみに、ノートパソコン搭載のGPUチップは、GT-540M。
Releaseでビルドして 650ms付近を中心に数10msばらつきます。
win32とx64で大差はなく、
win32の方が平均10ms程度速いような気がするという程度です。
CUDAの計算のところを1000回程度,繰り返すとかはできるでしょうか? そうすれば有意な差が出るかも知れません.
487 :
デフォルトの名無しさん :2012/02/13(月) 19:18:58.94
200回繰り返してみました。 OS:64bit CUDA:32bit VisualStudio:32bit → 131,671 ms OS:64bit CUDA:64bit VisualStudio:64bit → 132,146 ms 0.3%の差で、前者が速かったです。 100回繰り返しでも、0.3%差でした。
なるほど,有り難うございます. 有意な差が出たと言うことですね. では,後は分かる人,よろしく♪ ε=ε=ε=ε=ε=┌(; ・_・)┘
GPU側では64bitのメリットは使用可能なメモリが増えることだけで、その代わりにアドレス計算の命令数や使用レジスタ数が増えるデメリットがある。 前はGPU側だけ32bitモードで動かすという恐ろしい手法があったが、今使うのはなかなか難しい。
ptxオプションつけてptx出力を見れば見当がつくね。 64ビットアドレッシングだとアドレス計算の為にint32→int64の変換がしばしば行なわれる。
491 :
デフォルトの名無しさん :2012/02/14(火) 18:19:16.13
ざっくり捉えると、32bitと64bitでは、 CUDA自体はさほど変わらない。 CUDAと切り離されてコンパイルされるC++言語部分は、64bitの方がいくらか速い。 という感じでしょうか。
>>491 アドレス計算多用したり、レジスタ沢山使うようなカーネルだと
CUDA自体32bit版の方が無視できないほど
速くなる可能性はある。
493 :
デフォルトの名無しさん :2012/02/14(火) 18:53:10.54
doubleが速くなるっていう噂はどうなったの?
GPU上にあるコンスタントキャッシュメモリって、CUDAではどのようにして使われるのでしょうか? 計算の中にリテラルがあった場合(CPUでは命令の即値として埋め込まれますが)や、 constantを付けて宣言した定数が割り当てられたりするのでしょうか??
llvmで使えるようになったのなら boostも使える?
498 :
494 :2012/02/17(金) 00:11:19.70
>>496 試してみました。
const変数もリテラルもレジスタにマッピングされました。
ひょっとして、DirectXAPIなどのグラフィックスAPIを利用していた場合に
使われるだけ(要は旧来通りの使われ方)かも・・・
プロファイラを見ると、テクスチャキャッシュのヒット/ミスヒットが見れるようになっていますが
これなんかはグラフィックスAPI専用のような気が。
しかし同じようにコンスタントキャッシュの項目もあって然りと思うが在らず。
う〜ん、気にしないことにしますw
貴重な報告、ありがとうございます♪
GPUを演算器としてフルに使いたい場合って ビデオ出力をオンボードのものになるようBIOSで切り替えた方がいいのでしょうか linuxだったらxconfigいじるだけでいけるとか?
フルの定義は?
>>494 CPUでも即値として埋め込めない配列や文字列リテラルは
.rdataセクションのデータとして書き込まれるのと同様に、
GPUでも即値として使えないconstデータはコンスタント領域に置かれる。
で、コンスタント領域の値を直接オペランド指定できない
命令を使う場合は当然レジスタを介して使用することになる。
逆に一部の変数にデフォルトで定数が入っているなんてのは
ABIで決まったもの以外は普通無理だから、コンスタント境域などの
メモリに置いて読み込む以外やり様が無い。
>>500 普通はGPUカードに何もモニタを繋がなければ(≒出力する設定をしなければ)
GVRAMをほぼ全部使い切れるよ。
Linuxなら、Xを立ち上げなければGPUカードから出力していてもかなり使えるよ。
ついでに、3秒だか5秒だからの制限もなくなるからGPUカーネルで長時間計算させられるよ。
504 :
デフォルトの名無しさん :2012/02/17(金) 20:22:12.27
505 :
494 :2012/02/18(土) 16:32:22.87
>>502 ありがとうございます。
コンスタント変数の扱われ方、レジスタが必要な理由、理解しました。
書籍を読んでいたら、コンスタントメモリおよびテクスチャメモリについて書かれているのを見つけました。
コンスタントメモリは__constant__修飾子を付けて宣言することで割り当てられるようです。
物理的にはVRAMに置かれ、GPU内部のコンスタントキャッシュでキャッシュが効くものと思われます。
空間フィルタリングなどでフィルタ係数を格納しておくのに適しているのではと思いました。
(繰り返し参照されるのでキャッシュが効きそう)
テクスチャメモリはテンプレートクラス型のtexture<・・・>で宣言することで割り当てられるようです。
物理的にはVRAMに置かれ、GPU内部のテクスチャキャッシュでキャッシュが効くものと思われます。
さらに、グラフィック用途時の同様、テクスチャフィルタリングやサンプリング方式などのハードワイヤード回路機能を
利用できるよう、それらを指定するプロパティを保持しているようです。
複数カーネルを呼び出すんだけどカーネルのソースはそれぞれ別ファイル それぞれのカーネルから同じDevice関数を呼び出したいんだけどファイルまたがることできないようだし ソースごとに同じ内容の別名関数作るしかないのかな?
ファイルが別でも,プログラムとしてつながっていれば問題はないと思うが. 関数のプロトタイプ宣言をしっかりしておけば良いかと...
共通のdevice関数を別のヘッダに纏めて カーネルソースからincludeして使う。 device関数はどうせinline展開されるのだから 通常のプログラムでのinline関数と同様に扱えばいい。
>>508 サンクス
でも今やってみたんだけど全ファイルで展開されるんで
multiple definition of 関数名
ってエラーになっちゃう
実際やってみました?
カーネル起動時、1スレッドの起動につき1クロックかかり、128スレッドなら128クロックかかると書籍に書いてあるのですが、 これはCPU側でもかかるのでしょうか? それとも、CPU側はカーネル起動の合図を送ってすぐ制御を戻し、 GPU側の制御ブロックがクロックをかけてスレッドを起動するというかたちでしょうか?
511 :
デフォルトの名無しさん :2012/02/19(日) 01:11:11.40
教えて欲しいことがあります. CPU側で複数スレッドを立てて,互いに独立した処理を行っています. そして,それぞれスレッドにおいて,CUDAを使って画像処理をさせたいのですが,(たとえば,グレースケール化とか) この時,CUDAでの処理にテクスチャメモリを使いたい場合は,どのようなコードを書けばいいのでしょうか? テクスチャメモリを使う場合,グローバルで宣言しなけらばならないですよね? たとえば,5枚のRGB画像をグレースケール化するときに, CPU側で5スレッド立てて,各スレッドでCUDAを使ってグレースケールへの変換処理をしたいのですが, テクスチャメモリをグローバルで宣言するとおかしなことになるきがするんですが. どなたか教えて頂けないでしょうか? 宜しくお願い致します.
CUDA対応のGPUを5枚挿す。 またはGTX590などを3枚挿す。
513 :
デフォルトの名無しさん :2012/02/19(日) 01:26:39.83
>>512 回答ありがとうございます.
テクスチャメモリを使った処理の場合,マルチGPUじゃないと出来ない
とういことでしょうか?
514 :
デフォルトの名無しさん :2012/02/19(日) 01:44:29.68
公式マニュアル見れば分かること 試してみれば分かること これを聞く人が多すぎる
4-way SLI でも、ホスト上の4スレッドで利用する場合、 各々1GPUを独立して割り当てるため、4GPU間で通信して仕事をこなす本来的なSLIにはならない、って認識でイイんですよね? (SLIコネクタがなくても動く?)
517 :
デフォルトの名無しさん :2012/02/19(日) 02:47:03.16
>>516 回答ありがとうございます.
CPU側でスレッドを立てる段階で決めたいので,
できることなら,可変にしたいと考えています.
514さんは「これを聞く人が多すぎる 」と仰っていますが,
そもそも,このようなことをシングルGPUでやろうとするのが間違いなのでしょうか?
518 :
デフォルトの名無しさん :2012/02/19(日) 02:59:57.92
いやいや、まず試せば分かることじゃん。 その労力を2chに押しつけるのは、今この瞬間は楽では良いかもしれないけど、 自分の成長には繋がらないよ。
>>513 OS側のスレッド何本あろうとGPUが1つしかないんだから
順番に処理されるだけだろってことでしょ
520 :
デフォルトの名無しさん :2012/02/19(日) 05:43:18.09
concurrent kernel execution…
521 :
509 :2012/02/19(日) 07:12:00.07
>>508 失礼しました
インライン関数化で無事できました
ありがとうございました
522 :
デフォルトの名無しさん :2012/02/19(日) 11:42:52.58
>>518 心遣い感謝します.
>>519 >>520 回答ありがとうございます.
やっぱり,そうなんですね.
「concurrent kernel execution」ついて調べて確信しました.
回答・アドバイスして下さった皆様,ありがとうございました.
おかげで解決しました.
523 :
デフォルトの名無しさん :2012/02/19(日) 11:49:26.29
誰が質問者なのかわからないので謎な流れだが、とりあえずグレースケール化にテクスチャは不要だと思う今日この頃。
525 :
デフォルトの名無しさん :2012/02/19(日) 17:00:58.89
NPPで一発解決だよな
atomicCASのCASってどういう意味ですか??
Compare And Swap
528 :
526 :2012/02/20(月) 22:28:38.87
cudaの本には2次元配列より1次元の方が速いってあったけど実際どのくらい違うんですか? また、コピーとgpu上での演算の両方で遅くなるんですか?
530 :
デフォルトの名無しさん :2012/03/07(水) 03:39:04.89
なんでこのスレ突然止まってたの? 卒論修論シーズンが終わったから?
>>529 一行目: 速い
二行目: 遅くなる
馬鹿?
学会はこれからだというのに
> cudaの本には2次元配列より1次元の方が速いってあったけど実際どのくらい違うんですか? メモリーのアクセス時間の早さのこと? ちなみに2次元であっても,例えば f[i, j] で,先にjを変化させるか,i を変化させるのかで アクセス時間が違ったと思う. 要はメモリーが並んでいる順番にアクセスするのが早いはず.どの位,早くなるかは知りませぬ♪ > また、コピーとgpu上での演算の両方で遅くなるんですか? コピーとは? CPUとGPU間のメモリーの転送のことかな??
人いなくなったな 春になってGPUネタで研究する学生が増えるのを待つしかないか
二次元配列なんて存在しねぇ って考えると楽なのに
昔のオカルト本では「謎の四次元」「四次元失踪」とかあって、 「四次元」とか謎めいた雰囲気を感じた。 今なら四次元配列とか当たり前だw
>>536 そうそうw
四次元の神秘性が薄らいじゃうw
ハミルトンの四元数というのがあって近年計算機工学に応用されるようになって云々かんぬん
発見自体は1800年代だったっけ? たしか橋の上かなんかで思いついたとかw
HLSLとどう違うの?
四次元配列と四次元ベクトルは別物だろ 後者は要素数4の一次元配列
543 :
デフォルトの名無しさん :2012/03/11(日) 17:04:08.60
>>540 HLSLはShader Languageなんで整数とか扱えなかったと思う。
あとステップ数も制限があったような。
Compute Shaderを記述する言語も HLSLじゃ無かったっけ? リソースベースのHLSLと、ポインタ・配列ベースのCUDA
リソースベースって言い方、分かり易いね。 HLSLはまさにそんな感じだ。
誰かいますかね、三つほど質問が。 ■__syncthreads()だけでは共有メモリへの書き込みを保証できない? 1つのスレッドがグローバルメモリから共有メモリにデータを書き込み、 その後全てのスレッドがそのデータを使用して計算を行うような場合、 書き込み後に__syncthreads()だけではなく__threadfence_block()も必要なのでしょうか? 青木本には__threadfence_block()について特に言及ありませんでしたが・・・。 ■ブロック内の全スレッドからの同一グローバルメモリへのアクセス ブロック内で共通で使用する構造体などをグローバル→共有メモリに移す場合 全スレッドで行うよりもやはり if(threadIdx.x==0)・・・ のようにした方が良いでしょうか? ■カーネル内でのreturn文の使用悪影響あるか スレッドごとに計算を行うか判定をする場合、if文で囲っている例をよく見ますが これは if(条件)return; と書いてはいけないのでしょうか? 上のように書いてもとりあえず計算は流れたのですが何か悪影響はあるでしょうか?
>>548 ・取り敢えずsyncしか使ってないけど問題になったことはない。
・全スレッドから共有メモリへの書き込みを行なうのは多分遅くなるんじゃないかな?
・どちらで書いても同じこと。普通のCPUのような分岐とは違うことを判っていればOK。
550 :
548 :2012/03/15(木) 22:39:42.94
>>549 ありがとうございます!!
一個目の_syncthreads()と__threadfence_block()の件ですが、
syncだけだと今日うまくいかなかったもので。
ただ他のバグの影響なども考えられるのでもうちょっと調べてみます。
>>548 ・__syncthreads()は__threadfence_block()相当の処理を
含んでいた気がするけど気のせいかも。
・全スレッドで同じメモリにアクセスするのはたとえfermiでも遅くなるはず。
・カーネル内部で_syncthreads()使う必要があるなら
returnは使っちゃ駄目だろう。
いまだにアプリ開発環境すらまともに構築できてない・・・ visual studio 2008でやろうと思って一応ビルドは通ったけど 実行するとまずcutil32.dllがありませんって出た。 次にcutil.dllをデバッグ.exeと同じフォルダに置き実行!! CUDA version is insufficient for CUDART version. ってなる・・・orz まずなにからはじめるべきですか?
ちなみに.cuの中身は拾ってきたちょっと複雑なコード #include <stdio.h> #include <cutil.h> int main( int argc, char** argv ) { CUT_DEVICE_INIT(argc, argv); CUT_EXIT(argc, argv); return 0; } ・・・・・orz #include <stdio.h> #include <cutil.h> int main( int argc, char** argv ) { return 0; } これに書き換えると プログラムが完成し、エラーもなく実行もできる
GPUドライバのアップデート
>>554 ありがとうございました!!!!!!!!!!
動いた!!!!!!!!
おおw よかったな!
>LINK : /LTCG が指定されましたが、コードの生成は必要ありません。リンク コマンド ラインから /LTCG を削除し、リンカの性能を改善してください。 と表示されるのですがリンク コマンド ラインは固定されて編集できません。 解決方法はありますか?
>>577 補足:
開発環境はVisualStudio2008
cuda ver 2.3
windowsを窓から投げ捨てろ
560 :
509 :2012/03/17(土) 15:16:23.02
そんなことして道歩いてる人の頭に当たっちゃったら大変ですよ
角に当たったら痛そうだもんね・・・
>>557 リンカ -> 最適化 -> リンク時のコード生成 (/LTGG)
C/C++ -> 最適化 -> プログラム全体の最適化 (/GL)
Visual Studio 2008 の使い方なのでスレが違うかも。
>>562 ありがとうございます。
CUDA-Zの実行結果はどのように見たらいいですか?
日本のサイトが全然ないです。
コンスタントメモリキャッシュへのアクセスはバンクコンフリクトとかないんでしょうか??
>>565 そりゃキャッシュはバンクになってないからねー
567 :
565 :2012/03/21(水) 22:44:39.32
>>566 おお、やっぱり。
できるだけコンスタントメモリ使うようにしまつ。
アドレスが静的に解決できないというのが前提だけど 16ポートのSRAMなんてコスト的に不可能だからマルチバンク以外無いんじゃないの?
569 :
デフォルトの名無しさん :2012/03/22(木) 00:27:08.40
Fermi以前はコンスタントメモリ使う意味あったけど、 Fermi以降はL2キャッシュとあんまり変わらない印象
570 :
デフォルトの名無しさん :2012/03/22(木) 22:48:01.36
GTX680が発表されたけど、CUDA的には好ましくない方向の進化が多い。。
チップ名 GTX 680 GTX 580 GPC*1 4 4 SM*2 8 16 CUDAコア 1536 512 テクスチャーユニット 128 64 ROPユニット 32 48 ベースクロック*3 1.006GHz 772M/1.544GHz ブーストクロック 1.058GHz − メモリー転送レート 6.008Gbps 4.008Gbps メモリー容量 GDDR5 2048MB GDDR5 1536MB メモリーバス幅 256ビット 384ビット メモリー転送速度 192.26GB/秒 192.4GB/秒 製造プロセス 28nm 40nm 補助電源端子 6ピン×2 8ピン+6ピン 推奨電源ユニット出力 550W 600W TDP*4 195W 244W
GK104はミドルレンジだからGK110は全体的に上回ってくるでしょ
573 :
デフォルトの名無しさん :2012/03/22(木) 23:16:01.94
kepler誕生おめ! .o゜*。o /⌒ヽ*゜* ∧_∧ /ヽ )。*o ッパ (・ω・)丿゛ ̄ ̄' ゜ . ノ/ / ノ ̄ゝ
Keplerキタ━━━━━━(゚∀゚)━━━━━━ !!!!!
gen3じゃないんだっけ?
>>570 もともとはミドルレンジでグラフィック向けだったから仕方ない気もする。
予想以上にグラフィック方面に舵を切ったという感はあるけど。
このままグラフィック向けとGPGPU向けで大きく分かれていくのではないかという心配はあるかな。
1SM = 192コアか。おっそろしいなあ。
nVidia始まったな。
>>580 あれ?
コンスタントキャッシュって無くなった?
L1/L2キャッシュがその役割を担ってる?
ということはFermiからか・・・
582 :
デフォルトの名無しさん :2012/03/23(金) 01:06:15.11
48コアが192コアになったのに レジスタは2倍、 共有メモリは据え置き。 どーすんだこれ。。
レジスタ足りんくなりそうな。
Keplerはクロック落としてパイプラインを浅くする設計 演算器のレイテンシが小さくなるならレジスタの消費量は変わらない Fermiの18cycleは頭おかしすぎた これが例えば6cycleにになればレイテンシ隠蔽に必要なスレッド数が1/3になるから問題ない
x86 CPUと同じ道を辿ってるのか
586 :
デフォルトの名無しさん :2012/03/23(金) 15:07:43.87
誰か26次元計算してくれ、1000コアくらいじゃマジに足らんぞw
float a[1][1][1][1][1][1][1][1][1][1][1][1][1][1][1][1][1][1][1][1][1][1][1][1][1][1] = NULL; *a += 1;
588 :
デフォルトの名無しさん :2012/03/23(金) 17:47:27.55
こんにちは
国際暗号学会のプレプリントサーバにこんな論文があがってました
Usable assembly language for GPUs: a success story
Daniel J. Bernstein, et. al.
http://eprint.iacr.org/2012/137 GPUのことはさっぱりわかりませんが、なにかこのスレの足しにでもなれば幸いです
それでは
>584 グローバルメモリアクセスのレイテンシ隠匿とか、ループが遅いとかの情報が頭にあったんで 今まで深く考えず1024スレッド突っ込んでたんだけど、 スレッド減らしてループ回すような構造にしたほうがいい、って解釈でいいんだろうか?
>>588 PTXよりもっとネイティブ寄りのアセンブラ言語qhasm-cudasmを使って
パフォーマンスクリティカルな場面で力を発揮(nvccの148%)するよ!って話かな?
暗号学会で発表されるんだね。
メモリアクセスに対する演算の比率を上げないと、性能をフルに発揮できないことは分かったんですが、 具体的にどれくらいの比まで高めるべきかの目標はどうやって決めればイイでしょうか??
理論性能(カタログ値)がでるまで頑張れば良いのでは? それからグローバルメモリーのアクセス速度が,カタログ値の何%になっているのかも チェックすべきだと思う.
593 :
591 :2012/03/26(月) 18:59:54.70
>>592 ありがとうございます!
やはり、
理論性能が出ない → ボトルネックを割り出して改善 → 先頭に戻る
のループで追い込んでいくやり方ですね。
本を読んで、Visual Profilerを知ったのですが、 ひょっとして今はParallel Nsightで同じことができるでしょうか?
GTX 680駄目すぎるわ 死んだ
warp shuffle
VLIWの腐ったようなアーキテクチャになったくさいな
どのへんが?
斜め読みしたうえで完全にESPだが 命令は各スロットごとに別という点でVLIWでデータパスはSIMDみたいに独立とみた ソフトウェア的にはもちろん別スレッドとして書けるみたいな 全然違ったらごめんね
やっぱりGCNと同じでshuffle入れてきたな。
CUDA C Programming Guide Version 4.2 74p. Table 5-1. Throughput of Native Arithmetic Instructions (Operations per Clock Cycle per Multiprocessor) いろいろやばすぎるな。
>>589 いまさらだが
物理レジスタが足りてるなら同時に多数スレッドを保持しておけるが
複雑なカーネルだとレジスタは不足しがち
>604があるから時として64ビットアドレッシングより32ビットアドレッシングの方が有利なんだよね。
Visual studio2010、CUDA4.1、Windows7 64bitではじめようと思ったんだけど、ネットで拾ったプログラムとか動かすとcutil_inline.hが見つからないって出る。 これってどうすればいいの?
>>606 GPU Computing SDKをインストールしてパスを通す
スレタイがイイね。 くだすれくーだすれw
なんでGPGPUはみんな同じようなアプリしかつくらないん? 想像力が欠如してるから他人の猿真似ばかりしてんの? GTX680の性能を生かすアプリ教えろ
わりぃなぁ。そんじょそこらにはないアプリ作っているんだが公表できないんだわ。
恥ずかしいやつが湧いたな
>>609 BOINCでもやってな
なんぼぶん回しても次から次へ宿題出してくれるから
数値流体力学シンポジウムにでも参加すれば良いぞ♪
公表できないと言ってるけど、どうせCUDA ZONEに登録されてるようなものだろ? 外人の真似と負け惜しみしかできないの? 素人だからよくわからないけど、欧米に対して技術面で遅れているから この分野で日本に有名な人がいないでしょ? 自称一流の教授に俺の書き込み見せてあげて
だから数値流体力学シンポジウムにでも参加すれば? この分野では例えば東京工大の青木先生が有名だが
ちなみに 615
あれ? ちなみに615は大学生か??
大学生なら,物理や力学関係の学会に参加すれば,GPGPUを使った シミュレーションの研究結果が報告されていることがわかるはず♪ 高卒なら縁はないが...
大学生がこんな幼稚な文章書いてたら日本終っちまうぞw
GTX480である程度大きなサイズをホストからデバイスに転送するのに3回に2回ぐらいセグメンテーションエラーで落ちる。 うまく行くとなんの問題もなく実行できる。 デバイス側でメモリ確保ができてないみたいなんだが、こんなもんなのかね?
そのカードでモニターを表示させているとか?
ケプラーの倍精度計算は速くなったの? それとも以前と同じ?
HPC向けがでてみないと分からんけど GTX680じゃSPに対して1/24だよ
>>622 表示させている。
それがダメなのかな?
2枚さして,一枚はディスプレイ用 もう一枚は演算用にしないと,一枚では負荷に耐えられないのでは?
みんなすごいね せいぜい準備して 九九の掛け算一括処理くらいしかできないよ
629 :
営利利用に関するLR審議中@詳細は自治スレへ :2012/04/10(火) 00:59:53.96
CUDAを学ぼうとするからそうなるんだと思うよ。 何か問題があって、それをCUDAで解こう!って始めた方が早く習得できる。
>>624 それって今までよりも遅いってこと?
どこかに倍精度のベンチマークの比較はありませんか?
今研究室でGTX680か580のどっちかを買おうって話になってるんだけど CUDA的にはどっちがいいと思う? 一任されて困ってる・・・
floatかdoubleか。 あと整数演算も遅くなったらしい。 テスラ売るためとはいえ、なんかいやーんな感じ
> 631 両方買う. が,テスラの方が安定していると聞いているよ(速度は若干GTXより落ちるが)
>>631 迷わず580かと。
680はレジスタとか整数演算・論理演算のスループットとか色々と問題になりそう。
それに最新の正式版Toolkit 4.1のプログラミングガイド見てもKepler載っていないし・・・
>>632 グラフィック性能のワットパフォーマンスを上げるためというのが一番じゃないかな。
>>631 勉強用に1枚2枚買うって話ならGTX580だろうね。
1.資料がぜんぜん揃ってないKeplerを今買ってもしょうがない。
2.GPGPUとしての性能がGTX580のほうが「上」
(完全上位互換というわけではないし処理にもよるしスペック上FLOPSでは負けてはいるが)
GTX680のグラフィック・GPGPU性能を調べる ≪ dokumaru
http://dokumaru.wordpress.com/2012/03/27/gtx680-spec/ 10数枚買って研究室全体に大量導入…なら先生が決めるよね。
両方買ってもいい。でもそれならKeplerは対応するToolkitが出てからでも遅くはないかと。
あるいは一刻一秒を争うならなおさらKeplerは冒険かと。
今はまだ早いGK110をまて 5月にイベントあるから、そこでなんかあるかも
Teslaはメモリ容量が良いよね GTXだと計算領域が足りないよ・・・
638 :
デフォルトの名無しさん :2012/04/11(水) 23:21:32.32
FLOPS/MBで見ると、Teslaでも全然足りない。
何の計算??
640 :
デフォルトの名無しさん :2012/04/13(金) 04:19:22.20
倍精度計算が主なのですが、Ivyと680と580、どれがコストパフォーマンス的にお薦めですか?
IvyはCUDA動かないよ。
言語の問題じゃなくて、プログラムはこれから作るから倍精度計算をわんさかやろうと思うんだけどどれがいいかなあ? 程度の話じゃないかと。
そういうことか。 超並列に対応できるのであればGPUのほうがイイね。
えっ!? ここCUDAのスレだよね? てか、Ocelotとかどうなのかな。
倍精度計算が中心なら、CPUで最適化するのが一番。 例えば近似計算のようにGPUの単精度で近づけてから、 CPUの倍精度で収束させるとかならありだけど。
cuda zoneがメンテナス中・・・
toolkit4.2が来るのかな
>>640 GPU用の倍精度プログラムを書く気があるならTeslaにしとけ。
コストが厳しいならRadeonの最上位にしとけ。
sdk 4.1とtoolkit 4.1インストールしたんだけど アンインストールせずにそのまま sdk 2.3とtoolkit2.3をインストールしたらコンパイルやリンクの挙動とかおかしくなりますか?
自分でpathやMakefileなどを管理できるのなら無問題。
parallel nsightは甘え
>>651 甘えと言えるほどすごいのか。
今度使ってみようw
シングルGPUでもデバッグできるようになったのが凄くうれしい
654 :
デフォルトの名無しさん :2012/04/17(火) 03:52:59.70
680を手に入れたんだけど、ガッカリ性能だった ゲーム系は速くなってるんだけどね
何をやったの
PTX直接書いてプログラミングする人とかいるの?
一応いるけど。
GTX580が生産終了なんだとか
>658 在庫はけたからか。 しかしIntelのCPUはGPU載せてんのにGPGPUにはさっぱり対応しないからあんま意味ないな。 CUDA対応とかにしないのは戦略的判断なんだろうけど、なんとももったいない。
大丈夫 Intelも来週からGPGPU対応する
661 :
デフォルトの名無しさん :2012/04/18(水) 16:31:10.22
倍精度計算じゃまだインテルに分があるしね
ivyのeuどうなってのかね
> 654 今,それに触手を伸ばしているところだけど,どこがダメだった??
OpenCLならできなくもないのかな?
誰かHMPP使ったことある人いる?
ivyの影響だな
GPUのデメリットは同じ変数計算を毎回糞真面目に超高速で行うところ
メモリ読むより速いからな
>> 671
一つ一つの計算は超高速でもなんでもない
並列で行うので早くなるだけ
超高速になるか否かはプログラミングの問題
>>672 演算にはメモリーの読み書きを伴うので,演算が「メモリ読むより速い 」とはならないのでは?
>>671 意味が分からん。
アーキの概念の理解ができていないじゃねーか?
>>675 明日短いわかりやすいソースアップするからコンパイルして実行してみて
言いたいことがわかると思う。
CPUにはあってGPUにはない機能を使うことになる、まぁホントしょうもないことだけど・・・
???
言ったら悪いかも知れんけど単にアルゴリズムが悪いんじゃないのか。
>>671 アドレス計算とかまさにそれだよね。
普通のループなら+4で済むところが、
ptr + threadIdx.x*4 + threadIdx.y*hoge
とかになっちゃう。
それはGPUのデメリットじゃないな。 GPU(nvcc)でもループなら普通に書いたら普通に最適化してくれる。
>>679 これはデメリットと違う。
CPUでマルチスレッドでやれば同じように明示的にアドレス計算を行う必要がある。
シングルスレッドでの最適化が、そのままマルチスレッドに使えると思ってるなら、並列で組むのに向いてないな。 ひとつの処理として見たとき無駄でも、それで大多数の演算を同時に走らせることができるなら、 並列処理においてはそれこそが効率的なんだよ。
684 :
デフォルトの名無しさん :2012/05/02(水) 01:48:54.23
同じ世代のGPUでも生産地の違いで演算速度は全く違うからね もっと言うと転送速度が全く違う まあフラッシュメモリでも同じこと言えるけど
686 :
デフォルトの名無しさん :2012/05/02(水) 04:31:00.56
>683 ずいぶんニッチなところだな
>>684 マジレスすると、CUDAでやるメリットはない。
Sandyやivyの方がはるかに高速。
まあ、CUDA勉強するだけならいいが、もっと別のカードのほうがいいだろ。
>>684 メモリ転送が遅過ぎて4GBのメモリを活かしきれない悪寒。
SRAMを4GBつんでるカードはないのか?
>>667 >New Features
>Support for GK10x Kepler GPUs.
とりあえず、GK104対応にしました的か。
691 :
デフォルトの名無しさん :2012/05/02(水) 18:28:26.54
VRAM 4GB以上のカードって、ほとんどないんだね。 TeslaかQuadroしか見つからなかった。 お値段10万円越
4年後にはVRAM16GBが普通にでまわるんだよ
693 :
679 :2012/05/02(水) 22:18:41.81
>671は皮肉だろ。
ということにしたいのですね。
デメリットに感じる境地まで辿りついたんだよ、きっと 俺にはまだメリットにしか思えないんだけど・・・
>>692 今から4年前の2008年頃はG80世代で大体1GBだったが、4年後のGTX680でまだ2GBだから、
4年後はせいぜい4GBなんじゃないの?Tesla系で16GBにはなっていそうだけど。
そういやPCのDRAM搭載量に比べて、あんまり伸びないよね>ビデオカードのメモリ
GDDRは数が出ないからね。 DRAMメーカーがあんな状態だから尚更でしょう。
プロセスシュリンクが汎用DRAMと同じように進めば同じようにでかくなると思うんだけど。 だんだん引き離されてるってこと?
日本のメモリの会社が潰れたのはかなり痛いな・・・
ptxコード読まなきゃいけなくなったんだけど、typeの.predって何なのかいまいちわかってない
述部(predicate)だね。 ptxの場合は単に、比較などの結果を保持するだけのような希ガス。 で、そのレジスタの結果に依存してインストラクションの実行する、と。 例えば、 -- setp.gt.s32 %p1, %r5, %r7; @%p1 bra $Lt_0_12802; -- なら r5 > r7のときに分岐するし、 -- setp.lt.s32 %p2, %r9, %r11; @%p2 sub.s32 %r14, %r11, %r14; -- なら r9 < r11のときに引き算を行なう。
分岐マスクのためのレジスタは何本あるんだろ それとも汎用レジスタと共用なのか
705 :
702 :2012/05/11(金) 16:21:50.31
>>703 thx
そういう意味だったのか……
CUDAはC言語の延長だから大丈夫とか考えた三月の俺を叩きのめしたい
PTXコードの読み方って英語のやつしかないよねたぶん
>>705 私が書いたメモならあるよw
>>704 実験コードで見たところ、汎用レジスタと述語レジスタの合計で制限されてたかと。
述語レジスタだけでどこまで増やせるかは実験してない。
707 :
702 :2012/05/11(金) 17:26:43.70
>>706 恵んでください。
割と切実に。卒業したいので。
ISA的にはwarpあたり7本か6本じゃね。 3bitのどれかが常にalways扱いだったような。
709 :
デフォルトの名無しさん :2012/05/14(月) 15:19:34.27
初心者質問です。 お願いします。 cufftってcuda3.2でも使えるのでしょうか? cufftdestroyが未解決の外部シンボルだと言われてしまうのですが? ただ単に、リンクできてないだけなのでしょうか?
>>709 使えたと思うよ。
つーか、cufftdestroy()が未解決って、あんたの間違いだろ。
711 :
デフォルトの名無しさん :2012/05/15(火) 14:03:46.41
>>710 返信ありがとうございます
他の関数はコンパイルが通る(通っているように見えるだけ?)のに
cufftdestroy()
cufftExecZ2Z()
cufftPlan1d()
だけが未解決となっているのですが、
この関数だけ、他のライブラリが必要だなんてことがあるのでしょうか?
destroyはDestroy。 z2zは未実装。 Plan1dはしらね。 警告レベル引き上げれば?
cufft.hはインクルードしているのかな?
715 :
デフォルトの名無しさん :2012/05/17(木) 15:21:47.32
警告レベルって、デフォルトは最大なんですよね? Destroyに関しては、タイプミスです。 z2zは未実装っていうのが、よくわからないんですけど。。。。
GPGPU上でソケット通信とかって出来るかな
GPGPUの仮想マシン同士のn対n通信をシミュレートとかそういうのをイメージした
>>715 未実装: 実装されていないこと。
cufftのライブラリの中にz2zの関数そのものが存在していないのよ。
で、あんたがどんな環境で開発しているか判らんのに警告レベルがどうなっているかなんか判るかい。
そんなことは自分で調べなさいよ。
>>716 cuda5でLAN接続されているGPU同士で連携させる機能がつくらしいよ。
多次元配列を扱えないのは何でなんだろう. ブロックとスレッドインデックスで一次元化するの面倒なんだけど.
ピンメモリを確保すると、スワップによる退避を防げるのは分かったのですが、 実際はスワップ以外にも、メモリフラグメンテーション解消のためのコンパクションでも メモリアドレスの変化って起こり得ますよね? それもないようにするのがピンメモリですよね?
>>719 別に扱えなくはないぞ。普通にdata[blockIdx.x][thiredIdx.x]ってできると思う。
スレッド数を定数にしなくちゃならなくなるから却って煩わしいと思うけど。
つーか、面倒ったってオフセット計算する関数を作るだけじゃん。
>>719 多次元配列はサイズが大きくなるとポインタ分データが多くなるのでポインタ分だけでキャシュからの溢れる。
だから一次元でやる。
面倒ならマクロで定義すればいい。
>>723 >多次元配列はサイズが大きくなるとポインタ分データが多くなる
は?
>>722 いやさ、今は実装されているならいいんだけどね。
少なくとも暫く前のcufftにはcufft.hで宣言されていたのにライブラリにはなかった。
>>723 なんでオフセット計算をマクロでやるの?
gpu関数はインライン展開されるからマクロより悪くなることはないんだけど。
ここでいう多次元配列って静的に作ったやつのことか? ならインライン展開されるわな。
__device__の関数→インライン展開される。 静的な多次元配列→インライン展開される。 ∴動的な多次元配列をポインタ配列で実装した場合のみ、インライン展開されない。 つーか、ポインタ配列なんて面倒でキャッシュ効率の悪いデータ構造なんて使うなや。
>>721 CudaMalloc使ってるんだけど,
Data[Mx][My][Mz]の配列があったとして,Data[i][j][k]にアクセスするには,
Data[Mx*My*k+ My*i+ j]としないといけないじゃん.
配列要素を求める関数を作ればいいわけね,了解.
ありがとう.
>>729 __device__ static int offset(int p1, int p2, int p3, int w2, int w3)
{
return p3 + (p2 + p1 * w2) * w3;
}
731 :
デフォルトの名無しさん :2012/05/18(金) 20:25:07.43
プログラムによるけど、ループの早い段階でiやjを確定させてアドレス求めて、 最終要素だけポインタインクリメントの方がだいぶ速い。
それが一番コンパイラにやさしい。
tesla k20は24GBもメモリが載るみたいだな。 でもメッチャ高そう。 50万位かな?
735 :
デフォルトの名無しさん :2012/05/21(月) 15:46:37.20
>>712 解決しました。恥ずかしながら、ただ単にcufftライブラリにうまくリンクできてないだけでした。
全ての関数において、実装されているようです。
でも、なんで他のcufft関数においては、リンクエラーが出なかったのでしょうか?。。。。。
sampleが置かれたディレクトリでコンパイルすれば,そのままOKだし. sampleを別のディレクトリにコピーしてからコンパイルすると,設定し直しのところが多々あるが, そう言うことかい??
737 :
デフォルトの名無しさん :2012/05/23(水) 14:44:51.35
詳しい方に質問させてください。 Win7Pro64 + VC2010Pro + CUDA4.2で32bitのdll (Matlabのmex)を 作成して、カーネル(1回のみ呼び出し)の実行時間を計測しました。 Nsight Visual Studio Editionのプロファイルでは約3583msとなったのですが、 メイン関数側でカーネル呼び出しの時間を調べると12079msとなりました。 この差の原因がわからずに困っているのですが、心当たりのある方は いらっしゃいませんか? 足りない情報などあればツッコんでください。 よろしくお願いします。
>>737 詳細は実際のコードを見ないと判らないけれど、カーネル呼び出しは完了復帰じゃないよ。
逆に、それを巧く利用すればGPUとCPUで並列処理ができるって寸法だ。
739 :
737 :2012/05/23(水) 15:25:51.83
>>738 レスありがとうございます。
メイン関数側での測定は、カーネル呼び出し後に
cudaThreadSynchronize();
としてから測定していますが、これではダメですか?
これを入れ忘れると、メイン関数側での測定時間が極端に短くなる
という話はよく目にするのですが、今回は逆に、メイン関数側での
測定時間が極端に長くなっています・・・・
>>739 おっと、時間は見落としていた。
ドライバレベルで実装してないなら、カーネル呼び出しの陰に
いろんなものが呼ばれているんだと思うよ。
グリッド数が多ければカーネル呼び出しは複数回に分割されるし、
GPU側は純粋にカーネル時間しか計上されない筈。
しかし、12秒か。ほんとにそんなに掛かってるの?
741 :
737 :2012/05/23(水) 15:57:57.69
>>740 実際にmatlab側でdllの呼び出し時間を測っても12秒程度かかっています。
問題のサイズに応じてブロック数とスレッド数を変えても、無駄時間ごと
スケールしているような印象です・・・
12秒というのは、ブロック数512, スレッド数32で呼び出した時の結果です。
普通はどのくらいのラグがあるものなのでしょうか?
またグラボはGTX560Tiの1枚のみの環境なのですが、画面出力用に
別のグラボを挿したら改善する可能性などありますか?
>>741 SPをどれだけ使えるかは画面表示でどれだけ使っているかとトレードオフだから、
画面出力を止めると処理時間が短くなる可能性は充分あるね。
ブロック:スレッド=512:32でそれなら、256:64か128:128にしてみたら?
あー、問題の性質に対して固定なら難しいか。
できれば、問題の性質に関係なくスレッド数を決められるようにした方がいいのだけれど。
743 :
737 :2012/05/23(水) 16:24:24.92
>>742 ブロックとスレッドの積を揃えたときの実行時間も調べたのですが、
スレッド数が8,16,32,64の中では32が一番早かったです。
CUDAのドライバAPI(cuda*ではなくcu*の命令?)は使っていないのですが、
カーネル関数をただ呼び出すときでも陽に使うものなのでしょうか??
また、そのあたりにどのくらいの時間がかかっているかを調べるには
どのような方法がありますか?
>>743 私ゃ面倒だから使ってない。
カーネル内でスレッドに依存するようなコードでなければ、
スレッド数32にピークが来るのはなんか変だな。
スレッド数はもっと多く取れと書かれていたのを読んだ気がするし、
実際もっと多く突っ込んだ方が速くなった。
745 :
デフォルトの名無しさん :2012/05/24(木) 14:50:32.51
cudaSafeCall() Runtime API error : unknown error.ってなんだよ!? エラーの原因ぐらい解明してくれよ!! メモリ確保のあたりが、おかしいんですかね?
CUDAの関数は、完了復帰とは限らないのでご用心。 ソース見ないとなんとも言えないけど、その呼び出しの以前に発生したエラーが その呼び出しで露見して、エラーを返しているのかもしれない。
カーネル関数Aを実行して、カーネル関数Bを実行して、カーネル関数Cを実行して、 それをループ、という処理を考えてるんですが。 実行するカーネル関数を切り替えるのって、オーバーヘッドありますか?
>>747 そりゃぁ、ありますがな。できればカーネル一個でできるだけ時間稼ぐべきだけど、
問題の性質によっては分割数の関係でカーネルを分離せざるを得ないしね。
749 :
デフォルトの名無しさん :2012/05/25(金) 10:38:00.51
>>746 なるほどです。
直前のcuda関数からあらってみます。
750 :
737 :2012/05/25(金) 19:53:25.01
>>744 レスありがとうございます。
あれから色々調べてみたところ、カーネル呼び出しに無駄に時間が掛かるのは、
NVIDIA Nsight Visual Studio Edition (Win 64bit, 2.2.0.12132)が原因のようでした。
(これをアンインストールしたところ、呼び出し側で測定しても、カーネル側で測定
した時間と同じ程度の時間しかかかりませんでした。)
クリーンなWin7環境にVS2010とMatlabとCuda Tool Kitを導入した上で同じNsightを
インストールした場合には、特に遅延などの問題もなく呼出しができたので、
ご相談した件は、他のソフトウェアなどの環境とNsightとの複合的な要因による
かもしれません。
どうもお騒がせしました。
>>745 このエラー、たまにだけど確かにおきる。
デバッグのノウハウとかあったら教えてほしいです。
752 :
>24 ではないです :2012/06/01(金) 14:54:14.19
>>24 ,
>>26 で触れられてから9ヶ月経っていますが
NVidiaのコンシューマ向けGPUの計算結果の正しさをチェックする
スクリーニングソフトウェアって今でも公開されいないのでしょうか?
>>752 CPUとGPUの計算結果の比較を取ればいい。
CUDA5の正式版はいつ公開されるの
cudaをVM上からさわる方法ってないんだろうか
756 :
デフォルトの名無しさん :2012/06/01(金) 22:25:29.40
Linux系のVMでPCI IDを持った任意のデバイスをVM側のネイティブデバイスに見せかける 方法が5年くらい前はあったはず。
CUDAで計算した結果をDirect3Dで描画したいんですが、 同じプログラム内で同時に使っても問題ないでしょうか? 一つのGPUがCUDAとDirect3Dを切り替えて使うので、 切替のオーバーヘッドとか大きかったりするんでしょうか??
ここでどんな答えが得られたとしても 自分で実際にやってみないと無意味な事だと思わないか 聞いているヒマがあったら自分が納得のいくテストプログラムを作らないか
760 :
758 :2012/06/02(土) 02:54:57.22
できるのであればnVIDIAのハイエンドのGPUを買いたいです。 できないのであればRadeonを買いたいです。 その判断をしたいんです。
RadeonでCUDA使えるのか?意味が分からん
>>760 できるかできないかで言ったら、できる。
NVIDIA GPU Computing SDKの中にもCUDAによる計算結果をDirect3Dで描画しているサンプルもある
(An example of fluid simulation using CUDA and CUFFT, with Direct3D 9 rendering.)。
オーバーヘッドの大きい小さいは知らんけど。
763 :
758 :2012/06/02(土) 11:43:21.60
>>761 すみません、CUDA+Direct3DができないならCUDAは不要で、
純粋なグラフィックス用途(Direct3D)のみの観点でGPUを選ぶことになり、
その場合は個人的にRadeonのほうが良いと考えているからです。
>>762 ありがとうございます!
そのサンプルの情報を参考に検討してみます。
764 :
762 :2012/06/02(土) 12:51:16.13
ちなみにタイトルはFluids (Direct3D Version)。
意外とCUDA+描画はDirect3Dのサンプルってないんだなぁ。
このFluidsはD3D9を使ってはいるけど実質2Dの描画だから、
3Dのレンダリングをさせた場合とは負荷が違いそうなので注意。
しかしわざわざSDKに
Fluids (Direct3D Version)
Fluids (OpenGL Version)
とあるくらいなのでDirect3DとCUDAの併用って何か壁があるんだろうか。
CUDAは2Dどころかコンソールで済む使い方しかしたことないのでわからん。
サンプル内にCUDA+DirectXの組み合わせってなかなか見当たらないんだよね。
Ocean SimulationもN-Body SimulationもCUDA+OpenGLとDirectCompute+DirectXの2パターン。
ただ↓のリンク先にこう書かれてあるから可能は可能なのだろう。
>OpenGLおよびDirectXグラフィックスドライバとCUDAドライバを同時使用可能
AV関連機能編 1/2 | ビデオカード選びの新常識 | DOS/V POWER REPORT
http://www.dosv.jp/other/0907/09.htm
765 :
758 :2012/06/02(土) 12:58:25.72
>>764 とても参考になります。
そういえばDirectComputeがありましたね・・・
グラフィックスとの親和性という意味ではそちらのほうが適切かもしれないので検討してみます。
(ただDirect3D9しか知らないので、Direct3D11を覚えないといけませんが・・・)
>>764 すべての環境がDirect3Dが使えるわけじゃないからな。
Windows環境しかやらないならDirectComputeあたり使ってればいいじゃん?
カーネルを起動するときにグリッドとブロックの設定を、 コンパイル時に<<<16,128>>>としてビルドする場合と、 例えばプログラムの引数で動的に変えられるように、 <<<grid,block>>>のようにコンパイルすると、 後者のほうがパフォーマンスが悪いんだけど、 グリッド/ブロック指定によってインライン展開とかに影響があるのかな? 決め打ちだと、GPUの種類ごとにバイナリを用意しないとならんのが嫌なので、 引数で変えられるようにしたい。
その違いによってコンパイル時にどういう展開になるのかとかは知らんので 他の詳しい人に任せるとして、小手先の対症療法だけどパターンが限られているなら if(...) Kernel<<<16,128>>>(...); else if(...) Kernel<<<32,64>>>(...); とかは?これで決めうちに比べて速度低下しないのかどうかは知らんけど。
769 :
767 :2012/06/03(日) 17:18:27.98
>>767 すみません、単に自分のコードがミスってて、grid = 16, block = gridな感じで単なるバグでした。
スレ汚しすみませんでしたm(__)m
ぶw 原因わかってよかったね
突然すみません error LNK2001: 外部シンボル "_cudaD3D9SetDirect3DDevice@8" は未解決です。 cudaD3D9SetDirect3DDevice()という関数がうまくリンクされていないというのですが原因が全くわからないです。 このエラーを調べたのですが全然原因がわかりません。 誰かわかる方いませんか?
すいませんようやく解決しました。 Direct3D9との相互運用はなかなか難しい・・・
昔の経験では、D3D9連携でcudaD3D9RegisterResource()がかなり重かった cudaMemcpyでhostに転送してからGPUに書き込むのと速度が大差なかった データサイズが大きくなかったせいかもしれないが
もう一般向けのGPGPUの時代は終わったよ。。。もうね趣味の領域を出ないわ
金蔓になるのはしょぼい自作ユーザじゃなくて、テスラ買ってくれる方だろうね。ぬふぉも必死なんでしょ
tesla K10はいくらくらいなんだろ?
777 :
デフォルトの名無しさん :2012/06/08(金) 14:37:45.13
カーネル関数で乱数使いたいとき、みんなどのライブラリを使ってるの? curandとtinymtは使ったことがあるんだけど、他におすすめある?
その二つに何か問題があって質問してるの?
779 :
777 :2012/06/08(金) 15:22:52.38
>>778 標準でついてくるからcurandを最初に使ってみたんだけど、
constメモリやレジスタを食いつぶすうえに遅いので窓から投げ捨てた。
みんなほんとにこんなの使ってるのか?と思って聞いてみた。
curandのかわりに偶然見つけたtinymtを使ってみたんだけど、浮動小数の乱数を返す関数が
(区間(1,2]の一様乱数)- 1.0f
のような実装なんで、machine epsilon以下の値が出ないんじゃないかと気になってる。
ほとんどゼロみたいな確率だから、実用上影響ないってことなのかな?
こんな感じで気になることがあったんで、スレ民のおすすめを聞いてみたくなったんだ。
780 :
777 :2012/06/08(金) 15:28:07.77
間違えた。tinymt32_double12は区間(1,2]じゃなくて区間[1,2) の乱数を返すみたい。
781 :
777 :2012/06/08(金) 15:38:50.50
また間違えた… tinymt32_double12って何者だorz tinymt32_single12かtinymt64_double12に読み替えてね
CPU側のMTで生成した乱数をメモリに転送してるな たくさんいるならカーネルで生成したほうがいいのだろうが
>>779 一様乱数ならその実装の方が良いんじゃね。
784 :
777 :2012/06/08(金) 18:19:32.88
>>782 >CPU側のMTで生成した乱数をメモリに転送してるな
>たくさんいるならカーネルで生成したほうがいいのだろうが
ごめん、tinymtはcuda実装があるんで、そいつの話をしてるつもりだったんだ。
(www.math.sci.hiroshima-u.ac.jp/~m-mat/MT/TINYMT/CUDA/index-jp.html)
だからMTはCPU側じゃなくてGPU側で、thread毎に別のseedから生成しているつもり。
単にtinymtって言ったら元のc実装のことだと思っちゃうよね...すまん。
785 :
777 :2012/06/08(金) 19:19:16.31
>>783 curandの方は少し違う実装になってるみたいなんで、ちょっと気になったんだ。
[1.0 .. 2.0) というのは (rand() & 0x7fffff) | 0x3f800000 ということ。 これて通じないなら掲示板で説明するには分量が多すぎるから、そういうものだと思い込むか、疑似乱数スレ見たりするといいんじゃなかろうか。
数値計算の基礎ってことで行列ベクトル積とか行列の和とかCUDAで実装したいんだけどサンプルコードとか載せてるサイトってありますか?
789 :
777 :2012/06/09(土) 00:37:53.14
>>786 uint32の乱数の上位9bitをマスクして残りの仮数部だけを使えば[1,2)の乱数ができる、
ということなら分かってるつもり。
でも、ここから1.0fを引いたら、0とepsilonの間の値が取れないような気が....
curandの実装はその辺ちゃんとしてるのかな??
擬似乱数スレなんてものがあったとは気づかなかったんで、ちょっと覗いてきます。
epsilonってのは気にしても無駄な範囲だろ 一切計算する気ないのか
792 :
777 :2012/06/10(日) 23:28:03.14
擬似乱数スレかとおもってたら「疑」似乱数スレだった・・・
しかも5ヶ月以上止まってるしorz
>>790 >>791 一様と言うからには、epsilon以下の値についても、その浮動小数に丸められる区間の長さ?に応じて出現してくれないと、単精度で
if (rand() < 1e-10)
などとしたいときに困らない?
自分が今困ってるわけじゃないし、あってもレアなケースだとは思うんで、実用上問題なさそうなのはわかるんだけど、epsilon以下の値が帰ってくる実装もあるみたいなんで、ちょっと気になってただけなんだ。
有限精度な時点で飛び飛びの値を取るのは仕方が無い。 それでも、786みたいな実装なら乱数値として表れうる数は 等間隔で並んでいるから、表れうる値で任意の区間は 元の乱数が完全なら全て同じ濃度になる。 x / 2^32+1みたいな処理なら、無限精度があれば 上記と同じほぼ一様な擬似乱数となるだろうが 現実には浮動小数点の指数が変わり目で 濃度がグチャグチャになる。
>>792 そもそも浮動小数点数のことが分かっているのか?
例えば0.1から始まるfloatの値で、小数点以下7桁より下は表現できないわけだが
0-0.1の間だけで細かい精度が出るとかその方がおかしいだろ
あとepsilonというのはDBL_EPSILONやFLT_EPSILONのことだと思っているが違うのか
今から勉強するならOpenCLがいいのかな
OpenCLがGLレベルに普及するとは思えないなあ。 今からだったらOpenACCの方がいいかも。
C++ AMPもあるし正直まだわからんな
799 :
777 :2012/06/11(月) 18:05:58.78
>>793 最初の6行は理解してたつもりなんだが、
>現実には浮動小数点の指数が変わり目で
>濃度がグチャグチャになる。
これは実際に試さないと理解できないような気がするけど、ちょっと納得した。
どうもありがとう。
>>794 ,795
epsilonの意味するところはあってると思うけど、精度についてはなにか勘違いしてないかい?
スレ違いだけど、単精度浮動小数点の(絶対値の)最小値は 2^-126 ≒ 1.1755e-38 なんで、
FLT_EPSILON = 2^-23 ≒ 1.1921e-7 よりも小さい値を扱うことはできるよ。
wikiの浮動小数点の項目には詳しく書いてないけど、
ocw.kyushu-u.ac.jp/0009/0006/lecture/4.pdf の 10-12ページとか、
www-ics.acs.i.kyoto-u.ac.jp/~nagahara/nc_text/nc_text.pdf の9-12ページ(14-17枚目)
なんかに詳しく書いてあるみたい。
http://codepad.org/60zSd6wZ CUDAで行列ベクトル積を計算しようと思ったんですが
gemv.cu(38): error: argument of type "float *" is incompatible with parameter of type "float **"
のようにエラーが出ます。
配列の定義がおかしいというのはなんとなく分かるのですがどのように改善したいいのかアドバイスお願いします
>>800 エラーそのままじゃん
56行目のgemvの引数aの部分がおかしい
>>800 単にfloatポインタはfloatのポインタのポインタとの互換はありませんって言っているじゃない。
CUDAというよりもC言語の問題だぞそれ。
コード見てみたけど、ダブルポインタは必ずしも2次元配列を表すわけじゃないぞ。
cudaMallocHost((void **)&a,sizeof(float)*size);
<=これがおかしい。
やるなら、
cudaMallocHost((void **)&a,sizeof(float*)*size);
for(int i = 0; i < size; i++)
{
float *temp
cudaMallocHost((void **)&temp,sizeof(float)*size);
a[i] = temp;
}
とかやらないかん。メモリ管理はちゃんとせなアカンが。
>>800 a[i * size + j] = 1.0;
とか。
他にもg_ansはホスト側なのかデバイス側なのか、bを何回初期化するのか、
wkのところは+=ではないのか、この初期値で期待通りの動作をしているか確認できるのか、
とか色々と気になる点がある。
最新のおすすめの本ってありますか?
> 804 CUDAとOpenGLを組み合わせたい人には 先端グラフィックス言語入門 ~Open GL Ver.4 & CUDA~
GT440と640で性能大きくかわる?大差ないなら440買おうと思うんだけど
GT440の96coresに対してGT640は384coresで1:4か。他の条件が大体同じとするとどうなるんだろ。 Keplerは倍速シェーダじゃないから1:2相当。しかしcoreあたりの性能が違ってるから、 単精度演算メインならそのまま1:2くらいで、倍精度や整数演算とかが増えると1:1未満かな。 もし単精度メインで使う予定だったり新しいモノ好きなら640を、 用途はまだ不明で価格差がそれなりにあるなら440でいいんじゃないかな。
Xeon Phiが発表されたけど、K20の情報は出てこないな。
SCまで出ないんじゃね GTX685の情報の方がまずは気になるわ K20は当分触れる見込みないし
K20のどんな情報が知りたいの?
K20とXeon Phiはどっちの方がはえーんだろ? Xeon Phiの方がなんかすごいような感じがするが。 K20はメモリが多いのがメリットになるのかなあ。
220から470に乗り換えて倍精度計算を始めようと思ったらコンスタントメモリまわりでエラー。
理由を探していたら
>>580-581 で既出か。
Fermiでコンスタントメモリ無くなってたのか?
コンスタントメモリなくなったんか! ブロードキャストがなくなって性能落ちた!?
817 :
デフォルトの名無しさん :2012/07/16(月) 23:46:38.53
個人的感想だけど、Fermiになってキャッシュ効くから 小容量ならグローバルメモリもコンスタントメモリも 大差無い気がした。
>>817 キャッシュだとポート(バンク?)の取り合いが起こっちゃうよね?
コンスタントメモリは読み込み専用の強みを活かして、各コアに一斉にデータを供給することができるんだよね?
この差は大きくない??
サイトがハックされたってメールがNVIDIAから来たんだが、 このスレでは欠片も話題になっていないんだな
今時あわてるような話でもないと思ってた
GeforceでCUDAのコードを書いているんだが、 Linuxに比べWindows7は半分ぐらいのスピードにしかならない。 WDDMが問題みたいだけど、今は解決の方法があるのかな? NVIDIAのフォーラムはなんかよろしくない状況みたいだから情報が入ってこない。
Windows7ショボいな・・・
グラボ2枚挿して片方何も接続しなくても遅いの?
>>823 ダメ。
WDDMがすべてのメモリを占有してしまう。Tesla系ならTCCが使えるから問題はないんだが。
>>821 処理としてはなにをやってる?
半年前に480GTX@Windows7と580GTX@Ubuntuでやったけど
はじめWindows7のほうが速かった。
(微々たる差だけど)
なにか試行錯誤して最終的にはLinuxのほうが若干速くなった。
具体的になにをやったかは覚えてないけど。
>WDDMがすべてのメモリを占有してしまう。Tesla系ならTCCが使えるから問題はないんだが。
Physx専用にするとWDDMは関係ないはずだけど。
>>826 nVidiaコントロールパネルを開いたことあるか?w
>>825 Physx専用にしてもWDDMの管理下だぞ。
nvidia-smi.exeを実行してみ
最近CUDAの勉強を始めて、試しに画像処理アプリをGPU用に書き換えているのですが どうしても分からないことがあるので質問させてください double *a[5],*d_a[5]; int size=1280*1024; //実際には読み込んだ画像のサイズが入ります for(i=0;i<5;i++){ a[i]=(double *)calloc(size,sizeof(double)); for(j=0;j<size;j++)a[i][j]=i*j; //ここのi*jは実際には読み込んだ画像の画素情報が入ります } for(i=0;i<5;i++)cudaMalloc((void**)&d_a[i],sizeof(double)*size); for(i=0;i<5;i++)cudaMemcpy(d_a[i],a[i],sizeof(double)*size,cudaMemcpyHostToDevice); func<<< Dg, Db >>>(d_a); __global__ void func(double *f_a[5]); 上記のf_a[5]の、ホストにおけるa[i][j]にあたる値を参照するにはどうしたらよいでしょうか ホスト同様にf_a[i][j]とすると Warning: Cannot tell what pointer points to, assuming global memory space 上記の警告が出て上手くいきません 超初歩的な質問かと思いますが、アドバイス宜しくお願いします
>>829 d_aはホスト上の配列。そいつへのポインタを渡してもなんともならない。
d_aが固定長でいいなら、__global__ void func(doule * f_a0, doule * f_a1, doule * f_a2, doule * f_a3, doule * f_a4);
とするのが手っ取り早い。
d_aが可変長だったり巨大だったりするなら、
double ** dPtrArray, ** hPtrArray;
hPtrArray = malloc(sizeof(* hPtrArray) * length);
cudaMalloc((void **) & dPtrArray, sizeof(* dPtrArray) * length);
for (int ic = 0; ic < length; ++ic) {
cudaMalloc((void **) & hPtrArray[ic], sizeof(* tmp) * size);
cudaMemcpy(hPtrArray[ic], a[ic], sizeof(double) * size, cudaMemcpyHostToDevice);
}
cudaMemcpy(dPtrArray, hPtrArray, sizeof(* hPtrArray) * length, cudaMemcpyHostToDevice);
とでもしないと。
で、こんなの一々書くの面倒だからTrust使えよって話もあるな。
832 :
829 :2012/08/02(木) 17:59:10.35
>>830 早速の回答ありがとうございます
書き忘れましたが、a、及びd_aは実行対象次第で2~5の間で変動するので下のやり方が良さそうですね
この後少々用事があるためすぐに試すことが出来ないのですが、参考にさせて頂きます
ちなみにTrustとは何でしょう?
軽く検索をかけたところThrustと言うものが見つかりましたが、こちらのことですか?
833 :
830 :2012/08/02(木) 19:13:23.52
>>832 あーそうそう、Thrustだね。
画像サイズの上限を見込めるなら、固定長配列として確保した方が効率いいかもね。
>830には書いていないけれど、どうせサイズ情報も引き渡さないといけないのだろうし。
それはそうとして、画像処理なのにdoubleなの? floatでは精度足りないの?
834 :
829 :2012/08/03(金) 00:35:11.95
>>833 レスが遅くなってしまいすみません。
要素型については、他のアプリとの連携や扱う画像の種類の関係上doubleでないと都合が悪いようです。
Thrustは独力でマニュアルを読むしかなさそうで大変そうですが、余裕があったら使ってみます。
ちなみに教えて頂いた方法に関して、
カーネル関数の実行 → func<<< Dg, Db >>>(**dPtrArray);
関数の宣言 → __global__ func(double **f_d);
関数内部での値の参照 → f_d[i][j]
でよいのでしょうか?
ダブルポインタの扱いに不慣れなため、質問が多くてすみません。
835 :
デフォルトの名無しさん :2012/08/03(金) 07:57:43.31
3次元で偏微分方程式を解いています。 境界部分は別に解かなければいけませんが、どのようにするのが適切ですか? 通常のCだと、 for (int i=1;i<N-1;i++) のようにしますが、CUDAでやる場合は、どのようにしたらよいのでしょうか? ifで分けると、3次元の場合、膨大な数になってしまいます。
>>835 境界は計算しないでいいということなら、条件分けしないで計算してしまえばいい。
後でホスト側に転送する際に省略するか、転送してから捨ててもいいし。
境界は別の計算をするという場合も同じようにしておいて、ホスト側で計算してしまえばいい。
後は、データ構造で巧く工夫する位かな。
837 :
デフォルトの名無しさん :2012/08/03(金) 08:46:34.92
さっそくの返信ありがとうございます。 x[i]+x[i-1]+x[i+1] のような計算をした際に、配列が範囲外(i=0やi=N-1)の時、ほかの変数を書き換えたり、 暴走したりしてしまいますか? そうでなければ、ホスト側で計算してしまうという方法もつかえるのですが。 それともなにかもっとよい方法があるのでしょうか?
頭堅いなぁ。 例えば、メモリ確保量を2増やしておいてx[i + 1] + x[i] + x[i + 2]にアクセスすればいいじゃん。 後はそれらの応用。
839 :
デフォルトの名無しさん :2012/08/03(金) 09:09:03.87
それだとN-1の時だめのようです
841 :
829 :2012/08/03(金) 10:45:26.55
>>830 で教えて頂いた方法を参考に
>>834 の
>カーネル関数の実行 → func<<< Dg, Db >>>(**dPtrArray);
をfunc<<< Dg, Db >>>(dPtrArray);
で実行したところ上手く参照できました。ありがとうございます!
ただ、警告に関しては最初と同じものが出ます。
(Warning: Cannot tell what pointer points to, assuming global memory space )
一応画像サイズ上限まで確認しましたが、値の参照については特に問題は見受けられません。
警告は無視してこのままでも大丈夫なのでしょうか?
>>839 >838に激しく同意。ループカウンタのiを1から回す代わりに0から回せってことでしょ。
これ。 >for (int i=1;i<N-1;i++) Cudaでループじゃなく並列になる部分だとしても、同じことでしょ。
845 :
デフォルトの名無しさん :2012/08/04(土) 15:38:05.57
>>840 このスレにもKPOPに乗り遅れた
在日だらけのバーとかエイベっ糞とかの
芸能ヤクザ工作員が湧いてくるんだな
早く祖国に帰れば
846 :
デフォルトの名無しさん :2012/08/04(土) 15:39:41.19
【芸能】紅白K-POPゼロは、「芸能界の重鎮」某社長の意向?…桑田佳祐・KARA、
紅白出るか出ないか、舞台裏攻防戦(週刊新潮)[12/9]
1 :諸君、私はニュースが好きだφφ ★:2010/12/09(木) 13:33:00 ID:???
■サプライズ無し!?
思わぬ逆風が吹いたのは、アミューズだけではない。
「当初、6枠と予想されたジャニーズ勢は、結局、嵐、NYC、SMAP、TOKIOの昨年と同じ4組で、
芸能生活30周年の近藤真彦(マッチ)は落選。これも、昨年4枠に急増したジャニーズの拡大路線を
警戒する業界に配慮した結果でしょうね」
芸能デスクが言う。 「また、NHKの井上啓輔プロデューサーは、16組の演歌枠を半減させ、
『少女時代』『KARA』『BIGBANG』などのK-POPについて男女4枠を確保する意向を示していたが、
蓋を開けてみると韓国勢はゼロでした。
実は、紅白出場歌手については、発表の数日前に芸能界の重鎮であるバーニングの周防郁雄社長と
相談するのが慣例化している。今年はその折に、“演歌枠半減なんてとんでもない。
そもそも全体の出場枠が減っているのに韓国人歌手を出すことはない"という意向が示されたというのです。
これに対して、K-POP勢の日本での面倒を見ているプロダクション尾木の尾木徹社長らは怒り心頭。
これまでバーニングに公然と反旗を翻す事務所はなかっただけに業界の注視を集めているんです」
ところで今年の紅白、このまま目玉もサプライズも無し、で終わるのか?あるNHK関係者によれば、
「KARAは30、31日のスケジュールが空いているので要注目でしょう。桑田は状況がこじれすぎですが、
マッチの方は何らかのサプライズがあるかも」
あまり期待せず、待とう。
ソース(週刊新潮 12/16号 46〜47ページ)
http://www.shinchosha.co.jp/shukanshincho/newest/
847 :
デフォルトの名無しさん :2012/08/04(土) 15:40:34.00
463 :名無しさん@お腹いっぱい。:2011/06/21(火) 20:11:37.69 昨年末に勃発した日本音楽事業者協会(音事協)の会長である尾木徹氏(プロダクション尾木社長)と ユニバ ーサルミュージックの顧問塩田真弘氏が K-POP利権を独り占めしていたことを、 芸能界のドンことバーニングプロダクションの周防郁雄社長が激怒した件は知る人ぞ知る話。 その手打ちの一環として、ユニバーサル所属になることが決まっていた芦田愛菜の音楽出版権を 「塩田氏が周防氏に譲渡したそうです。それで大ブレークするんですから、ドンの威光はすさまじいですよね」 (中堅芸能プロマネジャー)
848 :
デフォルトの名無しさん :2012/08/04(土) 16:03:48.61
サウンドスキャン(サ)とオリコン(オ)の初動売り上げ枚数比較
http://www.morcoff.uu.dnsdojo.net/morcoff/ga-ldr/img/mc1851.jpg サウンドスキャンランキング
http://www.phileweb.com/ranking/cd-top20/ ウィンターマジック KARA サ 72084 オ 78148
GO GO サマー! KARA サ 105213 オ 113873
ジェットコースターラブ KARA サ 111058 オ 122820
ジャンピン KARA サ 48229 オ 54977
ガールズトーク KARA サ 102605 オ 107403
スーパーガール KARA サ 265,697 オ 275,206
MR.TAXI / Run Devil Run 少女時代 サ 96153 オ 100461
Gee 少女時代 サ 59575 オ 66203
GIRLS' GENERATION 少女時代 サ 223,595 オ 231,553
Bo Peep Bo Peep T-ARA サ 48487 オ 49712
Let me cry チャン・グンソク サ 117170 オ 119149
849 :
デフォルトの名無しさん :2012/08/04(土) 16:04:22.30
>この「オリコン独占体制」を崩したのが、95年に米国からやってきた「サウンドスキャン」です。
その情報源は、バーコード読みとり機が集計した「POSデータ」。カネが支払われたCDの数だけを、
ホストコンピューターから取得してカウントする。大統領選挙の出口調査から生まれたこのシステム、
一店でも買い占めと思しき不自然な数字が集計されると、警告マークが出て自動的に統計から除外してしまう。
人為的な操作が入り込むすき間がないのです。米国の音楽チャート誌「ビルボード」は
91年からこのデータを使っています。
>サウンドスキャンとの比較で初めて、オリコンには独特の「癖」があることがわかりました。
(省略)レコード会社の看板スターがアルバムを発売すると、第1週のオリコンでの数字はなぜか
サウンドスキャンより高めに出る。2週目から降下し、だいたい20週間の総売上数で
POSデータと合致するのだそうです。1週目のオリコン売上数はレコード会社が
宣伝材料に使いますから、なんだか勘ぐりたくもなります。
>かつては甘めの数字で満足していたレコード会社も、今やサウンドスキャンを無視できなくなった。
98年をピークにCDの売り上げが急落する昨今、甘めのチャートではじいたCDを出荷したあげく、
膨大な返品の山を抱えることになったからです。限られた売り場にどの商品を仕入れるべきか
頭を悩ませるCD店もまたしかり。
(AERA 2003年02月03日)
サウンドスキャンランキング
http://www.phileweb.com/ranking/cd-top20/ このヒットチャートは「サウンドスキャン・ジャパン(SSJ)」システムの1週間の推定売上データをもとに
作成したものです。同システムは、全国約4100店のソフト取扱店とEコマースのPOSシステムで
集計された実売データから市場推定したものです。
発売以来累計は2010年1月4日以降からEコマース分が含まれています。
850 :
デフォルトの名無しさん :2012/08/04(土) 16:04:56.53
>サウンドスキャンジャパン
サウンドスキャンジャパンは、音楽・映像ソフトの売り上げランキングを集計・発表するサービスである。
株式会社エス・アイ・ピーが、アメリカのニールセン・サウンドスキャン社からライセンスを受けサービスを
提供している。
>ビルボード・ジャパン
株式会社エス・アイ・ピーの提供する、リアルストア、Eコマース約4,100店舗における初回盤、
限定盤、通常盤がそれぞれ別集計されている推定売上枚数を、商品タイトルごとに合算した
アルバムセールスチャートです。
http://www.musicman-net.com/business/14955.html 音楽ソフトパッケージの各種マーケティング情報を取り扱うサウンドスキャンジャパン(SSJ)が、「2011年オーディオ・ソフト売上動向」を発表した。
2011年の「オーディオ・ソフト売上」は、前年比91.7%の2320億円(2010年は前年比91.4%の2530億円)となり、依然、音楽市場は減少傾向を続けている。
順位 アーティスト名 売上(単位:百万円)
1 AKB4 7,370
2 嵐 4,201
3 EXILE 4,014
4 KARA 3,907
5 少女時代 3,406
6 安室奈美恵 2,236
7 関ジャニ∞[エイト] 1,994
8 EXILE ATSUSHI 1,977
9 aiko 1,960
10 B'z 1,840
韓国ハンファグループ日本法人のハンファ・ジャパンは2日、
総合商社の丸紅が建設を計画している日本全域の太陽光発電所に、
向こう4年間で約50万キロワット分の太陽光モジュールを供給することで
丸紅側と合意したと明らかにした。近く本契約を締結する。
供給されるモジュールは全てハンファソーラーワンの製品で、
売上額は6000億ウォン(約416億円)に達すると見込まれる。
日本への太陽光モジュールの供給が本格化したのは、
東日本大震災が発生した昨年3月以降だ。
ハンファグループは震災後、鳩山由紀夫元首相の支援要請に応じて
太陽光発電システムなど10億ウォン相当の支援を実施。
これを機に金升淵(キム・スンヨン)グループ会長が日本を訪れ、
野田佳彦首相や丸紅の朝田照男社長と面会し、太陽光発電事業での提携を協議した。
http://japanese.yonhapnews.co.kr/headline/2012/08/02/0200000000AJP20120802003300882.HTML >ハンファグループは震災後、鳩山由紀夫元首相の支援要請に応じて
>太陽光発電システムなど10億ウォン相当の支援を実施。
>これを機に金升淵(キム・スンヨン)グループ会長が日本を訪れ、
>野田佳彦首相や丸紅の朝田照男社長と面会し、太陽光発電事業での提携を協議した。
またネトウヨの陰謀論妄想か
C++とCUDAを連携させたプログラムを作成しているのですが、どうしても解決出来ない問題が発生してしまったので教えて下さい。 C++のスレに質問するか迷ったのですが、エラーにはCUDAがメインで関わっていると思ったのでこちらに書き込みます。 C++のプログラムからCUDAのプログラム(既にexeになっているもの)を実行したいと思い、下記のようなプログラムを作成したのですが、 CUDAプログラム単体では動いているのにC++のプログラムから呼び出すと実行時にエラーが発生してしまいます。 SDK内のサンプルプログラムでも実行エラーになってしまうので、CUDAプログラム側のバグでは無いと思うのですが… system関数では実行出来ないのでしょうか?system関数以外でも良いのでなんとか実行する方法が知りたいです。 int _tmain(int argc, _TCHAR* argv[]) { int ret; printf("CUDAプログラム実行開始\n"); ret = system("\"D:\\Program File\\Test_Print.exe\""); if(ret == 0) { printf("プログラム実行成功\n"); } else { printf("プログラム実行失敗\n"); } }
そのTest_Print.exeと同じ場所に何でもいいのでC++で作った何か*.exeを置いて同じことを試してみなよ。 俺があてずっぽうで言うとたぶん同じようにうまくいかないつまりCUDAがどうこうって問題じゃない。
>>854 回答ありがとうございます。
とりあえず引数表示するだけのプログラムで同じ場所にexeを置いて実行してみました。
このプログラムは実行出来ました。引数:XYZ のとき 結果表示:XYZ
int _tmain(int argc, _TCHAR* argv[])
{
printf("%S \n",argv[1]);
Sleep(1000);
return 0;
}
>>855 そのCUDAのプログラム、何の準備もしていないDOSプロンプトから動かしてみたら何が起きる?
恐らく、それが答え。
>>853 system関数のマニュアル見てエラーの値をしらべたら?
858 :
デフォルトの名無しさん :2012/08/07(火) 16:21:01.57
2次元のグリッドやブロックはどのようなときに使うのでしょうか? 2次元は配列と同じように、1次元よりもメモリのアクセス速度が落ちるのでしょうか?
>>856 何の準備もしていないとは、環境変数等がセットされていないという意味でしょうか?
>>857 調べてやってみます
>>859 DOSプロンプトを立ち上げて、そのまますぐに実行できるかどうか。
それでエラーが起きるようなら、system()の引き数で起動できるわけがない。
Program File(s)ってオチじゃね。
>>857 errnoにエラーが格納されるとsystem関数のマニュアルに書いてあったので、
errno.h をインクルードしてsystem関数の直後にperror("ERROR:");を挿入してみました。
しかし、「No error」と表示されてしまいます。systemからの返り値が0で無く実行も出来ていないので何らかのエラーが発生していると思うのですが…
表示方法が間違っているのでしょうか?
>>860 Windowsのアクセサリからコマンドプロンプトを起動してCUDAプログラムの場所まで移動して、
実行してみたところ問題無く動作しました。
>>861 フォルダの場所を再度確認してみましたが、問題ありませんでした。
>>862 フォルダの位置じゃなくて、ファイルパスの文字列の取り扱いの事だと思うぞ。
すみません。自己解決しました。 動かしていたCUDAプログラムは対象のファイルに対して計算を行うプログラムだったのですが、 その対象ファイルを置く場所を勘違いしていました。 CUDAのプログラムが置いてある場所ではなく、Visual Studioのプロジェクトフォルダーの方に計算対象ファイルを置かないといけないんですね。 お騒がせして申し訳ありませんでした。
>>864 最初によくハマるケースだなw
カレントディレクトリの意識、大事に。
良い経験になったね。
>>865 まさかこんな事が原因とは思いませんでした。。。この機会にカレントディレクトリについても改めて勉強しなおそうと思います。
皆様アドバイスありがとうございました。
何か酷いな。
TeslaK20は30万だそうだ。
超基本的な質問をしたい。 「はじめてのCUDAプログラミング」を読んで色々動かしてるんだが、 ブロックとスレッドとSMとSPのそれぞれ関係がいまいちモヤッとしてる。 最初から順に読んでいて、57 ページの次の説明で躓いた。 > スレッド当たりのレジスタ数が > 「SM当たりのレジスタ数 / ブロック内のスレッド数」 より多い時は、「カーネル関数」を実行できなくなります。 SM や レジスタは物理的なもので、ブロックやスレッドは論理的なものなのに、 なんでごちゃ混ぜになってるのか・・・ その少し後のページも見たり、直前の Warp の説明を読んで 俺なりに考えた結果、次のような認識なのだが、これは間違ってる? ・1つの SP は一度に1つのスレッドを処理する ・1つの SM は一度に1つのブロックを処理する
870 :
869 :2012/08/11(土) 20:48:28.79
もう少し先まで読んだら、63 ページに
> 実は、これらのリソース使用量から、
> 「SM」内で同時に実行されるブロックの数が決まります。
と書いてある。
ひとつの SM 内で複数のブロックが同時に処理されるのか。
完全に俺の考え(
>>896 )が間違ってる。
この辺り、小出しにしないでズバッと解説したサイトや本は無いですか?
> 870 つ"NVIDIA CUDA Programming Guide"
872 :
869 :2012/08/11(土) 21:12:39.61
あぁ、公式のガイドを当たれということね。 もし可能なら日本語の情報源があると大変ありがたいのだが・・・ がんばって NVIDIA CUDA Programming Guide 読んでみるよ。 ありがと
874 :
869 :2012/08/11(土) 21:43:56.41
>>873 ありがたい
公式ガイドと併せて読んでみる。
CUBLASのmapping errorについて質問なんですが、 To correct: prior to the function call, unbind any previously bound textures. ↑これは結局どうればいいのでしょうか? 前に確保したテクスチャの開放するとかfunction callにpriority与えるとか 具体的にどうしたら良いのかよくわかりません。 馬鹿な質問でごめんなさい
>>875 プログラマがすべき具体的な作業は知らんが、
「関数呼び出しの前に、それ以前バインドしたあらゆるテクスチャをアンバインドしろ」
と言ってるのだから、function callにpriority与える、は違うと思うよ。
次のコードを ptx にコンパイルした。 __global__ void kf (float *ods, float *ids) { int i = blockDim.x * blockIdx.x + threadIdx.x; ods[i] = ids[i] * 2; } cuModuleLoad 関数でこの ptx ファイルをロードするところまではできたから、 cuModuleGetFunction 関数を使ってカーネルのハンドラを得たいんだが、 第3引数に "kf" を指定しても、エラー値 CUDA_ERROR_NOT_FOUND が帰ってくる。 コンパイルした ptx をテキストエディタで見てみると、 エントリー名(カーネル名)が kf ではなく _Z2kfPfS_ になってた。 (関数のオーバーロードを実現するため?) 試しに第3引数に "_Z2kfPfS_" を指定すると問題なくハンドラが得られた。 これ、cu ファイルに書いたカーネルの名前 fk を、 そのまま ptx ファイルのエントリー名にする方法は無いの? あるいは、fk が _Z2kfPfS_ に変わる変換規則が分かれば、それでもいいけど。
878 :
877 :2012/08/12(日) 22:54:40.95
>>877 分かると思うけど、訂正
> これ、cu ファイルに書いたカーネルの名前 fk を、
> そのまま ptx ファイルのエントリー名にする方法は無いの?
> あるいは、fk が _Z2kfPfS_ に変わる変換規則が分かれば、それでもいいけど。
これ、cu ファイルに書いたカーネルの名前 kf を、
そのまま ptx ファイルのエントリー名にする方法は無いの?
あるいは、kf が _Z2kfPfS_ に変わる変換規則が分かれば、それでもいいけど。
_Z
881 :
877 :2012/08/12(日) 23:00:45.44
すまん もしやと思ってカーネル関数に extern "C" をつけたら、あっさり解決した。
CUDA Programming Guide の Appendix G を参考に Driver API を使ってみた。 cuda のライブラリを cuda.lib 使って静的リンクした時はOKだけど、 LoadLobrary + GetProcAddress で nvcuda.dll 使って動的リンクした時はエラーが出る。 下記の関数を nvcuda.dll からロードして関数ポインタにした。 cuInit cuDeviceGetCount cuDeviceGet cuCtxCreate cuCtxDetach cuModuleLoad cuMemAlloc cuMemFree cuMemcpyHtoD cuMemcpyDtoH cuModuleGetFunction cuLaunchKernel 初期化して、コンテキストを作って・・・と進んで、 cuLaunchKernel でカーネル関数を実行したところまでは問題ない。 各関数の戻り値は正常。 で、計算結果をホスト側に戻そうと、cuMemcpyDtoH を読んだんだが、 戻り値が CUDA_ERROR_UNKNOWN になる。 (当然、ホスト側のメモリには何もコピーされない) 全く同じコードを、LoadLobrary + GetProcAddress を省いて、 cuda.lib を使うように切り替えると、全く問題なく結果がホスト側にコピーされる。 自分の書いたコードが悪いのかと思って、試しに $(NVSDKCOMPUTE_ROOT)\C\common\inc\cuda_drvapi_dynlink.c をインクルードしてみたが、 こちらでも同じところで同じエラーが出た。 動的リンクに成功した人いる? <環境> Windows7 64bit CUDA Driver & Runtime 4.2 Compute Capability 3.0
883 :
882 :2012/08/13(月) 13:09:00.00
どうも、cuda_drvapi_dynlink.c を利用した場合に限れば、 64ビットをターゲットにしてコンパイルする事で正常になった。 32ビットをターゲットにしてコンパイルするとエラー。 しかし、nvcuda.dll は System32 フォルダにも SySWOW64 フォルダにもあるわけだが、 自分の作ったラッパーだとどちらを明示的にパス指定してロードしても、エラーになる。 64ビットOS用の CUDA SDK やツールキットを入れて、 32ビットをターゲットにして nvcuda.dll を動的ロードするプログラムを コンパイルすることは無理なのかなぁ・・・
884 :
デフォルトの名無しさん :2012/08/15(水) 06:59:58.41
885 :
882 :2012/08/15(水) 12:00:15.21
OS 毎全て再インストールしたら、同じコードで、同じコンパイル方法で エラーもなく正しい結果が得られた。 訳分からんが、とにかくこれで実験がつづけられる。
886 :
869 :2012/08/15(水) 23:02:20.22
>>884 ありがと。
>>873 のページでだいぶ分かってきた。
メモリアクセスで待っているSPを遊ばせないようにという目的が分かると、
ブロックやスレッドなどの論理的な構造がある理由や、
SM や SP などの物理的な構造との関係が分かってくる。
>>884 のテキストでちゃんと理解できているか復習してみるよ。
887 :
デフォルトの名無しさん :2012/08/17(金) 18:51:51.24
韓国経済が急激に失速している。一番の原因は、これまで韓国経済を引っ張ってきた原動力である輸出に陰りが見え始めたためだ。
輸出に次いで内需の鈍化も懸念されている。住宅価格の下落とそれに伴う消費や投資の抑制がみられ、バブル崩壊以降日本が苦しんだ
「日本型長期デフレ」の兆候が現れている、との指摘も少なくない。
■好調に見えたのはヒュンダイやサムスンだけ
韓国の輸出は、2012年7月の通関ベースで前年同月に比べて8.8%減と大きく減った。マイナス幅は3年ぶりの高い数値だという。
これまで韓国の輸出をけん引してきた自動車の輸出が頭打ちになったのをはじめ、船舶や石油化学製品、携帯電話など主力製品の輸出が急速に落ち込んだ。
韓国の輸出額は国内総生産(GDP)対比で50%を超える。「輸出国」といわれる日本でもGDP対比では10%半ばだから、輸出依存度の高さは圧倒的だ。
そのため、輸出の不振は即韓国経済の失速に直結する。
なかでも韓国経済を支えてきたのが欧州連合(EU)向けの輸出。EUとは自由貿易協定(FTA)を結んでいる。そのEU向けが12年1〜6月期には前年同期に
比べて16.0%も減った。EU諸国の債務危機から発した景気低迷が影響した。
さらには中国向けも1.2%減った。第一生命経済研究所経済調査部の主任エコノミスト、西?徹氏は、「中国向け輸出の減少はボディブローのように効いています」と話す
。韓国の素材や部品メーカーは中国を介して、間接的にEU向け輸出を増やしてきたからだ。
http://headlines.yahoo.co.jp/hl?a=20120817-00000006-jct-bus_all
889 :
KPOPに乗り遅れた、在日が多くいる芸能ヤクザ工作員にはご注意を :2012/08/18(土) 02:07:28.32
もう、一斉にテレビ局の命綱といわれる 韓国ドラマをやめて、BS・CS畳んで テレビ局の給料体系や芸能人のギャラ体系が維持できなくなるから いったんキー局・在阪局・ローカル局を含めたテレビ局大リストラ 芸能人のギャラも大幅カットで テレビ業界・芸能界が一から出直した方がいいな しょせん、韓国ドラマをネット配信に取られたくないからという 批判しながらのダブルスタンダードが甘いんだよ もう、何もかもうんざりだわ
シェアードメモリって例えば48KBあったとして一つのブロックで48KB使えるものなの? それともカーネル全体で48KB? ブロック内でのデータのやり取りが出来るって言うから前者かと思ってたんだが、うまく動かなかった。 Streaming Multiprocessorごとに48KBだから48KB÷(ブロックの数÷SMの数)がひとつのブロックのシェアードメモリ使用量の限界なのかな?
ディフォルトは16KiBで、APIで切り替えないと48KiB使えないんじゃなかったっけ? それと、端から端まではフルに使えなかったかも。その場合は、コンパイルオプションで確認できたかと。
分っているかもしれないが、カーネル関数の引数で使う分も加えることを忘れずに
ちなみに、ホンオフェは魚を発酵させたもの。 ジュンタクは魚を牛糞に漬け込んだもの。 ホンタクは魚を人糞に漬け込んだもの。 ホンタクがもっとも高価で美味しく、珍味とされる。 韓国には昔から食糞の文化があり、体の調子をみるために糞を舐めたりしていた。 甘いと病気で苦いと健康とされていた。 両親の糞が甘かったので病気だとショックを受けたが、あまりの糞の甘さにオオ皿山盛りの糞を平らげて腹いっぱいで動けなくなったという、 ほほえましい昔話も韓国では伝えられている。
895 :
KPOPに乗り遅れた、在日が多くいる芸能ヤクザ工作員にはご注意を :2012/08/19(日) 21:25:35.26
>>894 嘘つくなボケ
2009-09-20
トンスル、その他の漢方薬関連の無知を基にした嫌韓差別都市伝説の考察
朝鮮史, 日本史, 嫌韓バカ
最近(でもないか)の嫌韓バカのお気に入りは「韓国人は食糞が文化」などというデマ。
年間250万人の旅行客が日韓間を行き来するこのご時世*1に、
そんなデマを信じる奴がいるのかねぇ・・・。とか思うのだが、いるようである。
今回はホンタク、トンスル、野人乾水、について調べてみた。
(珍説にも記載→珍説24・人糞を食するのは朝鮮の文化|誰かの妄想)
http://d.hatena.ne.jp/scopedog/20090920/1253462347
896 :
KPOPに乗り遅れた、在日が多くいる芸能ヤクザ工作員にはご注意を :2012/08/19(日) 21:26:14.99
ホンタク
以前は、ガンギエイという「エイの切り身を人糞に漬け込んで発酵させた料理」*2
というデマが流されていた。
さすがに最近は、まともな情報が出てきて嫌韓バカの主張がデマであることが
知られてきた*3。
で、ホンタクとは、ホンオフェ(洪魚膾・エイ刺身)とタクチュ(濁酒・マッコリ)の
組み合わせのこと。
ホンオフェはテレビでも臭い食べ物として紹介され*4よく知られるようになったためか、
なかなか嫌韓バカのデマは通じにくくなったのだろう*5。
さて、人糞説というデマがなぜ発生したかと言うと、ホンオフェの作り方に起因する。
ホンオフェはエイの刺身を醗酵させて作るのだが、その醗酵のために堆肥の熱を
利用している。最近は刺身を入れた壷を堆肥の中に入れるようだが、昔は藁に包んで
堆肥に入れていたようだ。つまり納豆と似た製法になっている。
ちなみに堆肥は人糞、牛糞などを使っているので同じことだ、と主張する嫌韓バカもいるが、
ホンタク用の堆肥は主に松葉や米糠を使っている*6。堆肥は排泄物を使うとは限らない。
さてこの堆肥だが、韓国語では「トン?」と言い漢字にすると「糞」になる。つまり韓国語の
「トン?(糞)」には、日本語と同じく「糞」という意味もあるが「堆肥」という意味も
あるわけだ*7。
と言うより「糞」という漢字の字源は「畑にばら撒くもの」なので、意味としてはむしろ
「堆肥」の方が近かったりする。
もう分かるだろうが、韓国語で「トン?」と書かれた堆肥のことを人糞と勘違いした
どっかのバカが広めたのが、人糞説の発端だったわけ。
嫌韓バカは韓国をバカにするための格好の材料として、このネタに飛びつき、ろくに検証も
せずにネット上に広めて回った。と、こういうことだろう。
参照
http://blog.goo.ne.jp/mpac/e/4aba204bd01800448deb3f509a492d75
897 :
KPOPに乗り遅れた、在日が多くいる芸能ヤクザ工作員にはご注意を :2012/08/19(日) 21:27:22.02
トンスル こちらはホンタクより最近のネタ。「スル」は酒のこと。「トン-スル」で想像つくだろうが 「糞酒」という薬のこと。嫌韓バカが触れ回っている”大便をふんだんに入れた酒” というのはデマ。 入り口を塞いだ竹筒を大便の中に浸し長時間置き、中に滲み出てきた液体をとって 酒に入れたものが「トンスル」である*8。 液体の成分が何なのかはよく分からないが、大便由来なのは間違いなかろう。 ここで、やはり韓国人は食糞の文化を持っている!とはしゃぐのは早計。 漢方に詳しい人は気づいたかもしれないが、この液体をとるまでの製法は「人中黄」 という漢方薬とほとんど同じなのだ。李時珍の「本草綱目」で「糞清」と呼ばれるものが この液体であろう*9。 「人中黄」については、トリビアの泉でも紹介されたので有名だと思うが、このように 説明されている。 No.117 漢方薬には人間のウンコから出来たものがある 〜確認VTR〜 ・人間のウンコから出来た漢方薬があるという 富山医科薬科大学 薬学博士 小松かよ子さんはこう語る Q.人間のウンコから出来た漢方薬がある? 「確かにあります。人糞を加工した薬を解熱、解毒剤にしています」 ・人間のウンコから出来た漢方薬「人中黄」 何故そんなものが体に効くのか? 人中黄にさらに詳しいという難波先生を紹介して頂いた。 ・ウンコから出来た漢方薬に詳しい 難波恒雄先生 Q.何故ウンコから出来た漢方薬が体に効く? 「これは1500年程前からずっと使ってきた経験で良いという事が分かっているので 本当に良いかどうかというのはまだ科学的には解明されていない。 我々は毎日お米を食べてる。じゃあお米の何が効いているのかというのと一緒ですよね、 分からない」
898 :
KPOPに乗り遅れた、在日が多くいる芸能ヤクザ工作員にはご注意を :2012/08/19(日) 21:27:56.47
〜補足トリビア〜
1)「人中黄」は「甘草」という薬草の粉末を人糞に浸して作られた漢方薬
2)VTR中にあった「人中黄」は難波先生が中国のウイグルで手に入れたもので無臭である。
3)漢方薬の中には「人中白」という小便の固まりを日干しにして作られる薬もある。
http://s.freepe.com/std.cgi?id=mudachishiki&pn=18 日本人はそんな漢方薬を使っていない、とか言い出すバカもいるかも知れないので、
日本の医学書を調べてみた。
江戸時代の医学書「用薬須知」の6巻(人ノ部)には「人中黄」が載っており、
「大便ノ汁ナリ」という説明がついている。
製法については、「用薬須知続編」の3巻に載っており、大竹を切って一方の節を残し、
大甘草を入れて円形の木片でこれを塞いで、隙間は蝋で埋める。
これを糞の中に入れて一ヶ月おいて取り出し、乾かす。こうして出来たものが
人中黄、とある。
「トンスル」の場合、大甘草に滲みこませて乾燥させるのではなく、
竹筒内にしみでた液体を酒に入れて飲むという方法を取っているだけで、
基本的には人中黄と同じ成分を取っていると言える。
ちなみに、李時珍の「本草綱目」には糞の薬としての使い方がたくさん載っており面白い。
以下のサイトに現代語訳と共に乗っているので興味のある人はどうぞ。
参照:
http://mayanagi.hum.ibaraki.ac.jp/LecRep/04/TextInterp2.html
899 :
KPOPに乗り遅れた、在日が多くいる芸能ヤクザ工作員にはご注意を :2012/08/19(日) 21:28:32.56
野人乾水
朝鮮国王が野人乾水という人糞を使った薬を服用したという記録がある。
中宗(チュンジョン)は高熱が出た際、解熱剤として「野人乾水」を食べた。
「野人乾」というのは人糞だ。ドラマ『大長今(宮廷女官チャングムの誓い)』の背景になった
中宗後半から景宗(キョンジョン)のときまでは御医が明の医学の影響を受けて
王の疾病治療の原因究明を重視するようにしたほか、特に壬辰の乱(文禄の役)後、
鍼灸術が発達し、内医院での治療に広く使われた。
http://japanese.joins.com/article/article.php?aid=66867&servcode=400&code=400 これ自体は事実だが、前述の通り、江戸時代日本でも人糞を使った薬を用いているので、
だから?としか言いようがない。
ちなみにこの「野人乾水」の製法も「用薬須知続編」に載っている*10。
野人乾水 後編獣部ニ出ス所ノモノト別ナリ 此養鶏ノ方ニ出ツ 道三家ニ野人乾水ハ
竹筒ニ糞水ヲ盛リ泥水中ニ置クコト久ヲ経テ中清水トナルヲ取用ユト
乾燥した人糞を水に溶き長期間置いて透明になった上澄みが「野人乾水」。
日本にもあった漢方薬で「本草綱目」にも似たような記載があります。
伸びているから何かと思えば……
901 :
KPOPに乗り遅れた、在日が多くいる芸能ヤクザ工作員にはご注意を :2012/08/19(日) 22:26:06.85
用薬須知続編に載ってる変な薬の事例 「本草綱目」とそう変わらないだろうが、日本でも現在の感覚から理解できないものを 薬として使っていたことを知っておくのは無益ではなかろう。 「用薬須知続編」2巻にある薬 「尿滓」(小便ノカスナリ) 「男子屎尖」(男ノ屎ノ尖リノ端) 「熱糞堆」(アツキ人ノ糞ノ重ナリタルモノ) 「焼人糞」(ヤキタル人ノ糞) ここまでが人の排泄物関連。 「頭垢」(アタマノフケ) 「交余」(交接ノ時拭タル布ナリ 一名精余ト云フ) フケとかナニとか。 「人龍」(人が吐き出した回虫のこと) 寄生虫まで薬にするってのはすごいなぁ、と思う。あと、名前だけ見るとかっこいい。 まとめ そもそも日中韓は同じような文化圏に属するわけで、韓国に変な文化があるぞ、 と笑っていると、日本にも似たような文化があったりする。笑った人は、 相手国の文化を理解できないのみならず自国の文化すら理解していない 無知をさらけ出しているわけで、結局は自分が笑われることになる。教訓じみてるなぁ。 時期の問題とか言いそうな人に対するおまけ 戦後、日本に進駐した米軍兵士が日本野菜をサラダにして食べたところ、 寄生虫に感染するものが多数発生したため、マッカーサーは激怒して、 「「寄生虫がいるなんてなんて不潔な国だ!今すぐ人糞肥料の使用を中止させろ!」と 日本政府に迫ったそうです。慌てた日本政府は「寄生虫予防会」を各市町村に作り、 人糞肥料から化学肥料へと一大転換が行われたのです。」
902 :
KPOPに乗り遅れた、在日が多くいる芸能ヤクザ工作員にはご注意を :2012/08/19(日) 22:28:35.26
慶尚北道金泉で生まれ育った小説家、金衍洙(キム・ヨンス)氏は、25歳のときガンギエイを初めて見た。
ソウル・仁寺洞の居酒屋で先輩の詩人に勧められるままに、ピンク色の魚を一切れ、何気なく食べてみた。
その瞬間、彼は「全校生が使用する便所を丸ごと飲み込んだような感じ」に圧倒された。
アンモニア臭が口から3日間消えなかった。
金氏は「人間はなぜ、このような食べ物を食べなければならないのか」と考えた。
人生の中でそれは「死ぬと分かっていながら死に向かって疾走する欲望」だった。
そのため、金氏はガンギエイを「大人の食べ物」と呼んだ。
ガンギエイを発酵させる過程を見ると、金衍洙氏が便所を連想したのも無理はない。
つぼに石を入れて、その上にワラを敷き、ガンギエイを置いて、繰り返して重ねていく。
そしてつぼの口をしっかり閉めて、暗い倉庫に置いておく。実際には腐らせるのと変わらない。
5、6日後につぼを開けると、息が詰まるほどのアンモニア臭があふれ出てくる。
昔、寒くてガンギエイがあまり発酵しない冬には、堆肥を腐らせる堆肥置き場にガンギエイを置いておくこともあったという。
それでも、ガンギエイが腐らずに発酵するのは、体内に多くの尿素があるからだ。
ガンギエイを発酵させると尿素が分解されてアンモニアが発生する。
タンパク質が腐敗してアンモニアを発生させるのとは異なる。アンモニアは臭いがきつい毒性物質だが、
ガンギエイから発生するアンモニアは体に害を及ぼすレベルではない。
むしろ、ガンギエイを強アルカリ性にし、腐敗細菌や食中毒細菌の繁殖を抑える。
キムチやチーズのような発酵食品と同様に、ガンギエイの味を覚えると、病みつきになる。
口の中や舌の皮がむけるほどピリッとする味に魅せられてしまう。
(つづく)
http://www.chosunonline.com/site/data/html_dir/2012/08/18/2012081800546.html http://www.chosunonline.com/site/data/img_dir/2012/08/18/2012081800530_0.jpg
韓国の新聞社がちゃんと証言してるんですけど? #昔、寒くてガンギエイがあまり発酵しない冬には、堆肥を腐らせる堆肥置き場にガンギエイを置いておくこともあったという。
905 :
KPOPに乗り遅れた、在日が多くいる芸能ヤクザ工作員にはご注意を :2012/08/20(月) 00:25:33.44
906 :
KPOPに乗り遅れた、在日が多くいる芸能ヤクザ工作員にはご注意を :2012/08/20(月) 00:29:39.30
韓流依存の低いテレ朝でも韓国ドラマなくなったら
倒産するらしい
他も推して知るべしといったところだな
テレビ局総倒れということだな
422+1 :名無しさん@お腹いっぱい。 [sage] :2012/08/18(土) 19:48:49.96
>>421 そんな事をやったら体力の無いテレ朝は潰れますからw
テレ朝の親玉・朝日新聞だってTBSと組んで韓流やってるのにw
新聞業界も将来どうなるか分からないし朝日新聞も潰れない保証は無い
もう、命綱といわれる韓国ドラマを一切やめて
一度全局倒産したうえですべてリセットした方がいいわ
批判しながら韓国ドラマはネット配信に取られたくないとかいう
ダブルスタンダードなんか甘いの一言なんだよ
もう、何もかもうんざりだ
テレビ局の給料なんて
持ち株会社化する前のフジテレビの
全社員の給料分でさえ
一日20時間平均で200万円台程度なんだから
こんなはした金じゃ給料カットしても
在日だらけの日本のテレビドラマの
放映権ぶんの足しにもならない
だいたい芸能人のギャラが高すぎるんだよ
テレビ局員の給料も芸能人のギャラもすべてリセットだ
,,,
( ゚д゚)つ┃
君たちはこのスレが何のスレかも分からないの?
なんのスレか以前に 日本最大の右翼掲示板ですから
ネトウヨの陰謀論掲示板
ひとつのスレッドがメモリアクセスする命令を複数回連続して発行したとします。 (それぞれのメモリアドレスは別々です) この場合、一つ目のメモリアクセス命令が完了するを待つ間に、 同じスレッドの次のメモリアクセス命令が実行を開始するなんてことはあるのでしょうか。 それとも、先のメモリアクセス命令が完了するまで、 次のメモリアクセス命令は実行されずにスレッド自体が止まるのでしょうか。 この辺りの挙動はグローバルメモリへのアクセスでも、 シェアードメモリへのアクセスでも同じでしょうか。
>>910 シェアードだろうがグローバルだろうが、両方ともあり得る
>>911 ということは、シェアードメモリにアクセスする命令を複数回連続して発行した時も、
バンク・コンフリクトは起こり得るのですか。
つまり最適化のためには、並列に実行される複数のスレッドからのアクセスだけでなく、
ひとつのスレッドからアクセスするシェアードメモリについても、
バンク・コンフリクトを意識しないといけないのですね。
まぁ要するに韓国が絡むとどこでも気持ち悪くなるという話だな。 スレ違いの話をいつまでも続けるんじゃねぇよカスが。
ネトウヨの妄想陰謀論は他でやれ
915 :
川田裕美とかいう基地外女狐アナがいるYTVには注意を :2012/08/23(木) 00:15:45.88
>>912 読みまちがえた、1つのスレッドが連続してメモリを読み書きする場合か。
それならシリアルに実行される
CPU→GPUのデータロードのボトルネック解消(?)のように思えるけど、 ボトルネックになってるのは大抵GPU内の計算かVRAM←→GPUのアクセスじゃなかったっけ?
つまりOut-Of-Orderのメモリ転送ってことなのかな? QってQueueを意味してたり。
1本当たりの帯域はへるけど最大32本までメモリ転送が平行でできるってこと のように見える。データ待ちを減らしてMPIのような全体協調アプリの性能を上げる というようなものかな
ひょっとして律速がGPU内の計算かVRAMアクセスかとかいうレベルじゃなくて ばんばん主記憶からデータを持ってこないといけないようなGPGPUプログラムに対する 向上策って感じなのかな?
925 :
924 :2012/08/27(月) 02:41:50.83
例えばDSPの代用みたいにリアルタイム処理で使うなら頻繁にCPU←→GPUがあるだろうから効果あるのかも。
いままで、非同期処理にストリームを使っていたところ、 結局ジョブのキューが1個しかないからいろいろ制限がかかってた。 それを、キューをコア数分だけ用意することによって・・・ という話かと考えた
927 :
ぼんじん :2012/08/28(火) 17:52:18.51
Geforce580からGeforce690に変更し, SDKのサンプルを動かしたのですが,動作が遅いです. そもそもCPUの動作も以前より格段に悪いです. なぜかわかる方いたら教えていただけないでしょうか? ちなみに ドライバーやその他の設定は580のときのままです. また,cuda4.0を使用しOSはwin7 64bitです.
>>927 CUDAの性能に限って言えば580が一番高いんじゃないの?
6xxシリーズは単精度の浮動小数点演算以外の性能は削りまくってる
みたいだし。命令ごとのクロック数の表をみたら一目瞭然でしょ。
自分は580 SLIだけど、Kepler 2が一般消費者向けに販売されるまで様子見してる。
>>929 マジかよ… なんじゃこりゃ…
I32加算:8/7、1、1
論理演算:24/17、1、1
>>929 うわぁ・・・。
ここまでヒドいとは・・・。
よく見たら680/690のGPGPU性能が芳しくないという話はこのスレでも出ているな。
>>635 には同じリンクが紹介されてる。
これ大事だから次スレのテンプレに入れておいたほうがいいな。
Kepler2だとこれ改善されるの? ↓ I32加算:8/7、1、1 論理演算:24/17、1、1
>>931 っつか、団子さんがCUDA関連のスレにいることに驚いたw
>>935 グラフィックに不要な部分をばっさり切ったのがkepler1なので当然改善するでしょう。
問題はマス向けのお値段でkepler2が出るのかどうか心配な点。。
939 :
ぼんじん :2012/08/29(水) 14:04:13.81
昨日の続きです. GTX690に変更したところ,580で動作していたプログラムが 動作しない(途中で落ちるor動作が鈍い)くなったんですが なぜだかわかりますか?? CUDA4.0とGTX690の相性が悪いのでしょうか??
CUDA 4.0で普通に動いていたプログラムが CUDA 4.2にしただけで速度が1/5ぐらいになったぞ。 ひょっとしてCUDA 4.2って地雷?
俺も4.0から4.2にした時に、 コンパイル時のメッセージで 「とりあえず浮動小数点は全部倍精度にしとくね!」 みたいなメッセージが出てた気がするけど 単精度メインの人だと処理速度に影響出たりするのかな? あんまり検証してなくてスマン
ほとんど整数演算だから一般的な使い方とは違うけど、それにしても遅すぎ! 4.0に戻して検証してみるよ。
遅くなったと思ったのは勘違いだった。スマソ。 でも4.0で何ら問題なくコンパイルできたプログラムが 4.2だとReleaseにするとコンパイラが落ちてビルドできなかったりと 挙動が不審すぎる。4.1でも4.2と変わらない。 …よくみたらマクロ使いまくった変態的プログラミングをしていた箇所で ciccが落ちてた。この箇所書きなおさないといけないなorz
944 :
ぼんじん :2012/08/29(水) 17:58:34.05
939の問題なのですが解決しました。 ただの接触不良みたいでした. もうひとつ質問なのですが, GPUに重い処理をさせたらOSごと落ちることありますよね?? あれってなぜなるのですが,また,GPUの性能とどのようにかかわりがあるのですか? すみません。よろしくお願いします。
熱排気がダメで高温落ちじゃないの
電力オーバーとか
947 :
ぼんじん :2012/08/29(水) 20:53:17.64
なるほど... 分かりました!! 調べてみます。 ありがとうございました。
やっぱりまだまだGPUプログラミングは厄介だね・・・
>>944 windows7だとGPU処理が一定時間返ってこないとディスプレイ落ちるが..
それとは違うんだな?
>>947 Windowsはドライバの応答がなくなると固まる
10秒以上かかりそうな処理は適当な間隔で戻してループしないとダメなようだ
951 :
デフォルトの名無しさん :2012/08/31(金) 03:37:51.28
その制限、解除出来るけどね
>>951 できねーよ
できるならここに詳細を書いてみろ(プゲラ
955 :
デフォルトの名無しさん :2012/08/31(金) 16:09:02.05
nVIDIAが出してるツール使えばレジストリエディタさえ使わずに制限解除できる。
ク ク || プ // ス ク ス | | │ // / ス | | ッ // ク ク ||. プ // / // ス ク ス _ | | │ // / ̄ ̄\ / ス ─ | | ッ // / _ノ .\ / // | ( >)(<) ____ . | ⌒(__人__) ./ ⌒ ⌒\ | ` Y⌒l / (>) (<)\ . | . 人__ ヽ / ::::::⌒(__人__)⌒ \ ヽ }| | | ` Y⌒ l__ | ヽ ノ、| | \ 人_ ヽ / . /^l / / ,─l ヽ \
竹島問題で 韓ドラとK−POPが消える!?
竹島問題を巡る日本と韓国の対立はテレビ界でも大きな影響が出始めている。韓流ドラマとK―POPが近いうちに
消滅するのではと囁かれているのだ。
この3年でテレビで放送された韓流ドラマは約500番組、日本でデビューしたK―POPのグループは30組以上
にも上る。テレビ局にとってはドル箱といっていいコンテンツだった。
「例えば、韓ドラはテレビ局には安く、そこそこ視聴率が稼げる番組として重宝がられました。買い付け価格は一時上がり
ましたが、最近はダウンし、初回購入の一番組の単価が500万円程度というものもあった。しかも、3回まで再放送が
認められるのでBSやCSが飛びついていた。一方、K―POPは音楽番組に出るのはプロモーションの一環で出演料が
かからなかったり、レコード会社が立て替えて支払ってくれるケースも多かったのでメリットが大きかった」(事情通)
だが、例の問題で状況が一変した。すでに日本テレビやテレビ朝日、TBSなどは年末年始の番組編成から韓ドラ枠を
大削減する方針を固めたという。音楽番組もK―POPのゲスト出演を極力減らしていく方針だという。また、
韓国びいきと批判されているフジテレビでさえも韓ドラからの撤退、K―POPの露出を減らしているという。
「李明博大統領が竹島上陸、ロンドン五輪でサッカー韓国代表メンバーが『独島は我が領土』というメッセージを掲げ、
韓流スターらが竹島に泳いで渡るという反日的ニュースが流れてから、BSとCSで視聴者の接触率が激減しました。
竹島問題で韓国の反日的なパフォーマンスが続く限り、来年4月の番組編成で韓ドラとK―POPは地上波から消滅する
可能性が高い。NHKも韓ドラをかなり流していますが、右に倣えだと思います」(民放編成マン)
テレビ局にとっては韓ドラに代わるコンテンツが見つからないのは頭が痛いが、K―POPに代わるダンスミュージック
を発掘中とか。KARAもチャン・グンソクも来春以降は見ることができなくなる?
http://gendai.net/articles/view/geino/138538
>>959 北チョン芸能ヤクザとべったりともっぱらの
日刊ゲンダイの記事かよ
北チョン芸能ヤクザ自身が書いているともいわれているな
>>957 北チョン芸能ヤクザとべったりともっぱらの
日刊ゲンダイの記事かよ
北チョン芸能ヤクザ自身が書いているともいわれているな
このスレは何故かネトウヨの陰謀論がよく沸くね
num = 0; for(i = 0;i < 100;i++) { if(条件) { list[num++] = i; } } のような感じである条件をみたすものだけを リストに追加、その個数をカウントしたいのですが CUDAで実装可能でしょうか。
可能は可能。極論言えばそのループを1スレッドで回せばいい 効率よくできるかどうかは詳細部分に依存する。
苦手な部類だな。 CPUの方がよっぽど手っ取り早くやってのけられる。 仮にiの最大値が100なんてしょぼい数じゃなく巨大なら、 スレッド数で分割してスレッドごとにlistを作って後で繋ぐこともできるけど。
リストに順に詰めていくことに意味があるのなら、それはつまり
i==nのときの処理の後でないとi==n+1のときの処理を行えないことを意味するので、
並列化に向いていない。
100じゃなくてもっと大きい数でなおかつ格納順を気にしなくていいなら
>>964 の言ってるとおり。
あとカウントはAtomicAddとかを使うよりトーナメント方式(?あの隣同士をlogN段階に分けて
足すやつ)のほうが早かった気がする。
966 :
962 :2012/09/11(火) 00:28:33.85
ありがとうございます。順序は関係ないので
>>964 さんの方法で分割しようと思います。
ただ、listを繋ぐために要素をコピーしていくのも
時間がかかりそうです(自分の理解が浅いかもしれませんが)
CPUが手っ取り早いかもしれません。
条件の真がint の1になるなら分岐せずに = (条件) * i にするとか
順序が関係ないなら、いろいろやりようがあるな。 メモリに余裕があれば、条件が真か偽かで1/0を入れるところまで 馬鹿並列で計算してしまって、それからリダクションみたいな操作 (たぶん上で言ってるトーナメント方式ってのと一緒)で詰めるとか
969 :
965 :2012/09/12(水) 02:35:53.57
>>962 prefix scanを使えばきれいに並列化できる
thrustのcopy_ifを参照
時代はOpenCLじゃないのか? 俺様のみのCUDAって先がないって感じじゃないか
事実上CUDAが標準じゃないの
openclの時代は来るだろうし誰も期待してない訳ではなかろ ただまだcuda脅かす感じじゃないよな 先はどうなんだろ、まだ見ぬkepler2の動的並列化が標準化される頃にはまた状況変わってるだろうしなあ xeon phiとか出てくるし
らくちんにプログラムしたけりゃCPUで、性能欲しけりゃCUDAで、って感じでOpenCLはその中間にある。 中間とはいっても、だいぶCUDA側に寄ってるわけで、 CPUプログラミングからわざわざやってくる人は性能を求めてるわけだからOpenCLを すっ飛ばしてCUDAやるんじゃないかっていう。 しかしコンシューマ用のKeplerを見る限りはなんかnVidiaは安心してサボってるような印象は受ける。 Geforce 7xxシリーズもGPGPUとしては使いにくいものになりそうなので MaxwellでFermiの上位互換的なものがくることに期待。 でもたぶんこない。コンシューマ向けのGPGPU体験/普及キャンペーンはFermiで終了したっぽい。
>>974 7xxはFermiの「上位互換」だと思ってたけど…
コアはTesla K20と同じGK110でしょ?
性能がほしいからGPU使うのにOpenCL使って性能でないとか馬鹿すぎるので みんなCUDA使ってるよ
とんがった用途に汎用環境とか、根本的に矛盾してるよね。 Radeon/ATI stream どうなった?
Occupancyを上げるためにPTXでコードを書き直してレジスタ数を削ったり、 48Kしかない共有メモリの使い方をああでもないこうでもないと工夫しているというのに、 NVIDIAのカードでOpenCLを使うなんて考えられない。
CUDAってFORTRANみたいな感じになって、デスクトップのiGPU付きCPUでGPUは 汎用環境用に進化していくんじゃないか。 デスクトップdGPUに重要なのはゲーム、それ以外(GPGPU用)は あんまり重要ではないから最新のデスクトップ用ケプラーでは GPGPUを軽視したものにしたんだろ。 今後はとんがった用途にデスクトップ用GPUなんて使うな、 GPGPU用GPUを使えになるんじゃないのか。 デスクトップPC,iGPUの性能向上でゲーム以外の用途では あんまりVGA要らなくなってきているってのが現状だし
APUはどこへ行くんだろうか
982 :
974 :2012/09/21(金) 01:37:47.43
>>975 GK110をそのままスケールダウンしたものがミドルやそれ以下にくればいいけど、
7xxの最上位以外はGK104/106/107のマイナーチェンジなんじゃないかな?
4xx→5xxがそうであったように。
それだとTeslaの名前が変わっただけで俺のような下々の人間にとっちゃ何も変わらんのよ。
寂しいけど
>>980 の言うようにこうなっちゃいそうなんだよね。
>今後はとんがった用途にデスクトップ用GPUなんて使うな、
>GPGPU用GPUを使えになるんじゃないのか。
>>982 個人的にはGK110が780に来てくれさえすればいいんだけど、
別にローエンドにこないと決まったわけじゃないでしょ。
とんがった用途って具体的に何を使って何をしているんだ?
今発売中のgeforceの中では、どれが一番GPGPUに適してます? 6xxシリーズはダメっぽいって感じでいいのかな?
>>985 でも単精度浮動小数点数だけでいいっていうんなら680/690もあり。
Tesla K10なんてのもあるぐらいだし、需要はあるんだろう。
OpenCLは全てが中途半端 OpenACCで学会の論文とかチラホラ出始めてるけど、当初の予想よりかなり良さそうなので、そのままOpenCLは駆逐される気がする。
>986,987 ふーむ。ありがとう
ここのひとはC++ AMPにも手を出してる? OpenCLみたいに中途半端なんだろうか。
今はCUDA触れる環境がないから実験で確認できないのだが、 Warpと__syncthreads関数との関係について質問したい。 「はじめてのCUDAプログラミング」という本には、 俺の記憶では次のような意味に取れる事が書かれていたような気がする。 条件分岐にさしかかった場合、 Warp内の「全てのスレッド」がまず一方の分岐を処理し、 本来そちらには行かないはずのスレッドは計算結果を捨てる。 それからもう一方の分岐をまた「全てのスレッド」が通り、 本来そちらには行かないはずのスレッドは計算結果を捨てる。 しかし「CUDA BY EXAMPLE」という本には、これとは違い、 分岐した各スレッドは分岐のもう一方の道は通らないと書かれている。 だから、ifの中に__syncthreadsがあってそこを通らないスレッドがあると、 全てのスレッドが__syncthreadsに到達するという条件が決して満たされず ハングアップするのだという。 前者の解釈ならifの中に__syncthreadsがあっても問題ない。 最後に計算結果を捨てるが、とりあえず全スレッドがそこを通る。 後者の解釈ならマズイことになる。 本当はどっちの解釈が正解? それとも実装依存で、環境によってどちらもあり得る?
>>991 if (...) {
....
__syncthreads();
...
} else {
....
__syncthreads();
....
}
は問題ない。昔エミュレータでエラーになることはあったけどGPUでは問題なく動作していた。
993 :
991 :2012/09/23(日) 18:41:53.63
>>992 if (...) {
....
__syncthreads();
...
} else {
....
}
でWarp内のスレッドが分かれるの場合の話。
「はじめてのCUDAプログラミング」の解釈だと問題ないと思われるけど、
「CUDA BY EXAMPLE」だとはっきりマズイと書かれている(日本語版73ページ)。
>>991 少なくともWarp内の全てのスレッドの分岐が同方向なら、
そちらには行かないはずでは。
__syncthreadsがWarp単位で数えていたか、
thread単位で数えていたかは分からんが、
少なくとも全Warpが同じ__syncthreadsを
通る必要はあるだろう。
>>992 MPIのbarrierと違って数だけ辻褄合わせても駄目で
ソースの同じ行にある__syncthreadsを通らなきゃいけないとか
制限無かったっけ?
>>991 CUDA4.2のCUDA_C_Programming_Guide.pdfの89ページには
以下の記述がある。
>__syncthreads() is allowed in conditional code but only if the conditional
>evaluates identically across the entire thread block, otherwise the code execution is
>likely to hang or produce unintended side effects.
仕様上は全スレッドで同じ__syncthreadsを実行しろと書いてあるから、
たとえ今のハードやコンパイラでたまたまOKでも
よろしくない書き方だろう。
>>995 でFAみたいだね。結局二つの本のどちらも間違ってなかったんじゃないかな?
っていうか
>>991 をよく見ると「はじめてのCUDAプログラミング」にしたって、
Warp内またはBlock内で分岐が割れるところに_syncthreads()を置いていい、とは書いてないみたいだし。
俺はその本持ってないけど、プログラムの動き方の説明あたりだと予想。
処理時間が片方はT1、片方はT2かかる分岐でなおかつWarp内で割れる場合だと、
処理時間はmax(T1,T2)じゃなくてT1+T2かかる的な話の。
次スレどうしよう。23:15まで誰も立てないようなら俺が立ててみようかな。
998 :
991 :2012/09/24(月) 21:22:27.56
みんなレスありがと。
>>995 俺も Programming_Guide 読んでみた。
確かに __syncthreads() を通るスレッドと通らないスレッドがあるとマズイと書かれてるね。
仕様なんで、ちゃんと従うようにするよ。
>>996 図書館で「はじめてのCUDAプログラミング」をもう一度確認したら、次のように書かれてた。
ちょっと長いけど引用する(改行は俺が勝手に入れた)。
-----( 52ページから引用 )-------------------------------
if (threadIdx.x & 1)
a = a + 2;
else
a = a + 1;
この例では「スレッド番号」が「奇数」の場合と「偶数」の場合で異なる処理をします。
どちらの方向に分岐するスレッドも、「ウォープ」内にあります。
このとき、「a=a+2」と「a=a+1」の両方の命令が実行されます。
単純に両方実行すると「a=a+3」と同じことになってしまいますが、
「a=a+2」のほうはスレッド番号が奇数のスレッドのみ有効に、
「a=a+1」のほうはスレッド番号が偶数のスレッドのみ有効にすることで実現されます。
ここで加算命令を実行した結果は有効なスレッドに対してのみ反映されます。
有効でないスレッドに関しては、命令を実行していないのと同じ状態になります。
--------------------------------------------------------
この「命令」や「実行」、「反映」、「実行していないのと同じ状態」
という言葉がやや曖昧ではっきりしない印象だけどね。
「a=a+2」の方を __syncthreads() に変えたら偶数番目のスレッドは実行される?
それとも命令を実行していないのと同じ状態になる?
この説明だけではよく分からん。
999 :
996 :2012/09/25(火) 01:49:08.14
>>998 「命令を実行していないのと同じ状態」だからつまりそういうことでいいんじゃないかな。
俺が前もって答え(
>>995 )を知ってるからそう思うのかもしれないけど。
クゥ〜〜・・・ ッダ!!
1001 :
1001 :
Over 1000 Thread このスレッドは1000を超えました。 もう書けないので、新しいスレッドを立ててくださいです。。。