SWDEV-362611 - Added hmax and hmin

Change-Id: I133a80e997e39357693df7ab969425c0d5585607

hipamd/include/hip/amd_detail/amd_hip_fp16.h

@@ -1250,6 +1250,38 @@ THE SOFTWARE.
             inline
             __HOST_DEVICE__
             bool __hbgtu2(__half2 x, __half2 y) { return __hbgt2(x, y); }
+            inline
+            __device__
+            __half __hmax(const __half x, const __half y) {
+              return __half_raw{__ocml_fmax_f16(static_cast<__half_raw>(x).data,
+                                   static_cast<__half_raw>(y).data)};
+            }
+            inline
+            __device__
+            __half __hmax_nan(const __half x, const __half y) {
+                if(__ocml_isnan_f16(x)) {
+                  return x;
+                } else if (__ocml_isnan_f16(y)) {
+                  return y;
+                }
+                return __hmax(x, y);
+            }
+            inline
+            __device__
+            __half __hmin(const __half x, const __half y) {
+              return __half_raw{__ocml_fmin_f16(static_cast<__half_raw>(x).data,
+                                   static_cast<__half_raw>(y).data)};
+            }
+            inline
+            __device__
+            __half __hmin_nan(const __half x, const __half y) {
+                if(__ocml_isnan_f16(x)) {
+                  return x;
+                } else if (__ocml_isnan_f16(y)) {
+                  return y;
+                }
+                return __hmin(x, y);
+            }
 
             // Arithmetic
             inline

hipamd/include/hip/amd_detail/hip_fp16_math_fwd.h

@@ -25,7 +25,6 @@ THE SOFTWARE.
 // /*
 // Half Math Functions
 // */
-
 #include "host_defines.h"
 #if !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__
 extern "C"
@@ -51,6 +50,8 @@ extern "C"
     __device__ _Float16 __ocml_sin_f16(_Float16);
     __device__ __attribute__((const)) _Float16 __ocml_sqrt_f16(_Float16);
     __device__ __attribute__((const)) _Float16 __ocml_trunc_f16(_Float16);
+    __device__ __attribute__((const)) _Float16 __ocml_fmax_f16(_Float16, _Float16);
+    __device__ __attribute__((const)) _Float16 __ocml_fmin_f16(_Float16, _Float16);
 
     typedef _Float16 __2f16 __attribute__((ext_vector_type(2)));
     typedef short __2i16 __attribute__((ext_vector_type(2)));
@@ -84,3 +85,8 @@ extern "C"
     __device__ __attribute__((const)) __2f16 __ocml_trunc_2f16(__2f16);
 }
 #endif // !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__
+//TODO: remove these after they get into clang header __clang_hip_libdevice_declares.h'
+extern "C" {
+    __device__ __attribute__((const)) _Float16 __ocml_fmax_f16(_Float16, _Float16);
+    __device__ __attribute__((const)) _Float16 __ocml_fmin_f16(_Float16, _Float16);
+}