     else {return max(a, b);}
 }
 
-#if __CUDACC_VER_MAJOR__ >= 9
+#if __CUDACC_VER_MAJOR__ >= 9 && (__CUDA_ARCH__ >= 530 || !defined(__CUDA_ARCH__))
 // specialization for half for handling NaNs
 template <>
 __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
 }
 #endif
 

#include "cupy_cub.h"
#include <stdexcept>

#if __CUDACC_VER_MAJOR__ >= 9
#include <cuda_fp16.h>
#endif


    else {return min(a, b);}
}

#if __CUDACC_VER_MAJOR__ >= 9 && (__CUDA_ARCH__ >= 530 || !defined(__CUDA_ARCH__))
// specialization for half for handling NaNs
template <>
__host__ __device__ __forceinline__ __half Min::operator()(const __half &a, const __half &b) const
{
#ifdef __CUDA_ARCH__
    // NumPy behavior: NaN is always chosen!
    if (__hisnan(a)) {return a;}
    else if (__hisnan(b)) {return b;}
    else {return CUB_MIN(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_MIN(__half2float(a), __half2float(b)));}
#endif
}
#endif

