From 84a028b9e1601129254f82ad4a7f00be91f0c564 Mon Sep 17 00:00:00 2001
From: Aaron Enye Shi <enye.shi@gmail.com>
Date: Fri, 11 Jan 2019 23:22:07 +0000
Subject: [PATCH 1/9] Add Dot functions as amd_mixed_dot function

Introduce the Dot functions which are available in the device library. Forward those prototypes, and introduce HIP API to expose the usage of the dot functions.
---
 .../include/hip/hcc_detail/math_functions.h   | 38 +++++++++++++++++++
 hipamd/include/hip/hcc_detail/math_fwd.h      | 24 ++++++++++++
 2 files changed, 62 insertions(+)

diff --git a/hipamd/include/hip/hcc_detail/math_functions.h b/hipamd/include/hip/hcc_detail/math_functions.h
index 63e48fab29..69c9f358c0 100644
--- a/hipamd/include/hip/hcc_detail/math_functions.h
+++ b/hipamd/include/hip/hcc_detail/math_functions.h
@@ -23,6 +23,7 @@ THE SOFTWARE.
 #pragma once
 
 #include "hip_fp16_math_fwd.h"
+#include "hip_vector_types.h"
 #include "math_fwd.h"
 
 #include <hip/hcc_detail/host_defines.h>
@@ -119,6 +120,43 @@ uint64_t __make_mantissa(const char* tagp)
     return __make_mantissa_base10(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);
+}
+__DEVICE__
+inline
+uint amd_mixed_dot(ushort2 a, ushort2 b, uint c, bool saturate) {
+    return __ockl_udot2(a, b, c, saturate);
+}
+__DEVICE__
+inline
+int amd_mixed_dot(char4 a, char4 b, int c, bool saturate) {
+    return __ockl_sdot4(a, b, c, saturate);
+}
+__DEVICE__
+inline
+uint amd_mixed_dot(uchar4 a, uchar4 b, uint c, bool saturate) {
+    return __ockl_udot4(a, b, c, saturate);
+}
+__DEVICE__
+inline
+int amd_mixed_dot(int a, int b, int c, bool saturate) {
+    return __ockl_sdot8(a, b, c, saturate);
+}
+__DEVICE__
+inline
+uint amd_mixed_dot(uint a, uint b, uint c, bool saturate) {
+    return __ockl_udot8(a, b, c, saturate);
+}
+
 // BEGIN FLOAT
 __DEVICE__
 inline
diff --git a/hipamd/include/hip/hcc_detail/math_fwd.h b/hipamd/include/hip/hcc_detail/math_fwd.h
index e5594924ba..df611dfe80 100644
--- a/hipamd/include/hip/hcc_detail/math_fwd.h
+++ b/hipamd/include/hip/hcc_detail/math_fwd.h
@@ -28,6 +28,30 @@ THE SOFTWARE.
     extern "C" {
 #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);
+__device__
+__attribute__((const))
+unsigned int __ockl_udot2(ushort2 a, ushort2 b, unsigned int c, bool s);
+__device__
+__attribute__((const))
+int __ockl_sdot4(char4 a, char4 b, int c, bool s);
+__device__
+__attribute__((const))
+unsigned int __ockl_udot4(uchar4 a, uchar4 b, unsigned int c, bool s);
+__device__
+__attribute__((const))
+int __ockl_sdot8(int a, int b, int c, bool s);
+__device__
+__attribute__((const))
+unsigned int __ockl_udot8(unsigned int a, unsigned int b, unsigned int c, bool s);
+
 // BEGIN FLOAT
 __device__
 __attribute__((const))

From 7ba7d9a0f65d1a3b1c97ee39065c6c4c4536e939 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 2/9] Move fp16 dot functions to hip_fp16

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

diff --git a/hipamd/include/hip/hcc_detail/hip_fp16.h b/hipamd/include/hip/hcc_detail/hip_fp16.h
index 5a6e650069..45390c8282 100644
--- a/hipamd/include/hip/hcc_detail/hip_fp16.h
+++ b/hipamd/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/hipamd/include/hip/hcc_detail/hip_fp16_math_fwd.h b/hipamd/include/hip/hcc_detail/hip_fp16_math_fwd.h
index 83fd1349b1..ebc048dcb7 100644
--- a/hipamd/include/hip/hcc_detail/hip_fp16_math_fwd.h
+++ b/hipamd/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/hipamd/include/hip/hcc_detail/math_functions.h b/hipamd/include/hip/hcc_detail/math_functions.h
index 69c9f358c0..96566eab4e 100644
--- a/hipamd/include/hip/hcc_detail/math_functions.h
+++ b/hipamd/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/hipamd/include/hip/hcc_detail/math_fwd.h b/hipamd/include/hip/hcc_detail/math_fwd.h
index df611dfe80..3636899618 100644
--- a/hipamd/include/hip/hcc_detail/math_fwd.h
+++ b/hipamd/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);

From bdb84f3d9cac41ed5ab482eb5aff8ba2541a1484 Mon Sep 17 00:00:00 2001
From: Aaron Enye Shi <enye.shi@gmail.com>
Date: Mon, 14 Jan 2019 23:14:49 +0000
Subject: [PATCH 3/9] Add backwards compatibility to dot funcs

---
 hipamd/include/hip/hcc_detail/hip_fp16.h          | 2 ++
 hipamd/include/hip/hcc_detail/hip_fp16_math_fwd.h | 2 ++
 hipamd/include/hip/hcc_detail/math_functions.h    | 2 ++
 hipamd/include/hip/hcc_detail/math_fwd.h          | 2 ++
 4 files changed, 8 insertions(+)

diff --git a/hipamd/include/hip/hcc_detail/hip_fp16.h b/hipamd/include/hip/hcc_detail/hip_fp16.h
index 45390c8282..103e9765e1 100644
--- a/hipamd/include/hip/hcc_detail/hip_fp16.h
+++ b/hipamd/include/hip/hcc_detail/hip_fp16.h
@@ -1403,6 +1403,7 @@ THE SOFTWARE.
             }
 
             // Math functions
+            #if (__hcc_workweek__ >= 19015)
             inline
             __device__
             float amd_mixed_dot(__half2 a, __half2 b, float c, bool saturate) {
@@ -1410,6 +1411,7 @@ THE SOFTWARE.
                                     static_cast<__half2_raw>(b).data,
                                     c, saturate);
             }
+            #endif
             inline
             __device__
             __half htrunc(__half x)
diff --git a/hipamd/include/hip/hcc_detail/hip_fp16_math_fwd.h b/hipamd/include/hip/hcc_detail/hip_fp16_math_fwd.h
index ebc048dcb7..2cb8556e63 100644
--- a/hipamd/include/hip/hcc_detail/hip_fp16_math_fwd.h
+++ b/hipamd/include/hip/hcc_detail/hip_fp16_math_fwd.h
@@ -53,7 +53,9 @@ extern "C"
     typedef _Float16 __2f16 __attribute__((ext_vector_type(2)));
     typedef short __2i16 __attribute__((ext_vector_type(2)));
 
+    #if (__hcc_workweek__ >= 19015)
     __device__ __attribute__((const)) float __ockl_fdot2(__2f16 a, __2f16 b, float c, bool s);
+    #endif
 
     __device__ __attribute__((const)) __2f16 __ocml_ceil_2f16(__2f16);
     __device__ __2f16 __ocml_cos_2f16(__2f16);
diff --git a/hipamd/include/hip/hcc_detail/math_functions.h b/hipamd/include/hip/hcc_detail/math_functions.h
index 96566eab4e..c5ebd9703d 100644
--- a/hipamd/include/hip/hcc_detail/math_functions.h
+++ b/hipamd/include/hip/hcc_detail/math_functions.h
@@ -121,6 +121,7 @@ uint64_t __make_mantissa(const char* tagp)
 }
 
 // DOT FUNCTIONS
+#if (__hcc_workweek__ >= 19015)
 __DEVICE__
 inline
 int amd_mixed_dot(short2 a, short2 b, int c, bool saturate) {
@@ -151,6 +152,7 @@ inline
 uint amd_mixed_dot(uint a, uint b, uint c, bool saturate) {
     return __ockl_udot8(a, b, c, saturate);
 }
+#endif
 
 // BEGIN FLOAT
 __DEVICE__
diff --git a/hipamd/include/hip/hcc_detail/math_fwd.h b/hipamd/include/hip/hcc_detail/math_fwd.h
index 3636899618..9d4d51cf36 100644
--- a/hipamd/include/hip/hcc_detail/math_fwd.h
+++ b/hipamd/include/hip/hcc_detail/math_fwd.h
@@ -29,6 +29,7 @@ THE SOFTWARE.
 #endif
 
 // DOT FUNCTIONS
+#if (__hcc_workweek__ >= 19015)
 __device__
 __attribute__((const))
 int __ockl_sdot2(short2 a, short2 b, int c, bool s);
@@ -47,6 +48,7 @@ int __ockl_sdot8(int a, int b, int c, bool s);
 __device__
 __attribute__((const))
 unsigned int __ockl_udot8(unsigned int a, unsigned int b, unsigned int c, bool s);
+#endif
 
 // BEGIN FLOAT
 __device__

From b0e0ec3c673d8e66920ccb03e4a13e40159f0a60 Mon Sep 17 00:00:00 2001
From: Aaron Enye Shi <enye.shi@gmail.com>
Date: Tue, 22 Jan 2019 16:21:16 +0000
Subject: [PATCH 4/9] Allow hip-clang to use amd_mixed_dot as well

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

diff --git a/hipamd/include/hip/hcc_detail/hip_fp16.h b/hipamd/include/hip/hcc_detail/hip_fp16.h
index 103e9765e1..a718678539 100644
--- a/hipamd/include/hip/hcc_detail/hip_fp16.h
+++ b/hipamd/include/hip/hcc_detail/hip_fp16.h
@@ -1403,7 +1403,7 @@ THE SOFTWARE.
             }
 
             // Math functions
-            #if (__hcc_workweek__ >= 19015)
+            #if (__hcc_workweek__ >= 19015) || defined(__HIP_CLANG_ONLY__)
             inline
             __device__
             float amd_mixed_dot(__half2 a, __half2 b, float c, bool saturate) {
diff --git a/hipamd/include/hip/hcc_detail/hip_fp16_math_fwd.h b/hipamd/include/hip/hcc_detail/hip_fp16_math_fwd.h
index 2cb8556e63..0c025db786 100644
--- a/hipamd/include/hip/hcc_detail/hip_fp16_math_fwd.h
+++ b/hipamd/include/hip/hcc_detail/hip_fp16_math_fwd.h
@@ -53,7 +53,7 @@ extern "C"
     typedef _Float16 __2f16 __attribute__((ext_vector_type(2)));
     typedef short __2i16 __attribute__((ext_vector_type(2)));
 
-    #if (__hcc_workweek__ >= 19015)
+    #if (__hcc_workweek__ >= 19015) || defined(__HIP_CLANG_ONLY__)
     __device__ __attribute__((const)) float __ockl_fdot2(__2f16 a, __2f16 b, float c, bool s);
     #endif
 
diff --git a/hipamd/include/hip/hcc_detail/math_functions.h b/hipamd/include/hip/hcc_detail/math_functions.h
index c5ebd9703d..c3a790751b 100644
--- a/hipamd/include/hip/hcc_detail/math_functions.h
+++ b/hipamd/include/hip/hcc_detail/math_functions.h
@@ -121,7 +121,7 @@ uint64_t __make_mantissa(const char* tagp)
 }
 
 // DOT FUNCTIONS
-#if (__hcc_workweek__ >= 19015)
+#if (__hcc_workweek__ >= 19015) || defined(__HIP_CLANG_ONLY__)
 __DEVICE__
 inline
 int amd_mixed_dot(short2 a, short2 b, int c, bool saturate) {
diff --git a/hipamd/include/hip/hcc_detail/math_fwd.h b/hipamd/include/hip/hcc_detail/math_fwd.h
index 9d4d51cf36..cdf8e78da4 100644
--- a/hipamd/include/hip/hcc_detail/math_fwd.h
+++ b/hipamd/include/hip/hcc_detail/math_fwd.h
@@ -29,7 +29,7 @@ THE SOFTWARE.
 #endif
 
 // DOT FUNCTIONS
-#if (__hcc_workweek__ >= 19015)
+#if (__hcc_workweek__ >= 19015) || defined(__HIP_CLANG_ONLY__)
 __device__
 __attribute__((const))
 int __ockl_sdot2(short2 a, short2 b, int c, bool s);

From b63597bb95db75d2d2dff0385bb90366e4fa5416 Mon Sep 17 00:00:00 2001
From: Aaron Enye Shi <enye.shi@gmail.com>
Date: Tue, 22 Jan 2019 22:31:19 +0000
Subject: [PATCH 5/9] Use correct OCKL native vector types

---
 .../include/hip/hcc_detail/math_functions.h   |  8 ++---
 hipamd/include/hip/hcc_detail/math_fwd.h      | 29 +++++++++++++++----
 2 files changed, 27 insertions(+), 10 deletions(-)

diff --git a/hipamd/include/hip/hcc_detail/math_functions.h b/hipamd/include/hip/hcc_detail/math_functions.h
index c3a790751b..24728550ad 100644
--- a/hipamd/include/hip/hcc_detail/math_functions.h
+++ b/hipamd/include/hip/hcc_detail/math_functions.h
@@ -125,22 +125,22 @@ uint64_t __make_mantissa(const char* tagp)
 __DEVICE__
 inline
 int amd_mixed_dot(short2 a, short2 b, int c, bool saturate) {
-    return __ockl_sdot2(a, b, c, saturate);
+    return __ockl_sdot2(a.data, b.data, c, saturate);
 }
 __DEVICE__
 inline
 uint amd_mixed_dot(ushort2 a, ushort2 b, uint c, bool saturate) {
-    return __ockl_udot2(a, b, c, saturate);
+    return __ockl_udot2(a.data, b.data, c, saturate);
 }
 __DEVICE__
 inline
 int amd_mixed_dot(char4 a, char4 b, int c, bool saturate) {
-    return __ockl_sdot4(a, b, c, saturate);
+    return __ockl_sdot4(a.data, b.data, c, saturate);
 }
 __DEVICE__
 inline
 uint amd_mixed_dot(uchar4 a, uchar4 b, uint c, bool saturate) {
-    return __ockl_udot4(a, b, c, saturate);
+    return __ockl_udot4(a.data, b.data, c, saturate);
 }
 __DEVICE__
 inline
diff --git a/hipamd/include/hip/hcc_detail/math_fwd.h b/hipamd/include/hip/hcc_detail/math_fwd.h
index cdf8e78da4..3021f92d13 100644
--- a/hipamd/include/hip/hcc_detail/math_fwd.h
+++ b/hipamd/include/hip/hcc_detail/math_fwd.h
@@ -32,22 +32,39 @@ THE SOFTWARE.
 #if (__hcc_workweek__ >= 19015) || defined(__HIP_CLANG_ONLY__)
 __device__
 __attribute__((const))
-int __ockl_sdot2(short2 a, short2 b, int c, bool s);
+int __ockl_sdot2(
+    HIP_vector_base<short, 2>::Native_vec_,
+    HIP_vector_base<short, 2>::Native_vec_,
+    int, bool);
+
 __device__
 __attribute__((const))
-unsigned int __ockl_udot2(ushort2 a, ushort2 b, unsigned int c, bool s);
+unsigned int __ockl_udot2(
+    HIP_vector_base<unsigned short, 2>::Native_vec_,
+    HIP_vector_base<unsigned short, 2>::Native_vec_,
+    unsigned int, bool);
+
 __device__
 __attribute__((const))
-int __ockl_sdot4(char4 a, char4 b, int c, bool s);
+int __ockl_sdot4(
+    HIP_vector_base<char, 4>::Native_vec_,
+    HIP_vector_base<char, 4>::Native_vec_,
+    int, bool);
+
 __device__
 __attribute__((const))
-unsigned int __ockl_udot4(uchar4 a, uchar4 b, unsigned int c, bool s);
+unsigned int __ockl_udot4(
+    HIP_vector_base<unsigned char, 4>::Native_vec_,
+    HIP_vector_base<unsigned char, 4>::Native_vec_,
+    unsigned int, bool);
+
 __device__
 __attribute__((const))
-int __ockl_sdot8(int a, int b, int c, bool s);
+int __ockl_sdot8(int, int, int, bool);
+
 __device__
 __attribute__((const))
-unsigned int __ockl_udot8(unsigned int a, unsigned int b, unsigned int c, bool s);
+unsigned int __ockl_udot8(unsigned int, unsigned int, unsigned int, bool);
 #endif
 
 // BEGIN FLOAT

From bbe5a0381f9635798a3122f49c4c45d4c5628f6c Mon Sep 17 00:00:00 2001
From: Aaron Enye Shi <enye.shi@gmail.com>
Date: Tue, 22 Jan 2019 22:32:28 +0000
Subject: [PATCH 6/9] Add tests for dot functions

---
 .../src/deviceLib/hipTestDotFunctions.cpp     | 69 +++++++++++++++++++
 .../tests/src/deviceLib/hipTestNativeHalf.cpp |  5 ++
 2 files changed, 74 insertions(+)
 create mode 100644 hipamd/tests/src/deviceLib/hipTestDotFunctions.cpp

diff --git a/hipamd/tests/src/deviceLib/hipTestDotFunctions.cpp b/hipamd/tests/src/deviceLib/hipTestDotFunctions.cpp
new file mode 100644
index 0000000000..72d7c2d26d
--- /dev/null
+++ b/hipamd/tests/src/deviceLib/hipTestDotFunctions.cpp
@@ -0,0 +1,69 @@
+/*
+Copyright (c) 2015-2019 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+/* HIT_START
+ * BUILD: %t %s ../test_common.cpp
+ * RUN: %t
+ * HIT_END
+ */
+
+#include <hip/hip_runtime.h>
+#include <hip/math_functions.h>
+#include "test_common.h"
+
+
+__global__ void DotFunctions(bool* result) {
+    #if (__hcc_workweek__ >= 19015) || defined(__HIP_CLANG_ONLY__)
+    // Dot Functions
+    short2 sa{1}, sb{1};
+    result[0] = amd_mixed_dot(sa, sb, 1, result[0]) && result[0];
+
+    ushort2 usa{1}, usb{1};
+    result[0] = amd_mixed_dot(usa, usb, (uint) 1, result[0]) && result[0];
+
+    char4 ca{1}, cb{1};
+    result[0] = amd_mixed_dot(ca, cb, 1, result[0]) && result[0];
+
+    uchar4 uca{1}, ucb{1};
+    result[0] = amd_mixed_dot(uca, ucb, (uint) 1, result[0]) && result[0];
+
+    int ia{1}, ib{1};
+    result[0] = amd_mixed_dot(ia, ib, 1, result[0]) && result[0];
+
+    uint ua{1}, ub{1};
+    result[0] = amd_mixed_dot(ua, ub, (uint) 1, result[0]) && result[0];
+    #endif
+}
+
+int main() {
+    bool* result{nullptr};
+    hipHostMalloc(&result, 1);
+    result[0] = true;
+
+    hipLaunchKernelGGL(DotFunctions, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, result);
+    hipDeviceSynchronize();
+    if (!result[0]) { failed("Failed dot tests."); }
+
+    hipHostFree(result);
+
+    passed();
+}
diff --git a/hipamd/tests/src/deviceLib/hipTestNativeHalf.cpp b/hipamd/tests/src/deviceLib/hipTestNativeHalf.cpp
index 6e618618f6..52fa0ba932 100644
--- a/hipamd/tests/src/deviceLib/hipTestNativeHalf.cpp
+++ b/hipamd/tests/src/deviceLib/hipTestNativeHalf.cpp
@@ -156,6 +156,11 @@ void __half2Test(bool* result, __half2 a) {
     result[0] = (a >= a) && result[0];
     result[0] = !(a < a) && result[0];
     result[0] = !(a > a) && result[0];
+
+    #if (__hcc_workweek__ >= 19015) || defined(__HIP_CLANG_ONLY__)
+    // Dot Functions
+    result[0] = to_bool(amd_mixed_dot(a, a, 1, 1)) && result[0];
+    #endif
 }
 
 #endif

From a014668d3a4ce719dd7c429a58c0a7c4ff983ce4 Mon Sep 17 00:00:00 2001
From: Aaron Enye Shi <enye.shi@gmail.com>
Date: Wed, 23 Jan 2019 16:39:25 +0000
Subject: [PATCH 7/9] Fix mixed dot for Jenkins rocm_head

---
 hipamd/tests/src/deviceLib/hipTestNativeHalf.cpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/hipamd/tests/src/deviceLib/hipTestNativeHalf.cpp b/hipamd/tests/src/deviceLib/hipTestNativeHalf.cpp
index 52fa0ba932..3169268dd7 100644
--- a/hipamd/tests/src/deviceLib/hipTestNativeHalf.cpp
+++ b/hipamd/tests/src/deviceLib/hipTestNativeHalf.cpp
@@ -159,7 +159,7 @@ void __half2Test(bool* result, __half2 a) {
 
     #if (__hcc_workweek__ >= 19015) || defined(__HIP_CLANG_ONLY__)
     // Dot Functions
-    result[0] = to_bool(amd_mixed_dot(a, a, 1, 1)) && result[0];
+    result[0] = amd_mixed_dot(a, a, 1, 1) && result[0];
     #endif
 }
 

From 77de87d0cd001cfc42dd73cb3b8c799805915bbe Mon Sep 17 00:00:00 2001
From: Aaron Enye Shi <enye.shi@gmail.com>
Date: Thu, 24 Jan 2019 21:00:45 +0000
Subject: [PATCH 8/9] Exclude mixed dot functions from nvcc path

---
 hipamd/tests/src/deviceLib/hipTestDotFunctions.cpp | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/hipamd/tests/src/deviceLib/hipTestDotFunctions.cpp b/hipamd/tests/src/deviceLib/hipTestDotFunctions.cpp
index 72d7c2d26d..4dbc12522d 100644
--- a/hipamd/tests/src/deviceLib/hipTestDotFunctions.cpp
+++ b/hipamd/tests/src/deviceLib/hipTestDotFunctions.cpp
@@ -21,8 +21,8 @@ THE SOFTWARE.
 */
 
 /* HIT_START
- * BUILD: %t %s ../test_common.cpp
- * RUN: %t
+ * BUILD: %t %s ../test_common.cpp NVCC_OPTIONS -std=c++11
+ * RUN: %t EXCLUDE_HIP_PLATFORM nvcc
  * HIT_END
  */
 

From 2ab6fd374c5325f6c20dc201e80ab615435f4bfd Mon Sep 17 00:00:00 2001
From: Aaron Enye Shi <enye.shi@gmail.com>
Date: Thu, 24 Jan 2019 22:48:35 +0000
Subject: [PATCH 9/9] Fix mixed dot for Jenkins rocm_2_0

---
 hipamd/include/hip/hcc_detail/hip_fp16.h           | 2 +-
 hipamd/include/hip/hcc_detail/hip_fp16_math_fwd.h  | 2 +-
 hipamd/include/hip/hcc_detail/math_functions.h     | 2 +-
 hipamd/include/hip/hcc_detail/math_fwd.h           | 2 +-
 hipamd/tests/src/deviceLib/hipTestDotFunctions.cpp | 2 +-
 hipamd/tests/src/deviceLib/hipTestNativeHalf.cpp   | 2 +-
 6 files changed, 6 insertions(+), 6 deletions(-)

diff --git a/hipamd/include/hip/hcc_detail/hip_fp16.h b/hipamd/include/hip/hcc_detail/hip_fp16.h
index a718678539..849b7278d8 100644
--- a/hipamd/include/hip/hcc_detail/hip_fp16.h
+++ b/hipamd/include/hip/hcc_detail/hip_fp16.h
@@ -1403,7 +1403,7 @@ THE SOFTWARE.
             }
 
             // Math functions
-            #if (__hcc_workweek__ >= 19015) || defined(__HIP_CLANG_ONLY__)
+            #if (__hcc_workweek__ >= 19015) || __HIP_CLANG_ONLY__
             inline
             __device__
             float amd_mixed_dot(__half2 a, __half2 b, float c, bool saturate) {
diff --git a/hipamd/include/hip/hcc_detail/hip_fp16_math_fwd.h b/hipamd/include/hip/hcc_detail/hip_fp16_math_fwd.h
index 0c025db786..eeb617c40b 100644
--- a/hipamd/include/hip/hcc_detail/hip_fp16_math_fwd.h
+++ b/hipamd/include/hip/hcc_detail/hip_fp16_math_fwd.h
@@ -53,7 +53,7 @@ extern "C"
     typedef _Float16 __2f16 __attribute__((ext_vector_type(2)));
     typedef short __2i16 __attribute__((ext_vector_type(2)));
 
-    #if (__hcc_workweek__ >= 19015) || defined(__HIP_CLANG_ONLY__)
+    #if (__hcc_workweek__ >= 19015) || __HIP_CLANG_ONLY__
     __device__ __attribute__((const)) float __ockl_fdot2(__2f16 a, __2f16 b, float c, bool s);
     #endif
 
diff --git a/hipamd/include/hip/hcc_detail/math_functions.h b/hipamd/include/hip/hcc_detail/math_functions.h
index 24728550ad..c1f4b7b1b3 100644
--- a/hipamd/include/hip/hcc_detail/math_functions.h
+++ b/hipamd/include/hip/hcc_detail/math_functions.h
@@ -121,7 +121,7 @@ uint64_t __make_mantissa(const char* tagp)
 }
 
 // DOT FUNCTIONS
-#if (__hcc_workweek__ >= 19015) || defined(__HIP_CLANG_ONLY__)
+#if (__hcc_workweek__ >= 19015) || __HIP_CLANG_ONLY__
 __DEVICE__
 inline
 int amd_mixed_dot(short2 a, short2 b, int c, bool saturate) {
diff --git a/hipamd/include/hip/hcc_detail/math_fwd.h b/hipamd/include/hip/hcc_detail/math_fwd.h
index 3021f92d13..64c5f114bb 100644
--- a/hipamd/include/hip/hcc_detail/math_fwd.h
+++ b/hipamd/include/hip/hcc_detail/math_fwd.h
@@ -29,7 +29,7 @@ THE SOFTWARE.
 #endif
 
 // DOT FUNCTIONS
-#if (__hcc_workweek__ >= 19015) || defined(__HIP_CLANG_ONLY__)
+#if (__hcc_workweek__ >= 19015) || __HIP_CLANG_ONLY__
 __device__
 __attribute__((const))
 int __ockl_sdot2(
diff --git a/hipamd/tests/src/deviceLib/hipTestDotFunctions.cpp b/hipamd/tests/src/deviceLib/hipTestDotFunctions.cpp
index 4dbc12522d..51e005f7fb 100644
--- a/hipamd/tests/src/deviceLib/hipTestDotFunctions.cpp
+++ b/hipamd/tests/src/deviceLib/hipTestDotFunctions.cpp
@@ -32,7 +32,7 @@ THE SOFTWARE.
 
 
 __global__ void DotFunctions(bool* result) {
-    #if (__hcc_workweek__ >= 19015) || defined(__HIP_CLANG_ONLY__)
+    #if (__hcc_workweek__ >= 19015) || __HIP_CLANG_ONLY__
     // Dot Functions
     short2 sa{1}, sb{1};
     result[0] = amd_mixed_dot(sa, sb, 1, result[0]) && result[0];
diff --git a/hipamd/tests/src/deviceLib/hipTestNativeHalf.cpp b/hipamd/tests/src/deviceLib/hipTestNativeHalf.cpp
index 3169268dd7..10a3f33a27 100644
--- a/hipamd/tests/src/deviceLib/hipTestNativeHalf.cpp
+++ b/hipamd/tests/src/deviceLib/hipTestNativeHalf.cpp
@@ -157,7 +157,7 @@ void __half2Test(bool* result, __half2 a) {
     result[0] = !(a < a) && result[0];
     result[0] = !(a > a) && result[0];
 
-    #if (__hcc_workweek__ >= 19015) || defined(__HIP_CLANG_ONLY__)
+    #if (__hcc_workweek__ >= 19015) || __HIP_CLANG_ONLY__
     // Dot Functions
     result[0] = amd_mixed_dot(a, a, 1, 1) && result[0];
     #endif