template <typename T>
__device__ void atomic_max(T* addr, T val){
  atomicMax(addr, val);
}

template <typename T>
__device__ void atomic_add(T *addr, T val)
{
    atomicAdd(addr, val);
}
template <typename T>
__device__ void atomic_min(T *addr, T val)
{
    atomicMin(addr, val);
}

__device__ __inline__ void atomic_max(int8_t* addr, int8_t val){
  if(*addr >= val)
    return;

  unsigned int* const addr_as_ull = (unsigned int*) addr;
  unsigned int old = *addr_as_ull;
  unsigned int assumed;
  do{
      assumed = old;
      if(reinterpret_cast<int8_t&>(assumed) >= val)
        break;
      old = atomicCAS(addr_as_ull, assumed, val);
  } while (assumed != old);
}

__device__ __inline__ void atomic_add(int8_t* addr, int8_t val){
  if(*addr >=val){
    return;
  }
  unsigned int* const addr_as_ull = (unsigned int*) addr;
  unsigned int old = *addr_as_ull;
  unsigned int assumed;
  do{
    assumed = old;
    old = atomicCAS(addr_as_ull, assumed, reinterpret_cast<int8_t&> (old) + val);
  }while(asuumed != old)
}