Sort declarations when generating Python bindings (#6701)

* Sort declarations when generating Python bindings

This helps resolve ambiguities in argument parsing according to
any rules we will need.

For now, this allows us to make scalar operations more conservarive
wrt. argument types, but makes them commutative again.

* Fix inconsistencies between mod with tensor and scalar

* Fix a stupid mistake

4 files changed, 45 insertions, 21 deletions

diff --git a/aten/src/TH/generic/THTensorMath.c b/aten/src/TH/generic/THTensorMath.c
index 692742e882..28a3bbb8cd 100644
--- a/aten/src/TH/generic/THTensorMath.c
+++ b/aten/src/TH/generic/THTensorMath.c
@@ -1051,8 +1051,9 @@ void THTensor_(fmod)(THTensor *r_, THTensor *t, real value)
   }
 }
 
-static inline bool has_different_sign(real a, real b) {
-  return (a < 0) != (b < 0);
+// Should wrap if the value (a) has a different sign than the divisor (b), but is not 0.
+static inline bool modulo_wrap(real a, real b) {
+  return (a != 0) && (a < 0) != (b < 0);
 }
 
 void THTensor_(remainder)(THTensor *r_, THTensor *t, real value)
@@ -1073,7 +1074,7 @@ void THTensor_(remainder)(THTensor *r_, THTensor *t, real value)
 #else
       // There is no NAN for integers
       rp[i] = tp[i] % value;
-      if (has_different_sign(rp[i], value))
+      if (modulo_wrap(rp[i], value))
         rp[i] += value;
 #endif
     }
@@ -1088,7 +1089,7 @@ void THTensor_(remainder)(THTensor *r_, THTensor *t, real value)
 #else
       // There is no NAN for integers
       TH_TENSOR_APPLY2_OMP(r_Size, r_Contig, tContig, real, r_, real, t, *r__data = *t_data % value;
-                                        if (has_different_sign(*r__data, value)) *r__data += value;);
+                                        if (modulo_wrap(*r__data, value)) *r__data += value;);
 #endif
     }
 #else
@@ -1101,7 +1102,7 @@ void THTensor_(remainder)(THTensor *r_, THTensor *t, real value)
 #else
     // There is no NAN for integers
     TH_TENSOR_APPLY2(real, r_, real, t, *r__data = *t_data % value;
-                                          if (has_different_sign(*r__data, value)) *r__data += value;);
+                                          if (modulo_wrap(*r__data, value)) *r__data += value;);
 #endif
   }
 }
@@ -1644,7 +1645,7 @@ void THTensor_(cremainder)(THTensor *r_, THTensor *t, THTensor *src)
 #else
         // There is no NAN for integers
         rp[i] = tp[i] % sp[i];
-        if (rp[i] * sp[i] < 0)
+        if (modulo_wrap(rp[i], sp[i]))
           rp[i] += sp[i];
 #endif
       }
@@ -1658,7 +1659,7 @@ void THTensor_(cremainder)(THTensor *r_, THTensor *t, THTensor *src)
         TH_TENSOR_APPLY3_OMP(r_Size, r_Contig, tContig, srcContig, real, r_, real, t, real, src, *r__data = (*src_data == 0)? NAN : *t_data - *src_data * floor(*t_data / *src_data););
 #else
         TH_TENSOR_APPLY3_OMP(r_Size, r_Contig, tContig, srcContig, real, r_, real, t, real, src, *r__data = *t_data % *src_data;
-                                                     if (*r__data * *src_data < 0) *r__data += *src_data;);
+                                                     if (modulo_wrap(*r__data, *src_data)) *r__data += *src_data;);
 #endif
       }
 #else
@@ -1674,7 +1675,7 @@ void THTensor_(cremainder)(THTensor *r_, THTensor *t, THTensor *src)
 #else
     // There is no NAN for integers
     TH_TENSOR_APPLY3(real, r_, real, t, real, src, *r__data = *t_data % *src_data;
-                                                     if (*r__data * *src_data < 0) *r__data += *src_data;);
+                                                     if (modulo_wrap(*r__data, *src_data)) *r__data += *src_data;);
 #endif
 
   }
diff --git a/aten/src/THC/THCNumerics.cuh b/aten/src/THC/THCNumerics.cuh
index cbc2743d9f..9be6d9e129 100644
--- a/aten/src/THC/THCNumerics.cuh
+++ b/aten/src/THC/THCNumerics.cuh
@@ -28,11 +28,6 @@ static inline __host__ __device__ scalar_t powi(scalar_t a, scalar_t b) {
   return result;
 }
 
-template <typename scalar_t>
-static inline __host__ __device__ bool has_different_sign(scalar_t a, scalar_t b) {
-  return (a < 0) != (b < 0);
-}
-
 template <>
 struct THCNumerics<uint8_t> {
   static inline __host__ __device__ uint8_t min() { return 0; }
diff --git a/aten/src/THC/THCTensorMathPairwise.cu b/aten/src/THC/THCTensorMathPairwise.cu
index f530d814d8..a208f29c2d 100644
--- a/aten/src/THC/THCTensorMathPairwise.cu
+++ b/aten/src/THC/THCTensorMathPairwise.cu
@@ -240,19 +240,33 @@ struct TensorDivConstantOp<half> {
 };
 #endif // CUDA_HALF_TENSOR
 
+template<typename T>
+static __device__ __forceinline__
+typename std::enable_if<std::is_signed<T>::value, bool>::type
+modulo_wrap(T a, T b) {
+  return (a != 0) && (a < 0) != (b < 0);
+}
+
+template<typename T>
+static __device__ __forceinline__
+typename std::enable_if<std::is_unsigned<T>::value, bool>::type
+modulo_wrap(T a, T b) {
+  return false;
+}
+
 template <typename T>
 struct TensorRemainderOp {
   TensorRemainderOp(T v) : val(v) {}
   __device__ __forceinline__ void operator()(T* out, T* in) {
     *out = *in % val;
-    if (has_different_sign<T>(*out, val)){
+    if (modulo_wrap<T>(*out, val)) {
       *out += val;
     }
   }
 
   __device__ __forceinline__ void operator()(T* v) {
     *v = *v % val;
-    if (has_different_sign<T>(*v, val)){
+    if (modulo_wrap<T>(*v, val)) {
       *v += val;
     }
   }
diff --git a/aten/src/THC/THCTensorMathPointwise.cuh b/aten/src/THC/THCTensorMathPointwise.cuh
index 0ff2f9220d..42b04862d9 100644
--- a/aten/src/THC/THCTensorMathPointwise.cuh
+++ b/aten/src/THC/THCTensorMathPointwise.cuh
@@ -440,11 +440,25 @@ struct TensorDivOp<half> {
 };
 #endif // CUDA_HALF_TENSOR
 
+template<typename T>
+static __device__ __forceinline__
+typename std::enable_if<std::is_signed<T>::value, bool>::type
+modulo_wrap(T a, T b) {
+  return (a != 0) && (a < 0) != (b < 0);
+}
+
+template<typename T>
+static __device__ __forceinline__
+typename std::enable_if<std::is_unsigned<T>::value, bool>::type
+modulo_wrap(T a, T b) {
+  return false;
+}
+
 template <typename T>
 struct TensorCRemainderOp {
   __device__ __forceinline__ void operator()(T* out, T* in) {
     T val =  *out % *in;
-    if ((val * *in)<0){
+    if (modulo_wrap(val, *in)) {
       val += *in;
     }
     *out = val;
@@ -452,7 +466,7 @@ struct TensorCRemainderOp {
 
   __device__ __forceinline__ void operator()(T* out, T* in1, T* in2) {
     T val = *in1 % *in2;
-    if ((val * *in2)<0){
+    if (modulo_wrap(val, *in2)) {
       val += *in2;
     }
     *out = val;
@@ -462,22 +476,22 @@ struct TensorCRemainderOp {
 template <>
 struct TensorCRemainderOp<float> {
   __device__ __forceinline__ void operator()(float* out, float* in) {
-    *out = *in != 0 ? *out - *in * floorf(*out / *in) : NAN;
+    *out = *in != 0.f ? *out - *in * floorf(*out / *in) : NAN;
   }
 
   __device__ __forceinline__ void operator()(float* out, float* in1, float* in2) {
-    *out = *in2 != 0 ? *in1 - *in2 * floorf(*in1 / *in2) : NAN;
+    *out = *in2 != 0.f ? *in1 - *in2 * floorf(*in1 / *in2) : NAN;
   }
 };
 
 template <>
 struct TensorCRemainderOp<double> {
   __device__ __forceinline__ void operator()(double* out, double* in) {
-    *out = *in != 0 ? *out - *in * floor(*out / *in) : NAN;
+    *out = *in != 0. ? *out - *in * floor(*out / *in) : NAN;
   }
 
   __device__ __forceinline__ void operator()(double* out, double* in1, double* in2) {
-    *out = *in2 != 0 ? *in1 - *in2 * floor(*in1 / *in2) : NAN;
+    *out = *in2 != 0. ? *in1 - *in2 * floor(*in1 / *in2) : NAN;
   }
 };