CUDA Programming atomic operation atomicadd reports error err:msb3721, return code 1

problem ： Atomic manipulation atomicAdd() Report errors err:MSB3721, Return code 1.

Problem description ： Today, I am writing and using cuda Accelerate the calculation of normal direction related code after 3D point cloud reconstruction , One step is to calculate the average point distance after counting the point cloud distance . This requires multiple threads to add the distance between different points and adjacent points , Then divide by the number of points .
   But without atomic operation atomicAdd() Words , The addition result will inevitably make mistakes . Be similar to OpenMP In parallel for The principle of loop locking the addition variable to ensure the correct addition of variables .
   But because of Dot distance is double type , Ignorant I use it directly atomicAdd() Adding dots causes compilation errors err:MSB3721, Return code 1, After searching the official documents, I learned , For double precision double Type of atomic operation atomicAdd(), Only in The computing power is greater than 6.0 On the machine that supports , For my computing power, only 3.5 Old age machine GT720, Unable to compile successfully . So far, the final reason has been found . Look at the Yellow characters .
Direct cause ： Double precision double The atomic operation of is done only when the computing power is 6.0 The above devices support , But my equipment is too LOW.
Solution ： But the government also gave us this kind of old GPU The way to live , Using the code in the figure, the computing power can be less than 6.0 The device performs double precision double Atomic manipulation .
【 Be careful ： Copy the following code , Don't use the official website code directly , Otherwise, the redefinition operator will report an error 】

#if define (__CUDA_ARCH__)||__CUDA_ARCH__ < 600
__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;

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

If this method can be used , Please pay attention to a favorite collection ！