Use correct OCKL native vector types

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

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