From 0e1860dd533daf8b91e463d29d5caf28da63c5cb 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] Add tests for dot functions

[ROCm/clr commit: bbe5a0381f9635798a3122f49c4c45d4c5628f6c]
---
 .../src/deviceLib/hipTestDotFunctions.cpp     | 69 +++++++++++++++++++
 .../tests/src/deviceLib/hipTestNativeHalf.cpp |  5 ++
 2 files changed, 74 insertions(+)
 create mode 100644 projects/clr/hipamd/tests/src/deviceLib/hipTestDotFunctions.cpp

diff --git a/projects/clr/hipamd/tests/src/deviceLib/hipTestDotFunctions.cpp b/projects/clr/hipamd/tests/src/deviceLib/hipTestDotFunctions.cpp
new file mode 100644
index 0000000000..72d7c2d26d
--- /dev/null
+++ b/projects/clr/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/projects/clr/hipamd/tests/src/deviceLib/hipTestNativeHalf.cpp b/projects/clr/hipamd/tests/src/deviceLib/hipTestNativeHalf.cpp
index 6e618618f6..52fa0ba932 100644
--- a/projects/clr/hipamd/tests/src/deviceLib/hipTestNativeHalf.cpp
+++ b/projects/clr/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