Ryuzのブログ

なかなか趣味の時間が取れない期間が続いていますが、GPU的なものの企ては続いています(数ヶ月ごとに数時間しか進捗してないですが)。

最後に考察した除算部分について、忘れ始めてきたので備忘録的に思い出しながら書いておきます。

射影変換ができてこそ3D-CGらしくなってくるわけですが、それには除算が必須です。
一方で単精度程度の除算や逆数演算はXilinxのIPコアだとレイテンシが長すぎて他の乗算などに隠すのが難しく、ボトルネックになりそうだというのが見えてきました。
そもそも画像画素値なら8bit、テクスチャ座標でも12bit程度の最終精度しかいらないので中間計算としても極端に厳密なものは要らないはずです。
そしてそもそも射影変換の場合、同じ除数での除算が続くので、先に逆数を求めておいて後は乗算のみとするのが効率的な気がします。

そこで、案として

• 平方の逆数を求めてから二乗する
• テーブルを使って直接逆数を求める

の２つの簡易逆数生成が考えられます。

平方の逆数は Fast inverse square root というテクニックがあるようで

https://en.wikipedia.org/wiki/Fast_inverse_square_root

http://homepage1.nifty.com/herumi/diary/1309.html

あたりが詳しい情報と解説があり、有難い限りです。

テーブルを使って直接逆数を求める方法は

http://www.kumikomi.net/archives/2010/10/ep32suc4.php?page=2

ディジタル数値演算回路の実用設計

に解説のある手法などがありそうです。

  ちなみに除算について、Intel CPU の実装としてAVX2のSIMD命令を  Intel(R) Intrinsics Guide で調べると、 _mm256_rsqrt_ps 命令 や __m256_mm256_rcp_ps 命令で

The maximum relative error for this approximation is less than 1.5*2^-12.

という記述があり、どちらも Latency 7 , Throughput 1 となっています。

  ちなみに、フル精度の除算を行う _mm256_div_ps は手元の Core i7-4770 (Haswell) でも Latency 17-21 , Throughput 13 となっていますので、Intel CPU であっても、まじめな除算は非常にコストの大きい課題であると言えます。

 

　ということで、まずは 平方の二乗の手法を試してみます。
  減算のみの簡単な実装(Newton法 0回)で実験したところ、 1.0～2.0の範囲の逆数で 0.0664284229 の最大誤差が出てしまいました。
   WikipediaにあるNewton法の項

   y  = y * (threehalfs - (x2 * y * y));

  を追加していくと

   １回：0.0033829212
   2回：0.0000087619
   3回：0.0000002980

  と、順次良くなっていきますが、乗算器使用コストが大きく、特に今回スループット１を目指しているので、この方法は一旦検討を保留します。

   で、次にテーブルの方法を検討して見たのですが

「ディジタル数値演算回路の実用設計」を取り寄せて見たところ、アドレス5bit(32個)のテーブルを用いているようです。
  Sprtan3 世代の想定(DistributedRAM 4bit を2個かな？)のようですが、現在 Zynq などの Xilinx  ７Seriase の場合だと、Distributed-RAM 1個で 6bit(64個)の幅ですのでそれにあわせて手を入れたほうがよさそうです。

  ありそうな設計選択肢として

• DistributedRAMで収まる範囲のテーブルにするか、BlockRAM を使うか
• DSPを1個で済ますか、2次誤差補正も含めて2個使うか

  といったリソースと精度のトレードオフになりそうです。

  2次誤差はまだですが、単純補間を軽くCで実装して実験したところ 1.0～2.0の範囲の逆数で実験できたので結果を張ると、

  5bit単純テーブル補間 ： 最大誤差 0.0002331734
  6bit単純テーブル補間 ： 最大誤差 0.0000596642
  9bit単純テーブル補間 ： 最大誤差 0.0000010133
10bit単純テーブル補間 ： 最大誤差 0.0000002980

  といったところでした。
  例によって実験用の適当な実装なのと、書籍とはテーブルの置き方や下位bitの幅を変えているのでがですが、5bitで書籍の精度と大きくは外していないのでたぶん目安程度の結果にはなっていると思います（あくまで「たぶん」ですが）。

  なので、シンプルに単純補間でいく場合、Distributed-RAMで6bit もしくは Block-RAMで 9bit or 10bit が、リソース的に効率の良い選択肢となりそうです。
  6bit でもそこそこいけそうな気もしますが、ここからさらに伸ばそうとすると、DSPをもう一個使うか、Block-RAMを使うかの選択肢になりそうです。DSPを使う場合、どうしても加算ステージが１段増えるのでレイテンシも増えそうだし、後で対応するとシェーダーの性能にも影響が出そうな気もします。Spartan3時代に比べると、BlockRAMやDSPの貴重さも変わってきているので選択肢としてはいろいろありそうです。

  一方でテーブル量だけなら後で変えてもリソース問題だけで、設計影響が少ないはずなので、なるべくシンプルなところに落としたいなとも悩み中で、6bit ～ 10bit の範囲組み換え可能なところで進めておくのが無難かなと思い始めています。

　さてさて、次の進展はいつになることやら。