CUDA 定义的atomicMin
函数不支持使用浮点数。参考文档,我们看到唯一可用的原型是、和(最后一个需要在计算能力为 3.5 或更高的 GPU 上编译和运行)。int
unsigned int
unsigned long long int
至少有2个选项。
您可以重构代码以用经典的并行减少替换原子。
如编程指南中所述,可以使用atomicCAS
(比较和交换)加上某种循环来创建“任意”原子。
这是一种可能的实现,对于double
:
__device__ double atomicMin_double(double* address, double val)
{
unsigned long long int* address_as_ull = (unsigned long long int*) address;
unsigned long long int old = *address_as_ull, assumed;
do {
assumed = old;
old = atomicCAS(address_as_ull, assumed,
__double_as_longlong(fmin(val, __longlong_as_double(assumed))));
} while (assumed != old);
return __longlong_as_double(old);
}
这个相关的问题和答案也可能很有趣,尽管它主要是float
考虑到的。
其他一些评论:
通过切换到float
而不是double
我相信可以简化atomicMin
(或atomicMax
)操作,如我在上面链接的答案中所示,可能有一些警告(例如,没有 NaN、INF 数据)。我相信 iee754float
遵循两个数量的排序规则A
,B
如果A > B
, then *reinterpret_cast<int*>(&A) > *reinterpret_cast<int*>(&B)
。我不确定是否double
遵循类似的规则long long
(可能其他人会说)。
在您的代码中,此循环可以首先对本地数量进行操作,然后在最后执行单个原子操作,如下所示:
double v = *g;
for( int u = 0; u<16; u++ )
{
v = min(v,d);
}
atomicMin_double(g, v);
我认为应该更快