私は CUDA 11.2 を使用しており、この__half型を使用して 16 ビット浮動小数点値を操作しています。
nvcc コンパイラーが融合乗算加算命令を適切に呼び出さないことに驚いています。
__half a,b,c;
...
__half x = a * b + c;
融合乗算加算を発行する代わりに、別個の mul 命令と加算命令を発行します。
mul.f16 %rs164,%rs1,%rs306;
add.f16 %rs167,%rs164,%rs65;
これは、--fmad=trueコンパイラ オプションを使用しているにもかかわらずです。
明示的なもの__hfma( a,b,c )は次のように出力されます。
fma.rn.f16 %rs164,%rs1,%rs300,%rs65;
明示的な組み込み関数を使用するために 16 ビット浮動小数点の乗加算を利用する唯一の方法はありますか?