__device__ static double atomicMul(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))); } while (assumed != old); return __longlong_as_double(old); }
__device__ static double atomicMin(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; if(__longlong_as_double(assumed) <= val) break; old = atomicCAS(address_as_ull, assumed, val); } while (assumed != old); return __longlong_as_double(old); }
__device__ static double atomicAdd(double* address, double val) { #if CV_CUDEV_ARCH >= 130 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))); } while (assumed != old); return __longlong_as_double(old); #else (void) address; (void) val; return 0.0; #endif }