From 52df82917adc42ef967495076e349bf28a251235 Mon Sep 17 00:00:00 2001
From: Tanguy Pruvot <tanguy.pruvot@gmail.com>
Date: Fri, 29 May 2015 14:25:44 +0200
Subject: [PATCH] cuda: fix uint2 substract operator

---
 cuda_helper.h | 43 ++++++++++++++++++++++++++++---------------
 1 file changed, 28 insertions(+), 15 deletions(-)

diff --git a/cuda_helper.h b/cuda_helper.h
index f6ff545..b46ef97 100644
--- a/cuda_helper.h
+++ b/cuda_helper.h
@@ -480,7 +480,7 @@ uint64_t SWAPDWORDS(uint64_t value)
 #endif
 }
 
-/* lyra2 - int2 operators */
+/* lyra2/bmw - uint2 vector's operators */
 
 __device__ __forceinline__
 void LOHI(uint32_t &lo, uint32_t &hi, uint64_t x) {
@@ -488,27 +488,41 @@ void LOHI(uint32_t &lo, uint32_t &hi, uint64_t x) {
 	asm("mov.b64 {%0,%1},%2; \n\t"
 		: "=r"(lo), "=r"(hi) : "l"(x));
 #else
-	lo = _LODWORD(x);
-	hi = _HIDWORD(x);
+	lo = (uint32_t)(x);
+	hi = (uint32_t)(x >> 32);
 #endif
 }
 
-static __device__ __forceinline__ uint2 vectorize(uint64_t v) {
+static __host__ __device__ __forceinline__ uint2 vectorize(uint64_t v) {
 	uint2 result;
-	LOHI(result.x, result.y, v);
+#ifdef __CUDA_ARCH__
+	asm("mov.b64 {%0,%1},%2; \n\t"
+		: "=r"(result.x), "=r"(result.y) : "l"(v));
+#else
+	result.x = (uint32_t)(v);
+	result.y = (uint32_t)(v >> 32);
+#endif
 	return result;
 }
 
-static __device__ __forceinline__ uint64_t devectorize(uint2 v) { return MAKE_ULONGLONG(v.x, v.y); }
+static __host__ __device__ __forceinline__ uint64_t devectorize(uint2 v) {
+#ifdef __CUDA_ARCH__
+	return MAKE_ULONGLONG(v.x, v.y);
+#else
+	return (((uint64_t)v.y) << 32) + v.x;
+#endif
+}
 
+/**
+ * uint2 direct ops by c++ operator definitions
+ */
 static __device__ __forceinline__ uint2 operator^ (uint2 a, uint2 b) { return make_uint2(a.x ^ b.x, a.y ^ b.y); }
 static __device__ __forceinline__ uint2 operator& (uint2 a, uint2 b) { return make_uint2(a.x & b.x, a.y & b.y); }
 static __device__ __forceinline__ uint2 operator| (uint2 a, uint2 b) { return make_uint2(a.x | b.x, a.y | b.y); }
 static __device__ __forceinline__ uint2 operator~ (uint2 a) { return make_uint2(~a.x, ~a.y); }
 static __device__ __forceinline__ void operator^= (uint2 &a, uint2 b) { a = a ^ b; }
 
-static __device__ __forceinline__ uint2 operator+ (uint2 a, uint2 b)
-{
+static __device__ __forceinline__ uint2 operator+ (uint2 a, uint2 b) {
 #ifdef __CUDA_ARCH__
 	uint2 result;
 	asm("{\n\t"
@@ -518,15 +532,14 @@ static __device__ __forceinline__ uint2 operator+ (uint2 a, uint2 b)
 		: "=r"(result.x), "=r"(result.y) : "r"(a.x), "r"(a.y), "r"(b.x), "r"(b.y));
 	return result;
 #else
-	// incorrect but unused host equiv
-	return make_uint2(a.x + b.x, a.y + b.y);
+	return vectorize(devectorize(a) + devectorize(b));
 #endif
 }
 static __device__ __forceinline__ void operator+= (uint2 &a, uint2 b) { a = a + b; }
 
-static __device__ __forceinline__ uint2 operator- (uint2 a, uint2 b)
-{
-#ifdef __CUDA_ARCH__
+
+static __device__ __forceinline__ uint2 operator- (uint2 a, uint2 b) {
+#if defined(__CUDA_ARCH__) && CUDA_VERSION < 7000
 	uint2 result;
 	asm("{\n\t"
 		"sub.cc.u32 %0,%2,%4; \n\t"
@@ -535,10 +548,10 @@ static __device__ __forceinline__ uint2 operator- (uint2 a, uint2 b)
 		: "=r"(result.x), "=r"(result.y) : "r"(a.x), "r"(a.y), "r"(b.x), "r"(b.y));
 	return result;
 #else
-	// incorrect but unused host equiv
-	return make_uint2(a.x - b.x, a.y - b.y);
+	return vectorize(devectorize(a) - devectorize(b));
 #endif
 }
+static __device__ __forceinline__ void operator-= (uint2 &a, uint2 b) { a = a - b; }
 
 /**
  * basic multiplication between 64bit no carry outside that range (ie mul.lo.b64(a*b))