[clang] b51ae6d - HIP: Directly call signbit builtins

[Matt Arsenault via cfe-commits]

Author: Matt Arsenault
Date: 2023-06-08T15:22:43-04:00
New Revision: b51ae6d31f86de2e47b1912eb045ea37a0c5de23

URL: https://github.com/llvm/llvm-project/commit/b51ae6d31f86de2e47b1912eb045ea37a0c5de23
DIFF: https://github.com/llvm/llvm-project/commit/b51ae6d31f86de2e47b1912eb045ea37a0c5de23.diff

LOG: HIP: Directly call signbit builtins

Added: 
    

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

Removed: 
    


################################################################################
diff  --git a/clang/lib/Headers/__clang_hip_math.h b/clang/lib/Headers/__clang_hip_math.h
index 7bfb98fba759f..46c44f29dec57 100644
--- a/clang/lib/Headers/__clang_hip_math.h
+++ b/clang/lib/Headers/__clang_hip_math.h
@@ -476,7 +476,7 @@ __DEVICE__
 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) {
@@ -1032,7 +1032,7 @@ __DEVICE__
 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); }

diff  --git a/clang/test/Headers/__clang_hip_math.hip b/clang/test/Headers/__clang_hip_math.hip
index 29ccfd9c09b2e..af829076e0e19 100644
--- a/clang/test/Headers/__clang_hip_math.hip
+++ b/clang/test/Headers/__clang_hip_math.hip
@@ -2798,37 +2798,22 @@ extern "C" __device__ double test_scalbn(double x, int y) {
   return scalbn(x, y);
 }
 
-// DEFAULT-LABEL: @test___signbitf(
-// DEFAULT-NEXT:  entry:
-// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call i32 @__ocml_signbit_f32(float noundef [[X:%.*]]) #[[ATTR14]]
-// DEFAULT-NEXT:    [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0
-// DEFAULT-NEXT:    [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32
-// DEFAULT-NEXT:    ret i32 [[CONV]]
-//
-// FINITEONLY-LABEL: @test___signbitf(
-// FINITEONLY-NEXT:  entry:
-// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call i32 @__ocml_signbit_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]]
-// FINITEONLY-NEXT:    [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0
-// FINITEONLY-NEXT:    [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32
-// FINITEONLY-NEXT:    ret i32 [[CONV]]
+// CHECK-LABEL: @test___signbitf(
+// CHECK-NEXT:  entry:
+// 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);
 }
 
-// DEFAULT-LABEL: @test___signbit(
-// DEFAULT-NEXT:  entry:
-// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call i32 @__ocml_signbit_f64(double noundef [[X:%.*]]) #[[ATTR14]]
-// DEFAULT-NEXT:    [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0
-// DEFAULT-NEXT:    [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32
-// DEFAULT-NEXT:    ret i32 [[CONV]]
-//
-// FINITEONLY-LABEL: @test___signbit(
-// FINITEONLY-NEXT:  entry:
-// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call i32 @__ocml_signbit_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]]
-// FINITEONLY-NEXT:    [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0
-// FINITEONLY-NEXT:    [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32
-// FINITEONLY-NEXT:    ret i32 [[CONV]]
+// CHECK-LABEL: @test___signbit(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[TMP0:%.*]] = bitcast double [[X:%.*]] to i64
+// CHECK-NEXT:    [[DOTLOBIT:%.*]] = lshr i64 [[TMP0]], 63
+// CHECK-NEXT:    [[CONV:%.*]] = trunc i64 [[DOTLOBIT]] to i32
+// CHECK-NEXT:    ret i32 [[CONV]]
 //
 extern "C" __device__ BOOL_TYPE test___signbit(double x) {
   return __signbit(x);