__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);
}
Example #3
0
__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
}