From 8be953f80ae9509e2a7cbde553809757b1a7063c Mon Sep 17 00:00:00 2001
From: Przemyslaw Tredak <ptredak@nvidia.com>
Date: Mon, 24 Aug 2020 19:55:20 -0700
Subject: [PATCH] Fix isnan usage in RTC (#18984)

* Fix isnan usage

* Add test
---
 src/common/cuda/rtc/forward_functions-inl.h | 50 ++++++++++-----------
 tests/python/unittest/test_numpy_op.py      | 10 +++--
 2 files changed, 31 insertions(+), 29 deletions(-)

diff --git a/src/common/cuda/rtc/forward_functions-inl.h b/src/common/cuda/rtc/forward_functions-inl.h
index 14ee83cd0759..3c64a92d2565 100644
--- a/src/common/cuda/rtc/forward_functions-inl.h
+++ b/src/common/cuda/rtc/forward_functions-inl.h
@@ -215,6 +215,31 @@ __device__ inline void store_add_index(const vector::VectorizedStorage<DType, nv
 const char function_definitions_binary[] = R"code(
 namespace op {
 
+template <typename DType>
+__device__ inline bool isnan(const DType val) {
+  return util::isnan(val);
+}
+
+template <typename DType>
+__device__ inline bool_t isinf(const DType val) {
+  return util::isinf(val);
+}
+
+template <typename DType>
+__device__ inline bool_t isposinf(const DType val) {
+  return util::isinf(val) && (val > 0);
+}
+
+template <typename DType>
+__device__ inline bool_t isneginf(const DType val) {
+  return util::isinf(val) && (val < 0);
+}
+
+template <typename DType>
+__device__ inline bool_t isfinite(const DType val) {
+  return !op::isnan(val) && !op::isinf(val);
+}
+
 template <typename DType, typename DType2>
 __device__ inline typename type_util::mixed_type<DType, DType2>::type
 add(const DType a, const DType2 b) {
@@ -867,31 +892,6 @@ __device__ inline bool_t np_logical_not(const DType val) {
   return !static_cast<bool>(val);
 }
 
-template <typename DType>
-__device__ inline bool isnan(const DType val) {
-  return util::isnan(val);
-}
-
-template <typename DType>
-__device__ inline bool_t isinf(const DType val) {
-  return util::isinf(val);
-}
-
-template <typename DType>
-__device__ inline bool_t isposinf(const DType val) {
-  return util::isinf(val) && (val > 0);
-}
-
-template <typename DType>
-__device__ inline bool_t isneginf(const DType val) {
-  return util::isinf(val) && (val < 0);
-}
-
-template <typename DType>
-__device__ inline bool_t isfinite(const DType val) {
-  return !op::isnan(val) && !op::isinf(val);
-}
-
 #undef DEFINE_UNARY_MATH_FUNC
 
 template <typename DType>
diff --git a/tests/python/unittest/test_numpy_op.py b/tests/python/unittest/test_numpy_op.py
index a45f973a92fa..cc97821fba94 100644
--- a/tests/python/unittest/test_numpy_op.py
+++ b/tests/python/unittest/test_numpy_op.py
@@ -3030,12 +3030,14 @@ def hybrid_forward(self, F, a, b, *args, **kwargs):
         'bitwise_and': (-100, 100, [None], None, [[_np.int32]]),
         'bitwise_xor': (-100, 100, [None], None, [[_np.int32]]),
         'bitwise_or': (-100, 100, [None], None, [[_np.int32]]),
-        'maximum': (-1, 1, [lambda y, x1, x2: _np.ones(y.shape) * (x1 >= x2)],
-                           [lambda y, x1, x2: _np.ones(y.shape) * (x1 < x2)]),
+        'maximum': (-10, 10, [lambda y, x1, x2: _np.ones(y.shape) * (x1 >= x2)],
+                             [lambda y, x1, x2: _np.ones(y.shape) * (x1 < x2)],
+                             [[_np.int32, _np.float16, _np.float32, _np.float64]]),
         'fmax': (-1, 1, [lambda y, x1, x2: _np.ones(y.shape) * (x1 >= x2)],
                         [lambda y, x1, x2: _np.ones(y.shape) * (x1 < x2)]),
-        'minimum': (-1, 1, [lambda y, x1, x2: _np.ones(y.shape) * (x1 <= x2)],
-                           [lambda y, x1, x2: _np.ones(y.shape) * (x1 > x2)]),
+        'minimum': (-10, 10, [lambda y, x1, x2: _np.ones(y.shape) * (x1 <= x2)],
+                             [lambda y, x1, x2: _np.ones(y.shape) * (x1 > x2)],
+                             [[_np.int32, _np.float16, _np.float32, _np.float64]]),
         'fmin': (-1, 1, [lambda y, x1, x2: _np.ones(y.shape) * (x1 <= x2)],
                         [lambda y, x1, x2: _np.ones(y.shape) * (x1 > x2)]),
         'copysign': (-1, 1,