Good day,
I am using the cuda function for implementing atomicAdd for double precision which, if I remember correctly, I got from the Nvidia site:
__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; //THIS IS LINE 16
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);
}
However when using cuda-memcheck I get a message saying there in an out of bounds error at line 16 (4 in this case) of the atomicAdd code:
========= CUDA-MEMCHECK
========= Invalid __global__ read of size 8
========= at 0x000000d8 in /home/carlos/CPP/mdChaSphere/cuda_ready/cuda/long_range/mean_f/v12/dev_functions.cu:16:atomicAdd(double*, double)
========= by thread (31,0,0) in block (15,0,0)
========= Address 0xb047d9ee8 is out of bounds
========= Device Frame:/home/carlos/CPP/mdChaSphere/cuda_ready/cuda/long_range/mean_f/v12/dev_functions.cu:947:correl_kernel(double const *, double*, double*, int const *, \
int, double const *, int, int, double) (correl_kernel(double const *, double*, double*, int const *, int, double const *, int, int, double) : 0x1660)
========= Saved host backtrace up to driver entry point at kernel launch time
========= Host Frame:/usr/lib/x86_64-linux-gnu/libcuda.so.1 (cuLaunchKernel_ptsz + 0x2c5) [0x1472e5]
========= Host Frame:/usr/local/cuda-7.5/lib64/libcudart.so.7.5 [0x14623]
========= Host Frame:/usr/local/cuda-7.5/lib64/libcudart.so.7.5 (cudaLaunch_ptsz + 0x154) [0x3d134]
========= Host Frame:./md_ch_sphere [0x52ba]
========= Host Frame:./md_ch_sphere [0x8b87]
========= Host Frame:./md_ch_sphere [0x8bec]
========= Host Frame:./md_ch_sphere [0x7e02]
========= Host Frame:./md_ch_sphere [0x3e7f]
========= Host Frame:/lib/x86_64-linux-gnu/libc.so.6 (__libc_start_main + 0xf5) [0x21ec5]
========= Host Frame:./md_ch_sphere [0x51cf]
I would appreciate if someone can help me understand why this could happen. Thank you!