Add backwards compatibility to dot funcs

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)

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);

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__

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__