Merge pull request #1 from PabloPalaciosAlonso/v2.x

Define explicitly atomic add of float2 numbers and double2 numbers
RaulPPelaez · Dec 4, 2023 · b2b96da · b2b96da
2 parents 64e82c5 + 3b838c2
commit b2b96da
Showing 1 changed file with 8 additions and 5 deletions.
diff --git a/src/utils/atomics.cuh b/src/utils/atomics.cuh
@@ -7,8 +7,6 @@ namespace uammd{
   template<class T>
   inline __device__ T atomicAdd(T* address, T val){ return ::atomicAdd(address, val);}
 
-
-#ifndef SINGLE_PRECISION
 #if !defined(__CUDA_ARCH__) || __CUDA_ARCH__ < 600
   inline __device__ double atomicAdd(double* address, double val){
     unsigned long long int* address_as_ull =
@@ -22,7 +20,6 @@ namespace uammd{
     } while (assumed != old);
     return __longlong_as_double(old);
   }
-#endif
 #endif
 
   inline __device__ real4 atomicAdd(real4* address, real4 val){
@@ -42,13 +39,19 @@ namespace uammd{
     return newval;
   }
 
-  inline __device__ real2 atomicAdd(real2* address, real2 val){
-    real2 newval;
+  inline __device__ float2 atomicAdd(float2* address, float2 val){
+    float2 newval;
     if(val.x) newval.x = atomicAdd(&(*address).x, val.x);
     if(val.y) newval.y = atomicAdd(&(*address).y, val.y);
     return newval;
   }
 
+  inline __device__ double2 atomicAdd(double2* address, double2 val){
+    double2 newval;
+    if(val.x) newval.x = atomicAdd(&(*address).x, val.x);
+    if(val.y) newval.y = atomicAdd(&(*address).y, val.y);
+    return newval;
+  }
 
   template<class T, class T2>
   inline __device__ T2 atomicAdd(T &ref, T2 val){