[PATCH] D155982: Partially revert "clang/HIP: Remove __llvm_amdgcn_* wrapper hacks"

[Matt Arsenault via Phabricator via cfe-commits]

arsenm updated this revision to Diff 543111.
arsenm marked an inline comment as done.

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D155982/new/

https://reviews.llvm.org/D155982

Files:
  clang/lib/Headers/__clang_hip_libdevice_declares.h
  clang/test/Headers/__clang_hip_math_deprecated.hip


Index: clang/test/Headers/__clang_hip_math_deprecated.hip
===================================================================
--- /dev/null
+++ 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);
+}
Index: clang/lib/Headers/__clang_hip_libdevice_declares.h
===================================================================
--- clang/lib/Headers/__clang_hip_libdevice_declares.h
+++ 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__((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);


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D155982.543111.patch
Type: text/x-patch
Size: 2938 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20230721/57a7ab31/attachment.bin>