cuda - CUDA点积

我正在尝试对双精度数组实现经典的点积内核，并通过原子计算跨各个块的最终和。如编程指南第116页所述，我使用了atomicAdd来实现双精度。可能是我做错了。每个块中线程的部分和被正确计算了，但后言原子操作似乎无法正常工作因为每次我用相同的数据运行内核时，我都会收到不同的结果。如果有人能发现错误或提供替代解决方案，我将不胜感激！
这是我的内核：

__global__ void cuda_dot_kernel(int *n,double *a, double *b, double *dot_res)
{
    __shared__ double cache[threadsPerBlock]; //thread shared memory
    int global_tid=threadIdx.x + blockIdx.x * blockDim.x;
    int i=0,cacheIndex=0;
    double temp = 0;
    cacheIndex = threadIdx.x;
    while (global_tid < (*n)) {
        temp += a[global_tid] * b[global_tid];
        global_tid += blockDim.x * gridDim.x;
    }
    cache[cacheIndex] = temp;
    __syncthreads();
    for (i=blockDim.x/2; i>0; i>>=1) {
        if (threadIdx.x < i) {
            cache[threadIdx.x] += cache[threadIdx.x + i];
        }
        __syncthreads();
    }
    __syncthreads();
    if (cacheIndex==0) {
        *dot_res=cuda_atomicAdd(dot_res,cache[0]);
    }
}

这是我的设备函数atomicAdd：

__device__ double cuda_atomicAdd(double *address, double val)
{
    double assumed,old=*address;
    do {
        assumed=old;
        old= __longlong_as_double(atomicCAS((unsigned long long int*)address,
                    __double_as_longlong(assumed),
                    __double_as_longlong(val+assumed)));
    }while (assumed!=old);

    return old;
}

最佳答案

您正在错误地使用cuda_atomicAdd功能。内核的这一部分：

if (cacheIndex==0) {
    *dot_res=cuda_atomicAdd(dot_res,cache[0]);
}

是罪魁祸首。在这里，您原子地添加到dot_res。然后非原子设置dot_res并返回结果。此函数的返回结果是原子更新的位置的先前值，并且仅提供给“信息”或调用方本地使用。您不将其分配给原子更新的对象，这完全违背了首先使用原子内存访问的目的。做这样的事情：

if (cacheIndex==0) {
    double result=cuda_atomicAdd(dot_res,cache[0]);
}