 
 template<typename T>
 struct AtomicAddIntegerImpl<T, 1> {
-  __device__ void operator()(T *address, T val) {
+  inline __device__ void operator()(T *address, T val) {
     unsigned int * address_as_ui =
         (unsigned int *) (address - ((size_t)address & 3));
     unsigned int old = *address_as_ui;


#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ < 600
// from CUDA C Programmic Guide
__device__  void 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;
  unsigned long long int assumed;


template<typename T>
struct AtomicAddIntegerImpl<T, 4> {
  inline __device__ void operator()(T *address, T val) {
    unsigned int * address_as_ui = (unsigned int *) (address);
    unsigned int old = *address_as_ui;
    unsigned int newval;
