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);