GPU 编程 CUDA C++ 使用 atomicCAS( ) 函数实现双精度浮点数的 atomicAdd( ) 函数

温柔的行子

已于 2023-04-01 09:59:21 修改

阅读量690

点赞数 1

文章标签： c++ 开发语言人工智能数据挖掘

于 2023-03-31 19:55:53 首次发布

本文链接：https://blog.csdn.net/qq_42536162/article/details/129887461

版权

该文章介绍了CUDA中一个特殊的功能函数atomicAdd_double，它能进行双精度浮点数的原子操作。通过将double类型转换为unsignedlonglong，利用atomicCAS进行比较和交换，确保‘读-改-写’操作的原子性，防止并发冲突。此函数在并行计算中对于保证数据一致性至关重要。

摘要由CSDN通过智能技术生成

本问题的关键是新创建的atomicAdd_double()函数不但能进行双精度浮点运算，还同时能继承原子函数的原子特性：“读-改-写”操作为最小操作单元，不可分割！

__device__ double atomicAdd_double(double* address, double val)
{

    unsigned long long* address_as_ull = (unsigned long long*)address;  // 将地址强制转换为unsigned long long类型的指针
    unsigned long long old = *address_as_ull, assumed;  // old保存原始值，assumed用于CAS操作
    
    do
    {
        
        assumed = old;  // 在每次循环开始时，将assumed设置为old的值
        old = atomicCAS
        (
            address_as_ull, assumed,
            __double_as_longlong(val + __longlong_as_double(assumed))
        );  // 使用atomicCAS()函数进行比较和交换操作
    } while (assumed != old);  // 如果assumed和old不相等，说明操作失败，需要继续循环            
    
    return __longlong_as_double(old);  // 将old转换为double类型并返回
}

其中，__double_as_longlong()和__longlong_as_double()是CUDA内置函数，用于将双精度浮点数和unsigned long long类型相互转换。需要注意的是，由于atomicCAS()函数只能操作32位或64位的数据，因此需要将双精度浮点数强制转换为unsigned long long类型。