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，常数直接来自常量缓存）。尽管四舍五入会让你远离精确的整数结果。