From f0b114ee1ccdaa3ae7f49dc782f36e757be4d985 Mon Sep 17 00:00:00 2001
From: Aaron Enye Shi <enye.shi@gmail.com>
Date: Mon, 14 Jan 2019 21:50:32 +0000
Subject: [PATCH] Move fp16 dot functions to hip_fp16

---
 include/hip/hcc_detail/hip_fp16.h          | 7 +++++++
 include/hip/hcc_detail/hip_fp16_math_fwd.h | 4 +++-
 include/hip/hcc_detail/math_functions.h    | 5 -----
 include/hip/hcc_detail/math_fwd.h          | 4 ----
 4 files changed, 10 insertions(+), 10 deletions(-)

diff --git a/include/hip/hcc_detail/hip_fp16.h b/include/hip/hcc_detail/hip_fp16.h
index 5a6e650069..45390c8282 100644
--- a/include/hip/hcc_detail/hip_fp16.h
+++ b/include/hip/hcc_detail/hip_fp16.h
@@ -1405,6 +1405,13 @@ THE SOFTWARE.
             // Math functions
             inline
             __device__
+            float amd_mixed_dot(__half2 a, __half2 b, float c, bool saturate) {
+                return __ockl_fdot2(static_cast<__half2_raw>(a).data,
+                                    static_cast<__half2_raw>(b).data,
+                                    c, saturate);
+            }
+            inline
+            __device__
             __half htrunc(__half x)
             {
                 return __half_raw{
diff --git a/include/hip/hcc_detail/hip_fp16_math_fwd.h b/include/hip/hcc_detail/hip_fp16_math_fwd.h
index 83fd1349b1..ebc048dcb7 100644
--- a/include/hip/hcc_detail/hip_fp16_math_fwd.h
+++ b/include/hip/hcc_detail/hip_fp16_math_fwd.h
@@ -53,6 +53,8 @@ extern "C"
     typedef _Float16 __2f16 __attribute__((ext_vector_type(2)));
     typedef short __2i16 __attribute__((ext_vector_type(2)));
 
+    __device__ __attribute__((const)) float __ockl_fdot2(__2f16 a, __2f16 b, float c, bool s);
+
     __device__ __attribute__((const)) __2f16 __ocml_ceil_2f16(__2f16);
     __device__ __2f16 __ocml_cos_2f16(__2f16);
     __device__ __attribute__((pure)) __2f16 __ocml_exp_2f16(__2f16);
@@ -75,4 +77,4 @@ extern "C"
     __device__ __2f16 __ocml_sin_2f16(__2f16);
     __device__ __attribute__((const)) __2f16 __ocml_sqrt_2f16(__2f16);
     __device__ __attribute__((const)) __2f16 __ocml_trunc_2f16(__2f16);
-}
\ No newline at end of file
+}
diff --git a/include/hip/hcc_detail/math_functions.h b/include/hip/hcc_detail/math_functions.h
index 69c9f358c0..96566eab4e 100644
--- a/include/hip/hcc_detail/math_functions.h
+++ b/include/hip/hcc_detail/math_functions.h
@@ -123,11 +123,6 @@ uint64_t __make_mantissa(const char* tagp)
 // DOT FUNCTIONS
 __DEVICE__
 inline
-float amd_mixed_dot(__2f16 a, __2f16 b, float c, bool saturate) {
-    return __ockl_fdot2(a, b, c, saturate);
-}
-__DEVICE__
-inline
 int amd_mixed_dot(short2 a, short2 b, int c, bool saturate) {
     return __ockl_sdot2(a, b, c, saturate);
 }
diff --git a/include/hip/hcc_detail/math_fwd.h b/include/hip/hcc_detail/math_fwd.h
index df611dfe80..3636899618 100644
--- a/include/hip/hcc_detail/math_fwd.h
+++ b/include/hip/hcc_detail/math_fwd.h
@@ -29,10 +29,6 @@ THE SOFTWARE.
 #endif
 
 // DOT FUNCTIONS
-typedef _Float16 __2f16 __attribute__((ext_vector_type(2)));
-__device__
-__attribute__((const))
-float __ockl_fdot2(__2f16 a, __2f16 b, float c, bool s);
 __device__
 __attribute__((const))
 int __ockl_sdot2(short2 a, short2 b, int c, bool s);