#ifdef FLOAT
  #define T float
#else
  #define T int
#endif

#ifdef FORUM
__device__ inline void atomicAdd(float *address, float val){
      int i_val = __float_as_int(val);
      int tmp0 = 0;
      int tmp1;
      while( (tmp1 = atomicCAS((int *)address, tmp0, i_val)) != tmp0)  {
              tmp0 = tmp1;
              i_val = __float_as_int(val + __int_as_float(tmp1));
      }
}
#else
__device__ inline float atomicAdd(float* address, float value){
  float old = value;  
  float ret=atomicExch(address, 0.0f);
  float new_old=ret+old;
  while ((old = atomicExch(address, new_old))!=0.0f){
    new_old = atomicExch(address, 0.0f);
    new_old += old;
  }
  return ret;
};
#endif

__device__ float fatomicMin(float *addr,float value){
          float old = *addr, assumed;
          if(old<=value) return old;
          do {
                     assumed = old;
                     old = atomicCAS((int*)addr, __float_as_int(assumed), __float_as_int(MIN(value, assumed)));
          }while(old!=assumed);
          return old;
};