Revert "SWDEV-362611 - Added hmax and hmin"

This reverts commit 84bd5f3. Reason for revert: This patch breaks release build #34 in rocprim. Change-Id: Iadba082e987c2e7564cd8f5a3aba49d4eb3ea250
ROCm · Nov 4, 2022 · dbcca1d · dbcca1d
1 parent 84bd5f3
commit dbcca1d
Show file tree

Hide file tree

Showing 2 changed files with 1 addition and 39 deletions.
diff --git a/include/hip/amd_detail/amd_hip_fp16.h b/include/hip/amd_detail/amd_hip_fp16.h
@@ -1250,38 +1250,6 @@ 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

diff --git a/include/hip/amd_detail/hip_fp16_math_fwd.h b/include/hip/amd_detail/hip_fp16_math_fwd.h
@@ -25,6 +25,7 @@ THE SOFTWARE.
 // /*
 // Half Math Functions
 // */
+
 #include "host_defines.h"
 #if !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__
 extern "C"
@@ -50,8 +51,6 @@ 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)));
@@ -85,8 +84,3 @@ 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);
-}