CUDA编程原子操作atomicAdd报错err:MSB3721,返回代码1_使用double类型atomicadd出错-CSDN博客

本文链接：https://blog.csdn.net/weixin_41966507/article/details/122899808

问题：原子操作atomicAdd()报错err:MSB3721,返回代码1。

直接原因： 双精度double的原子操作尽在仅在计算能力为6.0以上的设备支持，但我的设备太LOW。

问题描述： 今天在编写利用cuda加速三维点云重建后的计算法向相关代码，有一步骤需要对点云间距统计后求得平均点距。这就要求多个线程对不同点与邻点的距离进行相加，之后再除以点的个数。
但是如果没有原子操作atomicAdd()的话，相加结果会无法避免的出现错误。类似于OpenMP进行并行for循环锁住加和变量保证变量正确加和的原理。
但是由于点距是double类型，无知的我直接使用atomicAdd()进行点距相加导致编译出错err:MSB3721,返回代码1，经过搜索官方文档后得知，对于双精度double类型的原子操作atomicAdd()，仅在计算能力大于6.0的机器上支持，对于我这种计算能力只有3.5的老年机器GT720，无法成功编译。至此找到了最终原因。详看黄色字符。
在这里插入图片描述
解决方案： 但是官方也给了我们这种老GPU活路，使用图中的代码可以实现计算能力小于6.0的设备进行双精度double原子操作。
【注意：复制以下代码，不要直接使用官网代码，否则出现重定义操作符的报错】

#if define (__CUDA_ARCH__)||__CUDA_ARCH__ < 600
__device__ double atomicAdd(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(val +
                               __longlong_as_double(assumed)));

    // Note: uses integer comparison to avoid hang in case of NaN (since NaN != NaN)
    } while (assumed != old);

    return __longlong_as_double(old);
}
#endif