Точечный продукт CUDA

Я пытаюсь реализовать классическое ядро ​​скалярного произведения для массивов двойной точности с атомарным вычислением окончательной суммы по различным блокам. Я использовал atomicAdd для двойной точности, как указано на странице 116 руководства по программированию. Вероятно, я делаю что-то не так. Частичные суммы по потокам в каждом блоке вычисляются правильно, но после этого атомарная операция, похоже, не работает должным образом поскольку каждый раз, когда я запускаю ядро ​​с одними и теми же данными, я получаю разные результаты. Буду признателен, если кто-нибудь обнаружит ошибку или предложит альтернативное решение! Вот мое ядро:

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