Partially revert "clang/HIP: Remove __llvm_amdgcn_* wrapper hacks"

Revert part of f407a73.

Some of the HIP headers were using the f16 rcp inline, such that it
breaks compiling code against non-top-of-tree headers. Need to wait
for a few HIP releases to expire to fully remove these.

Fixes #63981

Author: arsenm

clang/lib/Headers/__clang_hip_libdevice_declares.h

@@ -10,6 +10,10 @@
 #ifndef __CLANG_HIP_LIBDEVICE_DECLARES_H__
 #define __CLANG_HIP_LIBDEVICE_DECLARES_H__
 
+#if !defined(__HIPCC_RTC__) && __has_include("hip/hip_version.h")
+#include "hip/hip_version.h"
+#endif // __has_include("hip/hip_version.h")
+
 #ifdef __cplusplus
 extern "C" {
 #endif
@@ -312,6 +316,29 @@ __device__ __attribute__((const)) __2i16 __ocml_isnan_2f16(__2f16);
 __device__ __attribute__((pure)) __2f16 __ocml_log_2f16(__2f16);
 __device__ __attribute__((pure)) __2f16 __ocml_log10_2f16(__2f16);
 __device__ __attribute__((pure)) __2f16 __ocml_log2_2f16(__2f16);
+
+#if HIP_VERSION_MAJOR * 100 + HIP_VERSION_MINOR >= 560 || 1
+#define __DEPRECATED_SINCE_HIP_560(X) __attribute__((deprecated(X)))
+#else
+#define __DEPRECATED_SINCE_HIP_560(X)
+#endif
+
+// Deprecated, should be removed when rocm releases using it are no longer
+// relevant.
+__DEPRECATED_SINCE_HIP_560("use ((_Float16)1.0) / ")
+__device__ inline _Float16 __llvm_amdgcn_rcp_f16(_Float16 x) {
+  return ((_Float16)1.0f) / x;
+}
+
+__DEPRECATED_SINCE_HIP_560("use ((__2f16)1.0) / ")
+__device__ inline __2f16
+__llvm_amdgcn_rcp_2f16(__2f16 __x)
+{
+  return ((__2f16)1.0f) / __x;
+}
+
+#undef __DEPRECATED_SINCE_HIP_560
+
 __device__ __attribute__((const)) __2f16 __ocml_rint_2f16(__2f16);
 __device__ __attribute__((const)) __2f16 __ocml_rsqrt_2f16(__2f16);
 __device__ __2f16 __ocml_sin_2f16(__2f16);

clang/test/Headers/__clang_hip_math_deprecated.hip

@@ -0,0 +1,29 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
+// REQUIRES: amdgpu-registered-target
+
+// RUN: %clang_cc1 -include __clang_hip_runtime_wrapper.h \
+// RUN:   -internal-isystem %S/../../lib/Headers/cuda_wrappers \
+// RUN:   -internal-isystem %S/Inputs/include \
+// RUN:   -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \
+// RUN:   -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -O1 -o - \
+// RUN:   -D__HIPCC_RTC__ | FileCheck %s
+
+// Test deprecated functions in the header that should be removed eventually
+
+// CHECK-LABEL: @test_rcpf16_wrapper(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[DIV_I:%.*]] = fdiv contract half 0xH3C00, [[X:%.*]]
+// CHECK-NEXT:    ret half [[DIV_I]]
+//
+extern "C" __device__ _Float16 test_rcpf16_wrapper(_Float16 x) {
+  return __llvm_amdgcn_rcp_f16(x);
+}
+
+// CHECK-LABEL: @test_rcp2f16_wrapper(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[DIV_I:%.*]] = fdiv contract <2 x half> <half 0xH3C00, half 0xH3C00>, [[X:%.*]]
+// CHECK-NEXT:    ret <2 x half> [[DIV_I]]
+//
+extern "C" __device__ __2f16 test_rcp2f16_wrapper(__2f16 x) {
+  return __llvm_amdgcn_rcp_2f16(x);
+}