/*-------------------------------------------------------------------
 * Function:    atomicAddf  (device code)
 * Purpose:     Atomically add the second argument to the first
 * In arg:      val
 * In/out arg:  address
 *
 * Note:        This is a slightly modified version of code
 *              found on p. 116 of the CUDA Programming Guide,
 *              Version 3.2.
 */
__device__  float atomicAddf(float* address, float val) {
   float old = *address, assumed;

   do {
      assumed = old;
      old = __int_as_float(
            atomicCAS((unsigned int*)address,
      __float_as_int(assumed),
      __float_as_int(val + assumed)));
   } while (assumed != old);

   return old;
} /* atomicAddf */