- 1 名前：デフォルトの名無しさん [2011/08/23(火) 22:08:06.09 ]
- このスレッドは、他のスレッドでは書き込めない超低レベル、
もしくは質問者自身何が何だが分からない質問を勇気を持って書き込むスレッドです。
CUDA使いが優しくコメントを返しますが、
お礼はCUDAの布教と初心者の救済をお願いします。
CUDA・HomePage
ttp://developer.nvidia.com/category/zone/cuda-zone
関連スレ
GPGPU#5
ttp://hibari.2ch.net/test/read.cgi/tech/1281876470/
前スレ
【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/
- 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%差でした。
- 488 名前：デフォルトの名無しさん mailto:sage [2012/02/13(月) 19:39:39.79 ]
- なるほど，有り難うございます．
有意な差が出たと言うことですね．
では，後は分かる人，よろしく♪
ε=ε=ε=ε=ε=┌(； ・＿・)┘
- 489 名前：デフォルトの名無しさん mailto:sage [2012/02/14(火) 10:35:58.19 ]
- GPU側では64bitのメリットは使用可能なメモリが増えることだけで、その代わりにアドレス計算の命令数や使用レジスタ数が増えるデメリットがある。
前はGPU側だけ32bitモードで動かすという恐ろしい手法があったが、今使うのはなかなか難しい。
- 490 名前：デフォルトの名無しさん mailto:sage [2012/02/14(火) 16:21:21.98 ]
- ptxオプションつけてptx出力を見れば見当がつくね。
64ビットアドレッシングだとアドレス計算の為にint32→int64の変換がしばしば行なわれる。
- 491 名前：デフォルトの名無しさん [2012/02/14(火) 18:19:16.13 ]
- ざっくり捉えると、32bitと64bitでは、
CUDA自体はさほど変わらない。
CUDAと切り離されてコンパイルされるC++言語部分は、64bitの方がいくらか速い。
という感じでしょうか。
- 492 名前：デフォルトの名無しさん mailto:sage [2012/02/14(火) 18:24:53.01 ]
- >>491
アドレス計算多用したり、レジスタ沢山使うようなカーネルだと
CUDA自体32bit版の方が無視できないほど
速くなる可能性はある。
- 493 名前：デフォルトの名無しさん [2012/02/14(火) 18:53:10.54 ]
- doubleが速くなるっていう噂はどうなったの？
- 494 名前：デフォルトの名無しさん mailto:sage [2012/02/14(火) 22:53:20.23 ]
- GPU上にあるコンスタントキャッシュメモリって、CUDAではどのようにして使われるのでしょうか？
計算の中にリテラルがあった場合（CPUでは命令の即値として埋め込まれますが）や、
constantを付けて宣言した定数が割り当てられたりするのでしょうか？？
- 495 名前：デフォルトの名無しさん mailto:sage [2012/02/16(木) 05:51:50.91 ]
- llvmで使えるようになったのなら
boostも使える？
- 496 名前：デフォルトの名無しさん mailto:sage [2012/02/16(木) 08:10:03.40 ]
- >>494
適当に数値を入れて，確かめてみるとか
- 497 名前：デフォルトの名無しさん mailto:sage [2012/02/16(木) 08:17:55.24 ]
- >>495
２年前にパッチが出ているようです．boostが何か知らんけど
https://svn.boost.org/trac/boost/ticket/3919
- 498 名前：494 mailto:sage [2012/02/17(金) 00:11:19.70 ]
- >>496
試してみました。
const変数もリテラルもレジスタにマッピングされました。
ひょっとして、DirectXAPIなどのグラフィックスAPIを利用していた場合に
使われるだけ（要は旧来通りの使われ方）かも･･･
プロファイラを見ると、テクスチャキャッシュのヒット/ミスヒットが見れるようになっていますが
これなんかはグラフィックスAPI専用のような気が。
しかし同じようにコンスタントキャッシュの項目もあって然りと思うが在らず。
う〜ん、気にしないことにしますｗ
- 499 名前：デフォルトの名無しさん mailto:sage [2012/02/17(金) 06:31:17.52 ]
- 貴重な報告、ありがとうございます♪
- 500 名前：デフォルトの名無しさん mailto:sage [2012/02/17(金) 07:54:24.51 ]
- GPUを演算器としてフルに使いたい場合って
ビデオ出力をオンボードのものになるようBIOSで切り替えた方がいいのでしょうか
linuxだったらxconfigいじるだけでいけるとか？
- 501 名前：デフォルトの名無しさん mailto:sage [2012/02/17(金) 07:57:38.45 ]
- フルの定義は？
- 502 名前：デフォルトの名無しさん mailto:sage [2012/02/17(金) 07:58:52.97 ]
- >>494
CPUでも即値として埋め込めない配列や文字列リテラルは
.rdataセクションのデータとして書き込まれるのと同様に、
GPUでも即値として使えないconstデータはコンスタント領域に置かれる。
で、コンスタント領域の値を直接オペランド指定できない
命令を使う場合は当然レジスタを介して使用することになる。
逆に一部の変数にデフォルトで定数が入っているなんてのは
ABIで決まったもの以外は普通無理だから、コンスタント境域などの
メモリに置いて読み込む以外やり様が無い。
- 503 名前：デフォルトの名無しさん mailto:sage [2012/02/17(金) 08:33:23.95 ]
- >>500
普通はGPUカードに何もモニタを繋がなければ(≒出力する設定をしなければ)
GVRAMをほぼ全部使い切れるよ。
Linuxなら、Xを立ち上げなければGPUカードから出力していてもかなり使えるよ。
ついでに、3秒だか5秒だからの制限もなくなるからGPUカーネルで長時間計算させられるよ。
- 504 名前：デフォルトの名無しさん [2012/02/17(金) 20:22:12.27 ]
- VRAM 4GBのボードが六千円台
akiba-pc.watch.impress.co.jp/hotline/20120218/etc_inno.html
- 505 名前：494 mailto:sage [2012/02/18(土) 16:32:22.87 ]
- >>502
ありがとうございます。
コンスタント変数の扱われ方、レジスタが必要な理由、理解しました。
書籍を読んでいたら、コンスタントメモリおよびテクスチャメモリについて書かれているのを見つけました。
コンスタントメモリは__constant__修飾子を付けて宣言することで割り当てられるようです。
物理的にはVRAMに置かれ、GPU内部のコンスタントキャッシュでキャッシュが効くものと思われます。
空間フィルタリングなどでフィルタ係数を格納しておくのに適しているのではと思いました。
（繰り返し参照されるのでキャッシュが効きそう）
テクスチャメモリはテンプレートクラス型のtexture<･･･>で宣言することで割り当てられるようです。
物理的にはVRAMに置かれ、GPU内部のテクスチャキャッシュでキャッシュが効くものと思われます。
さらに、グラフィック用途時の同様、テクスチャフィルタリングやサンプリング方式などのハードワイヤード回路機能を
利用できるよう、それらを指定するプロパティを保持しているようです。
- 506 名前：デフォルトの名無しさん mailto:sage [2012/02/18(土) 21:02:09.38 ]
- 複数カーネルを呼び出すんだけどカーネルのソースはそれぞれ別ファイル
それぞれのカーネルから同じDevice関数を呼び出したいんだけどファイルまたがることできないようだし
ソースごとに同じ内容の別名関数作るしかないのかな？
- 507 名前：デフォルトの名無しさん mailto:sage [2012/02/18(土) 22:43:38.30 ]
- ファイルが別でも，プログラムとしてつながっていれば問題はないと思うが．
関数のプロトタイプ宣言をしっかりしておけば良いかと．．．
- 508 名前：デフォルトの名無しさん mailto:sage [2012/02/18(土) 22:46:25.76 ]
- 共通のdevice関数を別のヘッダに纏めて
カーネルソースからincludeして使う。
device関数はどうせinline展開されるのだから
通常のプログラムでのinline関数と同様に扱えばいい。
- 509 名前：デフォルトの名無しさん mailto:sage [2012/02/18(土) 22:54:19.36 ]
- >>508
サンクス
でも今やってみたんだけど全ファイルで展開されるんで
multiple definition of 関数名
ってエラーになっちゃう
実際やってみました？
- 510 名前：デフォルトの名無しさん mailto:sage [2012/02/18(土) 23:22:56.53 ]
- カーネル起動時、1スレッドの起動につき1クロックかかり、128スレッドなら128クロックかかると書籍に書いてあるのですが、
これはCPU側でもかかるのでしょうか？
それとも、CPU側はカーネル起動の合図を送ってすぐ制御を戻し、
GPU側の制御ブロックがクロックをかけてスレッドを起動するというかたちでしょうか？
- 511 名前：デフォルトの名無しさん [2012/02/19(日) 01:11:11.40 ]
- 教えて欲しいことがあります．
CPU側で複数スレッドを立てて，互いに独立した処理を行っています．
そして，それぞれスレッドにおいて，CUDAを使って画像処理をさせたいのですが，（たとえば，グレースケール化とか）
この時，CUDAでの処理にテクスチャメモリを使いたい場合は，どのようなコードを書けばいいのでしょうか？
テクスチャメモリを使う場合，グローバルで宣言しなけらばならないですよね？
たとえば，5枚のRGB画像をグレースケール化するときに，
CPU側で5スレッド立てて，各スレッドでCUDAを使ってグレースケールへの変換処理をしたいのですが，
テクスチャメモリをグローバルで宣言するとおかしなことになるきがするんですが．
どなたか教えて頂けないでしょうか？
宜しくお願い致します．
- 512 名前：デフォルトの名無しさん mailto:sage [2012/02/19(日) 01:21:41.45 ]
- CUDA対応のGPUを5枚挿す。
またはGTX590などを3枚挿す。
- 513 名前：デフォルトの名無しさん [2012/02/19(日) 01:26:39.83 ]
- >>512
回答ありがとうございます．
テクスチャメモリを使った処理の場合，マルチGPUじゃないと出来ない
とういことでしょうか？
- 514 名前：デフォルトの名無しさん [2012/02/19(日) 01:44:29.68 ]
- 公式マニュアル見れば分かること
試してみれば分かること
これを聞く人が多すぎる
- 515 名前：デフォルトの名無しさん mailto:sage [2012/02/19(日) 01:53:02.78 ]
- 4-way SLI でも、ホスト上の4スレッドで利用する場合、
各々1GPUを独立して割り当てるため、4GPU間で通信して仕事をこなす本来的なSLIにはならない、って認識でイイんですよね？
（SLIコネクタがなくても動く？）
- 516 名前：デフォルトの名無しさん mailto:sage [2012/02/19(日) 02:32:56.76 ]
- >>511
テクスチャを5個宣言じゃダメなの？
- 517 名前：デフォルトの名無しさん [2012/02/19(日) 02:47:03.16 ]
- >>516
回答ありがとうございます．
CPU側でスレッドを立てる段階で決めたいので，
できることなら，可変にしたいと考えています．
514さんは「これを聞く人が多すぎる 」と仰っていますが，
そもそも，このようなことをシングルGPUでやろうとするのが間違いなのでしょうか？
- 518 名前：デフォルトの名無しさん [2012/02/19(日) 02:59:57.92 ]
- いやいや、まず試せば分かることじゃん。
その労力を2chに押しつけるのは、今この瞬間は楽では良いかもしれないけど、
自分の成長には繋がらないよ。
- 519 名前：デフォルトの名無しさん mailto:sage [2012/02/19(日) 05:31:26.29 ]
- >>513
OS側のスレッド何本あろうとGPUが1つしかないんだから
順番に処理されるだけだろってことでしょ
- 520 名前：デフォルトの名無しさん [2012/02/19(日) 05:43:18.09 ]
- concurrent kernel execution…
- 521 名前：509 mailto:sage [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 ]
- 次に同じ疑問を持たれた方のために，
参考資料のアドレスを貼っておきます．
ttp://www.nvidia.co.jp/docs/IO/81860/NVIDIA_Fermi_Architecture_Whitepaper_FINAL_J.pdf
上記アドレスにあるpdfの「コンカレントカーネル実行」に書かれていることが，参考になるかと思います．
- 524 名前：デフォルトの名無しさん mailto:sage [2012/02/19(日) 16:33:50.20 ]
- 誰が質問者なのかわからないので謎な流れだが、とりあえずグレースケール化にテクスチャは不要だと思う今日この頃。
- 525 名前：デフォルトの名無しさん [2012/02/19(日) 17:00:58.89 ]
- NPPで一発解決だよな
- 526 名前：デフォルトの名無しさん mailto:sage [2012/02/20(月) 22:07:59.44 ]
- atomicCASのCASってどういう意味ですか？？
- 527 名前：デフォルトの名無しさん mailto:sage [2012/02/20(月) 22:24:51.56 ]
- Compare And Swap
- 528 名前：526 mailto:sage [2012/02/20(月) 22:28:38.87 ]
- >>527
ありがとうございました！
- 529 名前：デフォルトの名無しさん mailto:sage [2012/03/07(水) 03:16:50.68 ]
- ｃｕｄａの本には2次元配列より1次元の方が速いってあったけど実際どのくらい違うんですか？
また、コピーとgpu上での演算の両方で遅くなるんですか？
- 530 名前：デフォルトの名無しさん [2012/03/07(水) 03:39:04.89 ]
- なんでこのスレ突然止まってたの？
卒論修論シーズンが終わったから？
- 531 名前：デフォルトの名無しさん mailto:sage [2012/03/07(水) 06:34:38.99 ]
- >>529
一行目: 速い
二行目: 遅くなる
馬鹿?
- 532 名前：デフォルトの名無しさん mailto:sage [2012/03/07(水) 16:56:25.77 ]
- 学会はこれからだというのに
- 533 名前：デフォルトの名無しさん mailto:sage [2012/03/07(水) 17:15:29.75 ]
- > ｃｕｄａの本には2次元配列より1次元の方が速いってあったけど実際どのくらい違うんですか？
メモリーのアクセス時間の早さのこと？
ちなみに２次元であっても，例えば f[i, j] で，先にｊを変化させるか，i を変化させるのかで
アクセス時間が違ったと思う．
要はメモリーが並んでいる順番にアクセスするのが早いはず．どの位，早くなるかは知りませぬ♪
> また、コピーとgpu上での演算の両方で遅くなるんですか？
コピーとは？ CPUとGPU間のメモリーの転送のことかな？？
- 534 名前：デフォルトの名無しさん mailto:sage [2012/03/08(木) 10:13:29.60 ]
- 人いなくなったな
春になってGPUネタで研究する学生が増えるのを待つしかないか
- 535 名前：デフォルトの名無しさん mailto:sage [2012/03/08(木) 11:44:54.65 ]
- 二次元配列なんて存在しねぇ
って考えると楽なのに
- 536 名前：デフォルトの名無しさん mailto:sage [2012/03/08(木) 19:50:20.17 ]
- 昔のオカルト本では「謎の四次元」「四次元失踪」とかあって、
「四次元」とか謎めいた雰囲気を感じた。
今なら四次元配列とか当たり前だｗ
- 537 名前：デフォルトの名無しさん mailto:sage [2012/03/09(金) 01:08:22.32 ]
- >>536
そうそうｗ
四次元の神秘性が薄らいじゃうｗ
- 538 名前： ◆QZaw55cn4c mailto:sage [2012/03/09(金) 22:10:56.52 ]
- ハミルトンの四元数というのがあって近年計算機工学に応用されるようになって云々かんぬん
- 539 名前：デフォルトの名無しさん mailto:sage [2012/03/09(金) 22:22:34.47 ]
- 発見自体は1800年代だったっけ？
たしか橋の上かなんかで思いついたとかｗ
- 540 名前：デフォルトの名無しさん mailto:sage [2012/03/11(日) 11:39:47.33 ]
- HLSLとどう違うの？
- 541 名前：デフォルトの名無しさん mailto:sage [2012/03/11(日) 11:49:19.59 ]
- >>540
C言語っぽくなってる。
- 542 名前：デフォルトの名無しさん mailto:sage [2012/03/11(日) 14:43:17.16 ]
- 四次元配列と四次元ベクトルは別物だろ
後者は要素数4の一次元配列
- 543 名前：デフォルトの名無しさん [2012/03/11(日) 17:04:08.60 ]
- >>540
HLSLはShader Languageなんで整数とか扱えなかったと思う。
あとステップ数も制限があったような。
- 544 名前：デフォルトの名無しさん mailto:sage [2012/03/11(日) 17:26:35.76 ]
- Compute Shaderを記述する言語も
HLSLじゃ無かったっけ？
リソースベースのHLSLと、ポインタ・配列ベースのCUDA
- 545 名前：デフォルトの名無しさん mailto:sage [2012/03/11(日) 17:38:23.47 ]
- リソースベースって言い方、分かり易いね。
HLSLはまさにそんな感じだ。
- 546 名前：デフォルトの名無しさん mailto:sage [2012/03/12(月) 13:07:18.89 ]
- >>541>>543
サンクス
自由度がまして打ちやすくなってるんだな
- 547 名前：デフォルトの名無しさん mailto:sage [2012/03/12(月) 13:35:35.95 ]
- >>544>>545
サンクス
- 548 名前：デフォルトの名無しさん mailto:sage [2012/03/15(木) 21:09:26.09 ]
- 誰かいますかね、三つほど質問が。
■__syncthreads()だけでは共有メモリへの書き込みを保証できない？
１つのスレッドがグローバルメモリから共有メモリにデータを書き込み、
その後全てのスレッドがそのデータを使用して計算を行うような場合、
書き込み後に__syncthreads()だけではなく__threadfence_block()も必要なのでしょうか?
青木本には__threadfence_block()について特に言及ありませんでしたが・・・。
■ブロック内の全スレッドからの同一グローバルメモリへのアクセス
ブロック内で共通で使用する構造体などをグローバル→共有メモリに移す場合
全スレッドで行うよりもやはり
if(threadIdx.x==0)・・・
のようにした方が良いでしょうか?
■カーネル内でのreturn文の使用悪影響あるか
スレッドごとに計算を行うか判定をする場合、if文で囲っている例をよく見ますが
これは
if(条件)return;
と書いてはいけないのでしょうか？
上のように書いてもとりあえず計算は流れたのですが何か悪影響はあるでしょうか？
- 549 名前：デフォルトの名無しさん mailto:sage [2012/03/15(木) 21:25:56.04 ]
- >>548
・取り敢えずsyncしか使ってないけど問題になったことはない。
・全スレッドから共有メモリへの書き込みを行なうのは多分遅くなるんじゃないかな?
・どちらで書いても同じこと。普通のCPUのような分岐とは違うことを判っていればOK。
- 550 名前：548 mailto:sage [2012/03/15(木) 22:39:42.94 ]
- >>549
ありがとうございます!!
一個目の_syncthreads()と__threadfence_block()の件ですが、
syncだけだと今日うまくいかなかったもので。
ただ他のバグの影響なども考えられるのでもうちょっと調べてみます。
- 551 名前：デフォルトの名無しさん mailto:sage [2012/03/16(金) 04:06:36.41 ]
- >>548
・__syncthreads()は__threadfence_block()相当の処理を
含んでいた気がするけど気のせいかも。
・全スレッドで同じメモリにアクセスするのはたとえfermiでも遅くなるはず。
・カーネル内部で_syncthreads()使う必要があるなら
returnは使っちゃ駄目だろう。
- 552 名前：デフォルトの名無しさん mailto:sage [2012/03/16(金) 20:10:45.37 ]
- いまだにアプリ開発環境すらまともに構築できてない・・・
visual studio 2008でやろうと思って一応ビルドは通ったけど
実行するとまずcutil32.dllがありませんって出た。
次にcutil.dllをデバッグ.exeと同じフォルダに置き実行！！
CUDA version is insufficient for CUDART version.
ってなる・・・orz
まずなにからはじめるべきですか？
- 553 名前：デフォルトの名無しさん mailto:sage [2012/03/16(金) 20:17:42.17 ]
- ちなみに.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;
}
これに書き換えると
プログラムが完成し、エラーもなく実行もできる
- 554 名前：デフォルトの名無しさん mailto:sage [2012/03/16(金) 20:51:20.83 ]
- GPUドライバのアップデート
- 555 名前：デフォルトの名無しさん mailto:sage [2012/03/17(土) 01:39:06.58 ]
- >>554
ありがとうございました！！！！！！！！！！
動いた！！！！！！！！
- 556 名前：デフォルトの名無しさん mailto:sage [2012/03/17(土) 08:54:13.85 ]
- おおｗ
よかったな！
- 557 名前：デフォルトの名無しさん mailto:sage [2012/03/17(土) 12:49:24.04 ]
- >LINK : /LTCG が指定されましたが、コードの生成は必要ありません。リンク コマンド ラインから /LTCG を削除し、リンカの性能を改善してください。
と表示されるのですがリンク コマンド ラインは固定されて編集できません。
解決方法はありますか？
- 558 名前：デフォルトの名無しさん mailto:sage [2012/03/17(土) 13:37:04.85 ]
- >>577
補足：
開発環境はVisualStudio2008
cuda ver 2.3
- 559 名前：デフォルトの名無しさん mailto:sage [2012/03/17(土) 14:59:10.69 ]
- windowsを窓から投げ捨てろ
- 560 名前：509 mailto:sage [2012/03/17(土) 15:16:23.02 ]
- そんなことして道歩いてる人の頭に当たっちゃったら大変ですよ
- 561 名前：デフォルトの名無しさん mailto:sage [2012/03/17(土) 15:44:28.84 ]
- 角に当たったら痛そうだもんね・・・
- 562 名前：デフォルトの名無しさん mailto:sage [2012/03/17(土) 18:01:09.23 ]
- >>557
リンカ -> 最適化 -> リンク時のコード生成 (/LTGG)
C/C++ -> 最適化 -> プログラム全体の最適化 (/GL)
Visual Studio 2008 の使い方なのでスレが違うかも。
- 563 名前：デフォルトの名無しさん mailto:sage [2012/03/18(日) 15:04:39.12 ]
- >>562
ありがとうございます。
CUDA-Zの実行結果はどのように見たらいいですか？
日本のサイトが全然ないです。
- 564 名前：デフォルトの名無しさん mailto:sage [2012/03/18(日) 15:07:18.71 ]
- >>563です
すみません。解決しました。
- 565 名前：デフォルトの名無しさん mailto:sage [2012/03/21(水) 21:06:39.97 ]
- コンスタントメモリキャッシュへのアクセスはバンクコンフリクトとかないんでしょうか？？
- 566 名前：デフォルトの名無しさん mailto:sage [2012/03/21(水) 22:20:52.26 ]
- >>565
そりゃキャッシュはバンクになってないからねー
- 567 名前：565 mailto:sage [2012/03/21(水) 22:44:39.32 ]
- >>566
おお、やっぱり。
できるだけコンスタントメモリ使うようにしまつ。
- 568 名前：デフォルトの名無しさん mailto:sage [2012/03/21(水) 23:02:34.40 ]
- アドレスが静的に解決できないというのが前提だけど
16ポートのSRAMなんてコスト的に不可能だからマルチバンク以外無いんじゃないの？
- 569 名前：デフォルトの名無しさん [2012/03/22(木) 00:27:08.40 ]
- Fermi以前はコンスタントメモリ使う意味あったけど、
Fermi以降はL2キャッシュとあんまり変わらない印象
- 570 名前：デフォルトの名無しさん [2012/03/22(木) 22:48:01.36 ]
- GTX680が発表されたけど、CUDA的には好ましくない方向の進化が多い。。
- 571 名前：デフォルトの名無しさん mailto:sage [2012/03/22(木) 23:03:25.80 ]
- チップ名 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
- 572 名前：デフォルトの名無しさん mailto:sage [2012/03/22(木) 23:06:16.47 ]
- GK104はミドルレンジだからGK110は全体的に上回ってくるでしょ
- 573 名前：デフォルトの名無しさん [2012/03/22(木) 23:16:01.94 ]
- kepler誕生おめ！
.o゜*。o
／⌒ヽ*゜*
∧_∧ ／ヽ ）｡*o ｯﾊﾟ
(・ω・)丿゛￣￣' ゜
. ノ/ /
ノ￣ゝ
- 574 名前：デフォルトの名無しさん mailto:sage [2012/03/22(木) 23:21:31.67 ]
- Keplerｷﾀ━━━━━━(ﾟ∀ﾟ)━━━━━━ !!!!!
- 575 名前：デフォルトの名無しさん mailto:sage [2012/03/22(木) 23:22:01.12 ]
- gen3じゃないんだっけ？
- 576 名前：デフォルトの名無しさん mailto:sage [2012/03/22(木) 23:42:00.14 ]
- >>570
もともとはミドルレンジでグラフィック向けだったから仕方ない気もする。
予想以上にグラフィック方面に舵を切ったという感はあるけど。
このままグラフィック向けとGPGPU向けで大きく分かれていくのではないかという心配はあるかな。
- 577 名前：デフォルトの名無しさん mailto:sage [2012/03/22(木) 23:53:32.45 ]
- 1SM = 192コアか。おっそろしいなあ。
- 578 名前：デフォルトの名無しさん mailto:sage [2012/03/23(金) 00:00:34.24 ]
- nVidia始まったな。
- 579 名前：デフォルトの名無しさん mailto:sage [2012/03/23(金) 00:17:06.29 ]
- >>577
warp の扱いどうなるんかな。。。
- 580 名前：デフォルトの名無しさん mailto:sage [2012/03/23(金) 00:54:42.91 ]
- >>579
ttp://pc.watch.impress.co.jp/docs/column/kaigai/20120322_520640.html
> 32スレッドのWARPに同じ命令を実行する、この基本は、Keplerでも変わっていない。
らしいから、変わらないんじゃないかな。
GF104/114のSMには48コアと2ワープスケジューラ、4ワープディスパッチャで
GK104のSMXには192コアと4ワープスケジューラ、8ワープディスパッチャになっている。
その上レジスタ数は倍、L1キャッシュ/シェアードメモリはそのままってことは
GF104/114よりさらにピーキーになっているのかな？
- 581 名前：デフォルトの名無しさん mailto:sage [2012/03/23(金) 01:01:32.84 ]
- >>580
あれ？
コンスタントキャッシュって無くなった？
L1/L2キャッシュがその役割を担ってる？
ということはFermiからか･･･
- 582 名前：デフォルトの名無しさん [2012/03/23(金) 01:06:15.11 ]
- 48コアが192コアになったのに
レジスタは2倍、
共有メモリは据え置き。
どーすんだこれ。。
- 583 名前：デフォルトの名無しさん mailto:sage [2012/03/23(金) 01:54:37.82 ]
- レジスタ足りんくなりそうな。
- 584 名前：デフォルトの名無しさん mailto:sage [2012/03/23(金) 04:09:37.63 ]
- Keplerはクロック落としてパイプラインを浅くする設計
演算器のレイテンシが小さくなるならレジスタの消費量は変わらない
Fermiの18cycleは頭おかしすぎた
これが例えば6cycleにになればレイテンシ隠蔽に必要なスレッド数が1/3になるから問題ない
- 585 名前：デフォルトの名無しさん mailto:sage [2012/03/23(金) 04:14:42.78 ]
- x86 CPUと同じ道を辿ってるのか
- 586 名前：デフォルトの名無しさん [2012/03/23(金) 15:07:43.87 ]
- 誰か２６次元計算してくれ、1000ｺｱくらいじゃマジに足らんぞｗ
- 587 名前：デフォルトの名無しさん mailto:sage [2012/03/23(金) 15:59:10.59 ]
- 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;
|
 |
