float2::x and float2::y may not be the same as float on ROCm (#35593)

xuhdev · facebook-github-bot · commit 788ef939d837 · 2020-03-30T13:51:32.000-07:00
Summary: This causes ambiguity and can be triggered sometimes (e.g., by #35217). Explicitly convert them to float. error: conditional expression is ambiguous; 'const hip_impl::Scalar_accessor<float, Native_vec_, 0>' can be converted to 'float' and vice versa Pull Request resolved: #35593 Differential Revision: D20735663 Pulled By: ezyang fbshipit-source-id: ae6a38a08e59821bae13eb0b9f9bdf21a008d5c0
diff --git a/caffe2/operators/relu_op.cu b/caffe2/operators/relu_op.cu
@@ -54,8 +54,13 @@ __global__ void ReluCUDAKernel<half2>(const int N, const half2* X, half2* Y) {
     Y[i] = __hmul2(__hgt2(__ldg(X + i), kZero), __ldg(X + i));
 #else
     const float2 xx = __half22float2(X[i]);
-    Y[i] =
-        __floats2half2_rn(xx.x > 0.0f ? xx.x : 0.0f, xx.y > 0.0f ? xx.y : 0.0f);
+    // There are explicit cast to float here, because it may otherwise cause ambiguity on ROCm and can be triggered
+    // sometimes:
+    //
+    //   error: conditional expression is ambiguous; 'const hip_impl::Scalar_accessor<float, Native_vec_, 0>' can be
+    //   converted to 'float' and vice versa
+    Y[i] = __floats2half2_rn(xx.x > 0.0f ? static_cast<float>(xx.x) : 0.0f,
+                             xx.y > 0.0f ? static_cast<float>(xx.y) : 0.0f);
 #endif
   }
 }
