[PATCH] D152312: HIP: Use frexp builtins in math headers

[Matt Arsenault via Phabricator via cfe-commits]

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

https://reviews.llvm.org/D152312

Files:
  clang/lib/Headers/__clang_hip_math.h
  clang/test/Headers/__clang_hip_math.hip


Index: clang/test/Headers/__clang_hip_math.hip
===================================================================
--- clang/test/Headers/__clang_hip_math.hip
+++ clang/test/Headers/__clang_hip_math.hip
@@ -1061,37 +1061,25 @@
   return fmod(x, y);
 }
 
-// DEFAULT-LABEL: @test_frexpf(
-// DEFAULT-NEXT:  entry:
-// DEFAULT-NEXT:    [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.frexp.exp.i32.f32(float [[X:%.*]])
-// DEFAULT-NEXT:    store i32 [[TMP0]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA11:![0-9]+]]
-// DEFAULT-NEXT:    [[TMP1:%.*]] = tail call contract float @llvm.amdgcn.frexp.mant.f32(float [[X]])
-// DEFAULT-NEXT:    ret float [[TMP1]]
-//
-// FINITEONLY-LABEL: @test_frexpf(
-// FINITEONLY-NEXT:  entry:
-// FINITEONLY-NEXT:    [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.frexp.exp.i32.f32(float [[X:%.*]])
-// FINITEONLY-NEXT:    store i32 [[TMP0]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA11:![0-9]+]]
-// FINITEONLY-NEXT:    [[TMP1:%.*]] = tail call nnan ninf contract float @llvm.amdgcn.frexp.mant.f32(float [[X]])
-// FINITEONLY-NEXT:    ret float [[TMP1]]
+// CHECK-LABEL: @test_frexpf(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[TMP0:%.*]] = tail call { float, i32 } @llvm.frexp.f32.i32(float [[X:%.*]])
+// CHECK-NEXT:    [[TMP1:%.*]] = extractvalue { float, i32 } [[TMP0]], 1
+// CHECK-NEXT:    store i32 [[TMP1]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA11:![0-9]+]]
+// CHECK-NEXT:    [[TMP2:%.*]] = extractvalue { float, i32 } [[TMP0]], 0
+// CHECK-NEXT:    ret float [[TMP2]]
 //
 extern "C" __device__ float test_frexpf(float x, int* y) {
   return frexpf(x, y);
 }
 
-// DEFAULT-LABEL: @test_frexp(
-// DEFAULT-NEXT:  entry:
-// DEFAULT-NEXT:    [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.frexp.exp.i32.f64(double [[X:%.*]])
-// DEFAULT-NEXT:    store i32 [[TMP0]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA11]]
-// DEFAULT-NEXT:    [[TMP1:%.*]] = tail call contract double @llvm.amdgcn.frexp.mant.f64(double [[X]])
-// DEFAULT-NEXT:    ret double [[TMP1]]
-//
-// FINITEONLY-LABEL: @test_frexp(
-// FINITEONLY-NEXT:  entry:
-// FINITEONLY-NEXT:    [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.frexp.exp.i32.f64(double [[X:%.*]])
-// FINITEONLY-NEXT:    store i32 [[TMP0]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA11]]
-// FINITEONLY-NEXT:    [[TMP1:%.*]] = tail call nnan ninf contract double @llvm.amdgcn.frexp.mant.f64(double [[X]])
-// FINITEONLY-NEXT:    ret double [[TMP1]]
+// CHECK-LABEL: @test_frexp(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[TMP0:%.*]] = tail call { double, i32 } @llvm.frexp.f64.i32(double [[X:%.*]])
+// CHECK-NEXT:    [[TMP1:%.*]] = extractvalue { double, i32 } [[TMP0]], 1
+// CHECK-NEXT:    store i32 [[TMP1]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA11]]
+// CHECK-NEXT:    [[TMP2:%.*]] = extractvalue { double, i32 } [[TMP0]], 0
+// CHECK-NEXT:    ret double [[TMP2]]
 //
 extern "C" __device__ double test_frexp(double x, int* y) {
   return frexp(x, y);
Index: clang/lib/Headers/__clang_hip_math.h
===================================================================
--- clang/lib/Headers/__clang_hip_math.h
+++ clang/lib/Headers/__clang_hip_math.h
@@ -257,8 +257,7 @@
 
 __DEVICE__
 float frexpf(float __x, int *__nptr) {
-  *__nptr = __builtin_amdgcn_frexp_expf(__x);
-  return __builtin_amdgcn_frexp_mantf(__x);
+  return __builtin_frexpf(__x, __nptr);
 }
 
 __DEVICE__
@@ -806,8 +805,7 @@
 
 __DEVICE__
 double frexp(double __x, int *__nptr) {
-  *__nptr = __builtin_amdgcn_frexp_exp(__x);
-  return __builtin_amdgcn_frexp_mant(__x);
+  return __builtin_frexp(__x, __nptr);
 }
 
 __DEVICE__


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D152312.529058.patch
Type: text/x-patch
Size: 3531 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20230606/99e33a1e/attachment-0001.bin>