fixed build issue of double definition of atomicAdd on modern GPUs

paddle/cuda/include/hl_device_functions.cuh

@@ -16,6 +16,8 @@ limitations under the License. */
 #ifndef HL_DEVICE_FUNCTIONS_CUH_
 #define HL_DEVICE_FUNCTIONS_CUH_
 
+#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ < 600
+
 namespace hppl {
 
 static __inline__ __device__ double atomicAdd(double* address, double val) {
@@ -40,4 +42,6 @@ static __inline__ __device__ double atomicAdd(double* address, double val) {
 using hppl::atomicAdd;
 #endif
 
+#endif
+
 #endif /* HL_DEVICE_FUNCTIONS_CUH_ */