Do relu in float32 instead of bhalf_t. Because bhalf_t is unsigned

include/ck/tensor_operation/gpu/element/element_wise_operation.hpp

@@ -150,7 +150,12 @@ struct Relu
 
     __host__ __device__ void operator()(half_t& y, const half_t& x) const { y = x > 0 ? x : 0; }
 
-    __host__ __device__ void operator()(bhalf_t& y, const bhalf_t& x) const { y = x > 0 ? x : 0; }
+    __host__ __device__ void operator()(bhalf_t& y, const bhalf_t& x) const
+    {
+        float x_f32 = ck::type_convert<float>(x);
+        float y_f32 = x_f32 > 0 ? x_f32 : 0;
+        y           = ck::type_convert<bhalf_t>(y_f32);
+    }
 
     __host__ __device__ void operator()(int32_t& y, const int32_t& x) const { y = x > 0 ? x : 0; }