From 3b838c2f030348f874198a571b80907cd0495048 Mon Sep 17 00:00:00 2001 From: PabloPalaciosAlonso Date: Fri, 1 Dec 2023 09:32:05 +0100 Subject: [PATCH] Define explicitly atomic add of float2 numbers and double2 numbers --- src/utils/atomics.cuh | 13 ++++++++----- 1 file changed, 8 insertions(+), 5 deletions(-) diff --git a/src/utils/atomics.cuh b/src/utils/atomics.cuh index fc592524..dba7e4bd 100644 --- a/src/utils/atomics.cuh +++ b/src/utils/atomics.cuh @@ -7,8 +7,6 @@ namespace uammd{ template 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 inline __device__ T2 atomicAdd(T &ref, T2 val){