原子函数的合理使用

参考书目： CUDA编程基础与实践

作者: 樊哲勇

章节： 第9章 原子函数的合理使用

原子函数

在数组归约的例子中,我们使用了BLOCK_SIZE长度的d_y的数组存每个block的和,然后再计算d_y的和(此过程为先copy到CPU上,然后串行计算数组和),如果要再GPU上计算这后一步,可以通过两种方法

• 将d_y再次进行归约
• 利用原子函数

本章介绍原子函数

如果我们不想用d_y存每个block的和,直接在计算时进行累加,代码可以初步构想为

if(tid==0)
{
    d_y += s_y[0];
}

但这样写是错的,因为在进行d_y += s_y[0];时需要读s_y[0],d_y的数据,然后相加重新赋予d_y,多个线程可能同时进行读写操作,比如第一个线程读取了d_y的数据还没有对其进行更新时,另一个线程也读取了d_y,得到的结果肯定不对,所以这也就引出了原子函数

原子函数可以保证每个线程中的操作一气呵成,不会被其他线程的原子函数干扰.此时d_y是一个长度为1的数组

if(tid==0)
{
    atomicAdd(&d_y[0], s_y[0]);//原子函数,不会被不同线程的原子操作干扰
}

利用atomicCAS实现atomicAdd的双精度,多线程的CAS锁原理

__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;//old会改变

    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);
}

这样就能看出原子操作只能保证操作完整,但无法保证不同线程的执行次序