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


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


template<typename T>
struct AtomicAddIntegerImpl<T, 1> {
  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;
