【GPGPU】くだすれCUDAスレ part7【NVIDIA】 [転載禁止]©2ch.net
レス数が1000を超えています。これ以上書き込みはできません。
このスレッドは、他のスレッドでは書き込めない超低レベル、
もしくは質問者自身何が何だが分からない質問を勇気を持って書き込むスレッドです。
CUDA使いが優しくコメントを返しますが、
お礼はCUDAの布教と初心者の救済をお願いします。
CUDA・HomePage
ttp://developer.nvidia.com/category/zone/cuda-zone
関連スレ
GPGPU#5
ttp://peace.2ch.net/test/read.cgi/tech/1281876470/l50
前スレ
【GPGPU】くだすれCUDAスレ【NVIDIA】
ttp://pc12.2ch.net/test/read.cgi/tech/1206152032/
【GPGPU】くだすれCUDAスレ pert2【NVIDIA】
ttp://pc12.2ch.net/test/read.cgi/tech/1254997777/
【GPGPU】くだすれCUDAスレ pert3【NVIDIA】
ttp://hibari.2ch.net/test/read.cgi/tech/1271587710/
【GPGPU】くだすれCUDAスレ pert4【NVIDIA】
ttp://hibari.2ch.net/test/read.cgi/tech/1291467433/
【GPGPU】くだすれCUDAスレ part5【NVIDIA】
http://toro.2ch.net/test/read.cgi/tech/1314104886/
【GPGPU】くだすれCUDAスレ part6【NVIDIA】
ttp://peace.2ch.net/test/read.cgi/tech/1348409867/ >>1
ああ、キミ!また会えたね。久しぶりだ。どうだいあの件は?どうなったか説明したまえな。
神戸市の東、芦屋西宮の知的障害者施設で未成年利用者に性的な行為をして淫行条例で逮捕された三田谷学園元職員の堂垣直人(西宮市老松町)は、結局どういう罪になったの?
被害者家族のケアを芦屋市役所と兵庫県警はちゃんとやったのか?
差別や虐待は環境を選べない子供には関係ない。
http://www.youtube.com/watch?v=JxMzW3ZlV4g&sns=em
まあ、こっちに座れよ。ゆっくり話そうじゃないか。 syncthreadsとthreadfence_blockの違いが分かりません
syncthreadsだけで十分な気がしますが、どういう時に使い分けるのでしょうか? >>6
競合防止なんて書いてる?
複数のスレッドが同じメモリに書き込む時はatomicをsyncthreadsやthreadfenceと同時に使わないといけない気がする(誰か教えてください) すみません、誤読しました。
「同じブロック内の全スレッドがこの命令にたどり着く」
まで待つかどうかの違いですね。 プログラムのカーネル部分がどうしても実行されません(サンプルプログラムでは実行されていました)
どなたか原因に心当たりはありませんか?(私はありません)
ブレークポイントで確認したところカーネルの上下にあるクロックは実行されており、
カーネルだけが実行されていませんでした(カーネルの中へ入って行かないという意味です)
以下が呼び出しで、dim3はグローバルで定義してあります
dim3 blocks((num + max - 1) / max, (num + max - 1) / max);
dim3 threads(max, (1024 + max - 1) / max, 1);
void calculation(void)
{
clock_t start, end;
start = clock();
cal<<<blocks, threads>>>(con, num, points, data);
end = clock();
cout << double(end - start) / CLOCKS_PER_SEC << "\n";
} >>11
dim3構造体のメンバ変数をプリントしたらどうなりますか? >>12
>11のcalculation()の最後の行にプリントの一文を入れてみましたが
ブロックが(128,128)、スレッドが(128,8)と想定通りでした
(スレッドは1ブロック当たり1024個まで配置可能なのでギリギリセーフなはずです)
忘れていましたが、呼び出し先です
この中にブレークポイントを配置してもプログラムが止まらないという魔の領域となっています
(もちろんnsightのcuda debuggingでデバッグしています)
__global__ void cal(double con, int num, a_data *points, b_data *data)
{・・・} >>13
カーネルが実行されているかどうかはどうやって確認していますか? >>14
>>13に書いた通り、ブレークポイントをカーネルの中に入れて実行されているかどうかを確認しています
描画をするプログラムなので図形が動くか動かないかでも判断できます おれもカーネルに入らないケース出たわ
原因調査中・・・ >>15
CUDAのデバッガ使った事がない(ひたすらprintf)ので
一般論的な事しか言えませんが、
怪しそうな処理をコメントアウトしていったらどうでしょう?
変なメモリアクセスで落ちるとか割とありがちな気が。 おれもカーネルに入らない
ただなぜかcygwinでコンパイルするとカーネルが起動する
なんでじゃ・・・
例のvisual studio2013もインストールしてみたいな〜 >>18だけど
osはwindows8.1
コンパイラはcuda6.5+vs2013 express
カーネルは担当する要素を+1するだけのもの
これをコマンドプロンプトでコンパイルしてもカーネルは動かなかった(?)
動かないと判断したのは結果をmemcpyしてホスト側に返しても+1されてなかったから
もしかしたら正常にmemcpyされてないだけかもしれない
カーネルに入らない人はcygwin使ってみるといいかもね
本質的な解決にはならないけど・・・ kernel実行後にcudaGetLastError()でRCを取得
そうするとkernel実行結果が分かる
RCの数値は自分で調べてね 21の情報は役だったかな?
自分の場合(LINUX)、RC=7(too many resources requested for launch)だったので
コンパイルオプションに -Xptxas -vを追加して使用レジスタ数を確認。
結果ハードウェアのレジスタ数を超過したためにカーネルの処理が行われなかったことが判明。
スレッドサイズを小さくして問題解決。
因みに使用レジスタはハードウェアによって変わる。 >>22
ちょっと興味があるんですけれど、
カーネルのサイズは動的に決めてるんですか?
前にソースコードにブロック数とスレッド数をべた書きで
大きいサイズを指定したらコンパイルの段階ではじかれた事があったんで。 >>21
ありがとうすごく役に立ったよ
ここ最近忙しくてpc触れなかったんだ
自分も同じく「error: too many resources requested for launch」だった
原因を調べてみるよ
あと、自動でエラー内容もだせるみたいだね
http://homepage2.nifty.com/takaaki024/tips/programs/gpgpu/cuda.html
>>20
ビジュアルプロファイラーも便利そうだから調べてみるよ x,y,zの3つの変数から成る構造体配列A,B(同じサイズ)があったとして
BからAへそれぞれ対応するデータを転送する場合
Ax,Ay,Az,Bx,By,Bzという同じサイズの構造体でない配列が6つあったとして
BからAへそれぞれ対応するデータを転送する場合(BxからAxなど)
前者と後者では後者の方が転送速度は上がりますか? >>25
一般論としてデータ量が同じなら一回にまとめてを転送した方が効率はよくなりますね。 kernelの中で使えるタイマー関数はありますか? 自分の持っていないGPUの共有メモリの量などの詳細を知ることはできますか? >>32
ありがとうございます
嬉しいことに5.0以降から容量が増えてるみたいですね なんでCCのバージョン3.5から5.0に飛んでるの?
SDKのバージョンと合わせたのか? 二つ以上のGPUでVBOを使用する場合、データの流れはどうなっているのでしょう?
やはり一旦ディスプレイに接続側のGPU出力データが集められ出力されるのですか? >>34
ゲフォの800番台がスルーされたからかも? >>36
少し気になって質問してみただけで今自分は一つしか持っていません!
もう一つGPUを購入した時の為にvisual profilerを使えるようになっておこうと思い
ビルドした実行ファイルでプロファイラを使用してみたところ「Warning: No CUDA application was profiled, exiting」とエラーが出てしまいました
6.5のツールキットを使用して新しいセッションを作成→ビルドした実行ファイルを選択→設定はデフォルト、としたのですが何がダメだったのでしょう?
nvidiaの説明書を見てもさっぱりです。ヒントだけでもいいので教えてください >>38
cudaDeviceReset() が必要のはず。
Visual Studioで「CUDA X.X Runtime」のプロジェクトを作ったら、
kernel.cu の return 0; の直前のコードに、以下のように書いてある。
// cudaDeviceReset must be called before exiting in order for profiling and
// tracing tools such as Nsight and Visual Profiler to show complete traces.
cudaStatus = cudaDeviceReset(); >>39
ループしているプログラムなのでエスケープキーを押すと後処理関数をatexit関数で呼び出して終了するようになっています
その後処理関数の中にcudaDeviceReset();を入れているのですがこれではダメなようです
それともcudaError_t cudaStatus = cudaDeviceReset();としてcudaStatusをどこかへ渡すのでしょうか? >>40
とりあえず、>>39を新規プロジェクト作ってプロファイラの
動作を確かめてから、あらためて自分のソースコードに反映すれば? >>41
新規プロジェクトでサンプルプログラムが生成されるのを忘れていました
おそらくサンプルのプロファイルに成功したので自分のプログラムに反映させようと思います
ありがとうございました 一応書いておきます
調べてみた結果、必要なものはcudaDeviceReset();を呼び出すことのみでした
自分のプログラムがプロファイル出来なかった原因は.dllが.exeと同じ場所に無かったからでした 並列化についての質問です
スレッドやブロックを増やしてもあまり計算速度に差が出ないのですがどのような理由が挙げられますか 最初は一つのスレッドにつき4回ほどループさせ計算をしていました
その後、スレッド数を2倍にしてループ数を半分の2回しました
計算速度は2倍になるだろうと予想していましたが、あまり変わりませんでした ハード的に同時に実行できるスレッドは有限なんだから、それ以上スレッドを増やしても
物理的に速くなりようがない。 プログラム上のスレッド数とハード上のスレッド数は違うと言うことですか スレッドが多ければ、メモリアクセスでスレッドが止まっている間
cudaコアは別のスレッドを実行できる 基本的にはcudaコアの数だけしか並列計算出来ないのですか
だけしかと言ってもコアは何百もありますが ある瞬間、実際に並列に処理されているということと、理論上並列に扱われるということは別の話ですよ
上のレスにもありますが、計算速度的には実際に処理を行うヤツが足りていなければそこで頭打ちになるのは当然かと cudaコアの数以上にスレッドを生成する利点は
メモリアクセスの遅延の隠蔽にある 550TIで使っていたプログラムを750TIで走らせて見たところ1.5倍ほど遅くなってしまいました
何故でしょうか? >>53
一度のカーネル実行で処理するデータ量を増やしたら改善しませんか? >>53
maxwellは倍精度がそーとーしょぼいので、
cuda-zかなんかで性能をチェックした方がよいかもしれません。
http://sourceforge.jp/projects/sfnet_cuda-z/ >>56
本当にしょぼかった。陽子ビームぶち込みたい。
どうやらマクスウェルさんは演算用には向いていないようですね コア数が3倍になっていることを考慮すればそれでも遅い気がしますね CUDA初学者です
cudaBindTexture2D()のpitchとoffsetは何を表しているのですか?
手元の書籍のサンプルから推測するに
pitchは一次元の配列を二次元のテクスチャに入れる場合の折り返し地点のようなもの
でしょうか?それならwidthとhighだけでも十分ではないかと言う疑問も出てきます。
そして一番の疑問がテクスチャメモリの存在です。
いくら二次元、三次元配列が使えるとは言え512バイトしか容量のないテクスチャメモリは64キロバイトもあるコンスタントメモリに劣るのではないでしょうか?
長々と失礼いたしましたm(_ _)m >>52
それじゃあストリームは何のためにあるのさ? >>60
どっちも使えるなら、実験して早い方を採用
同時に並んでいるスレッド数を増やしてcudaコアが遊ばない状況
を作り出すことが重要
ストリームの使い道は異なるカーネルの並列実行だと思ってる 970/980は確かにゲームのパフォーマンスは上がってるが帯域減ってるから
GPGPU用途では微妙になってしまったな Toolkit 6.0 + VS 2008から
Toolkit 6.5 + VS 2013に移行したら
遅くなっちゃったんだけど、そういう人ほかにいる? CUDAの日本語ページって4.0とかの古い情報ばっかりじゃね
6.0/6.5では全然仕様が違ってて全然使えない 970の影響でGPUメモリテストが流行っているね。 ここ何週間かデバッグを続けているのですが原因を突き止めることが出来ません
初学者がはまりやすいミスやデバッグのこつなんかを教えてもらえませんか?
明らかなバグなら原因を突き止めやすいのですが、かなり微妙なバグなのでなかなか見つけられず困っています printfとかで要素を表示してデバックしてみれば? syncthreadとか?
if文の中に書いてたりすると同期ずれが起こったりするなー
他には確保してないメモリへのアクセスとか?
>>70の通り、printfとかで、配列の添字とか値を表示するしかないのかな? >>70
>>71
ありがとうございます
1セットの計算量があまりにも多いのでprintfの方法は難しいです
シンクロや範囲外アクセスもありませんでした
原因が分かっちゃったかも知れないので質問です
中間計算結果→atomicAdd
中間計算結果→配列→atomicAdd
こんな風に同じ数値を使った計算でも一度配列を通してしまうとatomicAddによって追加された計算結果に差が出たりしますか? >>72
ちょっと面倒だけど、要素が多い場合は減らしてやってみるとか、どうだろうか?
何はともあれ、原因判明したみたいで、おめでとう >>74
それは意外な盲点でした
数を減らせば良かったのですね >>72
fpが(a+b)+c != a+(b+c)を知らないとかではないよね? volatile使うとか?
変数の宣言とか関数の引数の型の前にvolatileをいれると・・・ >>53
750TIでGPGPUって考えていたけど、750TIって2世代前の同ランクぐらいの550TIより性能悪いのか。
一般ゲーム用VGAではGPGPU能力ってたいして要らないから落としたのかな
いろいろなゲーム用VGAの単精度、倍精度の能力が載ったホームページ教えてください 質問☆
cudaってドライバインストして、画像表示をcuda設定にするだけでは
効果ない? >>78
FP32とFP64の一覧表ならこれとか。
ttp://www.geeks3d.com/20140305/amd-radeon-and-nvidia-geforce-fp32-fp64-gflops-table-computing/
ボトルネックになりうる点は他にもあるから、Compute Capability毎の仕様の違いも結構重要だと思う。 maxwellさん自体にに倍精度が無いようだから
一世代前のkeplerさんか次世代のpascalさんを選べば良いんじゃないかな >>80
有難う。750TiのFP64悪すぎだな。
なんか大衆向け用でGPGPUするならFP32よ、FP64は使わないでだな。 >>81
KeplerはMaxwell以上にピーキーだったような。
自分の用途がはっきりしていて、それがKeplerやMaxwellに向いているならありだろうけど。
64bit変数をほとんど使わなくても、不向きな処理ではGTX 680が570に惨敗したりする。
CUDAの開発環境とか情報量に魅力を感じて、あえて今から始めてみるという人に勧めるとすれば、個人的には
投げ売り続行中のGTX 570や580で、余裕があるなら型落ちCPU・マザボ・メモリのセット等と
組み合わせてCUDA専用マシンを用意かな。
>>82
それは差別化とか、グラフィック用途でのワットパフォーマンスとかで仕方がないかと。 keplerはinteger bit shiftが弱いGK110(tesla)以外は
maxwellはkepler比で2倍のスループットになってる 適当なプログラム作ってみても
maxwellの方がはやいね
shared memoryが倍になったのも大きいなぁ スレチなら申し訳ない
当方、モバイルでCUDAを使用したいけど
安い方法はどれが良いと思います?
(速度はそこそこで良く、外でテストして
パワーがいる場合はデスクトップを使用するつもり)
モバイル用は安く上げたいので
Chromebookかタブレットで探した方が良いですかね?
奇をてらってJetsonのtk1にACアダプタ用のバッテリーを積むとか
(可能かどうかわからないですが)
ちなみにゲームはやるつもりありません。 thinkpad w550sのquadro K620mはダメなの?
ノートパソコンだよ レスありがとうございます
安くあげたいので予算的に厳しいかと
最初だけ計算量は多いですけど
要所ごとに定数化すれば、その後は計算量がへるかと思っているので外での使用は少ないデータ量でプログラムチェックができればいいかなと考えています。
K1がのったタブレットも安いのでそこから考えてみようと思います
ありがとうございました 単精度が7TFLOPS、倍精度が0.2TFLOPSで良いんだよな??
詳しい事は良く分からないんだけど倍精度ってそんなに使わないものなの?
PhysXっていう物理エンジン使ったりするのに >>90
ゲームだと単精度で十分だったりするからじゃない?
もともとゲームのために作られたような物だから、倍精度の性能なんかあまりこだわってないと思う Dear Developer,
The CUDA? 7.0 Production Release is now available to the public. Run your application faster with this latest version of the CUDA Toolkit. It features 64-bit ARM support and the simplified programming model enabled by Unified Memory. Highlights include:
New cuSOLVER library
? Accelerates key LAPACK routines, 12x faster direct sparse solvers
New C++11 language features
? Increases productivity with lambdas, auto, and more
Runtime Compilation
? Enables highly optimized kernels to be generated at runtime
Download the CUDA 7 Production Release at www.nvidia.com/getcuda
Learn more about CUDA 7 by attending these webinars:
CUDA 7 Feature Review
Date/Time: Friday, April 10th at 10:00 AM PDT
Register: https://cc.readytalk.com/r/4b0lwgeqgzrk&eom
CUDA 7 Performance Overview
Date/Time: Wednesday, April 15th at 11:30 AM PDT
Register: https://cc.readytalk.com/r/empyu1qc65l6&eom
Best regards,
Nadeem Mohammad
NVIDIA Developer Relations Team CUDAがgcc4.8でうまく動かない問題ってもう解決したの? Driver API使ってコンパイル済みのptxを実行するプログラムを作ろうとしているんだけど、
64bitのホストコードから32bitのptxを実行することってできるんだっけ? nvidia-smiでutilizationやmemory usageを確認できると思うのですが、全てのパラメータを確認できるのはtesla、quadro、Titanです。
Geforce系のGPUで、utilization、各プロセスのmemory usage等を確認するにはどうすればいいでしょうか? 古いGeforceでもOpenCL1.1のプログラムなら動くのでしょうか 少なくともCUDA対応している必要があるが、基本的にGeForce8シリーズ以降で動くはず。
特定の機種について知りたいならGeeks3Dのデータベース検索してみるとか。
しかしスレチ。 thinkpadに入ってるquadro K620Mはcudaのサポート無いんだな
インストールしようとしてワラタ Linux(CentOS6)の環境でTITAN BLACKでCUDA(Ver.6.5 or 7.0)使ってるのだけど、計算走らせると
NVIDIA X Server SettingsのPerformance LevelsのLevelが3から2に落ちてMemory Transfer Rateが7000Mhzから6000Mhzに落ちてしまう。
で結果的にノーマルTITANと性能が同じになってしまう。このメモリクロックを固定させることはできませんかね?
ちなみにTelsa K20 での固定クロックの方法はnvidia-smiがTITANに対応してないらしくできなかった。
ちなみにPreferred Modeは「Prefer Maxmum Performance」になっていて、
CUDAで計算しなければLevel3でメモリクロックも7000Mhzとなっています。 >>102
何時間もフルで使ってるならまだしも、GPU叩いた瞬間に7000Mhzから6000Mhzに落ちるって、TITAN BLACKの公称メモリクロック「7GHz」ってのは詐欺にならないか?
これじゃノーマルTITANと同じ性能・・・ 環境がわからんからなんとも言えんが、
どうしても固定したければどっかからbiosを引っ張ってくるしかないね。 ん?TITANのクロック制御はマザーのBIOSが管理してるのか? pycudaで既存のコンパイル済みのオブジェクトとリンクして
実行させるにはどうすればいいのでしょうか cmakeの自動configがGTX900シリーズなのにsm_20とか言ってくるのなんとかならないの? 最近のマザーボードってビデオカード2枚刺して
両方に計算させることもできるの?
自分のcore2duoのパソコンだと1つしか刺す場所ないけど
最近のddr4を使うマザーみると
それらしき場所が3つぐらいあるから3並列計算できるってことかな? >>110
場合によっては、計算用のGPUと出力用のGPUで分けてたりする。少なくともNVIDIAはそういう方針 vexclのサンプルコードをcudaでコンパイルしようとしてるんだけど
エラーだらけ
何か依存パッケージが足りないのかな? .bashrcに設定する変数名が
CUDA_PATH
だったり
CUDA_ROOT
だったりみんな違う
なんで統一しないのか vexclについてくるexampleプログラムの中のベンチマークってプログラムが面白い
このプログラムのベンチマークによると
GT430とかいう古いビデオカードなのに最近のCPUの10倍の計算速度でてる thinkpad w550sのquadroでもcudaなんも問題なかった
やっと勉強できる 3D映像でパストレーシングやモンテカルロ法を行いたい場合、光線とシーン中
のあらゆるポリゴンとの交差を判定し、条件分岐を行う必要があります。
しかし、GPGPUは条件分岐が遅いらしいので、もしかするとCPUと役割分担を
行うべきなのでしょうか?
例えば、シーン中のポリゴンとの交差判定はCPUでやった方が良いとか? if
A
else
B
end
のような文なら、AとBは逐次の処理になる
if
A
end
なら、Aを行わないスレッドは待つだけだから問題ない
もし3項演算で解決できる条件分岐なら3項演算子を用いて書く >>117
>AとBは逐次の処理になる
これはどういう意味でしょう?
AとBが内部的には必ず両方とも実行されてしまうと言うことでしょうか? Intel系CPUの場合、確か cmov などという命令は、bool 値が1の時だけ
転送するのだったと思います。そういう命令は無いんでしょうか?
仮に何かあったとして、CUDA や OpenCL でどうやってそのような
「専用命令」を生成する事は可能でしょうか? 分岐命令の代わりに使えるプレディケート付き命令があるかという話なら
それは存在するし、分岐のあるコードではコンパイラさんがよしなにやってくれる。
と、ヘネパタ本には書いてあった。 >>120
ありがとうございます。
では、「>>118」の方はどうですか? >>121
warp divergenceの話じゃないの? 今まで一つのソースファイルで書いてきたのですが、いい加減見づらくなってきたのでソースファイルを分けてみたところエラーがでたので質問させてください
extern __constant__ unsigned short v[n]; //nは#defineで定義しています。
メインのソースファイルで
__constant__ unsigned short v[n];
と定義し、ビルドすると「メインの方で再定義されている」というエラーが出てしまいます
他にも、__device__ double atomicAdd の定義をサブのソースファイルに書くと
ptxas fatal : Unresolved extern function '_Z9atomicAddPdd'
とエラーが出てしまいます。どうか解決方法を教えてください。
環境は
NVIDIA Nsight Visual Studio Edition v.4.2
CUDA 7.0 Toolkit
のはずです CUDA5.0以降ではrdcを「はい」にするとファイルの分割が出来るという情報を得たのでさっそく試してみたところ
>>125のエラーは無くなったのですが
1>sub.cu.obj : error LNK2005: "long __cdecl genrand_int31(void)" (?genrand_int31@@YAJXZ) は既に main.cu.obj で定義されています。
上記のような二重定義エラーがいくつか出てきました
また、ヘッダファイルの拡張子を.cuhとしている方を見かけたのですが、.hとの違いはあるのでしょうか? デバイス関数ってインライン展開されるんじゃなかったっけ。
でもって、他のファイルからは使えなかったのでは? >>128
ファイルを分割出来ないとなるとすごく不便になので、何かしらの方法があるのでは?と思っています とりあえずCUDAの日本語wikiがあるから、そこの分割コンパイルの項目をみたらどうかな?
OSによっては解決できるかも cuファイルをcuファイルにインクルードしてみるとか。 レスありがとうございます
どうやら>>127のエラーはCUDAとは関係なさそうなのでC++のスレで聞いてきます >>127の方法で分割コンパイルに一応成功しました!
一応と言ったのはrdcを「はい」にした場合と「いいえ」にした場合とでプログラムの挙動が少し変わってしまうのです
今書いているプログラムは何億回もの計算結果を足し合わせていくものなので、誤差の蓄積による影響を受けていると思います
コンパイルの方法を変えた(?)事で誤差の蓄積が変化(?)し、プログラムの挙動が変わったということなのでしょうか?
コンパイルと誤差がどう関係するのかは分かりませんが >>「はい」にした場合と「いいえ」にした場合
ファイルを分割するともちろん「いいえ」では動かないので、一つのファイルで完結しているプログラムでの話です 最適化の仕方が変わるのかもね
誤差が出て困るなら、なるべく誤差が少なくなるような組み方を考えるしかない 誤差にも種類があるが・・・
浮動小数点演算等の誤差なのか
同期ミスによる誤差なのか・・・ >>135
かもね
今はatomicAddをそのまま使っているので、情報落ちをなんとかして誤差は減らす予定です
>>136
原因と結果の関係性が分からないのでなんとも言えないです
ちなみに同期ミスによる誤差とは何でしょうか? つうか最適化で結果が変わったら
それはバグ以外に無いけどねえ。 今年はGTC Japan開催のリリース出ないな
例年ならもう出ているのに
中止かもしくは時期がずれたのかな >>140
今年は9月18日に虎ノ門ヒルズで開催、だったかと。 Visual studioでCUDAプログラムを作成しています
Nsightでstart CUDA debuggingからプログラムを走らせた時だけプログラムが動かないんですが
どのような原因が考えられますか?
以下の条件ではすべて正常に動作・終了します
デバッグなしで開始(debug,release共に)
start CUDA debuggingでブレイクポイントを仕掛けて1スレッドずつ動かす 追加です
プログラムの詳細ですが
ある無向グラフに対して,与えられた2点が連結かどうか調べるプログラム
枝重み0と仮定してダイクストラのアルゴリズムを用いて判定しています.
1ブロックにつき16スレッドが同時に走ります.
その判定関数をfoo()とすると
if( blockIdx.x == tmp )
foo();
というように一つのブロックのみで扱うようにすると
tmpがどのような値でも正常に動作するのですが
2ブロック以上でfooに行くようにすると動かなくなります
エラーがでないことと,printfによるチェックですが
配列のアドレスも問題ありません >>143
何かスレッドの実行順序に依存した処理があるか、排他が必要な競合でもあるんじゃね? デバイスメモリの内容を一時的にホスト側からアクセスしたいんですが、OpenCLでいう
clEnqueueMapBufferに相当する機能はCUDAにあるんでしょうか?
ちょっと調べてみたんですが、cudaHostRegisterは方向が逆だし、cudaMallocManagedは
既にあるデバイスメモリにアクセスするものではなさそうなので。
cudaMemcpyで読み出して終わったらまた書き出すという方法しかないんでしょうか。 OpenCLを知らないので外してるかも知れないけど、
基本的にはcudamemcpyするしかない。
CUDA6.0からUnified Memoryっていう機能が入って、
GPU上のメモリをホストメモリと同じように読めて、
自動的に同期を取ってくれる機能。用途によっては使えるかも。 >>147
ありがとう。やっぱりcudaMemcpyするしかないのか。 今日はじめてCUDAを入れた者ですが、VS2013でサンプルをビルドするとerror MSB3721…\simpleTemplates.cu"" はコード 1 で終了しました
って出るんだが何か設定とか必要ですか?
あとインストールの確認(?)でコマンドプロンプトでnvcc -Vのコマンド入れてもアクセスが拒否されましたと出る…
これが原因ですかね?
いろいろ調べてみましたが、これといった解決方法はなかったので助けて下さい。
よろしくお願いします。 >>149です。
すいません自己解決いたしました。
ウィルスソフトが邪魔をしていたようです・・・
スレ汚し申し訳ございませんでした。 750Tiってあんまりいけてないのか…
CUDA試したくてあんまり調べずに衝動買いしてしまった
900台にした方が良かったのかなあ ゲームするならかなりいけてるが倍精度はしょぼいな
倍精度が欲しいならfermiかgk110コアのがいいのか?
ゲームしないならamdでopenCLの方がコスパ良さそうなんだがどうなの? CUDA7入れたのですが、
cublas.hが見つからないと出てしまいました。
cublas使うには他に何か必要なのでしょうか? >>154
C:> dir /s /b "%CUDA_PATH%\cublas.h"
C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v7.0\include\cublas.h >>152
>ゲームするならかなりいけてる
こういうのがイマイチよくわかりません
単純にFLOPSで性能を見ちゃいけないのでしょうか
ゲームやGPUプログラミングに向いてるかどうかってどうやって判断してるんでしょうか >>157
プロはGPGPUするのにゲーム用のGFやRadeonをはあんまり使わない
趣味レベルでOKのGPGPUならGPGPUがオマケであるゲーム用(ゲームドライバ)のVGAでも良いが。
Xeonが一般的な業務サーバーに普通のデスクトップPCをいっぱい使うなことは
しないのと同じ。あと、FLOPSは目的のGPGPUするのに必要なレベルあれあば良い。
低FLOPSでOKなGPGPU処理に超高FLOPSのものを使うって趣味だし プロだってコスパ考えるだろ。
倍精度演算性能と信頼性で妥協できるならゲーム用GPUの方が圧倒的に安いしな。 まあ中国の天河1号はFireStreamですらないRadeonのデュアルGPUカードを大量搭載したが
結局実用にならなくて次の更新で全部Teslaに差し替えたからね
額面のFLOPS数だけで実用性は評価できない 中国はスパコンにTESLAとか使えなくなったから、NVIDIAやインテルは大口顧客を失った。 >NVIDIAやインテルは大口顧客を失った。
その損失分は米国内の大型スパコンでの便宜をはかることで補填されてるから企業としては差し引きOKとなったはず windows7でchainer試そうとすると
pycudaでなんか不具合出る 並列計算でGTX780Ti使ってるんだけどGTX980Tiにしたら倍精度の計算って遅くなる? cudaってMFCのcppから呼べますか?
サンプルで.cuから.cppを呼んでるのは見つけたんですが。 >>165
単精度演算性能 コア数×クロック数×2
倍精度演算性能 単精度演算性能/24(Kepler), 単精度演算性能/32(Maxwell)
として計算すると、
780Tiが2880×875×2/24=205GFLOPS
980Tiが2816×1000×2/32=172GLOPS
だから、遅くなるんじゃね? CUDAって仕様上できないこともコンパイルとおるのか〜 Nvはゲームにゲーム用VGAには倍精度演算イラネから省くって流れだから
ゲーム用VGAで倍精度演算性能ほしいならRadeにしたほうがいいよな >>168
ありがとう
Teslaは高くて買えないからPascalまで待つか 全く値段見ないでカンで言うけど
中古の無印TITANならウンコに混じったコーン位の値段で売ってんじゃないの? CULAのエラーでpos 1022ってなんですか?
LAPACKドキュメントを見よとか書いてありますが、見当たりません 最近始めたんだけどversionって無理して最新にする必要無いのかな
7入れるのに結構苦労したんであんまり変えたくない
みなさんはガンガン上げてく感じですか? 分割コンパイルができないんだけど、どうしたらいいんだろ?
ビルドどころか1ファイルのコンパイルもエラーになる
vs2013ce, cuda7.0, fermi, debug,x64
D:\cuda_test\cuda_test>
"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v7.0\bin\nvcc.exe"
-gencode=arch=compute_20,code=\"sm_20,compute_20\" --use-local-env --cl-version 2013
-ccbin "C:\Program Files (x86)\Microsoft Visual Studio 12.0\VC\bin\x86_amd64" -rdc=true
-I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v7.0\lib\x64"
-I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v7.0\include"
-I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v7.0\include"
-G --keep-dir x64\Debug -maxrregcount=0 --machine 64
--compile -cudart static -g -DWIN32 -DWIN64 -D_DEBUG -D_CONSOLE -D_UNICODE -DUNICODE
-Xcompiler "/EHsc /W3 /nologo /Od /Zi /RTC1 /MDd "/wd 4819""
-o x64\Debug\cuda_body.cu.obj "D:\cuda_test\cuda_test\cuda_body.cu"
1> nvcc fatal : A single input file is required for a non-link phase when an outputfile is specified nvcc fatal : A single input file is required for a non-link phase when an outputfile is specified
コンパイルフェーズで出力ファイル指定するときは1ファイルでやれっていうけど、出力ファイル指定しないと何にも出力されないし、
cuda5.0以降は分割コンパイルは対応してるよね? rdcもdlinkも指定してるし何がいけないんだ 別にCUDAじゃなくても良いんですが、
GPUで直接JPEGを、DirectXのテクスチャに展開して、そのまま表示したいです。
そういったライブラリって無いですか?
かなり需要がありそうなんですが ヘッダファイルで定義した定数は__device__内では使えないのでしょうか? 少し分かったことがありました
グローバルで
const int i = 1;
const double d = 1;
と定義すると
__device__内でiは使えてdは使えませんでした
error: identifier "d" is undefined in device code
とエラーが出てしまいます
どうしてdouble型だとエラーになるのでしょうか? floatはダメでlong intはOK
だめならだめで引数とすれば良いんですけどね
整数だけというのはcudaの仕様ですかね? なんでだろうな。全く調べていないので100%想像だけど、
GPU上にstatic変数が確保されているとは思えないから、
(無いとは言い切れないが)
コンパイル時に定数展開しているんだろうな。
C++の言語仕様としてconst intは定数展開していいし、
(厳密な言い方だと間違いかもしれん)
クラスのstaticメンバにしても宣言のみで定義が不要ってのがあるから、
それに則った処理なんじゃなかろうか。 Visual Studioでエラーでた
似たような状況の人いたら教えてほしいです
環境
CPU core i5 4570
GPU GTX 980
Visual Studio 2010 Ultimate
CUDA v7.0
NSIGHT v5.0
使用言語はC
エラーが起こる状況↓
__device__ function1(){
int array1[ 1000 ];
function2( array1 );
}
__device__ function2( int *array1 ){
int array2[ 1000 ];
for( int k = 0 ; k < 1000 ; k++ ){
array1[ k ] = rand(); // rand()は何らかの乱数を返す関数
array2[ k ] = rand();
}
}
function1を呼び出したとき,時折だがarray1の値がarray2の値と同じになってしまう
エラーが起きた状況で以下のことは確認済み
・array1,2ともにグローバルメモリに確保されている
・array1とarray2のポインタは異なっている
・rand()が異なる値を返している
また,array2の宣言をfunction1で行い,function2にarray2のポインタを返すことで回避ができる 全然知らんけど、スタックオーバーフローじゃね?
static int array2[1000];にすれば解決 >>184
device修飾してないだけとか馬鹿なおちはないよな? VS2013 cuda7.0で
intelisenseがcudaの関数やthreadIdxなんかに赤線つけるんだけど
intelisenseを無効化させる方法ではなく、対応させられない? CUDAって余計な決まり事が多くてGLSLよりわかりにくくなってると思うんだけど >>191
#include<device_launch_parameters.h>
で消えない? cuda_header.hに以下をまとめて、これを.cuでインクルードしてるんだけど、インテリセンスがエラーとして表示したままなんだ。
また従来のcutil系の代わりのhelper関数のヘッダが開けないともなぜか出る。(プロジェクトの追加のインクルードDirには指定してあるんだが・・)
でもコンパイルはされるんだよ・・・。インテリセンスのキャッシュクリアしようと考えたんだが2013のメニューにはないみたいだし・・。
#include <cuda.h>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <device_functions.h>
#include <math_constants.h>
#include <helper_cuda.h>
#include <helper_math.h> 超初歩的な質問ですがよろしければお願いします
環境 windows7 64bit Visual Studio 2013 community CUDA7.5
CUDAインストール時にはVS2013をインストール済みと判定するものの
コマンドプロンプトからnvccでサンプルを実行させると
nvcc fatal : nvcc cannot find a supported version of microsoft visual studio
とエラーが発生します
パスなどの設定の問題なのでしょうか? コマンドプロンプトからcl.exe叩くのと同じだろう。
VSの開発者コマンドプロンプトからやってみれば? Communityでもいけるんだ?
少し前に断念してProにした記憶がある 質問です
OpenGLとCUDAを使って流体のシミュレーションをしようと考えています
OpenGLの描画とCUDAの計算を別スレッドで行ってCUDAの計算の完了を待たずに描画をしようとするとプログラムが落ちてしまいます
そもそもこういう使い方は間違いですか?
GPUを二つ使わないと不可能なのでしょうか? CUDAを待たずにどういったデータで描画してるの? 正確には
CUDAで計算
↓
計算用と描画用として二つのデータをメモリにコピー
↓ ↓
グラフィック描画 CUDAで計算 ←ここを並列化したい
↓
計算用と描画用として二つのデータをメモリにコピー
↓ ↓
グラフィック描画 CUDAで計算
↓
といった処理の流れです
描画と計算の処理の合計が少しでも速くなれば最終的には速くなる見込みなんですが、描画とCUDAは同時にはできないのでしょうか? >>198
そういうのは環境、sdk、cc、計算モード、HWによっても違うと思うんだけど。
大前提として計算してる最中でも、描画分のリソースは残してるのかね? すいません何故か解決しました
OpenGLによる描画をサブスレッドにしていたところメインスレッドにしたら問題なく実行できるようになりました AMD、GPGPUの向けの取り組み「ボルツマンイニシアチブ」を発表。CUDAとの互換性も提供 | スラド
http://srad.jp/story/15/11/18/0657225/ >>203
いまや一般向けプログラムですらGPUを使うことが普通になったから、
これがC++コンパイラーのデファクトになるだろうな。
DX12,VulkanでGPUするならAMDって流れにまってしまったからな え?
GPGPUの世界は圧倒的にNVIDIAなんだけど? 産業向けのガチのやつならNvidiaでくだ
カジュアルな一般向けならIntelをOpenCLでってとこかな 開発の手間で言えばCUDAが一番楽。NV限定にできるなら趣味でやるにしてもそれが一番手軽。
OpenCLならせめてAMD。IntelはOpenCLの経験を積んだ苦労と向け。 個人向けやオフィスPCでグラボ別にあるのって少数だから仕方ない。
数を優先するならむしろNvidia切ってIntelのみ対応という手もありうる。 慣れてるし開発ツール揃ってるからNが一番楽なんだけど。 HD GraphicsでOpenCLやってもCPUと大して性能変わらんぞ?
それに、オンボードグラフィックスのPC自体は多いだろうが、OpenCL1.0しか動かんとか
そもそもOpenCLが使えないのも少なくない。間口でいうならSSEでもやるほうがまし。
そうでなければAMDでOpenCL開発してNVとIntelで動作確認ってところだな。 >>211
CPUの代わりというよりは少しでも計算パワーがほしい場合もあってさ。
GPU使ってる間CPU使えるなら2倍になるじゃん?
あるいはBGでうごかしてフロントでは通常のアプリ使えるじゃない?
絶対的なパワーというよりはそういう所に使いでを見出してる。
確かにOpenCLはHWアーキ変更激しいし世代でがらりと変えてくるからやりにくそうなんだけど。 間違い。OpenCLではなくてIntelのGPUね AMD関連使って良かったと思った印象が一度も無いから使いたくねえな >>212
つまり、少しでもパフォーマンスは欲しいけどdGPUほどの性能は必要なくて、
数を優先するけどオンボードグラフィックスでOpenCLが動く環境に限定していいと。
最初の「カジュアル」からは想像つかないピンポイントな要件だなw カジュアルGPGPUなら内蔵GPUでOpenCL2確定
一方、外GPUを使う高性能GPGPUものは昔はNvidiaが圧倒していたけど、いまはAMDが圧倒しているんだろ? >>215
そうなのかなあ?自分の文章の書き方が良くないのかもしれないが
ほとんどは内蔵GPUしか期待できないから、
その範囲で少しでも計算パワー稼ごうと思ったらそれしかないように思うのだけど
dGPUほどの性能はいらないわけじゃなくて、ほとんど期待できないからしかたなくという感じ。 組み込みは知らないけど汎用PCでAMDが圧倒した話なんて聞いたことない bitcoin マイニングでは AMD 一択だったようだが deep learningではNVIDIA圧勝だろ? bitcoinは専用ASICが開発されてとっくにそっちが主流になりました
GPGPUでやったところで、最早電気代のもとを取ることすらできない
GPGPUが最適といわれるのはASIC以前にFPGA化して売り出すほどの人気もない
マイナー暗号通貨くらいでそんなニーズが世界的な需要をけん引することなんて
ありうるわけがないのだ(需要が少ないからこそGPGPUでやってるわけだし) >>217
一般に配布するアプリでなるべく多くの環境で動かしたいというならOpenCLで、
内蔵GPUしかない人は動けば御の字、もっと速く動かしたければちゃんとしたGPU
買ってね、でいいと思うが。
社内ツールとかで外付けがほとんど期待できないのなら、たいして速くならない
割りに動かないリスクばかり大きいんで素直にCPUのがマシ。 GPUばりばり使ってますってアピールしているプログラムで、avxもsseも使えないようなプログラミングしている例も多いからなぁ。 内蔵GPUはどのみちメインメモリの帯域で頭打ちになるのでたいていの問題では
AVX使ったほうが速いんだよね
どっかのA社のはCPUコア側のメモリ帯域にリミッタがかかってて内蔵GPU側も使わないと
フルに帯域を使えないうえCPU側のAVX性能が残念だから
そもそもそれ専用にプログラム書くほどの市場がない
CUDAは一握りでもちゃんと労力に見合った金を払ってくれる法人ユーザーが一定数いるから
ビジネスとして成立するんだよ。その点法人皆無で貧乏人しかいないA社ユーザーじゃ無理 ビットコインマイニングとか普通じゃない使い方で普及してるかどうかはどうでもよくないか 内臓GPUにこだわるのはそもそもGPUに興味もない人が増えていて
そんな人までプログラム、GPGPUに参加したり利用する可能性が高いからなんだよな
GPU自体ニッチ向けの商売になりつつあるし >>200
DX12でGeForce全般がAsync、非同期処理が苦手といった話があったけどそれとは関係ないよね?
自己解決ってことはやっぱ非同期で処理してたのが原因かもしれないけど
(サブで描画するのは普通にお行儀が悪いけど) 単にOpenGLの使い方の問題だろ。context作ったスレッドと別のスレッドからアクセスしようとした。 Win7PCに搭載されたQuadro4000とVC2012のC++AMPでfloatの計算をさせると「グラフィックスが応答しないのでリセットしました」とかメッセージをだしてプログラムが異常終了してしまう。
計算量はCPUを使った場合に8スレッド+AVX2で30分くらい。一枚のグラボで画面表示とGPGPU計算の両立は無理? CUDAを使えば可能なのだろうか? >>229
GPGPUではよくあること
そんなことがよくあるから一般PCでGPGPUが激活用されない
一般PCではGPGPUよりOSによる画面表示の方が大事だから、画面表示を激妨げるのものは殺す
GPGPUを必死したいならGPGPU用VGAと画面表示用VGAを搭載する
画面表示を激妨害しないようなゆとりGPGPUプログラムにする >>231
レスありがとう。
やっぱ、無理なんだ。
動くのは、サンプルだけか。 QuadroってTCCモードにできなかったっけか? Quadro自体結構前の世代のグラボじゃないっけ・・・ 2秒たったらハードウエア割り込みをかけて一瞬だけ画面表示させて、TDRが問題にならないように制御できないのかな? いや、むしろ向いてるかも
デスマーチを突破するにはそれぐらいのやっつけ仕事が出来なきゃねw 勢いでwindows10にしてからcudaが一切認識しなくなった・・・
version7.5、ドライバ最新版入れてGPU-ZでCUDAにチェック入っているのに初期化で失敗する。
付随してwaifu2x-caffeとかも動かなくなってるしもしかしてまだwin10には対応してなかったのか。 >>240
だよね・・・
最新ドライバや少し古いバージョンを入れてみるとGPU-Zでチェックに入ってるんだが
cudaを使うソフトがことごとく検出できませんでしたと出てサンプルプロジェクトも強制終了する。
NVEncのためにCUDAが使えなくなったとかいう記事見たけど、まさかこれが理由ってわけではないよね >>241
win10のvs2013で7.5問題なく使えてる。ついてくるサンプル、ビルド・実行ともに。 今のNvidiaはAMDに比べてかなり安定度低いからな >>241
GPU-ZでCUDAが見えるのにアプリケーションが動かないってのはRuntime DLLが
見つからないんじゃないのか?
どっちにしても、CUDAのどの関数を呼んだときに落ちるのか、あるいはそもそも
最初から起動できないのか、どこで落ちているかくらい調べることはできるだろ。 >>238
そんな態度でいるからデスマーチに陥るんじゃ?
いざとなった時の処世術は身につけておかなければいけないとは思うけど おかげさまでCUDAでchainer使えるようになりました。
爆速で進捗状況見てるだけで嬉しいです。
本当にありがとうございました。 教えていただきたいことがあります。
ubuntu14で、chainer1.5.1のmnistをGPUで計算できるか試行錯誤中です。
python train_mnist.py --gpu 0 #エラー
sudo python train_mnist.py --gpu 0 #エラー
python train_mnist.py --gpu 0 #うまく動いた
と、sudoをつけて外しただけで、うまく動いた理由を知りたいのです。自分の操作が正しいのか間違ってるのかよくわからなくなってきました…
ログのリンクも添付しています。2つは同じファイルです(dotupはアップローダーの都合上zip圧縮しています)
事前の処理として、vgaのドライバ、cudaとcudnnを入れてパスを通しています。
またchainerのmnistも、cpuの処理ならエラーは出ずに動いていました。
操作は
python train_mnist.pyを実行してcrl+cで止めた。
python train_mnist.py --gpu 0を実行してエラーになった(30-38行目と同じエラー)。
その後、再起動
1行目:VGAの番号を取得した(GTX750tiのIDが0)
28行目:enter押しただけ
29行目:python train_mnist.py --gpu 1 なんとなくid1で動かしてみた。エラー
39行目:id2で動かしてみた。内蔵gpuを使うかも?と思ったがエラー
49行目:id0で試すもエラー
59行目:sudoつけてみてエラー。パスを引き継いでないからnvccがない(=CUDAが見つからない)、cudnnが無いというエラーが出ているのだと思います。
91行目:sudo外したら普通に動いた
http://www.dotup.org/uploda/www.dotup.org702271.zip.html
http://www1.axfc.net/uploader/so/3602353 普通にsudoユーザー環境でcudaにパスが通ってないだけ
パスの通し方がわからなかったらこのスレじゃなくてLinuxの初心者スレで 素人な質問なんですがコア数×10スレッド立てたときとコア数と同じスレッド数で10ループした場合とでは前者の方が1.5倍くらい早かったんですけど何故でしょうか
物理的に真に並列で動作するのはコア数と同じスレッド数ではないのですか
最適なスレッド数が分からなくて困ってます Core数ではなくWarp数に合わせるべし
Maxwellの場合はコア数とWarp数は一致してるけどね
コア数の倍以上のスレッド数を立ち上げるメリットは、ロード・ストアを含む命令間の
レイテンシが隠蔽されるのでパイプラインの充填率が上げやすくなる。
半面、1スレッドあたりで使えるレジスタ本数が少なくなるので
ロード・ストアを頻発するほどのレジスタ不足になるのは本末転倒なので
並列化は程々に。
これに特に明確な基準はない。最適なスレッド数はトライ&エラーで決める。
慣れてくるとPTXとかCUBINを逆アセンブルして決める場合もある
(昔、理研で重力を自在に操る高貴なる男性研究者がメールで教えてくれた方法) >>251
仮想コア等の技術などで、本来あるコア数以上の並列処理をするのが最新のcpuやGPU、そしてライブラリ
そのへんはあまり考えずにリファレンス通りに投げてしまう方がいいと思う。
ライブラリによっても日々変わるので、最適なスレッド数というのはわからない 並列化のアレコレの時に気をつけるのは基本的にはオーバーヘッドやで
スレッド増やしたことで減るオーバーヘッドを考えてみなされ そんなもん処理にもよるわ
メモリアクセスがネックになることが多いから
物理的に真に並列で動作するなんてことの方がまれ >>252
時間のかかる処理を行っている間に別のワープへ命令を発行しているということなんですかね ECCにはこだわっても
マザーの品質にこだわらないならいいのではw 正直本格的な並列処理は非同期処理対応したGPUを待った方がよさそうなんだけどな RadeonはCUDA対応してないじゃん
nvidia製の非同期処理に対応したCPUは実質的にはまだ出てない >>260
>RadeonはCUDA対応してないじゃん
当たり前じゃんw
>nvidia製の非同期処理に対応したCPUは実質的にはまだ出てない
だからなに?
それまで待ってるの?
意味がわからない 1.メモリ読みにいく
2.大きな遅延が発生するからレジスタ待避
3.別なスレッド立ち上げる
を繰り返し
レジスタがどんどん減っていく >>263
これがあるから並列処理に制限がかかって
非同期処理はnvidiaがNG出す状況になってるんだろうなと思う >>264
お前さっきからなに訳のわからない事言ってんだよ 今あるもの使うしかないのに何が言いたいのかわからんぞ 非同期処理はゲームで使うもの
GPGPUには関係ないよ むしろNVにとっての痛手は倍精度のGPUを作っちゃうと「ゲーム用」が完全にAMD天下になっちゃう事
その最大の理由が、SCEに吹っかけすぎて、NV切りを起こされた事
GCN世代GPUがゲーム機を総獲りし、ローレベルAPIで組まれるAAAタイトルが全てGCNに最適化
これを奇貨としたAMDがDirectX12を提唱→DirectX史上初めての互換切り→AMDデファクトスタンダードの流れ
そのために科学技術演算と言う市場を捨てないといけなくなった
VulkanもDX12に倣って完全にAMDに合わせてるからな
だからディープディープ連呼するんだよ
半精度でディープが出来るからな カメラと連携してリアルタイムで判断をこなしながらってのは非同期処理無しでも大丈夫なのか? 多分非同期君はNV信者だと思う
そうでなければこんなところのこんなスレで五月蝿く連呼するか?
問題はなぜそうするのか、だ
俺なりに出した答えは、Pascal開発が暗礁に乗り上げてる所為だと思う
未だにモックしか出してないからな、Pascal VulkanとDirectX12は殆どAMDのためにあるようなものだからな
ゲーム機を支配したAMDはローレベルAPIを支配し、ゲームコードは全てAMD基準になった
これはNVにとっては痛恨
倍精度GPUなんか出す余裕はないだろう GPUのグローバルメモリのバンド幅計算で質問があります。
バンド幅(GB/s) = メモリのデータレート(MHz) * bit幅(bits) / 8(byte) / 1000
メモリのデータレート GDDR5の場合 clock * 4
DDR3の場合 clock * 2
上記のようになると思うのですが、
DDR3は立ち上がり立ち下がりで情報を伝達するのでクロックの2倍になるのはわかったのですが
GDDR5がDDR3に比べなぜ2倍になるのでしょうか? >>281
回答ありがとうございます。
NVIDIAのスペックを見ていたのですが仕様値と計算があわず、わからなくなってしまいました。
http://www.nvidia.co.jp/object/product_tesla_C2050_C2070_jp.html
これによると、C2050は384bit, 1500MHz, で144GB/sとあるのですが、計算してみると
384 * 1500 * 4 / 8 / 1000 = 288GB/s
になってしまうのです。仮に1500MHzがデータレートと仮定すると
384 * 1500 / 8 / 1000 = 72GB/s
どちらも仕様値と違ってしまいます。もしかしてバンド幅はメモリのリードとライト
合わせての値なのでしょうか? つまり上り72GB/s、下り72GB/sの合わせて144GB/sなのでしょうか? NVはもう倍精度のGPUは造らないよ
ローレベルAPIをAMDに奪われた現在、半精度GPUで描画特化をせざるを得ない GPGPUでは倍精度がものを言うが、NVはもう倍精度に手を出せないということ 単精度でも十分みたいな話じゃなかったの?
それでディープ・ラーニングを推し進めていると思ったんだけど 深層学習とビッグデータは半精度(でも可)
しかしこのスレの関心はあくまでも倍精度、出来るならば倍々精度だろ?
しかしこの倍精度は、CS機の競争環境でオマケがつくか削られるか都合で決まる
そう言うわけでVoltaはアメリカ政府との契約違反に問われないように出荷するだろうけど、
これはカスタム
AMDにCS機が獲られた現在、NVのGPUは描画/ワットパフォーマンス特化路線に軸足を置いた
DirectX11まではハイレベルAPIだったから倍精度GPUを出す余地があり、同時に科学技術演算市場を狙う二兎を追う路線を
追求できた
しかしDirectX12は完全にAMDのGCN世代GPUのハード直叩き規格だから普通に戦ったらfps/解像度共に負けてしまう
だからこそ半精度に活路を求めている DirectX12とかどうでもよろしい
ここはCUDAのスレ
AMDな方はOpenCL()のスレにお帰り下さい 倍精度倍精度って、使い道も無いくせに
滑稽だな
戦闘力みたいなイメージか 科学計算でも計算量が多くてそこまで精度が必要なものってどれだけあるんだ 科学技術計算では単精度や半精度で行う方がよっぽど珍しいと思うが 物理系じゃ大きい数+極小の数なんて普通にあるからな
倍精度でも足りなくて四倍精度必要になったりすることもざら Kerasってdeep learningフレームワークはbackend.pyてファイルを追加するだけでopenclでも動くようになってるよ。誰かbackendを書かないか? こうなったらこのスレの住人が声を上げて、PEZYに倍々精度のMPUを造ってもらおうよ
もうNVは金輪際倍精度に手を出さないよ 倍精度が必要な要件って
例えば何があるんだ?
2.0e+38を超えるスケールで
2.0e-38以下の誤差を必要とするプロジェクト?
そんなのあんの?
そりゃー誤差が小さくなれば
破綻する可能性は減るけどさ
おたくらが求めてるのってそれだけじゃねえの? >>296
でもシミュレーションにGPGPU使ってる人いるだろ
たいして速度でないのに倍精度でやってんのか
xeon phi使った方が良さそう >>300
指数分の大きさではなくて仮数部の大きさ(幅)だよね。
仮数部の幅分大きさの違う数値を加減算すると0を加減算したのと同じになると思うけど。 装置もんで数万枚/日のリアルタイム画像処理に使ってるけど、普通にdoubleが使えると楽なんだけどな。
速度も考慮すると大きい数と小さい数の計算前、画像処理の前後でdouble、float変換してるからCPU圧迫してくる。
でも、シンプルで意図した通りに動くからCUDAは組みやすい。
んー、この辺がジレンマだな。 それは倍精度が必要な理由になってないような
やっぱ盲信が多そうだな >>299
そこの製品一般人で買えるルートあんのかよ 大きな行列の計算、FFT、急峻なカーブの音声フィルター、レイトレーシング、物理シミュレーション、空間認識、...
倍精度以上の精度が必要な用途はいろいろとある
単精度のハードで倍精度相当の演算をすると、十数倍の時間がかかる
倍精度のハードで4倍精度相当の演算も同様に十数倍
多倍長演算も倍精度や4倍精度ハードがあれば、単精度よりもずっと速く計算が可能 簡単な画像フィルターや音声フィルターなら単精度でいいし、ディープラーニングなら半精度でも良かったりする
半精度〜4倍精度までのハードがあるのが一番なんだけど
倍精度演算ハードがあれば単精度も同じ速度で計算が可能だが、単精度のハードで倍精度相当の演算をするには十数倍の時間がかかる
ゲーム用途など、単精度を使うことがはっきりしてる時じゃなければ倍精度を積んでおくのが無難
4倍精度はさすがにGPUに積むのは規模が大きくなりすぎるが、CPUには積んで欲しい
半精度は、単精度との変換だけでとりあえずは良い http://www.4gamer.net/games/251/G025177/20160105002/
http://www.4gamer.net/games/251/G025177/20160105002/
>また,Huang氏が示したスライドで,第2世代Maxwellアーキテクチャベースの「GeForce GTX TITAN X」とDrive PX 2を比較すると,
>単精度浮動小数点演算性能は前者が7TFLOPSなのに対して後者は8TFLOPSと,
>劇的には高速化していない点も気になった。
↓その一方で…
>ちなみにHuang氏は,Drive PX 2の性能指標として,ディープラーニング処理の演算性能を示す
>「DL TOPS」を示しているが,これだとデュアルPascalはGeForce GTX TITAN Xの3倍以上というスコアになっている。
これ、完全にゲーム特化とビッグデータ特化GPUじゃん、Pascalwwwwww
俺の言ったとおりもう科学技術GPGPUはオワコンなんだよ
完全に半精度依存、完全にワットパフォーマンス/描画特化GPU路線だはwwwwww 👀
Rock54: Caution(BBR-MD5:0be15ced7fbdb9fdb4d0ce1929c1b82f) >>309
倍精度乗算器の回路コストは単精度の約4倍 >>311
まあ純粋に乗算器だけでいえばそうだけど...
>>310
半精度はさすがにゲームでも足りない
半精度専用GPUなんて出ない
少なくともPC用には >>313
業務っていろいろあるが
おれの範疇は単精度、倍精度、疑似4倍精度、多倍長といろいろ
多倍長も、筆算レベルからカラツバ、DFTレベルとさまざま
もちろん普通のレジスタサイズの整数演算も
特殊なジャンルであることは否定しない 半精度もIEEE754じゃない独自フォーマットなら使った事はあるが、使った事があるって程度 >>314
>半精度はさすがにゲームでも足りない
そんなこと分かるかよ
HDRで16bit浮動小数点フォーマットとか普通に使うだろ
メモリ帯域が半分で済むんだぞ
ゲームなんか正確さより速さ優先だから、色を表現するのに16bitで十分な場合がほとんどだろ だから違うよ
使ったことがあるかじゃねーよ
何でその精度が必要なのかってところよ
4倍精度とか軽々しく言うが
このスケールが本当にわかってるのか
疑わしいぜ 4倍精度とか専用ハード作れとしか言いようがない
汎用品に求めるものじゃない >>314
よく嫁
Pascalは基本混合精度
PK104は単精度と半精度の混合
PK100は倍精度の混合らしいが、多分これはテスラ新シリーズとして販売だと思う
もうグラボ価格の倍精度CUDA組はむりぽ
>>319
もうPEZYにお願いするしかないよ、倍々精度
これはマジでそう思ってる
逆を言えば理研=PEZYの特殊HPCはそこしか生き残りが出来ないのではないか
どうしてもというならARMカスタムとか
倍精度はNVはもうだめぽだから、AMDに期待するしかない 混合精度って単精度と半精度が両方載ってるって事だよな
仮に1個の演算器の回路規模が1:3として、演算器の回路比率を1:1にし、単精度は半精度の演算も出来るとすると、すべて単精度にした場合に比べて演算器の数は単精度が半分で単精度が2倍か
単精度2倍のインパクトよりも単精度半分のインパクトの方が大きい気がする 完全にグラフィックス描画効率/深層学習向けの構成だね
データ地獄になるから広帯域メモリも必要になるがHBM規格はAMDが握っているために、(HBM2の)開発が遅れた
GDDR5+を急遽導入する事になったのはそのため
このスレだとDX12とかVulkanとか一言口にするだけで「スレチガイだ馬鹿!」と怒鳴る奴がいるが、
これは全てが絡み合っている
DX12とVulkanという二つのローレベルAPIがAMD中心による策定である以上、これらのローレベルAPIへの依存を避ける
ワットパフォーマンスGPUに経営路線の舵を切るの必然 >>322
そのすごいAMDでお前はどんなGPGPUソフト/プログラム作っていいるんだ? 雑談は楽しいけど、ここCUDAスレやで
amdの話をするなってことじゃなく、スレチ話をいつまですんねんってこと Visual Studio 2015対応になるのはいつですかねぇ >>324 >>325
自演乙
わざわざ自作板から出張ご苦労様 TeslaでもXeon Phiでもなく一般ピープル向けGPUに倍精度性能が必要だって言ってるのは
ろくに科研費申請が通らない貧乏研究室くらいでしょ Maxwellは擬似的な倍精度も出来ないんだよ
それくらいしっておこうな 単精度と比べると凄い遅いだけで
倍精度演算自体はハードで実装している。
物凄い遅いといっても1/32で1/16のradeon furyと比べると半分だから大差ないと言えば無い。 Maxwellは倍精度の演算器が排除されただけで疑似倍精度は実装されている
そもそもソフトレベルでも実装できることだし FURYはMaxwellを参考にしてるからな
HBMの容量の少なさが祟って何じゃこりゃ状態だったけど、Polarisは一層Maxwell類似のワッパ路線がとられる
これでNVはもう倍精度を完全に捨てにかかる いや実際半精度が4倍速いほうが使いどころあると思うよ
ゲームで倍精度なんて使わないじゃん 深層学習がはやったらGPGPUももっと市民権を得るかな? 半精度を4倍速くしたら半精度専用になるな
そんなもの売れないぞ >>339
>ゲームで倍精度なんて使わないじゃん
団子!
お前はインテルスレとワッチョイIPの次世代AMDスレに帰れ
倍精度使わんGPGPUなんか炭酸のないコーラと同じ
ここは管スレ 倍精度を高速化したらどれだけ売り上げが見込めるのかをnvに分かるように説明したら。
>>341
頂点シェーダは単精度が必要でも、実行時間のかなりがフラグメントシェーダなんだから、別に半精度専用になって売れないとはならないだろう。
住みわけがある訳だろ。で、倍精度頑張ったのは期待の売り上げが出なかったわけだろ多分。
願望を語るスレであるより、アリモノでなんとかするスレであった方が建設的だと思うがね。 なぜ団子がここにいる?
まあ、CUDAもかなり広く使われてくるようになったから、
数値演算の枠組みだけじゃなくなってきているのはわかるが。 いやこのひとはかなり初期からCUDAつかってるの
それこそnvにgeforceくれてやるからCUDA使ってちょって
誘われるような時代から http://techon.nikkeibp.co.jp/atcl/event/15/020100028/020400019/?bpnet&d=1454641873634
CNN専用チップ、NVIDIAなどがGPUの改良で開発
http://techon.nikkeibp.co.jp/atcl/event/15/020100028/020400019/160204NVIDIAp0.jpg
Eyerissは、グラフィックス処理プロセッサー(GPU)と同様、SIMD(single instruction multiple data)などから成る
プロセッサー・エレメント(PE)を多数集積したICである。
ただし、GPUとはデータの移動を最小限に抑えた設計にした点が異なる。
具体的には、EyerissはSRAMから成るオンチップメモリー108Kバイトを備えており、
CNNの各層の演算結果をオンチップメモリーにバッファリングさせることで、
チップ外の主記憶に書き戻す頻度を大きく減らした。
また、PEに入力するデータ列に共通部分がある場合は、その共通部分を保持して差分だけを入力する。
さらに、複数のPEに同じデータを入力する場合は、各PEにそれぞれ送るのではなく、
途中まで1つのデータとして送り、そこからPEの個数分データをコピーして入力する。いわゆるマルチキャストである。
これらによって、データ処理時のデータの移動を最小限に抑え、消費電力を大幅に低減した。
加えて、データを主記憶に書き戻す時は圧縮し、読み出す時に伸長する。
CNNの畳み込み演算後のデータはフィルターを通すことで圧縮しやすい形になっていて、メモリー帯域の節約効果が高いという。
作製したEyerissの動作をNVIDIA社のGPU「TK1」と比較すると、
EyerissはデータのスループットこそTK1の約1/2だったが、消費電力はTK1の最小約1/37と大幅に低い。
主記憶に出し入れしたデータの帯域も、TK1の1120Mバイト/秒に対し、Eyerissは127Mバイト/秒だったとする。 >>346
だからお前はインテルスレに帰れっつってんだろ団子! >>350
倍精度イラネ、ここは管スレだぞ!と言ってたのはお前だったのか、団子! どうもこのスレには
CUDAをマシに使っている奴は
いないようにみえるな
ひたすらベンチマークしかやってない感じね 倍精度よりコンテキスト処理の自由化をだな
Oculasからも苦言言われてるのに バックグラウンド処理が作りにくいからこそ、Asyncやプログラマブルコンテキストは次のGPUでは絶対に搭載してほしい お前その会社好きだな
一般向けに売ってもいないのに 「Teslaを一桁安値で買わせろ」みたいなのは論外だろ
ゲーマーと倍精度が欲しい人は市場が全くかみ合わない
そんなのもう結論出てるだろ もともとTeslaとQuadro/GeForceはECCの有無で差別化してたが
世界最大級のGPGPUスパコンがNonECCで使ってるので
全く差別化になってない
だから倍精度演算器を物理的に廃することでの差別化を行った
それだけ。
Tesla買えばいいだけだろ?
どうせ予算通らないから買えないだけだろ
ろくな論文も出せない底辺大学の研究室あたりがいいご身分だな >世界最大級のGPGPUスパコンがNonECCで使ってる
これについて詳しく。 最初は影でこそこそディープディープ言ってて倍精度いらない、管スレでゲームコンソールを巡る政治的な思惑が
NVの倍精度切りに繋がった事実もスレチガイだ馬鹿野郎!って煽ってたのにな
流れが悪くなったのでコテの正体を出したんだろう
多分Async煽りも正体はこいつだ
ここのスレで倍精度がどうなるかは大学企業にとっても死活問題なのに
本当に陰険だね、団子は そこそこの分解能ならTeslaやTitanで十分だろ べつにphiでも使えばいいだけだが
teslaとphiの混合機も普通にあるし
とりあえずamdという選択肢は皆無 nvは高校生や高専向けのセミナーもやってるし
ことしのgtcもセッション数は500超える
やってきたことの重みが違いすぎる でも他GPUでは使えないんでしょ?ってのが重すぎる
GPGPUはnvidiaの私物ではない まともな研究機関ならTeslaは学割使えるんだが…
GeForceで倍精度が必要な理由なんてほんと世間知らずの戯言だな >>372
黎明期はデファクトスタンダードの取り合いだから当然の流れ >>372
言ってることが頓珍漢過ぎる
嫌ならopenCL使えばいいじゃん? dlといえばmitとnvが共同開発したeyerissはflotじゃなくfixedだったな
データ移動を極力控えた構造でcnnの場合データの再利用性が高いので
peからpeへの受け渡しが有効なのだとか
それでもエネルギーコストは
aluからレジスタファイルへの場合を1とすると
peへは3
内部のバッファーへは10
外部dramへは500のエネルギーコストがかかるんだそうな >ここのスレで倍精度がどうなるかは大学企業にとっても死活問題なのに
言うことがバカすぎて吹いた
まともな法人はこのスレに社運なんてかけませんから
Tesla相当の性能をGeForceにも実装しろという
悪質クレーマーを客だと思わない権利は
NVIDIAにもあるだろjk 米エネルギー省の評価ソフトがほとんど単精度で
倍精度使うのがlinpackしかいという 単精度じゃ足りないけど倍精度なら十分って領域は実は少ない。
単精度用演算器に少量の精度を追加して倍精度を1/4の速度でやるのが一番利に叶っている。
まあこれすら普及帯ではけちられて1/16になるのが倍精度の立ち位置
倍精度用の演算器積むくらいだったら単精度の演算器で可能なbit数で複数組み合わせると簡単に多倍長精度演算を実現できるような命令を用意してくれた方がよっぽどありがたい。 >ここのスレで倍精度がどうなるかは大学企業にとっても死活問題なのに
糞ワロタwww
このスレすごいんだな どっかのバKaveriは512SPで倍精度50GFLOPS未満
CPU1コアより遅いんだが 倍精度が高ければ許される
倍精度が低ければ、メモリが多くてもメモリ幅が多くても消費電力が低くても許されない 倍精度ハードフル実装マンセー
↓
トランジスタが増える
↓
ウエハ1枚あたりのチップ採れる数が減る(&歩止まり率も減る)
↓
製造コスト上がる
↓
(,,・´∀`・,,) 倍精度の話しないと死んじゃう人は点呼取ってくれる?
「倍精度でCUDA」スレと、「実在のGPUでCUDA」スレに分けた方がいいよね、多分。 件の赤い会社が帳簿の赤も解消しないのは、市場価格に対して製造原価が高すぎるからだよ
会社を潰れかけの状態にしてまで実装コストの高い倍精度ハードを実装するのは模範的な企業とは思えんがねえ
QuadroとTeslaのブランドが分かれてないようなもんだ まぁDLのおかげでteslaでも倍精度無しのものもあるけどね
ただ今後、DLはEyerissの方向で行きそうなんで
完全にDLは別ラインの製品でいくのか
GPU自体がEyerissの発展形になるか
GPUのデータ移動減らすとかレジスタファイルの階層化って研究はしてるようだけど 何とか倍精度不要論に持っていこうと必死な団子でした アホがいるね
特定のアプリケーションに特化して結果的に倍精度性能落とした製品は、その目的の用途から外れなきゃ何の問題もねーじゃん
ディープラーニングにもゲームにも倍精度の性能は必要ないから
倍精度が必要な用途にはK80やその後継あるいはPhiを買えばいい
科学技術演算用に設計されてないGPUに倍精度が必要だとかいうのはただのデムパ >>393
Teslaの単精度特化版は1スロット版のK8ならそうだけどK10なら発表は2012年で
ディープラーニング言い出す前の話
ダイ本体はGTX600番台そのものだからそれ専用に設計されてるわけではないしね
EyerissがよくわらんのだけどROPやテクスチャユニットも削ってるわけ? ここだけの話
K10はシェールガス探査特化モデルだよ 倍精度推しの池沼君がどんな世の中の役に立つソフト作ってて
LINPACKベンチ以外なんの実務に役立ってるのかは興味あるね 俺は必要だっていう人の言い分は理解してるし
それに合わせたチップを選べばいいと思うのだけどね
倍精度ハードウェアがTeslaはともかく一般消費者向けのGeForceに
載らないといけない理由はない
一般消費者が倍精度のスループットを見てGPUを選ぶトレンドを作りたいなら
まずてめえが倍精度載せる価値のあるソフト作って示せばいい 科学技術演算に倍精度は必要不可欠
はい論破完了
簡単だったね へー
それって
一般消費者向けの
GeForceに? 一般消費者が科学技術計算する必要あることがまったく示せていない TSMCのウェハ価格体系的に用途別に倍精度が削れる用途向けには削ったほうが
コストダウンになるからそうするでしょうよ
Maxwell世代のTeslaがないのは大型スパコン案件もないからっての大きいと思う Tesla M40「・・・」
Teala M4 「・・・」 普通に考えて一般向けGPUに倍精度はいらないよな
32bitと16bit浮動小数点しか使わねーよ
それでも贅沢なぐらい ほんとそれ
一般向けとしてはオーバースペックなくらい だからお前の発言のどこが論破なんだよ
論理破綻の略か? ここがCUDAスレだと何度言えば分かるんだ!団子
科学技術演算も出来ないGPUじゃ話にならんわ まあ、倍精度じゃないと、イタレーティブソルバーは収束しないわな。4倍精度ぐらいあれは、CG法で反復回数が極端に減るかもね。 今日の重力波観測にしたってあれは、理論屋が出したシミュレーションに従ってそれと同じデータを観測したわけだろ
そういうのが単精度/半精度で出せるわけねーだろ @No double precision, no life.
ADouble precision make me happy. ここではMaxwellやPascalはおよびじゃない CUDAをなんのために使うかだよなあ。
グラフィックスだけならOPENGLだけでええんのとちゃうか? >>427
だったら何
そんなのはGPUであつかう問題じゃないから GPUはグラフィックスを扱うのが本分なので
一般向けGPUは倍精度無しでも文句は言えない >>436
お前の文句なんかどうでもいい
GPGPUってのはこういうことのために使うんだぜ!?団子
-----------------------------------------------------------------------
645 : 名無しのひみつ@転載は禁止2016/02/12(金) 22:56:12.20 ID:oUDSwd87
>>626
別に2ちゃんだから適当に幾らでも変なことでも聞いていいんじゃね?
人間の頭で考えないと絶対にダメな部分が最も大きい
相対論も標準理論も量子力学も計算機だけでは絶対無理
計算機が答を出すというより、計算機は望遠鏡とか温度計のような「測定器の一種」という感覚が近いと思う
GPGPUは使いまくってるはずだけど、半ば問題に特化して設計するスパコンもあり
スパコンに関してはそれ自体がすごいというより、スパコン目当てに優秀な人が集まる効果のほうが大事という気がする、2020年のオリンピックみたいにスパコンという祭りに人が集まるから、良い方にも悪い方にも転ぶだろう
646 : 名無しのひみつ@転載は禁止2016/02/12(金) 22:57:56.72 ID:eXgf7vm3 結局何が言いたいの?
理想のものが欲しければ自分が計算機設計の専門家になればいいじゃん
手遅れかもしれないけど 自然言語処理や画像認識、いま産学連携で資金集まってるのはこういう分野
データサイエンス分野は文系の学者さんも研究に参画してるし応用範囲広いんだよね
何の金にもならない倍精度行列積で1位じゃないとダメなんですか?とかホントどうでもいい話 スパコンに倍精度演算性能が必要なのは誰も否定してないじゃん
個人のゲームやお絵かき用のパソコンにそれが必要かという話であって。
コピペで代弁してもらったつもりになってるのかもしれないけど
実質的に何の主張にもなってないわけよ >>442
>個人のゲームやお絵かき用のパソコンにそれが必要かという話であって。
ホントここが何スレか分かってない池沼
Async煽りをしてた団子が攻守入れ替わるとこれだからな > Async煽りをしてた団子が攻守入れ替わるとこれだからな
それこそお前の妄想じゃん
RadeonのGPGPUプログラミングの話をしたいなら超過疎のHSAスレへ > GPGPUは使いまくってるはずだけど、半ば問題に特化して設計するスパコンもあり
って引用しつつ、倍精度性能が不要な用途に特化する意義は認めないんだね ついでに攻められてるのがどっちなのかすら理解してない >>427の言ってるLIGOのリサーチに協力したのはうんこFireProではなくXeon Phiなんだよね
http://www.hpcwire.com/off-the-wire/taccs-stampede-supercomputer-assists-ligo-research/
重力演算は1格子点あたり小数点以下3桁程度の精度があれば十分といわれる分野
Stampedeに載ってるKNCは除算ユニットは遅く逆数・逆数平方根命令の
有効桁は14ビットだ。
まあこんだけあればニュートン法すら要らないわけよ。
倍精度で分解能を上げて遅い倍精度に単精度(実効精度14ビット)で格子点を
大量に捌くほうがより正しいシミュレーションができる。
ニュートン法不要ってのはニタドリ=サンの受け売りだが 14ビットは間違い。SkylakeXeonで使えるAVX-512Fのほうだね。
KNCのrcp23ps/rsqrt23psのスループットは23ビット。
つまり単精度の除算は乗算1回との合計2ステップで割り切りなしの十分な精度が得られる 倍精度必要論が否定されるから都合が悪いだけだろう?
重力演算のために長崎大がGeForce買いあさって作った激安スパコンだって
単精度特化仕様だったよ?
これもニタドリ=サンが協力している。古事記にも書いてある。 ゴードンベル賞で2度の受賞を経験したニタドリ=サンをして
天文学分野のシミュレーションは半精度程度の分解能があれば十分だって言ってるんだよ でもGRAPEは6倍精度実装してるじゃん
天文学と一纏めにしてるけど連星とかブラックホールが近接するような場合は高い精度が要求されるよね
逆に宇宙の大規模構造とかだと桁増やすより分解能上げた方が精度上がるのかもしれないけど 否定しない。
いずれにしてもどれも個人のパソコンに必要な要件ではない。
個人のパソコンのパーツを組み合わせて大規模演算ができるのはコストメリット的に重要だけど
個人のパソコンに倍精度が必要な理由の説明としては本末転倒です 結局自分じゃ使えもしないのに
倍精度倍精度連呼してるのが滑稽だってこった お前が行くべきスレ(HSAスレ)に行けば平和になるよ AMDな人が来てCUDAの話無・CUDA使っていない奴の雑談スレになったな
AMDのGPGPUスレ過疎すぎでこっちにきたんだろうが そもそも初代CUDAの時期に倍精度なんてサポートされてなかったし
メインストリームのGeForceが倍精度強かった時代なんて一度もないのに
いつから馬鹿の一つ覚えの倍精度連呼が始まったやら そもそもCUDAは科学技術演算を必要とする人のスレなのにな
いつの間にか糞団子はゲームに必要かどうかに論点をすり替え > そもそもCUDAは科学技術演算を必要とする人のスレなのにな
意味が分からない CUDAでハローワールドも書けないレベルのnoobが気取ってみても
頭の悪さは隠せないよ cudaで今nvidia的にイチオシの用途がdlで、half性能が効いてくるという奴じゃないの? と言うよりGPGPU自体がDLバブルに乗っかってるからな
ボリュームゾーンに合わせた製品仕様になるのは当然の流れ バイセイドマンセー「俺がボリュームゾーンだ!俺のためのGPUを作れ!」 DOE案件みたいな大口需要があるときだけ倍精度特化版作って
それ以外は民間需要重視でいくってのは間違いじゃないと思うよ
次のDOE案件はVoltaだ
逆に言うとそれまでは倍精度特化版は出ないということだな 認めろよ負け犬
お前らの需要は民間需要とみなされてないんだよ 残念だけどGeForceに倍精度が必要な理由を説くには
ゲームなどの目的で使って一般ユーザーが必要であることは絶対必要な条件だ
まさかそんな単純な理屈もわからないのか?
だからお前は俺ごときに言い負かされてるんだよ やっぱり次のスパコンはGPU入れてくるのかな。
CPUだけで今の100倍は難しいよね? Top500ランキングに使うLINPACKベンチは倍精度だからそれのためだけに必要なんだよ
ちなみにHPCGだと倍精度ユニットいくらあってももてあそぶから倍精度厨には逆風だね maxwellが発売されて2年たつが、pascalはいつごろ発売になるのだろうか。 pascalは今年
voltaは来年
pascalは大きな変更になり
voltaは小改良 倍精度の必要性すら認めない奴がどうしてこのスレにいる? dlはもうeyerissの系統でいくんだろうな
機会学習はもう大きな需要が在るのわかってるから
専用設計のチップでも儲かるのが見えてる
gpuはある種汎用的なHPC向けになるんだろうな
pascalで階層型のスケジューラやレジスタファイルが実装されてると
そのコントローラー(こんぱいら)ようにcpuが搭載され
レンダリングはタイルになりそうだけど そんな人は一人もいませんよ?
消費者向けGPUに倍精度が必要みたいなアホなこと言う負け犬君を馬鹿にしてるだけですよ Geforceで科学技術計算なんてことあるの?
2秒間の壁や1週間も100%負荷で連続使用したら過熱して壊れそう。 必要性もなにも倍精度演算自体は可能だろ。積和で単精度の1/32くらいの速度だけど。
それで不満な人間がCUDAにこだわる意味がわからん。
そもそもCUDA対応ハードで倍精度強かったのってfermiとkeplerの時期だけで、しかもコンシューマ向けで開放されたのもその極一部じゃない。 倍精度が必要な人は必要らしくマイノリティ税を払って4桁ドルのTeslaを使えばよろしい 防衛上の理由として
北朝鮮や中共がGeForceを買いあさって弾道計算に使われたら困るから
消費者向けGPUは倍精度を削るのは当然でしょう 倍精度演算を単精度の1/4より速く実行するのは基本的に無駄なはずなんだよね。
keplerは別回路で使わなければ完全に無駄だったし、fermiの1/2にもからくりがあって、整数演算器の方に少し回路を付け足せば単精度性能を倍にできるのを抑えた感じ。
納入先との契約で倍精度性能は単精度の1/2にすることみたいな縛りでもあったのかね。
1/4にしないでもっとけちれば次は1/16になるのが自然
なのにmaxwellは1/32 >>491
>防衛上の理由として
>
>北朝鮮や中共がGeForceを買いあさって弾道計算に使われたら困るから
>消費者向けGPUは倍精度を削るのは当然でしょう
中共とかなに時代錯誤なこと言ってんだよ
Nvidiaはファブレスで製造はTSMCだぞ
そもそも弾道計算にGPU要らんし 今存在している政権を時代錯誤と言ってみたり台湾と中国の区別がつかなかったり
いろいろ論外すぎる >>494
いやいや、中共が買い漁るとか言ってるのが時代錯誤なんだよ
ココム時代かよ
スパコンの現状見て今更「弾道計算に使われたら困るからGPUから倍精度削りました」とか寝言としてもありえ無いぞ >>494
あ、そうそう
TSMCに限らず台湾メーカーは当然中国工場持ってるからな
中台の政治はアレだが経済はズブズブなんだよ その寝言で米国政府が天河1号Aおよび2号のパーツ供給打ち切る命令を下してりゃ世話無いわ つーか今さらだよなぁ。
既に世界一のスパコン輸出した後で規制するとか、アメリカはアホなのか。 >>497
お前それ反論になって無いぞ
もしかしてGeForceに倍精度のせないのはマジで米政府が関与してるとか思ってるわけ?
お前正気か?
もう少しまともな奴だと思ってたが買い被り過ぎだったようだな 弾道計算wwww
不安定なGPUなんか使うわけがないwwww
そんなものは化石CPUで十分 2018年にかけて米国でknlとpascal,voltaでスパコンが更新されていくからな
天河はもう過去のものになる
日本はもう一番目指してないし 計算機なんて数増やせば良いだけだから単体性能よりコスパで1位を目指すべきだと思うんだが >>502
なんもわかってないな
ノード数の2乗に比例してインターコネクトは複雑になる
なんのために大きなスパコンが必要だと思う?
莫大なメモリ空間がないとできないことがあるからだ >>501
そーは言ってもあの国はコンシューマ向けGPUかき集めてでもスパコン作るから
現に初代天河1号でそれやったし
中国のPCゲーム市場は拡大しててGPUメーカーにとってもおいしい市場だ
軍事転用の懸念から輸出規制かけられたらNVIDIAだって大損なわけよ
だから最初から倍精度ハードは載せない
その理屈は最低限理解してほしいのだけど・・・ そもそも倍精度はいらないって持論言ってなかった?
それなら軍事用途でも倍精度なしでいいってことだよね?
そうすると、中国はいくらでもGeForceを軍事転用して作れるんじゃない?
倍精度載せないのが軍事転用を防ぐためってのは無理がないか? もちろんそれが一番の理由だなんて言う理由はない
どうせ倍精度なんて使えてもそれこそ豚に真珠を与えるがごとき日曜プログラマの道楽のために
NVIDIAは会社を傾ける気はないのは確かだろう ハイエンド以外はそうでしょうね。
でもハイエンドはTeslaとマスクを別にするほどメリットはないでしょう。
だからTeslaに搭載されている以上、ハイエンドにもつくだろうね。
ドライバや回路切断で無効化措置がとられるかもしれないけど。 GeForceに倍精度を載せないのはビジネス面からの理由で、
Teslaに倍精度が必要かどうかとはまた違う議論だよね
それからKeplar世代Teslaも最初1/12くらいだったのが1/3にまで上げられたのは
必要な人がいたからじゃないのかな。
4倍精度以上のライブラリがあることからも精度が必要な人はいると思うよ。
単品の精度はfloatでいいけど、累積誤差のせいで途中段階ではより高精度が必要という局面は多いんじゃないだろうか。
倍精度完全ゼロにしたらTeslaはこけると思う。 Nvidiaが倍精度どころか単精度ですら要らないんじゃね?路線になった理由は明快
(AMDの)MantleがDirectX12とVulkanの取り込まれたから
ゲーム機は完全にAMD一色になったから、もはやゲームはAMD中心に動いている
全てのAAAゲームはCS機のローレベルAPIで組まれ、その成果をPCに持っていく形が定着した
従ってPCゲームは今までのように互換性や異種GPU間の翻訳を行うのをやめた
DirectX12はそこの所が根本的に違う
しかし団子はそこの理解がないから最初っから倍精度不要論をぶっていればまだかわいげがあったが、
ゲームに倍精度は不要だ!と言い出したのは、Pascalの仕様が明らかになってからw
それまでの団子は、NVはPascalに倍精度路線を復活させる!(キリッ
の一点張りだった > GeForceに倍精度を載せないのはビジネス面からの理由で、
> Teslaに倍精度が必要かどうかとはまた違う議論だよね
それは誰も否定してる人はいない
> それからKeplar世代Teslaも最初1/12くらいだったのが1/3にまで上げられたのは
> 必要な人がいたからじゃないのかな。
大型スパコン案件があったから。
Maxwell世代にはなくて、民間需要を探したら結局単精度しか需要が無かった 団子>弾道計算には倍精度が必要!(キリッ
↑
マジでアホ 惑星探査機なんか発射時刻1秒のずれすら許されないレベル 弾道計算といったら大砲の弾やせいぜい弾道ミサイルといったイメージ
人工衛星は軌道計算というのが一般的だろう で
このhello worldも書けない倍精度バカは何が言いたいんだ スパコンは数増やせばいいとか言っちゃうレベルの超絶バカには理解できんだろう 呼ばれてもいないAMDファンボーイが勝手にわめいてるだけでしょ
Async Shader(笑)なんてスパコンでは全く需要がないんですよ。
繰り返すけど、需要がないんですよ。
何千並列のGPUが動く環境ではそもそも同じGPUで複数スレッドを
動かす必要性そのものが疑問だし本当に別々のスレッドを多並列で動かしたい用途には
60コア240スレッドのXeon Phi向けの問題だったりする。帯短襷長。 >>511
>大型スパコン案件があったから。
Titanとかだよね。ということはやっぱりむこうの研究者もより精度がほしいってことなんだな。
>Maxwell世代にはなくて、民間需要を探したら結局単精度しか需要が無かった
そもそも民間にGPGPUの需要があるのかも疑問だな。
大学以外でほとんど導入してる事例を見聞きしたことがない。
F1や飛行機の気流計算とかには大規模サーバに積んでシミュレーションしたりしてるのかな?
あとは建築における風の影響とかさ。
普通は世代ごとにプログラミングとテストし直す手間考えるとCPU増設するよって考えに至りそう。 >>510
pascalの仕様がいつ明らかになったんだ? >>514
>惑星探査機なんか発射時刻1秒のずれすら許されないレベル
その昔3ビットコンピュータの時代に米国は有人月探査をやってるんだが?
中共は弾道計算()にスパコンが必要なんだw 月がいつ惑星になったんだ
的が大きすぎてスイングバイのタイミング計算も必要ないぞ 月には本当に行ったのかね
よく帰ってこれたな
てかランディングよくできたな
よく月から離れられたな 核開発はおいそれと実験できないので
シミュレーション主体だろうけど >>523
弾道計算がいつの間にか惑星探査に成るよりマシ
スウィングバイの計算ですらスパコンいらないのに弾道計算にスパコンが必要な根拠を早く示せよw それは使用者がどの程度の精度を求めるかってだけじゃね
キロか
メートルか
センチか
ミリか
マイクロか
ナノか ipadのdawソフトだって64bit floatで合成できるし
サンプリンレートだって96khzがつかえる
で、聞いて違いがわかるのかって話だわ
さすがに32khz(昔のrolandのdtm音源とかね)と44.1khzはわかるけど 倍精度くんはDLで使われてる半精度の用途すら理解できてないと思うよ 科学技術計算で半精度が使えるのなんて、DLぐらいじゃねーの?
線形方程式を扱う問題は倍精度がないと使いもんになんないよ。
単精度で行けるのは、流体解析とかかな。 必要な部分だけ倍精度で
後は単精度って言う混合精度が普通ですけどね >>529
いいからお前は中共がスパコンで弾道計算()する根拠を示せよw
出来ないなら去れ そこいちいち突っ込むなよ
↓この池沼レスに突っ込んだやつ何人いるよ
427 :デフォルトの名無しさん [↓] :2016/02/12(金) 22:41:27.77 ID:10CGBYPs (2/2)
今日の重力波観測にしたってあれは、理論屋が出したシミュレーションに従ってそれと同じデータを観測したわけだろ
そういうのが単精度/半精度で出せるわけねーだろ 惑星探査機の打ち上げ時に1秒のズレが許されないのは、もし1秒のズレが
起こるとロケット側の打ち上げシーケンスはリアルタイムに変更出来ない以上、
生じたズレを探査機側でカバーする必要があり燃料を大量に消費するから
ロケット発射時の1秒分に相当するズレが軌道計算時の演算精度によって
生じると示されない限りは今回の話の流れで例として出す意味は無い >>533
話を逸らすな糞団子
GPUに倍精度が載らない理由としてお前が妄想で持ち上げた米政府陰謀論はどう始末すんだよ
いつの間にか惑星探査の話に誤魔化そうとしてドツボに嵌ってるし
どうやらお前には虚言癖があるようだな
それ精神疾患の初期症状だぞ こいついいかげんあたまわりーな
geforceにいらねぇってだけで
どこでgpuにいらねぇってかいてんだ
トットトキエロ無関係なamd信者 491 名前:,,・´∀`・,,)っ-○○○ [sage] :2016/02/14(日) 22:08:24.09 ID:vcWFNaoW
防衛上の理由として
北朝鮮や中共がGeForceを買いあさって弾道計算に使われたら困るから
消費者向けGPUは倍精度を削るのは当然でしょう >>537
てめえのケツは文体から丸見えなんだよ、糞団子! >>539
どこに弾道計算()って書いてあるんだ?
今更必死になって資料探してるのが笑える
糞団子www missile guidanceって書いてんですけど
本当にアホなんだ >>543
ミサイル誘導と弾道計算()をゴッチャにして無いか? お前らってCUDAスレでミサイルの話したいの?
うざいよ そもそも米国初のスパコンは300flopsで
弾道計算目的で作られたわけですけどね >>547
だからなに?
300flopsでいいならatomでも使えば?
弾道ミサイルと誘導ミサイルの区別もつかない知障が恥の上塗りしてどうする >>551
なんだよネタ切れか糞団子
早く弾道計算()の根拠を示せよwww ID:Rkm/CQ33
ひまだねこいつ
Teslaなんて買えるご身分じゃないことはわかったよ >>553
お前さぁ何で個人ユース前提で話進めてんの?貧乏F蘭学生か?
悪いけどプライベートでゴリゴリコード書かんから家ではオンボードチップのみだよ
んなこといいから早く弾道計算()のソース探して来いや AMDの株価がまたダダ下がりしてるよ
年内に潰れるんじゃないの?(煽 >>556
お前、俺がAMD信者だと思いたいわけ?
残念だが単なるツールに思い入れは無い
そもそもプライベートではインテルHDだしな
ぶっちゃけ、AMDでもNVでも倍精度でも単精度でもどうでもよろしい
糞団子がおかしな事言ってるから突っ込んでるんだよ
さぁ話を逸らさず、早く弾道計算のソース探して来いやwww 結局Asyncとか煽ってたのは団子だよ
そして株価ガーとまたぞろ論破された悔しさから関係ないバカッターぶり 糞団子は及びじゃねーんだよ
倍精度が嫌なら自作板帰れや 倍精度CUDAでAMDの倒産時期のシミュレーションをしようず この糞団子って奴はNV信者なのか?
かなりキショいな >>562
お前、先にやる事あんだろ
早く弾道計算()のソース探して来いや
糞団子 amd信者はスレのレベルを下げるの上手だよな
どこでも お前何なの?
団子団子団子って
この固ハン好きなの? 507 :,,・´∀`・,,)っ-○○○:2016/02/15(月) 20:06:18.36 ID:ZpISfAaX
もちろんそれが一番の理由だなんて言う理由はない
どうせ倍精度なんて使えてもそれこそ豚に真珠を与えるがごとき日曜プログラマの道楽のために
NVIDIAは会社を傾ける気はないのは確かだろう
団子いいこと言うな、その通りだよ
お前ごときが倍精度の存在意義なんかわかりゃしないんだから
そのままここから消えろ アンチの方が邪魔で迷惑というのは、嫌儲や嫌韓と変わらんな。 >>574
団子ははっきり言ってやらんと分からんアスペだから 確実に言えることは
こいつはCUDAより団子のほうが
興味があるようだってこと
CUDAスレなのにな
迷惑だから他でやって欲しい 団子はインテルNVマンセー馬鹿に過ぎない
Fermiの時代は倍精度の必要性を謳ってたのまた団子
現在倍精度不要論をぶってるのはこいつが工作員である証拠
くだスレにはまったく不要 >>575
団子はいくら言っても分からないようだから、そろそろ諦めてくれないか。
邪魔だし迷惑。
団子だけならNGできるんで、団子たたきはやめて、キッチリ完全スルーしてくれ。 >>579
>、団子たたきはやめて、キッチリ完全スルーしてくれ。
NGだけなら簡単だが、こいつは倍精度を否定してるんだよ
そして名無しで煽ってるから問題だろうが 俺がこのスレで名無しになったことなんて1度たりともない 一般消費者向けのGPUで倍精度の性能は必要ないとしか言ってないんだけどな俺は
なぜなら、一般消費者の用途で倍精度の性能が必要でないからだ
トートロジーを並べてるだけだけどこれは普遍の真理。
倍精度特化のTeslaが使いたい人は使えばいい
だからこそなおさらGeForceが倍精度に特化しなければいけない理由はない
いつ倍精度そのものを否定したんだ? >>583
はいはい
一般消費者向けのGPUから倍精度が削除されたのは米政府の陰謀なんだろw
お薬増やしおきますね〜 お前のGPGPUプロジェクトが
どのように許容誤差見積を出して
倍精度を必要とそっちを語った方が
反論になるだろう
出せ >>585
Teslaにも倍精度に特化してないモデルもある
倍精度が強いのはあくまで「スパコン用のTesla」だ
それの流用版のGeForceやQuadroも副次的に倍精度が強かったこともあったが
GeForceに倍精度の需要があったからではない
GeForceで倍精度が速くないと困るとか言ってる手合いはろくに研究予算もつかない
ゴミ研究者かなんかだと思ってるが
(そんなものに拘ってるからろくな論文が書けないんだろ) >>588
>GeForceで倍精度が速くないと困るとか言ってる手合いはろくに研究予算もつかない
>ゴミ研究者かなんかだと思ってるが
おいおい、GeForceは一般消費者向けが前提じゃなかったのか?
糞団子は論点ズレまくりでまともに議論が噛み合ってない だから学校から支給される予算じゃTeslaが買えないからGeForce使うしかないんでしょ あるいは赤ポスにつけなかったおちこぼれか
いずれにしても「一般」消費者でありはしない 所詮は自分が欲しいかどうかだけが基準点でもっと大きな市場のニーズを
理解できないアフォでしょう
大学院重点化の弊害ですね
就活失敗したようなレベルのゴミを研究室に引きこもらせたら
ますます世間知らずをこじらせるのも道理 論文博士という日本独自の制度もたいがいだがな
3年間も通って単位取得だけしかできず学校を放逐されたゴミにも
なまじ希望を与えるからタチが悪いというか >>589
糞団子は論点をずらすと言うより、負け犬になるといつも違う話をする&名無しでスレ汚しだから
NGだけでは駄目なんだよ
本当はDX12でNVがハブられた事実を認めたくないために架空のでっち上げ話をしてる
その証拠にMaxwell世代以前のNVはニタドリサン(笑い)の話をだして市販のNVGPUが大学研究機関
に納入された事を自慢してたからな
こいつは単なるアンチAMDのステマ部隊
454 :,,・´∀`・,,)っ-○○○:2016/02/13(土) 19:06:49.99 ID:dwcfM3wy
ゴードンベル賞で2度の受賞を経験したニタドリ=サンをして
天文学分野のシミュレーションは半精度程度の分解能があれば十分だって言ってるんだよ
453 :,,・´∀`・,,)っ-○○○:2016/02/13(土) 18:48:25.54 ID:dwcfM3wy
倍精度必要論が否定されるから都合が悪いだけだろう?
重力演算のために長崎大がGeForce買いあさって作った激安スパコンだって
単精度特化仕様だったよ?
これもニタドリ=サンが協力している。古事記にも書いてある。 >>588
特化の意味わかる?
単精度の1/3の速度じゃ特化じゃないから ただのAMDファンボーイだから
そもそもAMDはCUDAが使えないしCPUもAVXが半速のゴミしか出せないので
科学技術方面では全く役に立ちませんよ 価格の安いコモディティ製品本来の用途外に流用するのがGPGPUの始まり
用途外を前提に設計しろなんてのは片腹痛い話ですよ ちなみに仮数の単精度は24ビット、倍精度は53ビット
積和ユニットのの複雑度は仮数のビット数の二乗にほぼ比例するので
実際には倍精度積和ユニットの実装コストは単精度の5倍よりも多い >>597
倍精度特化じゃないCPUが1/2
イメージ的にはこれで普通
特化って言ったら単精度演算器を持たない、昔ながらのスーパーコンピューターみたいなやつ >>602
要は技術を知らないバカが無謀な要求をしていると自白しているわけだな。
そもそもCPUは倍精度特化で、単精度は搭載している回路の半分だけ利用しておまけでつけているくらいの勢い。
正確には倍精度以上から降りてきたという方が正しいような気もするが。 うん、じゃCPU使ってなよ病気の人
CUDA関係ないからもう来ないでね 逆に言うと倍精度をサポートしなければ同じトランジスタの実装コストで
5倍以上の単精度ユニットを実装できるということ
それがCPUとは別方向で進化してきたGPUの速さの根源
それを否定するのはGPUを否定するようなもの >>606
んな事はねぇよ。
回路を流用すれば1/4や1/16での
倍精度演算は殆ど追加コスト無く出来る。それを削ったからといって増やせる単精度性能は20%も無い。
まあ、これが大きいから1/4でなく1/16を選択するのが通常なのだけど。 >>607
>>606のいうサポートっていうのはフルスピードの倍精度っていう意味
もちろん単精度を複数回(or複数ユニット)回して倍精度を実現してもいいが
それが現状のMaxwellなどの在り様なわけで あったとしてもどういう役回りかわからんでしょ
ぶっちゃけると単精度で分割処理するとはいっても実際には
倍精度フォーマットをデノーマライズして複数の単精度データに変換するハードウェアと
演算結果を組み立てて倍精度フォーマットに書き戻すハードウェアも必要だからね
そのビットフィールドの変換をソフトで処理したら1/16すら出ない 本来遅いハードでもdpどころか64bit整数の乗算も1/9で済むはずなのにそこはサボって24bit乗算器を16bitとしてしか利用していないんだよな。 挙動の違いがあるとまずいから
同じユニット使ってるとか言う >>607
AMDのGCNアーキテクチャーのGPUは複数のコアを使うと倍精度の演算が出来るように設計してある
それがPS4/XBOに搭載されたものだから、ゲームは非常にGPGPUに強いんだよな
DX12とVulkanはGCNで最大性能を発揮できるように設計されたローレベルAPI
これが怖いからMaxwellは半精度FP16GPUになった
描画に特化してGPGPUが弱い >>614
ここはCUDA(言語)スレでGPGPUスレではない
AMDのGCNの宣伝したいなら、いまAMDはCUDAもある程度サポートしているらしいから、
AMDでCUDAしたときのベンチぐらい載せて宣伝してくれないか?
GPGPUのすばらしいぞ宣伝ならここでなくGPGPUスレでやれ HSA総合スレを盛り上げてやれよ
プログラムのプの字も書けないゴミには無理か? >>617
ここはディープラーニングスレじゃねーぞ糞団子 >>601
乗算の回路規模なんてぜんぜん大したこと無いって言ってなかった? >>618
プログラム書けないクソボケが居座る板でもないぞ
DQNのAMDファンボーイは自作PC板で妄想垂れ流してろクソボケ そもそもAMDのあれはCUDAを実行できるという理解は正しくない
CUDAのCソースをOpenCLに変換するトランスレータを作りましたという話で
PTXを実行できる環境を作るという話ではない >>622
>プログラム書けないクソボケが居座る板でもないぞ
ここお笑いですか?糞団子 倍精度連呼君は倍精度を使って何をやるのかすら提示してないよね >>624
AMDファンボーイはプログラムを書いたこともないから
AMDの開発環境の圧倒的なひどさを知らんのですよ
知識さえあればNVIDIA一択になる理由がわかる >>621
おまえの偽物か?
態度はおまえそのものだったけど >>625
いろいろと書かれてるのに、すべてお前が一般人は使わないと決めつけてスルーしてるだけ
プログラム板のCUDAスレッドなんだから一般人とは母集団が違うのは当たり前なんだが
お前が不要で、一般人の9割が不要でもこのスレ的には関係ない AMDのGPUは母集団そのものに入ってない
その点は理解してくれ だれも倍精度不要なんて書いてないのにな。
倍精度特化がコンシューマ向けには必要無いと言うだけで。 倍精度は必要な人は必要だけどAMDがこのスレに不要 >>632
てめえはすぐに論破されると単発で煽りやがる
こういう奴だからこそこっちは不必要にNGだけでなく相手をしないといけないんだよ馬鹿野郎!
コテならコテだけで通せ!糞団子!
それから、NVがこれからGPGPUを捨てる根拠となるソースだ
実際にここで科学技術演算をされてらっしゃる皆さんもご覧になった方がいい↓
http://camera-pc.com/?p=578
結局これもCS機をAMDが独占したためなんだよな 科学技術計算でAMDのGPUが上げた実績なんてないし
NVIDIAがHPCを重視する姿勢にも変わりはない
そもそもMaxwell世代に大型のスパコン案件が無かったから倍精度特化版を作ってないだけで ちなみにその個人ブログの主は科学技術計算に対してあまりに無知でお話にならない
ただカタログスペックのFLOPS数値を上げてヨシとするものではないのだわ
大規模システムはノード間のデータ転送帯域がものすごい重要で
NVLinkもその一環
AMDはカタログスペックありきで中身が無いし
実際2017〜18年の米国大型スパコン案件を何一つとして獲得できてない http://ascii.jp/elem/000/001/046/1046956/index-3.html
>Keplerの世代では倍精度浮動小数点のサポートの有無がこれにあたったが、
>Pascalはこの点ではMaxwellを継承しており、
>倍精度浮動小数点の性能は非常に低いと思われるので、ここでの差別化は難しい。
Pascalはここのスレ住人にとっては大変残念な事になるだろうな
その理由は先述したとおり、AMD準拠のDirectX12が策定された事でゲーム用GPUでNVは必死になって
電力効率の優位で喰らいつこうとしているためで描画特化でDX12世代に必須になるGPGPUを弱体化させた
トレードオフがある
倍精度とNVの関係は残念ながら離婚前の別居状態のようなさめた状態 相変わらず知恵おくれだねえ
その記事を読んでもPascalの倍精度特化版が存在しないということは読み取れないがね
大型案件がつけばMaxwellの倍精度特化版も出す用意はあったくらいだし
少なくともPascalには客がついている 倍精度の性能はスパコンでは今後も重要だけど
足回りの性能に見合わないほどのカタログスペックが必要かと問われれば
「必要ない」のが現実なわけで。
足回りの強化もなしにただ倍精度ユニットだけ並べておけばFLOPS数が稼げる
LINPACKベンチの実用性に疑問符が付くご時世ですからね
http://news.mynavi.jp/articles/2014/12/25/nvidia_coral/ AMDが無様な全敗を喫した2017-18年の案件はDOEではPre-Exa世代という位置づけで
Exaにむけたスケーラビリティのための足回りの
IntelのOmni-Pathしかり、選定されたメーカーの製品には選定されるだけの裏付けがある
AMDにはインターコネクトに対して2018年にただちに用意できる対案が何もない。
そのスパコンの選考がとっくに終わってからこれから倍精度強化だとか
何も理解してないというか既に周回遅れ。 病気の人は倍精度特化のcpu使うって宣言してんだから
関係のないcudaスレにあらしにくんなよ Polarisは結局スパコンのデザインウィンを未だ1件たりとも取れてないんだよな
全てが無駄だったね 野球ファンがどっちの贔屓チームの方が凄いか喧嘩してるようなアホらしさw たとえるなら阪神ファンがオフシーズンに東京ドームに居座って六甲颪を大合唱するレベルの
あつかましさ >>634
勝手にくそ団子と同一視するな。
不愉快だ。 米国のエクサスケールではcuda使わないんだけどね http://www.livevolume.com/?page_id=14&lang=en
このソフト動かしたいんだけど、nvidiaのカードが必要なんですか?
すいません、プログラムのプの字も知らない初心者です
pentium N3700でGPUがintel graphicなんですけど
やっぱりこの環境で動かすことはできないのでしょうか…
現状では「GLSL compile error」と出て描画ができません
もしスレ違いな質問だったら本当にごめんなさい >>651
Minimum system requirementsにNvidiaとあるし、
そして、実際に動かないとなるとnvidiaつけるしかないだろ 果たしてそのマシンにGPUなんてつけられるだろうか?
ノートならG-Tuneあたりのゲーミングノート買うのが一番手ごろじゃなかろうか
(メーカー製ノートでもGeForce搭載機はあんまり見かけない) スパコン案件の有無が倍精度版の有無を決めるというのは間違いないようで
NVIDIA Pascal based GP100 GPU to boast of 12TFLOPS SP & 4 TFLOPS DP performance
http://vr-zone.com/articles/nvidia-pascal-based-gp100-gpu-boast-12tflops-sp-4-tflops-dp-performance/106203.html
とりあえずTesla購入予定の方は何も心配いらないということです
悔しがるのはTeslaが買えない貧乏研究者とAMD厨だけでいいですよ 研究者しか倍精度を使わないと思ってるところが団子クオリティー GeForceは個人用のGPUであって倍精度で科学技術演算のためのものじゃないんですよ
倍精度の性能に正規の対価を払おうとしない悪質ユーザーの梯子を外されたからって
何が問題なんですかねえ 倍精度ユニットの実装コストはロハではないので
メインで使う単精度の演算性能向上や電力効率の改善に振ったほうが
大多数のユーザーのためなんですよ
実質TeslaをGeForceの価格で売れと言ってるような倍精度乞食の言葉に
耳を傾ける必要はないのデス NGにしてるから何かいてあるか知らないけど、どうせ相手にする価値がないことは>>659
から読み取れるw
このスレは研究者もヲチしてるだろうから言っておくけど、もうNVは倍精度は捨てるね
Felmi時代とは大違い
DirectX12(ウィンドウズ10)になるとAMDだけがローレベルAPIの恩恵に与れることを知ってるから、
Maxwell世代で素早くGPGPU切りの変わり身を演じた
今後は半精度と単精度の混合GPUばかりでCUDAはどんどん腐ってくるw >>661
てめえをNGにして非難合戦のスレ汚しをしたくないのにノコノコ単発で出てくるな!
糞団子! >>660
なんかAMDが勝ち組でNVが負け組の様に書いてあるけど、実際は真逆だよ
AMDはゲーム機でなんとか持ちこたえてる状況で、NVはDL特需で歴史的な好況
GPGPUはCUDAがデファクトになりつつあるからAMDはほぼ詰んでる
まぁ、先の事は判らないけどね。 >>662
お前みたいな負け犬アム厨が自作PC板から出てこなきゃ平和なんだけど >>663 >>663
まだ単発自演で俺をやり込めようってか?糞団子!
NVはもはや半精度でしか生き残りが出来ないんだよ!
このスレは倍精度の行方が一番大事なんだからてめえがいることが一番スレを汚すんだ馬鹿野郎! 倍精度の行方なら団子がソース付きで書いたけどね
結局NVIDIAがHPCから撤退なんてのは出鱈目な言いがかりと ほれ
NVIDIA Pascal based GP100 GPU to boast of 12TFLOPS SP & 4 TFLOPS DP performance
http://vr-zone.com/articles/nvidia-pascal-based-gp100-gpu-boast-12tflops-sp-4-tflops-dp-performance/106203.html
そらあんだけのシェアとエンプラ向けの大躍進あればコンシューマ向けとHPC向けで
それぞれに合わせた演算ユニット構成のカスタマイズ費用くらい捻出できるだろう
そしてここはAMDについて語るスレでもNVIDIAをネガキャンするスレでもない
誰が場所をわきまえていないのか自覚したら? >>665
おいおい糞団子と一緒にするなよ
ここはCUDAスレだぞ
倍精度の行方とかどうでもよろしい
お前自分が糞団子レベルになってる事に気付けよ 浮動小数点数の精度なんて
コード組む側からすればどれも一緒だよ
倍精度だと収束するが
単精度だと発散するとか
それ単にコードの質が悪いだけだしね
何がそんなに気になるのかわからん FireProがHPC向けでまったといっていいほど使われないのは
分岐粒度が荒過ぎて実効でスペックの半分も出ないとか
特定用途向けの商用ライブラリが用意されてないのか
それ以前の問題としてランタイムが頻繁にクラッシュするとか
結局のところトータルの問題なんだよね
単精度特化の一般のGPUで多倍精度使う場合は遅い倍精度を使うより
単精度DFTのほうがはるかに速い 中小企業や個人開発のプログラムにもCUDA開発の無償支援してくれる
そういう実績があるからCUDAが鉄板になってるわけで
そういうことをやったためしもなければ大赤字でやれる予算も割けないAMDに
期待できることは何一つとしてない、これが現実
うちの商品はほぼ整数しか使わないだから倍精度実数が速いかどうかなんて
微塵も関係ない 糞団子はコテ使ってると相手してくれないから寂しくて単発コテはずし作戦に出たのか?
NVはどう考えても半精度中心のグラフィックス/深層学習中心に経営資源をシフトしてるんだがw
http://ascii.jp/elem/000/001/046/1046956/index-3.html
>Pascalはこの点ではMaxwellを継承しており、倍精度浮動小数点の性能は非常に低いと思われるので
>Pascalはこの点ではMaxwellを継承しており、倍精度浮動小数点の性能は非常に低いと思われるので
↑
大事な事だから二度言うぞ
>つまりPK100は半精度浮動小数点をサポートするが、PK104はこれをサポートしない(あるいは回路を持たない)可能性がある。
↑
プロセスルール次第だが半精度に重心を置いていることは明白だな half2floatをPascalでは強化していくから、結局AMD対NVの戦いは「half2float」の描画特化対
「Async」の戦いになっていくんだろうね 今後この馬鹿に一切かまうな
かまうから喜んで荒らし続ける どっかのニュースのコピペばかりで
うんざりするよ
自分の引き出しは空だと言っているようなものだ >>674
AMDなんか興味ない
お前は糞団子と同レベルの荒らしだ
失せろ 一般利用でCUDAが何に使えるのか
ネタ出せよ
ゲームで使えんの? >>677
お前はAMDに一番興味あるだろうが糞団子!
お前が去ればこのスレは一気にクリーンになるんだよボケ
お前が一番危機感抱いてるのは倍精度をなくす方向のNVに対して非難が集中する事だけだ
だからこそお前が一番イラネーんだよ屑 >>681
お前はいったい何と戦ってるんだ?
日和ったの意味も違うしな
ここのみんなはAMDにもNVにも思い入れはないから
迷惑だから糞団子と一緒に消えてくれ 団子よりヤバイ真性のキチガイが団子以外も団子に見えているらしい。 >>683
ほんと真性基地だよな。日中でも必死カキコっていうのもな
この人、プログラムに関する話がでないことからして、プログラムやったことないんだろうな >>671
分岐粒度が荒すぎってCPUとの比較なら未だしも、NVIDIAとの比較だったら普通そんなに差は出ないぞ。
差が出るのは64のうち前半もしくは後半の32がほぼ埋まっているのに、反対の32が完全に空みたいな状況が頻発する場合だけ。
まあCUDAからのベタ移植ならあり得ないこともないのだろうが。 プログラムの話しようぜ
そうすれば団子その他プログラム書けない人が消えるだろうから >>686
まあそうだな
今のnVidia32、AMD64の設定は回路効率と性能効率を考えて美味しいところなんだろうか?
ただの惰性? 将来的に4にするっていこうだが
実行レベルでは
warpは32のままだろうけど >>680-683 >>685
糞団子は黙ってろ!
文句があるならコテ付けろやコラ!
テメエが鬱陶しいんじゃ、ボケ!
>>686
それも糞団子の自演だよ
NVIDIAの優位性は3dfxを買収した事で得られたソフトノウハウだな
自作板でもそうだけど、よくよく考えるとおかしな文面でごまかすのが団子
たいてい論破されるとどうでもいい専門用語を並べてイカスミ作戦に出る >>690
あの程度の文面を理解できないほど技術に疎いのだな。
そもそも団子認定連呼廚のスタンスがよくわからん。 団子既に消えてるじゃん?
イカスミが煽るから戻って来るんだよ nvidia ってC++AMPって力いれてるの?
CUDAが主力なのはわかるんだが。
OpenMPのようにC++AMPがいいところまでいってからtoolkitにマージするつもりなのか
それともC++AMPが自然消滅するのを待っているのかわからん。 C++ AMPは事実上終了したよ
MSはいまNVIDIA ThrustをベースにしたParrarel STLに注力中 >>693
NVはC++AMPとか言うより倍精度をボリュームゾーンから外しにかかってるな
Pascalには全く期待できんよ
IBMの場合は大口契約だから安く調達するだろうけど、それだけ
いずれGPGPU用(テスラ)は値段を上げざるを得ないだろうなあ それなら死んでないでしょう。
スラストがCUDA用のライブラリと同様に
C++AMPで使うためのものだろうから。 Pallarel STLのバックエンドはC++AMPに限らずOpenACCやOpenMP(4.0), TBBでもいい
NVIDIAが推してるのはOpenACC, IntelはOpenMP >>692
俺3dfxを知らないからググッタよ。イカスミ、かなりの歳なひとなんだな
C++AMPを使っている奴どれぐらいいるんだ? >>702
ミス指摘してさらにひどくミスすんなって >>703
最初の頃のテスラは円高もあり、20万位だったが、いまじゃ80万くらいになってもうた。まあA30とかは120万くらいしたんだけど。 >>705
そうだOpenACCだ。勘違いしてたよ。
C++AMPではなかった。
MSだけだとC++AMPはきびしいだろうね。
まだOpenACCのほうがめはあると思う >>709
その頃とくらべると開発環境がかくだんに良くなってる。
HWのコストだけでなく、そういうの含んでの価格だろうね。 >>690
後半はまったくその通りだな団子
どこかで調べた用語を使うが、意味を良く分からずに使うからおかしな日本語になる nvはpgi買収してopenaccのコンパイラも力入れてるよ
doeのエクサスケールで使うのもcudaじゃなくてopenaccだし 「,,・´∀`・,,)っ-○○○」って鳥屋氏でしょ。
悪いけどここに同じレベルの人はいないんじゃない。 NVの工作員
GPGPUはVLIW時代のATIGPUを持ち出してネガキャン 俺がいてもいなくてもCUDAプログラミングの話ができない無能集団のスレ 囲碁ソフトにディープラーニングが使われ始めてるんだが、
CUDA対応ソフトもそろそろ来るかもしれん。
一般人にも身近な技術になるかもね。 そのうち教えられたことだけしかコーディングできないクソPGもAIにとって代わられるんだろうな 鳥屋ってただのサンデープログラマーだからな
知ったか決めてりゃ嫌われる >>724
鳥屋氏の書いたコード見たことないでしょw
申し訳ないけど、ここにいる人たちとは次元が違うよ。
頭が良すぎて理解されないタイプ。 >>719
,,・´∀`・,,)っ-○◎●
と持っている団子が違うが同じ人なのか? 君はその名前欄に●入れてそのまま表示されるか試したことがあるかい? 別人なのか?
,,・´∀`・,,)っ-○◎●
が持っているのは焼き鳥
で
,,・´∀`・,,)っ-○○○
が持っているのは団子
ということか 鳥屋氏が両方つかってるの見たことあるから本人なんじゃない? なら、次は,,・´∀`・,,)っ-○◎● で書き込んでくれるな っていうか●が名前欄に使えなくなった経緯も知らないんじゃ、鳥屋氏のことを知らなくても無理は無いね。 >>733
俺もだが団子と言っている奴は鳥屋時代のこと知らないだろ
>>725の言うような奴ならすごいCUDAコード披露して欲しいが
でも、雑談に必死なコテの荒らしって感じだからな。
とりあえず、鳥屋時代のコードどっかにないのか? 見てみたい気がする 糞団子がどんなコード書くか知らんけど、弾道計算の件でシドロモドロになった事実は消えんぞ
所詮その程度の知的レベル >>713
速度自体は生でcudaやったほうが速いらしいから
これは開発効率を優先してPGIなのかね?
あるいは現時点ではいまいちでも今後の成長を鑑みてとかさ
さすがにdod,doeでも世代ごとのコード書き換えは負担なんだろうか キャッシュの増量やシャッフルIA、HyperQ,DP等を全部こみこみで考えると
1コアあたりの実効性能(flaot)はfermi, keplar, maxwellと上がってると考えていいのかね?
flops/core/clockではなくてflops/coreという意味で、かつ理論値での比較ではなくていかに実測値がでるかという意味で。 >>738
あれ、中村氏もそのAAつかってなかったっけ?
とにかく鳥屋氏(=中村氏)はプログラミングの技術も問題に対するアプローチの仕方も異次元の存在だったな。 やっぱり団子って複数人いるよね?
すごく詳しい話をするとおもったら、なんか急に低レベルだったり、
てにをはが変わったりする。 > あれ、中村氏もそのAAつかってなかったっけ?
なにそれ見たことない 結局鳥屋が馬鹿団子の振りしてるだけかもだが、詐欺の天才は結局馬鹿より有害
NVのやる事は何でもマンセーではスレ汚しとしか言いようがない
CS機をAMDが独占したからNVは倍精度に手を出せなくなった、と言うといつも発狂
これでは冷静な議論が出来ない おまえはcuda関係ないんだからくるなよ
かまってちゃん そっち界隈しばらくチェックしてなかったけど海外のハカー面白いことやってるね。
ちょうどCUDAに特化した話題だ。
MaxwellのLOP3.LUT命令を使ったビット演算の高速化テクニックについて興味ある人は
一度この講演ビデオ観ておくといいと思う。
http://www.irongeek.com/i.php?page=videos/passwordscon2014/bitslice-des-with-lop3lut-steve-thomas
実際問題AVX-512のvpternlogもLOP3.LUTと引数互換なんだわ
ぶっちゃけソース置換するだけで移植可能 >>745
鳥屋=糞団子!テメエこそCUDAの何が分かるんだボケ! 単発のお時間がきまちゅたね〜
鳥屋=糞団子が自作板からこっちのスレの巡回時間が来たみたいw
本当に鳥屋が使える天才なら2ちゃん談義なんかしねえな
10年前ならゲーム業界か、少し前なら京に関わるか
或いはアメリカでイラン向けバックドアを仕掛けるお仕事か…
どっちにしろヒマを持て余すようではお里は知れてる https://community.amd.com/thread/197348
AMDのGPUにはLOP3.LUT相当の命令がないことにユーザーからも不満が噴出しているようです 単発でも見事につられまちゅね〜wwwww
団子=鳥屋ちゃんピンチでちゅね〜wwwwww なんだかすごく質問しにくい流れなんですがいいですか?
バイトニックソートのプログラムを作ってみまして、
要素数512を境にGPUの方が速くなったのは良かったんですが、
要素数8192くらいからGPUの時間増加の傾きがCPUより悪化してしまって、
1048576ともなるとほとんど大差ないくらいになってしまいました。
これが普通なのか、おかしいのか、
おかしいなら原因として何がありそうか教えてほしいです。
一応可能なところはshared memory使ってますが、要素数256以降は
比較の組がとびとびになるはずなのでglobal memoryを直接読んでます。
以下、ソートを10000回繰り返した際の1回あたりの所要時間です。
要素数 CPU[s] GPU[s] CPU/GPU
2 0.00000 0.00004 0.004
4 0.00000 0.00005 0.008
8 0.00000 0.00005 0.011
16 0.00000 0.00006 0.020
32 0.00000 0.00007 0.036
64 0.00001 0.00008 0.099
128 0.00002 0.00009 0.229
256 0.00005 0.00010 0.522
512 0.00013 0.00013 1.027
1024 0.00031 0.00016 1.909
2048 0.00076 0.00025 2.975
4096 0.00179 0.00036 4.919
8192 0.00423 0.00068 6.248
16384 0.00933 0.00202 4.627
32768 0.02142 0.00637 3.363
65536 0.04903 0.01776 2.761
131072 0.10829 0.05003 2.165
262144 0.24928 0.13078 1.906
524288 0.54765 0.32437 1.688
1048576 1.22586 0.77101 1.590 >>755
バードショップ氏と話しがしたいなら @chapuni に凸ってきたら? バイトニックソートはキャッシュメモリで収まる範囲内の小さなデータ集合に対して有効
サイズが大きくなるとGPUのローカルメモリから外れて大幅にパフォーマンスが落ちる
CPUのほうが低レイテンシ・大容量のキャッシュ積んでる分だけ有利になるわけだ
CPUのキャッシュに収まらない量になれば再びGPUのほうが上回るはずだ
まあそのレベルに達する段階でO(n log n)のソートの検討の余地があるが >>754
Nvidiaならおかしいことはよくある。気にするな
GPGPUでおかしいことをさけたいなら倍精度の自慢のAMDを使う 質問者がシェアードメモリって言ってるのにキャッシュとかいい加減なこと言ってる
シェアードメモリに言及しないなんてあり得ない
そもそもCUDA使ったことないんだろうな
やっぱ糞団子は低レベルだな いったい何年前の話をしてるんだ
Fermiから階層型キャッシュ導入してるが
http://www.gdep.jp/column/view/2 GPUについてはどっちにも解せる「ローカルメモリ」と言い換えてる
ちゃんと読めない子は呼吸を止めて1秒あなた新鮮だけど煮魚 >>759
結局団子を買いかぶってた奴は団子自身だろwwwwwwwwww
糞っぷりを言い触らされたくないチンケなプライドが虚勢を張りたくなる理由なんだろう GMにマッピングされたメモリは一定の局所性があればキャッシュにマッピングされる
常識が通じないバカが勝手に自爆してるな >>757
キャッシュメモリ……
すいません、当方キャッシュについては不勉強でして、よくわかってません。
多分L1とかL2とか書いてあるやつだと思うのですが、手元の参考書が古い(6年前)もので……。
DeviceQueryだとL2 Cache Size: 524288 bytes って出るのでこれでしょうか。
Total amount of shared memory per block: 49152 bytes とか
Total number of registers available per block: 65536 とかもあります。
英語版Wikiの平均計算時間O(log(n)^2)parallel timeもよくわかってなかったりします。
log(n^2)なのか(log n)^2なのかよくわからんとです。
>>758
倍精度だと性能がうんたらとややこしかったので、単精度整数のソートなんです。
正確には同サイズの単精度整数配列2つを作り、
片方を基準にソートして、もう片方を一緒に並び替えてます。
構造体にしてもいいんですが、一緒に並び替えたほうだけが後々必要なデータでして。
最終的にはソートした結果を使って色々処理するためのプログラムなんですが、
現状だとこのソートで全体の8割くらい時間食っちゃってるんですよね。
ソートしなかったらもっと遅い(というかO(n^2)になる)ので無意味ってわけでもないんですが。
忘れてましたがGeForce GTX 760を使ってます。
以前機会があったので、別のGPU(GTX570だった気がする)でほぼ同じ計算させてみたこともあるんですが、
そっちだともっとマシでした。CPUより確実に速い、と胸を張って言えるくらい。
スペックの違いを加味しても、うちのGPUだけ下手したらCPUより遅いってのも変だなあ、と。 >>766
参照先読みました。キャッシュの話だったので自信はありませんが、
Fermi(GX570)世代のキャッシュは容量超えると(多分)通信の速いL1D等に退避されてたけど、
Kepler(GTX760)世代だと通信の遅いglobalに退避されるようになった、ってことですか?
さすがにアーキテクチャの仕様だと自分のスキルじゃ対応できません……。
>>767
答えになってないというか、私の知らない範囲の回答だったんでよくわからなかった感じです。
言ってることは766さんの参照先と多分同じなんだと思います。自信はありませんけど。
で、みなさん改善する方法はないでしょうか……。
大規模計算にウンタラ使うとカンタラがヤバいから避ける、みたいなことありませんか?
オプションにホニャララって追加すれば設定が変わる、とか(実はオプションもよくわかってない)。
大規模の時、キャッシュ代わりにあえてsharedを定義しても意味ないのかな…普通はないよな… そもそもキャッシュはシェアードメモリの退避先ではなくグローバルメモリの一部の読み書きを高速化するためのもの
シェアードメモリに収まらないデータのソートで遅くなる問題なのでシェアードメモリは全く関係ない
アホが自爆すんなよ データ全体がキャッシュに収まらない場合はキャッシュローカリティを意識した演算順序の最適化をする
これはCPUでもGPUでも理屈は変わらない 自分の答えを示さず謂れのないイチャモン付けに必死になる奴ほど
スレにとって邪魔な存在はないね >>770-771
キャッシュローカリティって何さ…となったので、Google先生に勧められた
wikiの『参照の局所性』を読んで、ちょっとわかった気がします。
自分が作ったバイトニックソートは大規模データの場合、
各threadがほとんど毎回前回と異なる&近傍でもないメモリ参照することになるので
時間的にも空間的にも最適化ができてないみたいです。
ない知恵絞ってif文減らして実質1つにした結果がこれだよ!
あ、いや、thread単位はそうでも、warpとかblock単位で見れば近傍データなのか…?これはバンクの話か…?
実装のややこしさを飲み込めばシェアソートの方が参照の局所性は大きくしやすいのかな…?
まだよくわかってないことが1つありまして、
プログラム上では何万何十万の計算を並列に扱ったとしても、
物理的にはハード側の制限があるわけじゃないですか。
GTX760ならCUDAコア1152基なわけで、一万を並列化しても
実際には同時じゃなくて9回くらいに分けて順に処理する、みたいな。
となると、その9回で各CUDAコアは別のメモリを参照するわけなんですが、
この時前回のキャッシュって残るんでしょうか? 残らない…ですよね?
それともカーネルとしての処理が終わるまでは残るんでしょうか?
はたまたプログラムが終わるまでは念のため残しておく、とか?
キャッシュの利用自体が悪いんじゃなくて、
キャッシュにも収まらないデータ量がよくないってのはわかってるんですが、
なんかこのあたりの理解がまだ曖昧なんです。 だらだら長文書いたわりに今自分の気になることからは逆に離れた気がする。
手っ取り早く要約すると、CPU↔メモリ(キャッシュ)の速度は
CUDAコア↔各種メモリの速度のどの当たりに相当するのか、でしょうか?
GPU内ならレジスタ>shared(キャッシュ)>globalの順で速いとかは聞きますけど、
もう一歩踏み込んで『この不等号の列に割り込ませるなら』CPUの速度はこのあたり、という指標が見つからないんですよね。
もちろん機種によって千差万別だとは思いますが、
傾向としてはこのくらいかなあ、みたいな話でもいいので教えていただけたら幸いです。
あ、もちろん>>774の質問も答えていただけると助かります。
キャッシュがよくわかってないのも確かなので。 後続の処理が走ってもデータが保持されるかどうか?
これについては実装依存としかいいようがない
なにせGPUはCPUと比べてもキャッシュの容量は小さいので
いつまでも今処理してないデータを保持し続ける理由も乏しい。
処理順序の並べ替えが可能なら、データアクセスの局所性に合わせて並べ替えることをお勧めする。 この会話を分かりやすく噛み砕くとこういう感じか
754
GPUの並列化の最適化を教えてください
団子
シェアードメモリーガー
754
あの…ですからデータの並列化を…
団子
データアクセスの局所性ガー
754
・・・・(その並び替えが分からないんだけどなー)・・・・まあいいです >>775
うーん、速さといってもデータ転送なのかレイテンシなのかで答えようが無いね。
レイテンシに関してはコアに近いところだとCPUの方がクロックが高いから圧倒的にCPUだね。 >>761
いやいや、CUDAユーザーならローカルメモリなんて言い方はしない
そもそもローカルではないし
SRAMはキャッシュとシェアードだけ
ボロ出しまくり糞団子 >>776
なるほど。基本的に残ってないと考えた方がよさそうですね。
バイトニックソートは綺麗に並列化できちゃうので、並び替えは無理っぽいです。
>>778
Occupancy確認しました。ビンゴでした。4%とかなにそれクソ使えてない。
というか、よくよく確認したらsharedもまだまだ余裕があったので、
要素数256でsharedの利用を切り替えていたのを2048まで拡張しました。
8192越えたあたりで悪化し始める、という性質こそ変わりませんでしたが
sharedの恩恵を受けられる範囲が増えたおかげか悪化しつつも
約100万要素でもCPUより7倍強速くすることができました。これでOccupancyは50%。
100%にするにはsharedのサイズの設定を変えなきゃならないけどやり方わかんない…。
しかし、こんなアホみたいな理由で遅くなっていたなんて。
みなさまにはご迷惑おかけしました……。
>>780
やっぱりCPUですか。GPUの性能をフルに発揮させるには
それなりに大きい規模を並列化して誤魔化さないといけないわけですね。 名前入れ忘れたりsage忘れたりしてますが>>782は754です。
IDで察してくださいすみません。 迷アドヴァイスを見るに見兼ねたスレ住民のおかげで助かりましたとさめでたしめでたし >>782
シェアードをローカルとか言う奴の事は聞かないように
GPUはシェアードメモリを使ってなんぼだから っていうかgpgpuでソートが
そもそも間違ってるよなあ てかFermi以降のSMはキャッシュを構成するSRAMの一部を
GMとは独立した空間のスクラッチパッドメモリとして割り当ててるだけなんだけどね
SMXが実質的なコアとみなすならシェアードメモリって別にシェアードでもなんでもないぞ ますます支離滅裂だな
シェアードメモリがシェアードじゃないとかCUDAの概念を全否定する発言じゃねぇか
つまりお前はCUDAの事なんか全く判ってない
ここはCUDAのスレだぞ
そろそろ出て行けや糞団子 まったく技術論になってないぞ
NVIDIAの用語がインチキ用語なのは今に始まったことじゃない
AMDのBulldozerのコア詐称以前からのコア数水増しやってるのがNVIDIA
命令ストリームレベルで見ればSMXが本来のコアで
CUDA Coreと呼んでるものはSIMDユニットの1エレメントにすぎない 確かにCUDAというかNVIDIAの定義しているシェアードメモリは一般的なシェアードメモリとは違うね。どちらかというとローカルメモリといった方が正しいかもしれん。 同じコアだからこそメモリを介さずレジスタ上でシフト・シャッフルができる(Kepler以降)
CUDAのSIMTとは単純に32要素1024ビットのSIMDを32スレッドと読み替えてるだけ
この点はG8x世代から変わらないね 自演会話すれば糞団子が味噌団子に変化するじゃないぞ?団子さんよお 自演認定しても頭がよくなりませんよ、他所者君
CUDAをちゃんと勉強してる人にはみんな知ってる事実だ また夜中の自演が始まったか
バレてないと思うところが糞団子だね
何度も言うがここはCUDAのスレ
僕の考えたシェアードメモリ論ならチラシの裏に書いとけ CUDAでいうシェアードメモリはCPUでいう一次キャッシュだし、CELLで言うところのLSだから、SMメモリとかコアメモリとかにしておけばよかったのにね。
CUDAの黎明期はGPUをスパコンって呼ぶ戦略があったから、1スレッドを一コアに見立て、共有メモリだと言っていたのかもしれないね。 >>795
それを言うならオンチップメモリ
少なくともソフトレイヤーではキャッシュとシェアードは別物
そもそもMaxwellではデータバスも異なる
んでここはCUDAのスレだ
スペオタと糞団子は去れ
>>792
cudaDeviceSetCacheConfigは試してみたのかな? >>792
俺に聞かれてもなw
Occupancy確認してからスムースになってきたみたいだし、使わなくても順調ならそれでいいという
考えもある ソフトレベルって言ってもなぁ
.cuやPTXのレベルならそうだろうがcubinレベルでみて言ってるんじゃないんだろ?
実際Sharedなんて言ってもSMX(実質上のコア)の中でしか参照できないんじゃ名前に偽りありとしか
SPMDをオレオレ用語で置き換えた概念で説明すればSharedなんだろうが物理的にはCellのSPUのLSと同じトポロジにすぎん 逆に複数のsmx間で共有されてたら
くそ遅いだろう
ローカルだから他のsmx気にしなくていいわけで >>798
相変わらず頓珍漢なこと言ってるな
お前がどう思うかは全く興味がないから
ここはCUDAスレだ
CUDA使った事ない糞団子は去れ >>796
CUDA_Occupancy_Calculator.xlsで確認してたんですが、
CC3.0だと48kb/16kbと16kb/48kbの他に32kb/32kbがあるっぽくて、
32kbなら100%になる、みたいな謎結果が出てたんです。
で、うちの古い参考書にはその設定のやり方が載ってなかったんです。
だけどそのあと.xlsファイルを開き直したらなぜが100%になってました。
なんでやねん。 >>800
お前がCUDA使ったことないゴミだと自己紹介してることはわかったよ
俺は初代の8系から触ってるんだがな
墓穴ほり過ぎ また低脳の団子が自演かよ!?
いい加減にして欲しい
コイツがいるとスレが腐るんだけど(怒) >>802
初代から使ってそのレベルならお前には向いてないからもう諦めろ
結局質問者には何ひとつ有益な情報を与えられない低レベル糞団子 >>801 (>>754)
ああいう大人(団子)みたいになったらだめですよ(笑い 文句言うだけの奴より団子のほうがこのスレの役に立っている。 【親米涙目】 ★ ロシアがアメリカを一発KO ★ 【反中逃亡】
宇宙人側からの申し入れは、とにかく核の利用と戦争をやめなさい、もう一つは宇宙人の存在を公表しなさい。
つけ足したのが、60年の猶予を与えましよう、2014年には発表しなければならなかったんですね、宇宙人の存在を。
ロシアという大国の首相がね、あれは冗談だよでは済まないですね、しかも2回も言ってるんだからね。
https://www.youtube.com/watch?v=FIRXKetUkq8
【親米屈辱】 ★ 宇宙人いない派のおバカさん残念www ★ 【反中赤面】
NASAは、UFOに対して長年取ってきた態度のために、無用な組織とされることを恐れています。
マイトレーヤが公に現れるにつれて、UFOが姿を表すでしょう。
世界平和の脅威は、イスラエル、イラン、アメリカです。イスラエルの役割は跪いて、パレスチナに許しを請うことです。
アメリカによる他国の虐待に反対の声を上げなければなりません。
彼らは今世紀(21世紀)をこの帝国が出来上がるアメリカの世紀と呼ぶ。しかし、そうはならないだろう。
彼らが世界中に‘民主的’制度を確立したいという衝動をコントロールするのは、マイトレーヤの任務である。
Q 経済崩壊が2015年から、テレパシー世界演説が2017年の初めなら、2016年に大戦でしょうか。A 大戦は起こりません。
【親米命乞い】 ★ マイトレーヤのテレパシーによる世界演説は英国BBCが放送 ★ 【反中土下座】
https://www.youtube.com/watch?v=6cOvo6n7NOk
マイトレーヤが世界に向かって話をする準備は良好に進行している。
25分か35分くらいかもしれませんが、歴史上で初めて、世界的規模のテレパシーによる接触が起こるのです。
テレビ中継はこの出来事のために存在するのであり、この時、初めてご自身の本当の身分を明らかにされます。 はい、単発来ましたw
糞団子のどこが有益なんだ?
シェアードメモリの割り当て変更にも答えられずに、ローカルメモリとか嘘の情報垂れ流してる奴だぞ?
まともなCUDAユーザーならローカルメモリとシェアードメモリが別物なんて事は常識以前の問題だ
糞団子の自論なんか誰も聞きたくない
嫌われ者の糞団子は去れ! 他人のレスで我が物顔してるだけじゃん
無益な厨房だな つーか自称シェアードメモリが実質的なLocalStorageでしかないことは知っておかないとハマる事例もあるんだけどね
SIMTはAVXユニットを8コア、Xeon Phiの1コアを16コアって読み替えるトンデモ抽象化 別に何も間違ったことは書いてないんだが理解できないバカがいる シェアードメモリをローカルメモリとか言う嘘を垂れ流してる事はスルーなんだな
簡単な初期化関数すらアドバイス出来ない低レベル糞団子
ちなみにシェアードメモリは自称でもなんでもなく公式な名称だ
気に入らないなら自分で会社起こして俺の考えた最強のGPUでも作ればいいさ
まぁ、お前のレベルじゃディスクリートすら無理な事は明らかだがな
早く去れよ糞団子 さっき、Anacondaの最新バージョンである2.5.0(64-bit python3.5.1)を入れたの
だけど、それにコマンドプロンプトがデフォルトで入っていないので、pip installが
出来なくて困っています。ipythonやインタラクティヴシェルは入っています。
前のAnacondaのバージョン(python2.7で使っている)であれば、ipythonや
対話シェルからq()やquit()で、普通に、コマンドプロンプトに戻れて、そこから
pip installなどが出来たのになぜなのでしょうか。
普段使っているコマンドプロンプトはpython2.7用になっているので、そこから
python3.5.1用のpip installも出来ません。コマンドプロンプトをpythonのヴァージョン
ごとに切り替えることができるのかもしれないけど、その方法も分かりません。
どなたか、知っている人がいたら教えて下さい。
よろしくお願いします。m(__)m >>815
恐ろしくスレ違いなんだけど
python2系と3系を共存させるならパスを適切に設定しないとダメだよ
詳しくはググってくれ
取り敢えずコマンドプロンプトでpyって打ってからpipしてみたら? >>816
パスは、最新のAnacondaをインストールする時に、パスを自動設定するなら
チェックを入れよという欄があって、チェックを入れてインストールを終了
すると、新しい環境設定が入っています。
念のためシステムで環境設定でパスをチェックしたけど、Anaconda3で、
新しいのが入っている。Anacondaは、3ヴァージョン使っている
のだけど、新しくなるほど使いにくくなっている感じですね。
q()でコマンドプロンプトに戻れると>>815で書いたのは、間違えだった。
それは、R言語のやり方だった。 nv用語だとローカルにないものが
ローカルメモリなんだなこりゃややこしい 今更の問題だけど、pythonはやはり、2系と3系の互換・移植性が
悪いのが問題かな。今日買った、
基礎 Python (IMPRESS KISO SERIES) 単行本(ソフトカバー) – 2016/3/4
大津 真 (著)
ttp://www.amazon.co.jp/%E5%9F%BA%E7%A4%8E-Python-%E5%A4%A7%E6%B4%A5-%E7%9C%9F/dp/484438015X
が、3系で、ちょっと変わった雰囲気のコードの使い方があったので、面白いの
だけど、2系と併用だと>>815みたいになって、pip installも出来ないし、
ipythonでしか、3系が使えないのが困りもの。2系で使える3系のものも
あるけれど…。
2系を全部消せば3系を使えるのは分かっているのだけど、2系の方が
出来ることがまだまだ多い感じがするからね。まだ2系が多数派でしょ。
あ、仮想環境ヴァージョンを切り替えてやればいいのかな。 あ、pythonのクダスレだと、勘違いしてました。
失礼いたしました。 >>821
何を今更
だからこそ嘘を垂れ流す糞団子はここから去れよ >>821
もっとも、スピルされたレジスタのデータが真っ先に格納されるのはGDDR*ではなく
キャッシュだけどな
なぜならキャッシュはDRAMの読み書きを高速化するためにあるから >>825
何が嘘なのか具体的に示せよ
俺は具体的に示したろ
糞団子は去れ >>826
>なぜならキャッシュはDRAMの読み書きを高速化するためにあるから
そんなパソコン教室の初日に教わる様な内容を語るスレじゃねぇよ
糞団子は去れ! 団子ちゃん悔しいだろホントはおともだちをつくってワイワイ話したいんだろ?
じゃあ自作板帰ろうな 自分がやってるからって
他人もそうだと思うなよ
病気の人 まだやってんのかよっw。
シェアードメモリはローカルメモリだよwww。 方言のない一般名称でいうと「スクラッチパッドメモリ」か ぼろ出してもコテでい続けるのは正直苦痛だろ?
な、団子 名無しばっかりだと不安で団子認定に必死な病気の癖に いい加減通常時に戻ってください
正直どっちもどっちなんだよね
どっちも邪魔 コテの人はいい加減ブログでもなんでも始めたらいいよ
荒らし行為だよ
コテ煽る人もやめてよ 団子さん
お願いです。
あなたが来てからこのスレは機能しなくなりました。
ここは数少ないCUDAの情報交換の場です。
もう、ここに来ないで下さい。 こいつ一人が団子団子連呼してて
すげー気持ちわりーよ 最近は傷心の団子をおもちゃにするうんこごっこが流行ってるのか? いいえ
精神病の人をどうバカにするかが流行ってます つーかまともな話しようとすると団子認定されるのは書き込む気を無くさせるためなのだろうか。
技術的な話をすると内容がどうでも団子にされるから、そういう話が理解できないのと、ある意味団子崇拝が酷いのだろうな。 >>849
団子はある程度(正しいかは別にして)技術的話が出来るけど、
その人はAMDは倍精度がすごいよ、GPGPUの勝者はAMDで決まり!と
I love 団子だけだからな。 あんまり団子をおもちゃにするなって
ウンコの匂いが充満するじゃないかw こんにちは。
CUDA では、__device__ 修飾された関数は、GPU内部から呼び出せますが、
再帰コールが出来ないのはなぜですか?
以下はできますよね?
__device__ int aaa(int x) {
return x * 2;
}
__device__ int bbb(int y) {
return aaa(y) + 3;
}
では、以下は駄目ですか?
__device__ int kaijou(int y) {
if ( y == 1 ) {
return y;
}
else {
return y * kaijou(y-1);
}
} ハード的に絶対できないと言うことも無いのだろうが、スタックどれだけ使うか計算できないってのが大きいんじゃね。 最適化の邪魔になるから最初から再帰はやめてねーんってことかな。 >>858
おまえが正解答えてやれば
あれずに済むだろ >>861
俺ならその無駄レス書いてる暇で書けるぞ>>858は無能だから理解できないだけで
>>859
CUDAの自称コアが要素毎プレディケートつきのSIMDユニットを
マルチコアって言い張ってるだけから本物の独立したコアでないから
で終わりなんだけどね。
命令ポインタは「自称コア」ごとに独立してないんだ
「分岐」も本物の分岐ではなく、if-else節を全部一列のコードに展開してプレディケート
しているだけ。プレディケートビットの全要素が0あるいは1になって初めて関数を
飛び越えることができる。
SIMDの要素(自称)ごとに命令ポインタと関数コールスタックを管理できて別々の
命令ステートを独立して持てるならそれは既にSIMDじゃない ベクトルを隠蔽しない、1ワープを1つのスレッドとして処理できるプログラミングモデルが
提供されるなら再帰は理屈上可能
それはCPUのSIMDの在り方と同じってことだけど 独立したコアという前提だから、コア毎のレジスタで表現できる範囲の処理しか出来ないって事だと思ってたけど違うのか。 基本はSIMDプロセッサでの本当の分岐だよ。
簡単な処理ならpredicate実行で済ますけどこれはモダンなプロセッサではSIMT、SIMD、に限らず一般的な手法
状態保存のメモリを最小化するためにdivergent pathを深さ優先でスキャンしていくから、どこかのdivergent pathで無限ループになるとwarp内の別threadには何時までたっても制御が帰ってこないけどね。 再帰だけならCell SPUすら出来てるからね
命令ポインタの制約のある偽りのコアであること以上の理由はないでしょう
末尾再帰をループに展開するオプションすら用意してないのは処理系の実装の手抜きだろうと思うけど 可能な限り手を抜くのがハードにとってのGPGPUという物かも 実害は具体的にはバックトラッキング型の探索アルゴリズムが使えない
(実装方法に制限がある)とかかな cudaは性能出すために最適化しだすと、結局SIMD構造を意識して記述することになる。
で結局、SIMTでプログラミングするメリットを感じなくなる。
scatter,gatherも便利なのだけど、それしかないから、性能出すためにアラインメント気にしながら連続アドレスにアクセスするように記述していると、普通のSIMDのようにアドレス1つ計算すればすむのに無駄だなぁと。
で、複数のアドレスを受け取ったハードウェアが同じキャッシュラインに乗っていることを検出してメモリアクセス回数を最小化するとか、バカじゃねと。
まあGCNはそこら辺の無駄に対応していて美しいのだが、いかんせんソフトやサポート体制がクソ過ぎて。 >>865
>状態保存のメモリを最小化するためにdivergent pathを深さ優先でスキャンしていく
横からですが参考になります。
再帰がまずいっていうより結果として条件分岐をたくさん通るのがまずいって感じですかね。 最近は、5,000円くらいの低価格GPUでも、Keplerコアを使っていて、
Fermiコアの上位互換だから、再帰コールも出来るのですか?
でも、関数のローカル変数は、C/C++と全く同じで、スレッド独立だが、
関数引数だけは、スレッド共通だと聞きましたが、これだと再帰コール
できなさそうですが。 あーあまた糞団子が答えちゃった
このスレの事情を知らん奴が見ると余計に迷うぞ 答えられないクズは存在する価値ないでしょ
とっとと失せろよゴミ つーか俺も含めてまだ正解出してないんだけどな
んで正解はこれ、Kepler以降なら再帰に【一応は】対応
http://developer.download.nvidia.com/assets/cuda/files/CUDADownloads/TechBrief_Dynamic_Parallelism_in_CUDA.pdf
誰かツッコミ入れると思ったけど期待外れだったなあ
まあ、俺の説明も何も間違ってないのだけど
(>>862も「出来ない」という説明をしてるのではなく制約を説明しているだけ) 確かにこの団子は鳥屋氏じゃないな。
鳥屋氏はこういう物言いはしない。 お前は、中村ちゃぷにちゃんが中身の団子が存在すると思ってるんだ
そこから既に間違ってるよ ちゃぷにさんはだんごやさんを嫌ってるはずなので彼がその顔文字を使う理由はないんだよ
自分でNGキーワードに入れてる宣言してたし 更に言うと彼は精神が壊れて自作PC板で機械翻訳コピペを貼るだけの人に成り下がってしまった 団子=鳥屋でも別にいいじゃねーか
糞撒き散らすだけならどっちもNGだ >>879
鳥屋氏に何かあったの? そっちが気になる。 >>882
だから本人に直接聞いて来いよTwitterアカウントあるだろ 別に本人と直接連絡とれるけどなんか聞きにくいじゃん。
おしえてよ。 ちなみに彼は10年以上前から初心者の質問板で「メンヘル屋」と呼ばれてたぞ 日曜日の真昼間に2ch張り付きかよ
さすが糞団子だな マジで氏ねよ!ここ糞団子!!
お前のせいでスレがめちゃくちゃ! 2chやっているおっさんなら起きてから寝るまで2chが普通だろ
いまはスマホで仕事サボって2chしているおっさんが多いからな >>898
今やちょっとゆる会社だと、トイレが満室になって10分待ちだからな。 そんなにトイレが好きならドアを接着剤で止めてやろうかと思った 平日仕事サボって2chなら、まぁわからなくも無い
理解できないのは日曜日の朝から晩まで2ch張り付きの糞団子 スマホゲーの体力ゲージ消化してるだけじゃね?
2chなんてとっくにブーム終わってるだろ >>901
話し相手が欲しいんだよ
スレをウンコ臭くしておいて 糞団子が急に大人しくなって笑える
かなり効いた様だなw コテ団子をNGにして華麗にスルーしないとね
名無しで煽る時は怒るんじゃなくてファブリーズでも置いておきましょう ここはプ板のかまってちゃんが集うところです
Nvidiaユーザーってへんなやつ多いよね。どう? 糞団子が荒らして過疎ったな
ホントろくでもねぇ奴だったな では素人が質問してもいいですか
オプションのCUDA RuntimeをStatic CUDA runtime library (-cudart static)にしているのにもかかわらずcudart32_75.dllが生成されるんですけど不要ですよね?
消しても動くのでいらなさそうではあるんですけど一応残してます >>918
CUDAを使ったことがない連中がたくさん来場してたからな
CUDAは映像・画像処理の一般プログラムではなく専門性が非常に高い分野(機械学習や車の自動運転とか)で使うって感じに
なってきたからな。 レジスタファイルがダイ全体で14Mもあるとかホント化け物ですね
アクセスレイテンシはでかそう 以前エクサスケールへむけての改良点てやってたけど
warpサイズが4になったりレジスタファイルの階層化とか
やってるのかな Pascalは忘れられてるけど、倍精度も強化されてるんだよね
単精度を徹底的に削った思い切った方針には目を見張った 大型スパコン案件に特化して設計したものをGPUに流用してるだけで
それが全部だと言い張られても困りますが GPGPUは深層学習のブレークで今後一気に需要が高まる? 高まるだろうね
倍精度と半精度がちょうどいい具合
そうやってGPGPU市場を頑張ってもらって12.1の妨害行為は止めて欲しいわ その需要の大半はnVIDIAが発表したEyerissみたいな専用チップに喰われるのでは あれはCNNに特化したものでGPUほどのプログラマビリティはない
ディープラーニングって一言で言っても1個のASICで賄えるほど狭い分野じゃない 疎行列に特化したマトリクス演算回路とそれなりの帯域のリングバスネットワークが主な構成要素
って覚えておけば大体間違いない
実はGPUの特性と正反対(GPUは密行列なら得意) fermi世代ってdouble⇒floatで倍速なはずだよね?
ちっとも速くならないんだが。 大分前だけど構造体でデータを渡すとdouble扱いになったことはある
原因は不明 >>937
thx。
調べてみた感じだと、倍精度が無いカードで暗黙のdouble→floatキャストがあったり、
floatへの数値代入で接尾語Fを付けないとdoubleにキャストされたりもあるっぽいね。
fermiでarch=compute_20,sm_20指定してるから基本的には宣言通りだと思ったのだが、実際どうなってるのか・・。
全部floatに置き換えてもわずかしかスピードアップしないのは、バイトコピー分が浮いてるだけな気がする。 匿名通信(Tor、i2p等)ができるファイル共有ソフトBitComet(ビットコメット)みたいな、
BitTorrentがオープンソースで開発されています
言語は何でも大丈夫だそうなので、P2P書きたい!って人居ませんか?
Covenantの作者(Lyrise)がそういう人と話したいそうなので、よろしければツイートお願いします
https://twitter.com/Lyrise_al
ちなみにオイラはCovenantの完成が待ち遠しいプログラミングできないアスペルガーw
The Covenant Project
概要
Covenantは、純粋P2Pのファイル共有ソフトです
目的
インターネットにおける権力による抑圧を排除することが最終的な目標です。 そのためにCovenantでは、中央に依存しない、高効率で検索能力の高いファイル共有の機能をユーザーに提供します
特徴
Covenant = Bittorrent + Abstract Network + DHT + (Search = WoT + PoW)
接続は抽象化されているので、I2P, Tor, TCP, Proxy, その他を利用可能です
DHTにはKademlia + コネクションプールを使用します
UPnPによってポートを解放することができますが、Port0でも利用可能です(接続数は少なくなります)
検索リクエスト、アップロード、ダウンロードなどのすべての通信はDHT的に分散され、特定のサーバーに依存しません
h eyerissは大量のデーターが複数のフィルター処理されるのに向いてる 質問があります
環境はwindows 10 | arch linux, cuda 7.5, gcc | visual c++ 2013です
#include <stdio.h>
__global__ void a(char *A){
}
int main(){
a<<1, 1>>(); // (1)
}
(1)で「expression must have integral or enum type」とエラーが出て困っているのですが
ホストコンパイラを呼び出す際に何か間違えているのでしょうか
コンパイルに使用したコマンドは
windows
nvcc -ccbin "C:\Program Files (x86)\Microsoft Visual Studio 12.0\VC\bin" -c -DWIN32 -D_CONSOLE -D_MBCS -Xcompiler /EHsc,/W3,/nologo,/Wp64,/O2,/Zi,/MT test.cu
linux
(失念したが大体同じ)
です >>941
≪≫のところじゃない?
二重カッコじゃなくて、三重カッコにすべきかな
二重カッコだと別の演算子になって、意味が変わってくる ああ、3重括弧でしたか
間抜けなミスで騒いでしまって申し訳ないです
ありがとうございます グーグルがTensolFlow用のTPUとかいうの作ってたらしいがどうなの? CUDA fortranで一つの配列に対して2GB以上のメモリを
アロケートすることが出来ず、
PGIのサイト見たら「2GBの壁」とかって説明されているのですが、
これってどうにか克服出来ないですかね? >>945
fortranなんか使っているからだよ。 CU_LIMIT_MALLOC_HEAP_SIZE か cudaLimitMallocHeapSize を拡げりゃいいんじゃないかな。 ttps://developer.nvidia.com/cudnn
>Supported on Windows, Linux and MacOS systems with Pascal, Kepler, Maxwell, Tegra K1 or Tegra X1 GPUs.
のTegra K1の文字を見て、サポート復活を喜んでいたのですが、
Tegra K1は32bitでして、
必要なCUDA7.5は64bit必須
もしかして、cuDNNとCUDAは無関係で独立して使えたりしますか? CUDA APIで1番始めに呼んだAPIは時間がかかるとのことですが、
時間計測に影響が出ないように事前にどんな関数を呼んでますか?
慣習だったり、よく使われる関数があれば教えてください。 工学じゃなくて物理系でもcuda覚える価値ってある? 安く実験したいなら
N体や粒子法なんかは普通のCPUじゃたいしたことはできない この部分がエラー吐くんだが
cuda7.5
struct cuComplex {
float r;
float i;
cuComplex( float a, float b ) : r(a), i(b) {}
__device__ float magnitude2( void ) {
return r * r + i * i;
}
__device__ cuComplex operator*(const cuComplex& a) {
return cuComplex(r*a.r - i*a.i, i*a.r + r*a.i);
}
__device__ cuComplex operator+(const cuComplex& a) {
return cuComplex(r+a.r, i+a.i);
}
}; CUDA by EXAMPLEのサンプル?
<cuComplex.h>の定義とぶつかってるんじゃね? >>957
そうそう
エラーはこんな感じ
Documents/Visual Studio 2013/Projects/template/cuda_by_example/chapter04/julia_gpu.cu(43): error
: calling a __host__ function("cu_Complex::cu_Complex") from a __device__ function("julia") is not allowed 試してないけど、コンストラクタに__device__付けてないのが悪いと言っているような。 615 :名前は開発中のものです。:02/06/10 21:48 ID:???
>将来的にはスカートの頂点を触ってヒラヒラさせたりしたいので、どうせシステムメモリーに
ボーンでやる気なのか?つらそー。
616 :◆CKmHzWpc:02/06/11 03:48 ID:zXaLQtp.
みなさんありがとうございます
1.モデルを分割するのが嫌→頂点シェーダ−を使うorシステムメモリに置いてソフトでやる
2.モデルを分割してもいい→ハードのジオメトリブレンドを使ってやる・・
というのが考えられる・・・んですよね
色々組み合わせを試して、高速なのを選びたいと思います
>>615
これは、スカートの頂点をCPUで直接操作しようと思っています
ボーンでやらずに、簡単なクロスシュミレーションのような事を目指しています
どうせモデルの一部をシステムメモリに置くなら、もう全部システムメモリにし
てもいいか、と思ったのです pascal titan xでcudaを動かしている人に質問ですが高負荷時にクロックダウンは起きませんか a.exe: a.cu
nvcc -Xcompiler "/wd 4819" a.cu -o a.exe
ウインドーズでメイクファイルで上みたいにやってクダーCコンパイラーでコンパイルしたんですけど
コマンドラインにちゃんと
nvcc -Xcompiler "/wd 4819" a.cu -o a.exe
と表示されたあと
wd
4819
c1xx: fatal error C1083: ソース ファイルを開けません。'略':No such file or directory
a.cu
c1xx: fatal error C1083: ソース ファイルを開けません。'4819':No such file or directory
と出ます。たぶん"に囲まれた中の/がメタ文字じゃないのになぜか消されているとおもうんですけど
なぜエラーになるのか教えてください。
/を//あるいは-に置き換えたらちゃんとコンパイルできました
よろしくお願いします。
" /wd 4819"と"と/の間にスペースを入れても出来ました。 このスレッドは1000を超えました。
新しいスレッドを立ててください。
life time: 1065日 22時間 39分 42秒 5ちゃんねるの運営はプレミアム会員の皆さまに支えられています。
運営にご協力お願いいたします。
───────────────────
《プレミアム会員の主な特典》
★ 5ちゃんねる専用ブラウザからの広告除去
★ 5ちゃんねるの過去ログを取得
★ 書き込み規制の緩和
───────────────────
会員登録には個人情報は一切必要ありません。
月300円から匿名でご購入いただけます。
▼ プレミアム会員登録はこちら ▼
https://premium.5ch.net/
▼ 浪人ログインはこちら ▼
https://login.5ch.net/login.php レス数が1000を超えています。これ以上書き込みはできません。