From d26ab68485fe32da1b696d0ce3a273a8c43ae207 Mon Sep 17 00:00:00 2001 From: Adam Paszke Date: Thu, 19 Apr 2018 03:51:35 +0200 Subject: 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 --- aten/src/TH/generic/THTensorMath.c | 17 +++++++++-------- aten/src/THC/THCNumerics.cuh | 5 ----- aten/src/THC/THCTensorMathPairwise.cu | 18 ++++++++++++++++-- aten/src/THC/THCTensorMathPointwise.cuh | 26 ++++++++++++++++++++------ 4 files changed, 45 insertions(+), 21 deletions(-) (limited to 'aten') 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 -static inline __host__ __device__ bool has_different_sign(scalar_t a, scalar_t b) { - return (a < 0) != (b < 0); -} - template <> struct THCNumerics { 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 { }; #endif // CUDA_HALF_TENSOR +template +static __device__ __forceinline__ +typename std::enable_if::value, bool>::type +modulo_wrap(T a, T b) { + return (a != 0) && (a < 0) != (b < 0); +} + +template +static __device__ __forceinline__ +typename std::enable_if::value, bool>::type +modulo_wrap(T a, T b) { + return false; +} + template struct TensorRemainderOp { TensorRemainderOp(T v) : val(v) {} __device__ __forceinline__ void operator()(T* out, T* in) { *out = *in % val; - if (has_different_sign(*out, val)){ + if (modulo_wrap(*out, val)) { *out += val; } } __device__ __forceinline__ void operator()(T* v) { *v = *v % val; - if (has_different_sign(*v, val)){ + if (modulo_wrap(*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 { }; #endif // CUDA_HALF_TENSOR +template +static __device__ __forceinline__ +typename std::enable_if::value, bool>::type +modulo_wrap(T a, T b) { + return (a != 0) && (a < 0) != (b < 0); +} + +template +static __device__ __forceinline__ +typename std::enable_if::value, bool>::type +modulo_wrap(T a, T b) { + return false; +} + template 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 { __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 { __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; } }; -- cgit v1.2.3