typos

torch_scatter/kernel/THCAtomics.cuh

@@ -58,16 +58,17 @@ struct AtomicMaxIntegerImpl<T, 4> {
 template<typename T>
 struct AtomicMaxIntegerImpl<T, 8> {
   inline __device__ void operator()(T *address, T val) {
-    unsigned long long * address_as_ui = (unsigned long long *) (address);
-    unsigned long long old = *address_as_ui;
-    unsigned long long newval;
-    unsigned long long assumed;
-
-    do {
-      assumed = old;
-      newval = val +  (T)old;
-      old = atomicCAS(address_as_ui, assumed, newval);
-    } while (assumed != old);
+    int *address_as_ull = (int*) (address);
+    int newval = *address_as_ull;
+    atomicMax(address_as_ull, newval);
+    /* unsigned long long newval; */
+    /* unsigned long long assumed; */
+
+    /* do { */
+    /*   assumed = old; */
+    /*   newval = val +  (T)old; */
+    /*   old = atomicCAS(address_as_ui, assumed, newval); */
+    /* } while (assumed != old); */
   }
 };
 
@@ -77,7 +78,9 @@ static inline  __device__ void atomicMax(int8_t *address, int8_t val) {}
 
 static inline __device__ void atomicMax(int16_t *address, int16_t val) {}
 
-static inline __device__ void atomicMax(int64_t *address, int64_t val) {}
+static inline __device__ void atomicMax(int64_t *address, int64_t val) {
+  AtomicMaxIntegerImpl<int64_t, sizeof(int64_t)>()(address, val);
+}
 
 #ifdef CUDA_HALF_TENSOR
 static inline __device__ void atomicMax(half *address, half val) {}