 __host__ __device__ __forceinline__ __half Max::operator()(const __half &a, const __half &b) const
 {
     // NumPy behavior: NaN is always chosen!
-#ifdef __CUDA_ARCH__
-    if (__hisnan(a)) {return a;}
-    else if (__hisnan(b)) {return b;}
-    else {return CUB_MAX(a, b);}
-#else
-    // TODO: avoid cast to float
-    if (isnan(__half2float(a))) {return a;}
-    else if (isnan(__half2float(b))) {return b;}
-    else {return __float2half(CUB_MAX(__half2float(a), __half2float(b)));}
-#endif
+    if (half_isnan(a)) {return a;}
+    else if (half_isnan(b)) {return b;}
+    else if (half_greater_than(a, b)) {return b;}
+    else {return a;}
 }
 #endif
 

    const KeyValuePair<int, __half> &a,
    const KeyValuePair<int, __half> &b) const
{
#ifdef __CUDA_ARCH__
    if (__hisnan(a.value))
        return a;
    else if (__hisnan(b.value))
        return b;
    else if ((b.value < a.value) || ((a.value == b.value) && (b.key < a.key)))
        return b;
    else
        return a;
#else
    // TODO: avoid cast to float
    if (isnan(__half2float(a.value)))
        return a;
    else if (isnan(__half2float(b.value)))
        return b;
    else if ((__half2float(b.value) < __half2float(a.value)) ||
            ((__half2float(a.value) == __half2float(b.value)) && (b.key < a.key)))
        return b;
    else
        return a;
#endif
}
#endif
/* ------------------------------------ End of "patches" ------------------------------------ */

template <>
__host__ __device__ __forceinline__ __half Min::operator()(const __half &a, const __half &b) const
{
    // NumPy behavior: NaN is always chosen!
    if (half_isnan(a)) {return a;}
    else if (half_isnan(b)) {return b;}
    else if (half_greater_than(a, b)) {return a;}
    else {return a;}
}
#endif

