From dbcca1d7a796a47cd5c78574d7cb186cb39f03c4 Mon Sep 17 00:00:00 2001 From: Zhongyu Zhang Date: Fri, 4 Nov 2022 02:01:11 -0400 Subject: [PATCH] Revert "SWDEV-362611 - Added hmax and hmin" This reverts commit 84bd5f3af44d91d7b2014ec1d5f83f6e60d45fa5. Reason for revert: This patch breaks release build #34 in rocprim. Change-Id: Iadba082e987c2e7564cd8f5a3aba49d4eb3ea250 --- include/hip/amd_detail/amd_hip_fp16.h | 32 ---------------------- include/hip/amd_detail/hip_fp16_math_fwd.h | 8 +----- 2 files changed, 1 insertion(+), 39 deletions(-) diff --git a/include/hip/amd_detail/amd_hip_fp16.h b/include/hip/amd_detail/amd_hip_fp16.h index 9a72bf1b..78010364 100644 --- 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 index 8b35a4ba..a5b8342d 100644 --- 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); -}