09/12/06 10:18:05
>>313
仰る通りだったようです。shared[125*32]をshared[32]としたら劇的に速くなりました。
バンクコンフリクトを疑うあまりブロックあたりのスレッド数を32にしてたのも不味かったようですね。まだ未確認。
>あとpxがアドレス参照でローカルメモリ（VRAM上)に行ってないか心配だ。
*((uint1 *)&px) = shared[i] を px.x = shared[i].x;
と書き換えたところ速度アップしたのでその可能性大です。最適化で同じコードになるんじゃね？と勝手に思ってました。
>>314
>vram2[i] がレジスタのってたりしないかな。
このコードは質問のために単純化したものだったのですが、単純化前にでptxファイルを出力して比較したときは
ld.global.u32のところがld.shared.u32に変化しただけでしたのでレジスタにのったり最適化で消えたりはしてなさそうです。
とはいえptxの書式の資料を見つけられなかったので自信ありません。
参照を i+1の意味はちょっとわかりませんでした。すみません。
>>313-314
おかげさまでグローバルメモリより共有メモリが遅いという現象は解決しました。ありがとうございます。
ブロック専用なのにextern __shared__ の構文がなぜあるのか不思議に思っていたのですがこういう理由だったのかと。
しかし大きな共有メモリで遅くなるのは厄介な仕様ですね。