clang/HIP: Remove deprecated rcp pseudo-intrinsics

[arsenm]

These shouldn't have been used by external users in the first place,
but have also been marked as deprecated for a number of releases.

[arsenm]

• clang/HIP: Remove deprecated rcp pseudo-intrinsics #171123 👈 (View in Graphite)
• main

This stack of pull requests is managed by Graphite. Learn more about stacking.

[llvmbot]

@llvm/pr-subscribers-backend-x86

@llvm/pr-subscribers-clang

Author: Matt Arsenault (arsenm)

Changes

These shouldn't have been used by external users in the first place,
but have also been marked as deprecated for a number of releases.

---

Full diff: https://github.com/llvm/llvm-project/pull/171123.diff

2 Files Affected:

• (modified) clang/lib/Headers/__clang_hip_libdevice_declares.h (-22)
• (removed) clang/test/Headers/__clang_hip_math_deprecated.hip (-29)

diff --git a/clang/lib/Headers/__clang_hip_libdevice_declares.h b/clang/lib/Headers/__clang_hip_libdevice_declares.h
index fad9c6ca7ffc5..565c68334023a 100644
--- a/clang/lib/Headers/__clang_hip_libdevice_declares.h
+++ b/clang/lib/Headers/__clang_hip_libdevice_declares.h
@@ -309,28 +309,6 @@ __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(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);
diff --git a/clang/test/Headers/__clang_hip_math_deprecated.hip b/clang/test/Headers/__clang_hip_math_deprecated.hip
deleted file mode 100644
index caba3e9ad83d1..0000000000000
--- a/clang/test/Headers/__clang_hip_math_deprecated.hip
+++ /dev/null
@@ -1,29 +0,0 @@
-// 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> splat (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);
-}