     // 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 b;}
+    else if (half_less(a, b)) {return b;}
     else {return a;}
 }
 #endif

    // 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

#endif
}

__host__ __device__ __forceinline__ bool half_less(const __half& l, const __half& r) {
#ifdef __CUDA_ARCH__
    return l < r;
#else
