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 常量直接来自常量缓存）。不过，四舍五入会使您离精确的整数结果更远。

GPU 上的廉价近似整数除法

Cheap approximate integer division on a GPU

cuda

integer-division

approximate