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

[arsenm]

Summary

Revert part of f407a7399575a6821940973c54754d42e72dd9ce.

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

[yaxunl]

clang/lib/Headers/__clang_hip_libdevice_declares.h

322

Can we add the deprecated attribute to urge people not to use them?

327

same as above

[arsenm]

clang/lib/Headers/__clang_hip_libdevice_declares.h

322

I initially added those, but thought it was a bit aggressive to put in always included headers. Is there an established practice for deprecating builtin header functions?

[arsenm]

clang/lib/Headers/__clang_hip_libdevice_declares.h

322

Is there a hip header version macro I could guard this on?

[yaxunl]

clang/lib/Headers/__clang_hip_libdevice_declares.h

322

I initially added those, but thought it was a bit aggressive to put in always included headers. Is there an established practice for deprecating builtin header functions?

I saw it used by libstdc++ headers and pthread headers to deprecate APIs.

I think it is common practice to use deprecate attibute.

We'd also better mention it in clang release note and give a timeline to remove it, e.g. next clang major release.

322

Is there a hip header version macro I could guard this on?

Yes like https://github.com/llvm/llvm-project/blob/main/clang/lib/Headers/__clang_hip_runtime_wrapper.h#L80

[arsenm]

Add versioned deprecated macro

[yaxunl]

clang/lib/Headers/__clang_hip_libdevice_declares.h

13–15

pls condition this with

#if !defined(__HIPCC_RTC__)
#endif

Since for hiprtc we concatenate header files linearly into one header file.

321

How about providing a replacement hint:

#define __DEPRECATED_SINCE_HIP_560(replacement) __attribute__((deprecated("replaced by clang builtin", #replacement)))

[yaxunl]

LGTM. Thanks.

[arsenm]

9b2dfff57a382b757c358b43ee1df7591cb480ee

clang/lib/Headers/__clang_hip_libdevice_declares.h

 /*===---- __clang_hip_libdevice_declares.h - HIP device library decls -------===
  *
  * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
  * See https://llvm.org/LICENSE.txt for license information.
  * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
  *
  *===-----------------------------------------------------------------------===
  */
 
 #ifndef __CLANG_HIP_LIBDEVICE_DECLARES_H__
 #define __CLANG_HIP_LIBDEVICE_DECLARES_H__
 
+#if __has_include("hip/hip_version.h")
+#include "hip/hip_version.h"
+#endif // __has_include("hip/hip_version.h")
+
 #ifdef __cplusplus
 extern "C" {
 #endif
 
 // BEGIN FLOAT
 __device__ __attribute__((const)) float __ocml_acos_f32(float);
 __device__ __attribute__((pure)) float __ocml_acosh_f32(float);
 __device__ __attribute__((const)) float __ocml_asin_f32(float);

 __device__ __attribute__((const)) __2f16 __ocml_floor_2f16(__2f16);
 __device__ __attribute__((const))
 __2f16 __ocml_fma_2f16(__2f16, __2f16, __2f16);
 __device__ __attribute__((const)) __2i16 __ocml_isinf_2f16(__2f16);
 __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
+#define __DEPRECATED_SINCE_HIP_560 __attribute__((deprecated))
+#else
+#define __DEPRECATED_SINCE_HIP_560
+#endif
+
+// Deprecated, should be removed when rocm releases using it are no longer
+// relevant.
+__DEPRECATED_SINCE_HIP_560
+__device__ inline _Float16 __llvm_amdgcn_rcp_f16(_Float16 x) {
+  return ((_Float16)1.0f) / x;
+}
+
+__DEPRECATED_SINCE_HIP_560
+__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);
 __device__ __attribute__((const)) __2f16 __ocml_sqrt_2f16(__2f16);
 __device__ __attribute__((const)) __2f16 __ocml_trunc_2f16(__2f16);
 __device__ __attribute__((const)) __2f16 __ocml_pown_2f16(__2f16, __2i16);
 
 #ifdef __cplusplus
 } // extern "C"
 #endif
 
 #endif // __CLANG_HIP_LIBDEVICE_DECLARES_H__

clang/test/Headers/__clang_hip_math_deprecated.hip

+// 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);
+}