[PATCH] D138396: HIP: Directly call signbit builtins

[Matt Arsenault via Phabricator via cfe-commits]

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

https://reviews.llvm.org/D138396

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
@@ -2724,10 +2724,9 @@
 
 // CHECK-LABEL: @test___signbitf(
 // CHECK-NEXT:  entry:
-// CHECK-NEXT:    [[CALL_I:%.*]] = tail call i32 @__ocml_signbit_f32(float noundef [[X:%.*]]) #[[ATTR15]]
-// CHECK-NEXT:    [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0
-// CHECK-NEXT:    [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32
-// CHECK-NEXT:    ret i32 [[CONV]]
+// CHECK-NEXT:    [[TMP0:%.*]] = bitcast float [[X:%.*]] to i32
+// CHECK-NEXT:    [[DOTLOBIT:%.*]] = lshr i32 [[TMP0]], 31
+// CHECK-NEXT:    ret i32 [[DOTLOBIT]]
 //
 extern "C" __device__ BOOL_TYPE test___signbitf(float x) {
   return __signbitf(x);
@@ -2735,10 +2734,10 @@
 
 // CHECK-LABEL: @test___signbit(
 // CHECK-NEXT:  entry:
-// CHECK-NEXT:    [[CALL_I:%.*]] = tail call i32 @__ocml_signbit_f64(double noundef [[X:%.*]]) #[[ATTR15]]
-// CHECK-NEXT:    [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0
-// CHECK-NEXT:    [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32
-// CHECK-NEXT:    ret i32 [[CONV]]
+// CHECK-NEXT:    [[TMP0:%.*]] = bitcast double [[X:%.*]] to i64
+// CHECK-NEXT:    [[DOTLOBIT:%.*]] = lshr i64 [[TMP0]], 63
+// CHECK-NEXT:    [[TMP1:%.*]] = trunc i64 [[DOTLOBIT]] to i32
+// CHECK-NEXT:    ret i32 [[TMP1]]
 //
 extern "C" __device__ BOOL_TYPE test___signbit(double x) {
   return __signbit(x);
Index: clang/lib/Headers/__clang_hip_math.h
===================================================================
--- clang/lib/Headers/__clang_hip_math.h
+++ clang/lib/Headers/__clang_hip_math.h
@@ -495,7 +495,7 @@
 float scalbnf(float __x, int __n) { return __ocml_scalbn_f32(__x, __n); }
 
 __DEVICE__
-__RETURN_TYPE __signbitf(float __x) { return __ocml_signbit_f32(__x); }
+__RETURN_TYPE __signbitf(float __x) { return __builtin_signbitf(__x); }
 
 __DEVICE__
 void sincosf(float __x, float *__sinptr, float *__cosptr) {
@@ -1057,7 +1057,7 @@
 double scalbn(double __x, int __n) { return __ocml_scalbn_f64(__x, __n); }
 
 __DEVICE__
-__RETURN_TYPE __signbit(double __x) { return __ocml_signbit_f64(__x); }
+__RETURN_TYPE __signbit(double __x) { return __builtin_signbit(__x); }
 
 __DEVICE__
 double sin(double __x) { return __ocml_sin_f64(__x); }


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D138396.476801.patch
Type: text/x-patch
Size: 2329 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20221121/9eb552f8/attachment-0001.bin>