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

[Matt Arsenault via Phabricator via cfe-commits]

arsenm created this revision.
arsenm added a reviewer: yaxunl.
Herald added a project: All.
arsenm requested review of this revision.
Herald added a subscriber: wdng.

Revert part of f407a7399575a6821940973c54754d42e72dd9ce <https://reviews.llvm.org/rGf407a7399575a6821940973c54754d42e72dd9ce>.

      

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


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
@@ -279,6 +279,7 @@
 __device__ __attribute__((pure)) _Float16 __ocml_log_f16(_Float16);
 __device__ __attribute__((pure)) _Float16 __ocml_log10_f16(_Float16);
 __device__ __attribute__((pure)) _Float16 __ocml_log2_f16(_Float16);
+
 __device__ __attribute__((const)) _Float16 __ocml_rint_f16(_Float16);
 __device__ __attribute__((const)) _Float16 __ocml_rsqrt_f16(_Float16);
 __device__ _Float16 __ocml_sin_f16(_Float16);
@@ -312,6 +313,18 @@
 __device__ __attribute__((pure)) __2f16 __ocml_log_2f16(__2f16);
 __device__ __attribute__((pure)) __2f16 __ocml_log10_2f16(__2f16);
 __device__ __attribute__((pure)) __2f16 __ocml_log2_2f16(__2f16);
+
+// Deprecated, should be removed when rocm releases using it are no longer
+// relevant.
+__device__ inline _Float16 __llvm_amdgcn_rcp_f16(_Float16 x) {
+  return ((_Float16)1.0f) / x;
+}
+
+__device__ inline __2f16
+__llvm_amdgcn_rcp_2f16(__2f16 __x)
+{
+  return ((__2f16)1.0f) / __x;
+}
 __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.542986.patch
Type: text/x-patch
Size: 2732 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20230721/9513e745/attachment.bin>