消えたCUDA関連の旧ブログ記事を復元するひとり Advent Calendar 2024の記事です。

なんの話か

CUDAの単精度浮動小数点数（binary32）用の除算命令には近似計算命令があります。
この近似除算命令の精度をちょっと見てみました、という話です。

除算近似命令

近似命令というのはdiv.approxです。
CUDA Toolkit Documentation 9.7.3.8. Floating Point Instructions: divによると、a / bをa * (1 / b)と計算するとのことです。
で、デフォルトの除算命令より速いとのことです。
逆数計算は近似逆数平方根計算などで行うのでしょうか？
ただしこの命令はbが[2^{-126}, 2^{126}]という制約があります。
binary32の指数部の最大値は127なので、指数部が126で仮数部が非零の場合と指数部が127の場合は計算できないということとなります。
で、この制約をスケーリングにより外したものがdiv.fullです（full-rangeのfullです）。
また、これらの近似除算命令では丸めの指定はできません。

精度評価

div.approx、div.full、div.rn（IEEE準拠で丸めにRNを用いる除算命令）の除算計算精度を倍精度で計算した場合と比較したのが下の図です。
1/3をb_FP32（横軸）で割った精度を倍精度計算との相対誤差（縦軸）で表示しています。

やはり近似計算の方が若干精度が悪いことがわかります。
また、div.approxでb_FP32が大きいときにerrorが飛び出ているのはdiv.approxのbの対応範囲外のためです。
近似計算か否かに関わらず、b_FP32が大きいときは精度が少し悪くなることも確認できます。

評価コード

template <>
__device__ float div<approx_div>(const float a, const float b) {
	float r;
	asm(
			R"(
{
div.approx.f32 %0, %1, %2;
}
)": "=f"(r) : "f"(a), "f"(b)
			);
	return r;
}

除算部分は以下のようにinline asmを使って記述しました。
Fullなコードはこちら。
enp1s0/cuda-fiv-approx - GitHub