7 Replies Latest reply on Oct 12, 2011 3:20 PM by MicahVillmow

    can I do atomic add on float using compare and swap

    shingoxlf

      In the cuda manual, attomic add can be implmented using compare and swap like this:

       

       

      __device__ double atomicAdd(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);
      }
      So I tried to implement atomic add in Opencl like this:
      double atomicDoubleAdd(__global double* address, double val)
      {
      ulong * address_as_ull =(ulong *)address;
      ulong old = *address_as_ull, assumed;
      do {
      assumed = old;
      old = atom_cmpxchg(address_as_ull, assumed,as_ulong(val +as_double(assumed)));
      } while (assumed != old);
      return as_double(old);
      }
      however, it fails when I tried to build the program, so if anybody knows how to do this?