問題タブ [fma]

計算シェーダーが SPIR-V (特に Vulkan) を介して変換される API を使用して GPU の計算コードを作成する場合、浮動小数点除算の ULP エラーが最大 3 になることが保証されます。その他の基本的な算術演算 (加算、乗算) は、正しく丸めます。

これらの状況で正しく丸められた除算を効率的に実装するにはどうすればよいですか? FMA が利用可能で、正しく丸められていると仮定しましょう。

デノーマルで何が起こるかは、基盤となるハードウェアによって異なります。Vulkan API を使用すると、デバイスがデノーマルを保持できるかどうか、およびデノーマルをゼロにフラッシュできるかどうかを照会できます (したがって、デノーマルを完全にサポートしていない GPU は「canPreserve: false, canFlush: true」になります)。さらに、GPU がデノーマルをゼロにフラッシュせずに生成および処理できると仮定しましょう (そうでなければ、サブノーマルである正しく丸められた結果を生成しようとしても無駄に思えます)。

私は CUDA 11.2 を使用しており、この__half型を使用して 16 ビット浮動小数点値を操作しています。

nvcc コンパイラーが融合乗算加算命令を適切に呼び出さないことに驚いています。

融合乗算加算を発行する代わりに、別個の mul 命令と加算命令を発行します。

これは、--fmad=trueコンパイラ オプションを使用しているにもかかわらずです。

明示的なもの__hfma( a,b,c )は次のように出力されます。

明示的な組み込み関数を使用するために 16 ビット浮動小数点の乗加算を利用する唯一の方法はありますか?

次の計算の実装を検討していました。ここで、divisorはゼロではなく、2 のべき乗ではありません

これは、64 ビットの整数および浮動小数点命令を欠いているが、32 ビットの融合乗加算を備えている可能性があるプロセッサ (同様に除算がない GPU など) にとって効率的な方法です。

この計算は、除数の最適化に関係する「魔法の乗数」を見つけるのに役立ちます。これは、除数が事前にわかっている場合に、上位乗算命令とそれに続くビットごとのシフトです。コンパイラで使用されるコードやlibdivideの参照コードとは異なり、そのような乗数の最大値を検出します。

もう 1 つのひねりは、私が見ていたアプリケーションでは、divisorほとんどの場合、float型で表現できると予想していたことです。したがって、これらの除数を処理する効率的な「高速パス」と、残りを処理するサイズが最適化された「低速パス」を用意することは理にかなっています。