From 7b3b82bd0d320c00820b020d954174e32cdea45a Mon Sep 17 00:00:00 2001
From: "Yaxun (Sam) Liu" <yaxun.liu@amd.com>
Date: Mon, 30 Sep 2019 01:10:42 -0400
Subject: [PATCH] Fix cast of __half for HIP-clang (#1475)

---
 hipamd/include/hip/hcc_detail/hip_fp16.h   | 1 +
 hipamd/tests/src/deviceLib/hipTestHalf.cpp | 5 +++++
 2 files changed, 6 insertions(+)

diff --git a/hipamd/include/hip/hcc_detail/hip_fp16.h b/hipamd/include/hip/hcc_detail/hip_fp16.h
index 11e1b6be11..18ad566121 100644
--- a/hipamd/include/hip/hcc_detail/hip_fp16.h
+++ b/hipamd/include/hip/hcc_detail/hip_fp16.h
@@ -223,6 +223,7 @@ THE SOFTWARE.
                     Enable_if_t<
                         std::is_floating_point<T>{} &&
                         !std::is_same<T, double>{}>* = nullptr>
+                __host__ __device__
                 operator T() const { return data; }
             #endif
             __host__ __device__
diff --git a/hipamd/tests/src/deviceLib/hipTestHalf.cpp b/hipamd/tests/src/deviceLib/hipTestHalf.cpp
index d48ce9b4f6..2056cffa91 100644
--- a/hipamd/tests/src/deviceLib/hipTestHalf.cpp
+++ b/hipamd/tests/src/deviceLib/hipTestHalf.cpp
@@ -30,6 +30,11 @@ THE SOFTWARE.
 
 #if __HIP_ARCH_GFX803__ || __HIP_ARCH_GFX900__ || __HIP_ARCH_GFX906__ || __HIP_ARCH_GFX908__
 
+__device__ void test_convert() {
+  __half x;
+  float y = (float)x;
+}
+
 __global__
 void __halfMath(bool* result, __half a) {
   result[0] = __heq(__hadd(a, __half{1}), __half{2});