https://github.com/arsenm created https://github.com/llvm/llvm-project/pull/171123
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. >From 6bc0d37d008fcfb2d17b61e8250348fa65c1caaf Mon Sep 17 00:00:00 2001 From: Matt Arsenault <[email protected]> Date: Mon, 8 Dec 2025 14:36:33 +0100 Subject: [PATCH] clang/HIP: Remove deprecated rcp pseudo-intrinsics 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. --- .../Headers/__clang_hip_libdevice_declares.h | 22 -------------- .../Headers/__clang_hip_math_deprecated.hip | 29 ------------------- 2 files changed, 51 deletions(-) delete mode 100644 clang/test/Headers/__clang_hip_math_deprecated.hip 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); -} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
