message: adjustments to atomic - adds flexibility
diff: @@ -37,11 +37,6 @@ void cuWrap(const int N, DBODY &&d_body)
  const cudaError_t last = cudaGetLastError();
  MFEM_ASSERT(last == cudaSuccess, cudaGetErrorString(last));
  }
  -template<typename T>
  -__host__ __device__ inline T AtomicAdd(T* address, T val)
  -{
  - return atomicAdd(address, val);
  -}
  #else // ***********************************************************************
  #define __host__
  #define __device__
  @@ -51,12 +46,26 @@ typedef int CUcontext;
  typedef void* CUstream;
  template <int BLOCKS, typename DBODY>
  void cuWrap(const int N, DBODY &&d_body) {}
  +#endif // __NVCC__
  +
  +#if defined(__CUDA_ARCH__)
  +template<typename T>
  +__device__ inline T AtomicAdd(T* address, T val)
  +{
  + return atomicAdd(address, val);
  +}
  +#else
  template<typename T> inline T AtomicAdd(T* address, T val)
  {
  - return *address += val;
  +#if defined(_OPENMP)
  +#pragma omp atomic
  + *address += val;
  +#else
  + *address += val;
  +#endif
  + return *address;
  }
  -#endif // __NVCC__
  -
  +#endif //__CUDA_ARCH__
  // *****************************************************************************
  namespace mfem
  {
  