GPU 上廉价的近似整数除法

Question

所以，我想在 GPU 上除以一些 32 位无符号整数，我不关心得到一个精确的结果。事实上，让我们宽容一点，假设我愿意接受高达 2 的乘法误差因子，即如果 q = x/y 我愿意接受介于 0.5*q 和 2*q 之间的任何值。

我还没有测量任何东西，但在我看来，这样的东西（CUDA 代码）应该很有用：

__device__ unsigned cheap_approximate_division(unsigned dividend, unsigned divisor)
{
    return 1u << (__clz(dividend) - __clz(divisor));
}

它使用 "find first (bit) set" integer intrinsic 作为廉价的以 2 为底的对数函数。

注意：我可以使这个非 32 位特定，但我必须使用模板使代码复杂化，用模板函数包装 __clz() 以使用 __clzl() 和__clzll()等

问题：

• 就时钟周期而言，这种近似除法是否有更好的方法？也许有稍微不同的约束？
• 如果我想要更高的精度，我应该使用整数还是只使用浮点算术？

Answer 1

通过浮点运算可以得到更精确的结果，在大多数架构上指令数略低，并且可能会提高吞吐量：

__device__ unsigned cheap_approximate_division(unsigned dividend, unsigned divisor)
{
   return (unsigned)(__fdividef(dividend, divisor) /*+0.5f*/ );
}

注释中的+0.5f 应表明您还可以将float->int 转换为适当的舍入，除了更高的能耗（它将fmul 转换为fmad 与常数直接来自常量缓存）。不过，四舍五入会使您远离确切的整数结果。