- 1 名前：デフォルトの名無しさん mailto:sage [2007/09/17(月) 14:54:28 ]
- 本家
developer.nvidia.com/object/cuda.html
- 231 名前：デフォルトの名無しさん mailto:sage [2008/03/13(木) 03:10:17 ]
- warpサイズって簡単に言うと何？
- 232 名前：デフォルトの名無しさん mailto:sage [2008/03/13(木) 06:48:58 ]
- 同一インストラクションが同時に走る、演算ユニットの集まり?
- 233 名前：デフォルトの名無しさん mailto:sage [2008/03/13(木) 13:50:16 ]
- CUDA初心者です。
CUDA難しいよ
難しいよCUDA
・<<<Dg, Db, sizeof(float)*num_threads*4>>> て書いたら動かないんすかね。
int sharemem_size = sizeof(float)*num_threads*4;
<<<Dg, Db, sharemem_size >>> て書いたら動いた。
・畳み込み演算させたら、CPUでやるより1.7倍しか速くないの(´・ω・`)
・動いてる間、システム全部固まってる。最初焦った。
CUDAの道険しいﾅﾘ
- 234 名前：デフォルトの名無しさん mailto:sage [2008/03/13(木) 15:37:14 ]
- 初歩的な質問かとは思いますが、各ストリームプロセッサにそれぞれ1つずつスレッドが割り当てられるのですか？
- 235 名前：デフォルトの名無しさん mailto:sage [2008/03/13(木) 18:22:14 ]
- __syncthreads()ではブロック内でしか同期できませんが、全ブロックを同期したいときは
カーネルを抜けていったんCPU側に戻すしかないのでしょうか？
- 236 名前：デフォルトの名無しさん mailto:sage [2008/03/14(金) 02:24:24 ]
- cudaMallocHostとCのmallocの違いを教えてください
- 237 名前：デフォルトの名無しさん mailto:sage [2008/03/14(金) 05:10:46 ]
- >>235
確かカーネルは非同期で実行されたはずだから、
カーネルを連続で呼んだら同期されない・・・と思う。
<<<カーネルの呼び出し>>>
cudaThreadSynchronize();
<<<カーネルの呼び出し>>>
ならいけるんじゃないの？自信ないけど。
- 238 名前：デフォルトの名無しさん mailto:sage [2008/03/14(金) 11:06:04 ]
- >>236
cudaMallocHostでメモリ確保すると、ページングなしスワッピングなし
物理メモリに連続して常駐、のメモリ領域が確保されるはずなのだ！
と思ってcudaMallocHost使ってます。mallocとmlockの合せ技？
- 239 名前：デフォルトの名無しさん mailto:sage [2008/03/14(金) 16:02:07 ]
- Windows, VC2005なんですが、PTXを出すためのオプションはどこに
書けばいいんどすか( ・ω・`)
苦労してPhenomの6倍速いまで持ってきたんですが、さすがに
この先はアセンブラ見てみないと何をどうすればいいか分からんどす。
モノは信号処理用の畳み込み(4096)どす。
- 240 名前：デフォルトの名無しさん mailto:sage [2008/03/14(金) 16:41:32 ]
- >>239
プロパティでカスタムビルドであることをチェックできない?
オプションは言うまでもなく -ptx
- 241 名前：デフォルトの名無しさん mailto:sage [2008/03/14(金) 16:56:28 ]
- グラフィックカードを2枚挿した状態で、メインメモリからGPU側に転送する際に
どのようにしてGPUを指定すればよいのでしょうか？
カーネルから見られるのは、1枚のGPUだけという認識でよろしいのでしょうか？
- 242 名前：デフォルトの名無しさん mailto:sage [2008/03/14(金) 23:48:11 ]
- >>241
cudaGetDeviceCount()からcudaChooseDevice()あたりでは。プログラミング
ガイドのD.1。二枚刺しﾃﾗｳﾗﾔﾏｼｽ
>>240
プロジェクトのプロパティのカスタムビルドステップ、等にもコマンドライン
指定が何もなかとです。次はVSディレクトリに潜って.cuのRuleファイルを
探してみるです。
- 243 名前：デフォルトの名無しさん [2008/03/15(土) 03:32:18 ]
- RuntimeAPIとDriverAPIの違いを教えてください。
- 244 名前：デフォルトの名無しさん mailto:sage [2008/03/15(土) 12:17:00 ]
- CUDA素晴らしすぎです。ようやくグリッドとブロックとスレッドと
globalとsharedとconstantを理解して、今「Xeonより100倍速いぜ！」
な速度を叩き出しています。もっとチューニングがんばってみよう。
>>243
RuntimeよりDriverの方が低レベルとして書いてありますけど、
似たようなの一杯有るし、違い良く分からないですよね。
- 245 名前：デフォルトの名無しさん mailto:sage [2008/03/15(土) 13:20:14 ]
- 恐らくは、Driverの方はC++で使うことを想定してない。
実際問題として、nvccはC++としてコンパイルするのでRuntimeAPIの方が使いやすいと思う。
- 246 名前：デフォルトの名無しさん [2008/03/15(土) 17:14:57 ]
- >>244
segdmm のC800では120GFlops程度と言われていますが
4coreのXeonでも75GFlopsです。
どのくらい違いますか教えてください。
- 247 名前：244 mailto:sage [2008/03/16(日) 00:22:24 ]
- >>246
やってみたのが4096単位データ×128個に対して4096のフィルタで
コンボリューション掛ける、なんですが、Xeon2GHzの1スレッド
だと250M操作/秒位しか出来ないですよ。キャッシュミス連発？
GeForce8800ですと、28G操作/秒で動きます。
- 248 名前：244 mailto:sage [2008/03/16(日) 04:08:51 ]
- あ、そんで、Xeonでやるとこうだってのはもれじゃ無くて既存の
なので、弄れないし、作りはよく分からんのです。とりあえず
「このルーチン速くするライブラリ作れんか？CUDAとか言う奴で」
と言われてやってみたら100倍速いので怒りが有頂天になった所です。
でも冷静に考えると、CPU側もカリカリに書けばあと10倍位速くなる
可能性はありますわな。たぶんSSEとか使ってないし。
でもそこから10倍にするには、つまりもう何台かパソコン追加しないと
ならんわけなので、やっぱりCUDA可愛いよ可愛いよCUDAｗ
- 249 名前：デフォルトの名無しさん mailto:sage [2008/03/16(日) 08:32:33 ]
- 私のところじゃ、Xeonの3-10倍だなぁ。単体テストで。
# 他の演算と組み合わせると、1.5-2倍が限度（TT
- 250 名前：デフォルトの名無しさん mailto:sage [2008/03/16(日) 09:07:29 ]
- 「Xeon」てみなさんが言ってるのは、DPマシン8コアのこと？それとも1コアのこと？
1コアに負けるんじゃ確かにちょっと悲しいよね。
- 251 名前：デフォルトの名無しさん mailto:sage [2008/03/16(日) 09:14:14 ]
- 1スッドレ、って書いてるじゃん
- 252 名前：デフォルトの名無しさん mailto:sage [2008/03/16(日) 09:34:49 ]
- つまりこういう暑苦しい感じが最強だと。
喧嘩してる場合じゃねぇ！Quadコア×デュアルCPU＋8800GTXの2枚刺し、
の友情で最強を目指そうじゃないかお前ら！夕日に向かって海まで走れ！
敵は巨大行列のLU分解だ！手ごわいぞO(N^3)！
- 253 名前：デフォルトの名無しさん mailto:sage [2008/03/16(日) 12:57:56 ]
- グリッド数、ブロック数、スレッド数の決め方がよくわかりません。
例えば独立な300個のスレッドがあった場合、
一つのブロックに300のスレッドを割り当てるのがよいのでしょうか？
それともスレッド数は1つで300のブロックを作成するのがよいのでしょうか？
ワープ数などはどのように考慮すればよいのでしょうか？
- 254 名前：デフォルトの名無しさん mailto:sage [2008/03/16(日) 13:54:06 ]
- >>253
あくまでも漏れの場合ですが、
ブロック数＝16。これは使っているカードの「Multi Processor」に合せて。
違うグリッドに属するデータのやり取りは出来ない事に注意。
スレッド数＝512。目安で。
で、カーネル関数でループで処理を回す場合、最外では
for (i = 0; i < maxcalccount; i=i+スレッド数×ブロック数)
のようにします。これで8192単位での並列化になります。
ループ内部では、グローバルメモリからたとえば１ワード読むなら
sharedメモリ[スレッドID]
＝グローバルメモリ[スレッド数×ブロック数＋ブロックID×スレッド数＋スレッドID]
となります。
floatを使う場合、スレッドが512なら、各スレッドで最大8ワードもてますが、
sharedメモリは使い切らないほうが良いようです。
- 255 名前：デフォルトの名無しさん mailto:sage [2008/03/16(日) 14:57:39 ]
- ブロック数は、どうせ多く割り当ててもCUDA内部で直列に並べるだけだから
非同期で少しでもCPUと並列にしたい場合を除けば大目に割り当ててOK。
スレッド数についても多い分はどうせ別のワープに割り当てられるから多めでOK。
但し、同期を取る場合には多過ぎるとダメ。
手元のデバイス関数の場合、ブロック数*スレッド数は少なくとも1024か2048以上必要(8800GTXの実測で)。
これらを踏まえると、スレッド数が32ならブロック数は64以上、スレッド数が64ならブロック数は32以上くらいか。
ブロック数の上限は、実測しながら適当に調整するとして、大体1024を超えるといくつでも変わらないと思う。
# これも、具体的なテスト用のサンプル用意したいところだね。
- 256 名前：デフォルトの名無しさん mailto:sage [2008/03/16(日) 16:42:50 ]
- >>188みたいな話が載っている本で使っている本があれば教えてほしいなあ。
入門書の理屈を素直になぞっているような人間には、普通のcpuとgpuとの差を出せなさそう。
- 257 名前：デフォルトの名無しさん mailto:sage [2008/03/16(日) 17:01:32 ]
- >>256
>188は統計学を知らんと自分で導けないし、統計学学ぶ香具師がプログラマになるとも限らんしね。
Webでも知識があることを前提としているか、律儀に平均値との差を求める方式しか載ってなかったり。
私の場合は、電卓の統計機能が個別のデータと平均値の差を使わずにどうやって標準偏差を得るのか
不思議に思って調べて以来の知識かな。
ってことで、統計学を知らない私も>188みたいな話の解説があるなら見たい希ガス。
- 258 名前：デフォルトの名無しさん mailto:sage [2008/03/16(日) 17:05:10 ]
- CUDAというかGeforceってSIMDなの？
SSEは確かに1つの命令で4つとかの足し算が行われるけど、
CUDAの場合、どれがそれに当てはまるの？
- 259 名前：253 mailto:sage [2008/03/16(日) 17:17:24 ]
- >>254>>255
早速のご回答ありがとうございます。
基本的にはスレッド数は目一杯であとはブロックすうを調整すると言うことでいいみたいですね。
ところで、処理する数がスレッド数×ブロック数で割り切れない場合、最後のループはどうするのがいいのでしょうか？
やはりifでやるしかないのでしょうか？
- 260 名前：255 mailto:sage [2008/03/16(日) 17:38:35 ]
- >>259
んにゃ、私はブロック数もスレッド数も実測で決定している。
で、割り切れない場合はループ数で制御するのがいいみたい。
# やっぱり説明するためにはさんぷるが必要かw
>>258
GPUの場合は、SIMDじゃなくてMIMDということになるのかな。
CUDAで考えると、Warp単位で同じインストラクションを走らせて、
複数のデータスレッドを処理するイメージ。
だから、条件分岐をCUDAで書いても、Warp単位で同じインストラクションなので
データスレッドごとに条件が違うと無駄が生じてしまう罠。
詳細は手ごろなサンプルがないので割愛w
- 261 名前：253 mailto:sage [2008/03/16(日) 18:08:52 ]
- >>260
やっぱり基本は実測なのですね。
ところでループ数で制御するというのはどういうことなのでしょうか？
処理する数/(スレッド数×ブロック数)=ループ数
だと思うのですが・・・
- 262 名前：255 mailto:sage [2008/03/16(日) 18:27:15 ]
- 例えばこんな感じ。
__global__ void func(float const * a1, float * a2, unsigned num)
{ // 実際には一時変数を使った方がいい希ガス
for (unsigned idx = blockIdx.x * blockDim.x + threadIdx.x; idx < num; idx += gridDim.x * blockDim.x) {
a2[idx] = a1[idx] * a1[idx];
}
num=ブロック数*スレッド数ならどのスレッドでもループは一回だけ回る。
num<ブロック数*スレッド数なら、暇なスレッドが発生する(から効率は宜しくない)。
num/(ブロック数*スレッド数)が11と12の中間なら一部のスレッドは12回回って残りは11回回る。
このループ回数が1や2じゃなければ、暇なスレッドの割合が相対的に少なくなる寸法。
# これが理由で、ブロック数は無闇に増やせばいいというわけでもないということになる。
あー説明が下手な私（：；
# サンプル作ったら誰か買わない?w
こういう回し方をすることによって、近くのスレッドが近くのメモリをアクセスする状態のままループが進行する。
つまり、ProgrammingGuideで言うところの"coalesced"。これをプログラミングガイドでは「結合した」と訳しているけど……
- 263 名前：253 mailto:sage [2008/03/16(日) 20:27:25 ]
- >>262
ありがとうございます。
なるほど非常にうまい手ですね。
プログラミングガイドはなんか日本語訳がめちゃくちゃで非常に読みにくいですが、
がんばって解読したいと思います。
- 264 名前：デフォルトの名無しさん [2008/03/16(日) 20:55:07 ]
- shared memoryを使ったら、global memoryを直接使うより遅くなってしまいました。
どうも、カーネルを実行する際のshared memoryの確保に時間がかかっているようなのですが、
カーネルを実行するたびに毎回確保するのではなく、常時確保するような方法はないでしょうか？
- 265 名前：255 mailto:sage [2008/03/16(日) 21:20:30 ]
- >>264
その現象は確認していないので不明です。
が、保持する方法は多分ないでしょう。
実は一つ問題があって、shared memoryは16バイト同時アクセスができないので
その点でglobal memoryより遅いのは確認済みです。
# ptx出力を見れば見当がつきますが。
- 266 名前：デフォルトの名無しさん mailto:sage [2008/03/17(月) 00:53:35 ]
- >>261
もれは、CUDAに投入する前に、CPU側で後詰めゼロとかやって、
「CUDA内では一切条件判断は不要」を原則にしています。
>>264
global、sharedへのアクセスがcoalescedじゃないと、めがっさ遅くなります。
それと、sharedに確保した後で、ブロック内の各スレッドでデータを
shareしない場合には、そもそもsharedを使う理由無いです。レジスタ
に順番にグローバルから読んで処理すればいいです。
- 267 名前：デフォルトの名無しさん mailto:sage [2008/03/17(月) 06:10:01 ]
- ブロック間の同期をとる方法を教えてください
- 268 名前：デフォルトの名無しさん mailto:sage [2008/03/17(月) 09:17:39 ]
- 「ブロック」は、そもそも同時にブロック001とブロック002が動いているか
どうかも分からないんで、同期出来ないんでは？
CPU側で、カーネル関数その１呼び出して、cudaThreadSyncronizeして、
カーネル関数その２呼び出す。ではないかと。
- 269 名前：255 mailto:sage [2008/03/17(月) 09:22:38 ]
- 一応、__syncthreads()という手はありますよ。
但し、全ブロックが同時に動いているわけではない点には注意。
つまり、上の方に書いた分割方法で書いたような大きなブロック数の時には
恐らく巧くいかないかかなり待たされることになるかも知れず。
私の手元のロジックでは、ブロック数が64の時には巧いこと同期が取れています。
# C870, 8800GTX, GTS, GTでしか確認してないけど。
- 270 名前：デフォルトの名無しさん [2008/03/17(月) 14:41:06 ]
- >>269
それって各ブロック内のスレッドの同期じゃないの？
- 271 名前：255 mailto:sage [2008/03/17(月) 14:55:21 ]
- あー、そうそう。失敬、その通りです。
>269は全て、「ブロック」を「スレッド」に読み替えてください。
# どっかでブロック数に切り替わってしまった……_/￣|○
と言うことで、やっぱりブロック間の同期は取れないのでした。
- 272 名前：デフォルトの名無しさん [2008/03/17(月) 15:29:11 ]
- >>268
カーネル2はカーネル1が終了してから実行されるんじゃないの？
- 273 名前：デフォルトの名無しさん mailto:sage [2008/03/17(月) 15:56:55 ]
- >>272
そこが良く分からんのです。
サンプルではcudaThreadSyncronizeでカーネル呼び出しを挟んでるんで、
一応もれも呼び出すことにしておるのですが。
- 274 名前：デフォルトの名無しさん mailto:sage [2008/03/17(月) 16:04:25 ]
- >>273
今、挟んだのと挟んでないのを比べてみたけど、全く同じ結果だった。
ifがあるので、各ブロックはばらばらのタイミングで終わるはずなのだが。
- 275 名前：255 mailto:sage [2008/03/17(月) 16:10:18 ]
- 前にも書いたかもしれないけれど、カーネル終了を明示的に知る手段はないので注意。
つまり、<<<>>>による呼び出しから復帰しても、GPUの処理が終了したわけではないので。
なので、次のような呼び出しをした場合、k1()とk2()は並列に動く可能性があるわけ。
--
k1<<<1, 8>>>();
k2<<<1, 8>>>();
--
終了を保証するには、cudaThreadSynchronize()を呼ぶかcudaMemcopy*()を呼ぶかすればOK。
- 276 名前：デフォルトの名無しさん mailto:sage [2008/03/17(月) 19:44:31 ]
- カーネルのループの段数を増やしたら、スレッド数が512だとカーネルの実行ができなくなった。
256ならできる。257はだめ。なんでだろう・・・
- 277 名前：デフォルトの名無しさん mailto:sage [2008/03/17(月) 19:51:18 ]
- 1BlockのThread数の制限512ってのと関係があるのでは？
- 278 名前：デフォルトの名無しさん mailto:sage [2008/03/17(月) 19:59:35 ]
- ちょっと舌足らずだったかもしれませんが、カーネル内のループの段数を増やしたところ、そのカーネルが実行できなくなってしまいました。
そこで、スレッド数を512から256に減らしてみたところうまくいきました。
スタックか何かがあふれてしまったのではないかと疑っているのですが・・・調べる方法がわからず行き詰まってしまいました。
- 279 名前：デフォルトの名無しさん mailto:sage [2008/03/17(月) 20:10:15 ]
- 実行できないってことは何かエラーが出るの？
- 280 名前：デフォルトの名無しさん mailto:sage [2008/03/17(月) 23:59:13 ]
- 使っているレジスタの総数が8192を超えたとか。
unrollしすぎだとか。
- 281 名前：デフォルトの名無しさん mailto:sage [2008/03/18(火) 05:17:35 ]
- カーネルを2つ同時に実行できますか？
2枚のカードでそれぞれ別のタスクを割り当てて、並列で実行したいのですが。
- 282 名前：デフォルトの名無しさん mailto:sage [2008/03/18(火) 07:49:07 ]
- ブロック数が少なければ、勝手に同時実行するな。
- 283 名前：デフォルトの名無しさん [2008/03/18(火) 08:14:14 ]
- Maximum memory pitchって何？
- 284 名前：デフォルトの名無しさん mailto:sage [2008/03/18(火) 08:33:44 ]
- VCのnmakeって使えないの？
cudaをコンパイルできるサンプルMakefileがあったら教えてください。
- 285 名前：デフォルトの名無しさん mailto:sage [2008/03/18(火) 11:56:28 ]
- >>281
サンプルのMultiGPU、8800GTS＋8400で試したら動いたよ。
単にこれ、WINAPIのCreateThreadして中でカーネル関数呼んでるだけですな。
でWaitForMultipleObjectsでWindowsスレッドの終了待っている。
>>283
2次元配列でGPUのメモリを使うときの幅じゃないすか。
cudaMemcpy2Dとか使うときに気にするんだと思われ。
ごめん二次元配列使ってないので良く分からない。
8800GTSも8400も262144bytesです。
- 286 名前：デフォルトの名無しさん mailto:sage [2008/03/18(火) 14:42:22 ]
- 演算性能がボトルネックになっているのか、バンド幅がボトルネックになっているのか知る良い方法はないものか
- 287 名前：デフォルトの名無しさん mailto:sage [2008/03/18(火) 15:15:30 ]
- >>286
演算抜きでcudaMemcpyのHostToDevice、DeviceToHostだけやって
時間を測って見る、そして演算入りの時間と比べる、でどうだろう。
- 288 名前：デフォルトの名無しさん mailto:sage [2008/03/18(火) 16:42:29 ]
- __constant__ floatとconst floatの違いって何？
__constant__の方が速いの？実測してみたけど、全然違いが出ないのだけど。
ちなみに何故か定数を直接、式の中に書いたら唯一5％ほど遅くなった。
- 289 名前：デフォルトの名無しさん mailto:sage [2008/03/18(火) 18:11:38 ]
- プログラミングガイドの5章が全然わかんねー
結局のところcoalescedって何よ
Figure 5-2の右側は何でダメなの？？？？
教えて偉い人
- 290 名前：デフォルトの名無しさん mailto:sage [2008/03/18(火) 23:54:19 ]
- >>288
__constant__は全部で64Kバイトある定数格納用メモリ、ホストから
memcpyすることが出来る。const floatは8192個のレジスタを使って
処理される。なので、普通のコード内定数はconst、ホストからどさ
っと書き込みたい固定データは__constant__でいいと思い。
>>289
5-2の右側は、132=33×4、だから？ 各スレッドは、
16×sizeof()×何か＋スレッドID、にアライメントされた
アドレスにアクセスしないとcoalescedしてくれないと。
- 291 名前：289 mailto:sage [2008/03/19(水) 04:35:23 ]
- >>290
coalescedを辞書で調べても「合体した」とかしか出てこないので、いまいち意味がよくわからないのですが、
結局0番スレッドは
a[0]
a[16]
a[32]
:
１番スレッドは
a[1]
a[17]
a[33]
:
以外のアドレスにアクセスしないようにすればよいと言うことでいいのでしょうか？
- 292 名前：デフォルトの名無しさん mailto:sage [2008/03/19(水) 04:51:32 ]
- 要は、配列のアクセスはarray[m * blockDim.x + threadIdx.x]のような形で
blockDim.xは16の倍数になればよろしって話なんじゃないかと。
- 293 名前：デフォルトの名無しさん mailto:sage [2008/03/19(水) 11:08:25 ]
- cuMemGetInfoフリーメモリと合計メモリを取得したのだけど、４MBにしかならない。
使い方は
unsigned int f,m;
cuMemGetInfo(&f,&m)
でいいんだよね？
- 294 名前：デフォルトの名無しさん mailto:sage [2008/03/19(水) 11:35:28 ]
- >>291
つまり「coalesced」は・・・もれの理解では・・・
・sharedメモリには入出力が32bit×8ポートずつあるかも。
でも開始アドレスは一個かも。
・各スレッドからのアクセス要求アドレスが「0,4,8,12,16,20,24,28」の場合、
GPU側で「ではsharedメモリの出力を綺麗にストリーミングプロセッサに
一個ずつ割り当てればいいのであるな、開始アドレスは0だな」と判断し
てくれるのかも。
>>293
もれの場合は、debugでは動くけどReleaseでは動かないのね。
debugでは FreeMem 465529088 TotalMem 536543232
と取れた。8800GTS/512MB。理由分かったら教えてください。
- 295 名前：デフォルトの名無しさん mailto:sage [2008/03/19(水) 19:26:24 ]
- >>290
__constant__って書き換えられるの？
- 296 名前：デフォルトの名無しさん mailto:sage [2008/03/19(水) 22:06:01 ]
- >>295
__constant__は定数メモリと考えればいい。
つまり、GPUからは書き換えられないのでCPUから書き換えると。
# constな変数は初期化はできるけど代入できないのと似てるかもしれない。
>>294>>293
>294のdebugでの数値は多分正しい値が出ているみたいだね。
>293の方は、もしかしたらエラーが返っていないかな?
私のところでも、整数にして201だったかのエラーが返ってきていたのだけれど、
その状態では値をセットしてくれないようだった。
つまり、もしかしたら>293の4MBとというのはfとmが未初期化で偶々入った不定値かも。
# 気になるなら、0で初期化してみればいい。
>>294>>291
coalescedなアクセスパターンにすることは、sharedだけでなくglobalでも重要なので
慣れておくことをお勧め。まぁ、例えば>262に書いたようにすればいいだけだけど。
- 297 名前：デフォルトの名無しさん mailto:sage [2008/03/19(水) 23:54:01 ]
- >>295
cudaMemcpyToSymbol()でCPUからセット出来るよ！
速いし64KBあるんで結構便利です。
- 298 名前：デフォルトの名無しさん mailto:sage [2008/03/20(木) 00:01:09 ]
- 問題は、ベクタアクセスできる__global__と違って1データずつのアクセスになるから
逆に遅くなるケースもあることだな。
- 299 名前：デフォルトの名無しさん [2008/03/20(木) 12:55:05 ]
- ややこしさはCellとあんまり変わらなそうだな。
- 300 名前：デフォルトの名無しさん mailto:sage [2008/03/20(木) 14:25:54 ]
- >>299
とにかくメモリのコピーがヤヤコシイんだわ。ホスト側、GPUの
グローバルメモリ、GPUのチップ内メモリで転送しまくらないと
いかんで。やり方間違えると全然性能出ないし。
Cellはその辺どうなの？
- 301 名前：デフォルトの名無しさん mailto:sage [2008/03/20(木) 14:28:32 ]
- 256kしかないLSでやりくりするのが大変って聞いたな
- 302 名前：デフォルトの名無しさん mailto:sage [2008/03/20(木) 15:10:07 ]
- 両方試した私に言わせて貰えば、どっちもどっち。
確かにCBEは256KiBの壁がねぇ。GPUも64KiBの壁やcoalescedの沼があるけど。
超越関数を使える点ではCUDAが有利。ホストの性能でもPPEじゃ結構泣けるし。
- 303 名前：デフォルトの名無しさん mailto:sage [2008/03/20(木) 16:02:10 ]
- PPEはひどいよな。ホステスにC2D使ってるけど、ifがいっぱいあるような場合はC2Dの方が速いしね。
- 304 名前：デフォルトの名無しさん mailto:sage [2008/03/20(木) 16:14:13 ]
- >>298
それ実験してみたんだけど、変わらないみたい。
・global→shared(行列多数)と__constant__に置いた定行列で行列積
・global→shared(行列多数)とsharedの一部に置いた定行列で行列積
で、後者が1%遅いくらいだった。リードオンリーなだけで、コアとの距離や
所要クロック数はconstantもshared・レジスタも同じなのかもと？
- 305 名前：デフォルトの名無しさん mailto:sage [2008/03/20(木) 16:28:58 ]
- CUDAで逆行列を求める方法を教えてください
- 306 名前：デフォルトの名無しさん mailto:sage [2008/03/20(木) 16:49:55 ]
- page-lockedってどういう状態を表すのでしょうか？
- 307 名前：298 mailto:sage [2008/03/20(木) 17:42:11 ]
- >>304
-ptxでニモニックを出力してみれば違いが判るかと。
例えば、ld.global.v2.f32はあるけどld.const.v2.f32はないからld.const.f32が2回になってしまう。
coalescedなアクセスができるglobalは4クロックでアクセスできるからsharedやconstと変わらないわけで。
尤も、一旦そのパターンから外れるとglobalは数百クロックだそうだから途端に劇遅になるけど。
- 308 名前：298 mailto:sage [2008/03/20(木) 17:48:04 ]
- >>305
先ずは、逆行列を求めるプログラムをCで書いてみてください。
それの、ループの部分を分割する方向でGPU関数を作るのが第一歩になります。
つーかさぁ、一行質問する人達ってなんなの?
情報交換する気がなくて、単にくれくれの精神なんだったら勘弁して欲しいんだけど。
# なんか回答者が数人とそれより少し多いレス要員だけで持ってる気がするよ……
>>304
あーそうそう、書き忘れたけどレジスタアクセスは1クロックじゃなかったカナ。
確実にメモリよりも速いみたい。
- 309 名前：デフォルトの名無しさん mailto:sage [2008/03/20(木) 19:43:17 ]
- nvccにはO2とかの最適化オプションは無いのでしょうか？
- 310 名前：デフォルトの名無しさん mailto:sage [2008/03/20(木) 20:26:00 ]
- -Wallが使えないみたいね
たぶん無さそうな予感
- 311 名前：298 mailto:sage [2008/03/20(木) 21:19:57 ]
- 一部のオプションはgccと共通ですよ。
例えば、-O3(恐らく-O2なども)や-pgは使えます。
- 312 名前：デフォルトの名無しさん mailto:sage [2008/03/20(木) 23:15:10 ]
- >>299
どっちもどっちで、用途や慣れの問題だと思う。302がいうように、SFUの有無は
大きいが、それも用途しだい。
自分的には比較的情報がオープンなCellのほうがいろいろいじりやすいと思うが、
最近人気なくて心配。たぶんそのうちLarrabeeが全部かっさらってくんじゃないかな。
- 313 名前：デフォルトの名無しさん mailto:sage [2008/03/22(土) 18:23:49 ]
- GeForce 9800 GX2で現バージョンのCUDAは動きますか？
- 314 名前：デフォルトの名無しさん mailto:sage [2008/03/22(土) 18:28:28 ]
- ストリーム番号の同じカーネルやメモリコピーは順番に、違うものは並列に実行される？？の認識でいい？？？
- 315 名前：デフォルトの名無しさん mailto:sage [2008/03/22(土) 18:43:27 ]
- テクスチャメモリの容量はどのようにすれば調べられるのでしょうか？
どこかに資料があるのでしょうか？それとも何かAPIがあるのでしょうか？
- 316 名前：デフォルトの名無しさん mailto:sage [2008/03/22(土) 23:25:42 ]
- いつも適当に答えてるおじさんです。まことにもうしわけない。
>>313
もれ、8800GTS(G92)で動いてるから、×２で二個見えるんじゃないかと
予想しますけど、一応どっちも対応外だから。漢は度胸でひとつ人柱に。
>>314
ストリームが二つあっても、実行順は関係ないのではと予想。ストリーム
１の半分実行した後ストリーム２の半分実行、もありえるかと。
cudaStreamSynchronizeでストリームの実行完了を待たないといけないです。
>>315
テクスチャメモリはグローバルメモリのキャッシュですが、実サイズが
不明ですね。cudaGetDevicePropertiesでは「alignment=256Bytes」と
読めます。これ以下と以上で実行時間を調べてみると、キャッシュミス
してるかしてないかわかるのではないでしょうかと。
- 317 名前：デフォルトの名無しさん mailto:sage [2008/03/23(日) 13:19:47 ]
- openSUSE 10.3でCUDAは動きますか？
- 318 名前：デフォルトの名無しさん mailto:sage [2008/03/23(日) 13:33:44 ]
- 動くらしい
- 319 名前：デフォルトの名無しさん mailto:sage [2008/03/23(日) 17:02:07 ]
- 次のバージョンっていつ頃出るの？
新しいカードも出ているわけだしそろそろ、SDKの方もアップデートしてほしいものなのだが
- 320 名前：デフォルトの名無しさん mailto:sage [2008/03/23(日) 23:52:41 ]
- Vista64と9600,9800対応版が欲しいすね。
早く出してくれないとAMD/ATIに浮気したくなりますよ。
- 321 名前：デフォルトの名無しさん mailto:sage [2008/03/24(月) 10:55:41 ]
- サンプル作ると言いつつなかなか暇が取れない私が来ましたよ。
>>313
ドライバが対応しているかどうかが問題。ドライバが認識さえすれば、動くとは思いますが。
>>314
ストリームのご利用は計画的に。ちなみに、CUDA1.0世代の(要はG8xかな)GPUだと
処理の隠蔽ができないらしく、全く速くならないという噂もあります。
>>317
>27-30
荒らしじゃないんだったら、スレくらい検索してほしいと思う今日子の頃。
# いや、荒らしなら検索如何に関わらずご遠慮願いますが。
>>319
取り敢えず、待つしか。下手すると夏の新製品までお預けの可能性も。
>>320
え〜、待っていれば倍精度も来るのに。
- 322 名前：デフォルトの名無しさん [2008/03/26(水) 23:04:49 ]
- 英wikiでは、9600GTでもcudaできることになっている
悪い子と漢は居ね゛ーがー
- 323 名前：デフォルトの名無しさん mailto:sage [2008/03/26(水) 23:38:52 ]
- >>322
もれ9600GT買おうと思ってる。
ARCTICのS1Rev2が付く、という報告待ち・・・。ファンレスで使いたくて。
- 324 名前：デフォルトの名無しさん [2008/03/31(月) 21:17:49 ]
- www.nvidia.com/object/cuda_learn_products.html
9800GX2 CUDA OKだとでてるよん
- 325 名前：デフォルトの名無しさん mailto:sage [2008/04/01(火) 00:04:20 ]
- おー、更新されている。情報THX。
ドライバは変わっていないから、チップ依存なんですかね。
- 326 名前：デフォルトの名無しさん [2008/04/04(金) 22:15:16 ]
- ML115祭りに参加してOpetron使用をポチりました。
で、nbody.exeを実行すると・・・・ 160〜170GFLOPSという値が。
うぅ・・速過ぎる。初めて見るそのパワーに感動ものですた。
- 327 名前：デフォルトの名無しさん [2008/04/04(金) 22:15:54 ]
- 忘れてた GPUは8800GTですぅ。
- 328 名前：デフォルトの名無しさん mailto:sage [2008/04/04(金) 22:24:35 ]
- そろそろまじめに実用的実装に入らないのかな？
圧縮解凍ソフトに組み込むとか
暗号化ソフトに組み込むとか
一番簡単なのは画像処理ソフトだろうけど
- 329 名前：デフォルトの名無しさん mailto:sage [2008/04/05(土) 02:08:09 ]
- 大真面目に実用的な実装で悪戦苦闘していますが何か。
- 330 名前：デフォルトの名無しさん mailto:sage [2008/04/05(土) 09:56:17 ]
- 画面の解像度によって使えるメモリ量は変わりますか？
- 331 名前：デフォルトの名無しさん mailto:sage [2008/04/05(土) 12:08:56 ]
- CUDA使うアプリをインストーラで配るにはどうしたらいいもんだろうか？
というところが。Vista対応してないし。
|
 |
