         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
 

        return a;
}

#if __CUDACC_VER_MAJOR__ >= 9
// specialization for half for handling NaNs
template <>
__host__ __device__ __forceinline__ KeyValuePair<int, __half> ArgMin::operator()(
    const KeyValuePair<int, __half> &a,
    const KeyValuePair<int, __half> &b) const
{
    if (__hisnan(a.value))
        return a;
    else if (__hisnan(b.value))

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