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));
}

它使用固有的“查找第一个（位）集”整数作为便宜的基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注释中的in表示您还可以将float-> int转换为除基本能耗外几乎没有任何成本的适当舍入（将a fmul转换为an fmad，而常量直接来自常量缓存）。四舍五入会使您远离精确的整数结果。