Author: Matt Arsenault Date: 2023-06-08T15:25:01-04:00 New Revision: d14ac1d11a5cb3994c63ecaaa5c950636a289fa1
URL: https://github.com/llvm/llvm-project/commit/d14ac1d11a5cb3994c63ecaaa5c950636a289fa1 DIFF: https://github.com/llvm/llvm-project/commit/d14ac1d11a5cb3994c63ecaaa5c950636a289fa1.diff LOG: HIP: Directly call isinf 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 46c44f29dec57..b92d7b74f25cf 100644 --- a/clang/lib/Headers/__clang_hip_math.h +++ b/clang/lib/Headers/__clang_hip_math.h @@ -271,7 +271,7 @@ __DEVICE__ __RETURN_TYPE __finitef(float __x) { return __ocml_isfinite_f32(__x); } __DEVICE__ -__RETURN_TYPE __isinff(float __x) { return __ocml_isinf_f32(__x); } +__RETURN_TYPE __isinff(float __x) { return __builtin_isinf(__x); } __DEVICE__ __RETURN_TYPE __isnanf(float __x) { return __ocml_isnan_f32(__x); } @@ -820,7 +820,7 @@ __DEVICE__ __RETURN_TYPE __finite(double __x) { return __ocml_isfinite_f64(__x); } __DEVICE__ -__RETURN_TYPE __isinf(double __x) { return __ocml_isinf_f64(__x); } +__RETURN_TYPE __isinf(double __x) { return __builtin_isinf(__x); } __DEVICE__ __RETURN_TYPE __isnan(double __x) { return __ocml_isnan_f64(__x); } diff --git a/clang/test/Headers/__clang_hip_math.hip b/clang/test/Headers/__clang_hip_math.hip index af829076e0e19..d099691fcfd8f 100644 --- a/clang/test/Headers/__clang_hip_math.hip +++ b/clang/test/Headers/__clang_hip_math.hip @@ -1191,17 +1191,14 @@ extern "C" __device__ BOOL_TYPE test___finite(double x) { // DEFAULT-LABEL: @test___isinff( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isinf_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: [[TMP0:%.*]] = tail call contract float @llvm.fabs.f32(float [[X:%.*]]) #[[ATTR17:[0-9]+]] +// DEFAULT-NEXT: [[CMPINF_I:%.*]] = fcmp contract oeq float [[TMP0]], 0x7FF0000000000000 +// DEFAULT-NEXT: [[CONV:%.*]] = zext i1 [[CMPINF_I]] to i32 // DEFAULT-NEXT: ret i32 [[CONV]] // // FINITEONLY-LABEL: @test___isinff( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isinf_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]] +// FINITEONLY-NEXT: ret i32 0 // extern "C" __device__ BOOL_TYPE test___isinff(float x) { return __isinff(x); @@ -1209,17 +1206,14 @@ extern "C" __device__ BOOL_TYPE test___isinff(float x) { // DEFAULT-LABEL: @test___isinf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isinf_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: [[TMP0:%.*]] = tail call contract double @llvm.fabs.f64(double [[X:%.*]]) #[[ATTR17]] +// DEFAULT-NEXT: [[CMPINF_I:%.*]] = fcmp contract oeq double [[TMP0]], 0x7FF0000000000000 +// DEFAULT-NEXT: [[CONV:%.*]] = zext i1 [[CMPINF_I]] to i32 // DEFAULT-NEXT: ret i32 [[CONV]] // // FINITEONLY-LABEL: @test___isinf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isinf_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]] +// FINITEONLY-NEXT: ret i32 0 // extern "C" __device__ BOOL_TYPE test___isinf(double x) { return __isinf(x); @@ -1756,11 +1750,11 @@ extern "C" __device__ long int test_lround(double x) { // DEFAULT-LABEL: @test_modff( // DEFAULT-NEXT: entry: // DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5) -// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17:[0-9]+]] +// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18:[0-9]+]] // DEFAULT-NEXT: [[CALL_I:%.*]] = call contract float @__ocml_modf_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]] // DEFAULT-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA15:![0-9]+]] // DEFAULT-NEXT: store float [[TMP0]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA15]] -// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_modff( @@ -1780,11 +1774,11 @@ extern "C" __device__ float test_modff(float x, float* y) { // DEFAULT-LABEL: @test_modf( // DEFAULT-NEXT: entry: // DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca double, align 8, addrspace(5) -// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]] // DEFAULT-NEXT: [[CALL_I:%.*]] = call contract double @__ocml_modf_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]] // DEFAULT-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA17:![0-9]+]] // DEFAULT-NEXT: store double [[TMP0]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA17]] -// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_modf( @@ -2401,11 +2395,11 @@ extern "C" __device__ double test_remainder(double x, double y) { // DEFAULT-LABEL: @test_remquof( // DEFAULT-NEXT: entry: // DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca i32, align 4, addrspace(5) -// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]] // DEFAULT-NEXT: [[CALL_I:%.*]] = call contract float @__ocml_remquo_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]] // DEFAULT-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA11]] // DEFAULT-NEXT: store i32 [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA11]] -// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_remquof( @@ -2425,11 +2419,11 @@ extern "C" __device__ float test_remquof(float x, float y, int* z) { // DEFAULT-LABEL: @test_remquo( // DEFAULT-NEXT: entry: // DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca i32, align 4, addrspace(5) -// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]] // DEFAULT-NEXT: [[CALL_I:%.*]] = call contract double @__ocml_remquo_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]] // DEFAULT-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA11]] // DEFAULT-NEXT: store i32 [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA11]] -// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_remquo( @@ -2822,12 +2816,12 @@ extern "C" __device__ BOOL_TYPE test___signbit(double x) { // DEFAULT-LABEL: @test_sincosf( // DEFAULT-NEXT: entry: // DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5) -// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]] // DEFAULT-NEXT: [[CALL_I:%.*]] = call contract float @__ocml_sincos_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]] // DEFAULT-NEXT: store float [[CALL_I]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA15]] // DEFAULT-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA15]] // DEFAULT-NEXT: store float [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA15]] -// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]] // DEFAULT-NEXT: ret void // // FINITEONLY-LABEL: @test_sincosf( @@ -2848,12 +2842,12 @@ extern "C" __device__ void test_sincosf(float x, float *y, float *z) { // DEFAULT-LABEL: @test_sincos( // DEFAULT-NEXT: entry: // DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca double, align 8, addrspace(5) -// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]] // DEFAULT-NEXT: [[CALL_I:%.*]] = call contract double @__ocml_sincos_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]] // DEFAULT-NEXT: store double [[CALL_I]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA17]] // DEFAULT-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA17]] // DEFAULT-NEXT: store double [[TMP0]], ptr [[Z:%.*]], align 8, !tbaa [[TBAA17]] -// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]] // DEFAULT-NEXT: ret void // // FINITEONLY-LABEL: @test_sincos( @@ -2874,12 +2868,12 @@ extern "C" __device__ void test_sincos(double x, double *y, double *z) { // DEFAULT-LABEL: @test_sincospif( // DEFAULT-NEXT: entry: // DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5) -// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]] // DEFAULT-NEXT: [[CALL_I:%.*]] = call contract float @__ocml_sincospi_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]] // DEFAULT-NEXT: store float [[CALL_I]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA15]] // DEFAULT-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA15]] // DEFAULT-NEXT: store float [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA15]] -// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]] // DEFAULT-NEXT: ret void // // FINITEONLY-LABEL: @test_sincospif( @@ -2900,12 +2894,12 @@ extern "C" __device__ void test_sincospif(float x, float *y, float *z) { // DEFAULT-LABEL: @test_sincospi( // DEFAULT-NEXT: entry: // DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca double, align 8, addrspace(5) -// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]] // DEFAULT-NEXT: [[CALL_I:%.*]] = call contract double @__ocml_sincospi_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]] // DEFAULT-NEXT: store double [[CALL_I]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA17]] // DEFAULT-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA17]] // DEFAULT-NEXT: store double [[TMP0]], ptr [[Z:%.*]], align 8, !tbaa [[TBAA17]] -// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]] // DEFAULT-NEXT: ret void // // FINITEONLY-LABEL: @test_sincospi( _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
