This is an archive of the discontinued LLVM Phabricator instance.

Differential D138399

HIP: Directly call isinf builtins
ClosedPublic

Authored by arsenm on Nov 20 2022, 9:02 PM.

Details

Reviewers
b-sumner
JonChesterfield
yaxunl

Diff Detail

Event Timeline

arsenm created this revision.Nov 20 2022, 9:02 PM
Herald added a project: Restricted Project. · View Herald TranscriptNov 20 2022, 9:02 PM
arsenm requested review of this revision.Nov 20 2022, 9:02 PM
Herald added a subscriber: wdng. · View Herald TranscriptNov 20 2022, 9:02 PM
arsenm added a parent revision: D138391: clang/HIP: Add new header test for math IR gen.Nov 20 2022, 9:03 PM
Harbormaster completed remote builds in B198720: Diff 476803.Nov 21 2022, 9:12 AM
arsenm added a comment.Jun 6 2023, 3:04 PM

ping

yaxunl accepted this revision.Jun 7 2023, 5:54 AM

LGTM. Thanks

This revision is now accepted and ready to land.Jun 7 2023, 5:54 AM
arsenm closed this revision.Jun 8 2023, 12:27 PM

d14ac1d11a5cb3994c63ecaaa5c950636a289fa1

Revision Contents

Diff 476803

clang/lib/Headers/__clang_hip_math.h

Show First 20 Lines • Show All 284 Lines • ▼ Show 20 Lines
__DEVICE__ __DEVICE__
int ilogbf(float __x) { return __ocml_ilogb_f32(__x); } int ilogbf(float __x) { return __ocml_ilogb_f32(__x); }
__DEVICE__ __DEVICE__
__RETURN_TYPE __finitef(float __x) { return __ocml_isfinite_f32(__x); } __RETURN_TYPE __finitef(float __x) { return __ocml_isfinite_f32(__x); }
__DEVICE__ __DEVICE__
__RETURN_TYPE __isinff(float __x) { return __ocml_isinf_f32(__x); } __RETURN_TYPE __isinff(float __x) { return __builtin_isinf(__x); }
__DEVICE__ __DEVICE__
__RETURN_TYPE __isnanf(float __x) { return __ocml_isnan_f32(__x); } __RETURN_TYPE __isnanf(float __x) { return __ocml_isnan_f32(__x); }
__DEVICE__ __DEVICE__
float j0f(float __x) { return __ocml_j0_f32(__x); } float j0f(float __x) { return __ocml_j0_f32(__x); }
__DEVICE__ __DEVICE__
▲ Show 20 Lines • Show All 538 Lines • ▼ Show 20 Lines
__DEVICE__ __DEVICE__
int ilogb(double __x) { return __ocml_ilogb_f64(__x); } int ilogb(double __x) { return __ocml_ilogb_f64(__x); }
__DEVICE__ __DEVICE__
__RETURN_TYPE __finite(double __x) { return __ocml_isfinite_f64(__x); } __RETURN_TYPE __finite(double __x) { return __ocml_isfinite_f64(__x); }
__DEVICE__ __DEVICE__
__RETURN_TYPE __isinf(double __x) { return __ocml_isinf_f64(__x); } __RETURN_TYPE __isinf(double __x) { return __builtin_isinf(__x); }
__DEVICE__ __DEVICE__
__RETURN_TYPE __isnan(double __x) { return __ocml_isnan_f64(__x); } __RETURN_TYPE __isnan(double __x) { return __ocml_isnan_f64(__x); }
__DEVICE__ __DEVICE__
double j0(double __x) { return __ocml_j0_f64(__x); } double j0(double __x) { return __ocml_j0_f64(__x); }
__DEVICE__ __DEVICE__
▲ Show 20 Lines • Show All 474 Lines • Show Last 20 Lines

clang/test/Headers/__clang_hip_math.hip

Show First 20 Lines • Show All 1,129 Lines • ▼ Show 20 Lines
// CHECK-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0 // CHECK-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0
// CHECK-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32 // CHECK-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32
// CHECK-NEXT: ret i32 [[CONV]] // CHECK-NEXT: ret i32 [[CONV]]
// //
extern "C" __device__ BOOL_TYPE test___finite(double x) { extern "C" __device__ BOOL_TYPE test___finite(double x) {
return __finite(x); return __finite(x);
} }
// CHECK-LABEL: @test___isinff( // DEFAULT-LABEL: @test___isinff(
// CHECK-NEXT: entry: // DEFAULT-NEXT: entry:
// CHECK-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isinf_f32(float noundef [[X:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.fabs.f32(float [[X:%.*]]) #[[ATTR18:[0-9]+]]
// CHECK-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0 // DEFAULT-NEXT: [[CMPINF_I:%.*]] = fcmp contract oeq float [[TMP0]], 0x7FF0000000000000
// CHECK-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32 // DEFAULT-NEXT: [[CONV:%.*]] = zext i1 [[CMPINF_I]] to i32
// CHECK-NEXT: ret i32 [[CONV]] // DEFAULT-NEXT: ret i32 [[CONV]]
//
// FINITEONLY-LABEL: @test___isinff(
// FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: ret i32 0
// //
extern "C" __device__ BOOL_TYPE test___isinff(float x) { extern "C" __device__ BOOL_TYPE test___isinff(float x) {
return __isinff(x); return __isinff(x);
} }
// CHECK-LABEL: @test___isinf( // DEFAULT-LABEL: @test___isinf(
// CHECK-NEXT: entry: // DEFAULT-NEXT: entry:
// CHECK-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isinf_f64(double noundef [[X:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.fabs.f64(double [[X:%.*]]) #[[ATTR18]]
// CHECK-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0 // DEFAULT-NEXT: [[CMPINF_I:%.*]] = fcmp contract oeq double [[TMP0]], 0x7FF0000000000000
// CHECK-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32 // DEFAULT-NEXT: [[CONV:%.*]] = zext i1 [[CMPINF_I]] to i32
// CHECK-NEXT: ret i32 [[CONV]] // DEFAULT-NEXT: ret i32 [[CONV]]
//
// FINITEONLY-LABEL: @test___isinf(
// FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: ret i32 0
// //
extern "C" __device__ BOOL_TYPE test___isinf(double x) { extern "C" __device__ BOOL_TYPE test___isinf(double x) {
return __isinf(x); return __isinf(x);
} }
// CHECK-LABEL: @test___isnanf( // CHECK-LABEL: @test___isnanf(
// CHECK-NEXT: entry: // CHECK-NEXT: entry:
// CHECK-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isnan_f32(float noundef [[X:%.*]]) #[[ATTR15]] // CHECK-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isnan_f32(float noundef [[X:%.*]]) #[[ATTR15]]
▲ Show 20 Lines • Show All 506 Lines • ▼ Show 20 Lines
// //
extern "C" __device__ long int test_lround(double x) { extern "C" __device__ long int test_lround(double x) {
return lround(x); return lround(x);
} }
// DEFAULT-LABEL: @test_modff( // DEFAULT-LABEL: @test_modff(
// DEFAULT-NEXT: entry: // DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5) // 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]]) #[[ATTR18:[0-9]+]] // DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR19:[0-9]+]]
// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract float @__ocml_modf_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]] // DEFAULT-NEXT: [[CALL_I:%.*]] = call contract float @__ocml_modf_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]]
// DEFAULT-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA13:![0-9]+]] // DEFAULT-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA13:![0-9]+]]
// DEFAULT-NEXT: store float [[TMP0]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA13]] // DEFAULT-NEXT: store float [[TMP0]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA13]]
// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]] // DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR19]]
// DEFAULT-NEXT: ret float [[CALL_I]] // DEFAULT-NEXT: ret float [[CALL_I]]
// //
// FINITEONLY-LABEL: @test_modff( // FINITEONLY-LABEL: @test_modff(
// FINITEONLY-NEXT: entry: // FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5) // FINITEONLY-NEXT: [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5)
// FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18:[0-9]+]] // FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18:[0-9]+]]
// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract float @__ocml_modf_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]] // FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract float @__ocml_modf_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]]
// FINITEONLY-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA13:![0-9]+]] // FINITEONLY-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA13:![0-9]+]]
// FINITEONLY-NEXT: store float [[TMP0]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA13]] // FINITEONLY-NEXT: store float [[TMP0]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA13]]
// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]] // FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
// FINITEONLY-NEXT: ret float [[CALL_I]] // FINITEONLY-NEXT: ret float [[CALL_I]]
// //
extern "C" __device__ float test_modff(float x, float* y) { extern "C" __device__ float test_modff(float x, float* y) {
return modff(x, y); return modff(x, y);
} }
// DEFAULT-LABEL: @test_modf( // DEFAULT-LABEL: @test_modf(
// DEFAULT-NEXT: entry: // DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca double, align 8, addrspace(5) // 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]]) #[[ATTR18]] // DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR19]]
// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract double @__ocml_modf_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]] // DEFAULT-NEXT: [[CALL_I:%.*]] = call contract double @__ocml_modf_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]]
// DEFAULT-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA15:![0-9]+]] // DEFAULT-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA15:![0-9]+]]
// DEFAULT-NEXT: store double [[TMP0]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA15]] // DEFAULT-NEXT: store double [[TMP0]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA15]]
// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]] // DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR19]]
// DEFAULT-NEXT: ret double [[CALL_I]] // DEFAULT-NEXT: ret double [[CALL_I]]
// //
// FINITEONLY-LABEL: @test_modf( // FINITEONLY-LABEL: @test_modf(
// FINITEONLY-NEXT: entry: // FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[__TMP_I:%.*]] = alloca double, align 8, addrspace(5) // FINITEONLY-NEXT: [[__TMP_I:%.*]] = alloca double, align 8, addrspace(5)
// FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]] // FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract double @__ocml_modf_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]] // FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract double @__ocml_modf_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]]
// FINITEONLY-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA15:![0-9]+]] // FINITEONLY-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA15:![0-9]+]]
▲ Show 20 Lines • Show All 606 Lines • ▼ Show 20 Lines
// //
extern "C" __device__ double test_remainder(double x, double y) { extern "C" __device__ double test_remainder(double x, double y) {
return remainder(x, y); return remainder(x, y);
} }
// DEFAULT-LABEL: @test_remquof( // DEFAULT-LABEL: @test_remquof(
// DEFAULT-NEXT: entry: // DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca i32, align 4, addrspace(5) // 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]]) #[[ATTR18]] // DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR19]]
// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract float @__ocml_remquo_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]] // DEFAULT-NEXT: [[CALL_I:%.*]] = call contract float @__ocml_remquo_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]]
// DEFAULT-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA19:![0-9]+]] // DEFAULT-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA19:![0-9]+]]
// DEFAULT-NEXT: store i32 [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA19]] // DEFAULT-NEXT: store i32 [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA19]]
// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]] // DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR19]]
// DEFAULT-NEXT: ret float [[CALL_I]] // DEFAULT-NEXT: ret float [[CALL_I]]
// //
// FINITEONLY-LABEL: @test_remquof( // FINITEONLY-LABEL: @test_remquof(
// FINITEONLY-NEXT: entry: // FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[__TMP_I:%.*]] = alloca i32, align 4, addrspace(5) // FINITEONLY-NEXT: [[__TMP_I:%.*]] = alloca i32, align 4, addrspace(5)
// FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]] // FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract float @__ocml_remquo_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]] // FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract float @__ocml_remquo_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]]
// FINITEONLY-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA19:![0-9]+]] // FINITEONLY-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA19:![0-9]+]]
// FINITEONLY-NEXT: store i32 [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA19]] // FINITEONLY-NEXT: store i32 [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA19]]
// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]] // FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
// FINITEONLY-NEXT: ret float [[CALL_I]] // FINITEONLY-NEXT: ret float [[CALL_I]]
// //
extern "C" __device__ float test_remquof(float x, float y, int* z) { extern "C" __device__ float test_remquof(float x, float y, int* z) {
return remquof(x, y, z); return remquof(x, y, z);
} }
// DEFAULT-LABEL: @test_remquo( // DEFAULT-LABEL: @test_remquo(
// DEFAULT-NEXT: entry: // DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca i32, align 4, addrspace(5) // 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]]) #[[ATTR18]] // DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR19]]
// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract double @__ocml_remquo_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]] // DEFAULT-NEXT: [[CALL_I:%.*]] = call contract double @__ocml_remquo_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]]
// DEFAULT-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA19]] // DEFAULT-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA19]]
// DEFAULT-NEXT: store i32 [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA19]] // DEFAULT-NEXT: store i32 [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA19]]
// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]] // DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR19]]
// DEFAULT-NEXT: ret double [[CALL_I]] // DEFAULT-NEXT: ret double [[CALL_I]]
// //
// FINITEONLY-LABEL: @test_remquo( // FINITEONLY-LABEL: @test_remquo(
// FINITEONLY-NEXT: entry: // FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[__TMP_I:%.*]] = alloca i32, align 4, addrspace(5) // FINITEONLY-NEXT: [[__TMP_I:%.*]] = alloca i32, align 4, addrspace(5)
// FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]] // FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract double @__ocml_remquo_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]] // FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract double @__ocml_remquo_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]]
// FINITEONLY-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA19]] // FINITEONLY-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA19]]
▲ Show 20 Lines • Show All 376 Lines • ▼ Show 20 Lines
// //
extern "C" __device__ BOOL_TYPE test___signbit(double x) { extern "C" __device__ BOOL_TYPE test___signbit(double x) {
return __signbit(x); return __signbit(x);
} }
// DEFAULT-LABEL: @test_sincosf( // DEFAULT-LABEL: @test_sincosf(
// DEFAULT-NEXT: entry: // DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5) // 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]]) #[[ATTR18]] // DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR19]]
// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract float @__ocml_sincos_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]] // DEFAULT-NEXT: [[CALL_I:%.*]] = call contract float @__ocml_sincos_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]]
// DEFAULT-NEXT: store float [[CALL_I]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA13]] // DEFAULT-NEXT: store float [[CALL_I]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA13]]
// DEFAULT-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA13]] // DEFAULT-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA13]]
// DEFAULT-NEXT: store float [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA13]] // DEFAULT-NEXT: store float [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA13]]
// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]] // DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR19]]
// DEFAULT-NEXT: ret void // DEFAULT-NEXT: ret void
// //
// FINITEONLY-LABEL: @test_sincosf( // FINITEONLY-LABEL: @test_sincosf(
// FINITEONLY-NEXT: entry: // FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5) // FINITEONLY-NEXT: [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5)
// FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]] // FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract float @__ocml_sincos_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]] // FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract float @__ocml_sincos_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]]
// FINITEONLY-NEXT: store float [[CALL_I]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA13]] // FINITEONLY-NEXT: store float [[CALL_I]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA13]]
// FINITEONLY-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA13]] // FINITEONLY-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA13]]
// FINITEONLY-NEXT: store float [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA13]] // FINITEONLY-NEXT: store float [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA13]]
// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]] // FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
// FINITEONLY-NEXT: ret void // FINITEONLY-NEXT: ret void
// //
extern "C" __device__ void test_sincosf(float x, float *y, float *z) { extern "C" __device__ void test_sincosf(float x, float *y, float *z) {
sincosf(x, y, z); sincosf(x, y, z);
} }
// DEFAULT-LABEL: @test_sincos( // DEFAULT-LABEL: @test_sincos(
// DEFAULT-NEXT: entry: // DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca double, align 8, addrspace(5) // 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]]) #[[ATTR18]] // DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR19]]
// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract double @__ocml_sincos_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]] // DEFAULT-NEXT: [[CALL_I:%.*]] = call contract double @__ocml_sincos_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]]
// DEFAULT-NEXT: store double [[CALL_I]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA15]] // DEFAULT-NEXT: store double [[CALL_I]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA15]]
// DEFAULT-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA15]] // DEFAULT-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA15]]
// DEFAULT-NEXT: store double [[TMP0]], ptr [[Z:%.*]], align 8, !tbaa [[TBAA15]] // DEFAULT-NEXT: store double [[TMP0]], ptr [[Z:%.*]], align 8, !tbaa [[TBAA15]]
// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]] // DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR19]]
// DEFAULT-NEXT: ret void // DEFAULT-NEXT: ret void
// //
// FINITEONLY-LABEL: @test_sincos( // FINITEONLY-LABEL: @test_sincos(
// FINITEONLY-NEXT: entry: // FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[__TMP_I:%.*]] = alloca double, align 8, addrspace(5) // FINITEONLY-NEXT: [[__TMP_I:%.*]] = alloca double, align 8, addrspace(5)
// FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]] // FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract double @__ocml_sincos_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]] // FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract double @__ocml_sincos_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]]
// FINITEONLY-NEXT: store double [[CALL_I]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA15]] // FINITEONLY-NEXT: store double [[CALL_I]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA15]]
// FINITEONLY-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA15]] // FINITEONLY-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA15]]
// FINITEONLY-NEXT: store double [[TMP0]], ptr [[Z:%.*]], align 8, !tbaa [[TBAA15]] // FINITEONLY-NEXT: store double [[TMP0]], ptr [[Z:%.*]], align 8, !tbaa [[TBAA15]]
// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]] // FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
// FINITEONLY-NEXT: ret void // FINITEONLY-NEXT: ret void
// //
extern "C" __device__ void test_sincos(double x, double *y, double *z) { extern "C" __device__ void test_sincos(double x, double *y, double *z) {
sincos(x, y, z); sincos(x, y, z);
} }
// DEFAULT-LABEL: @test_sincospif( // DEFAULT-LABEL: @test_sincospif(
// DEFAULT-NEXT: entry: // DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5) // 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]]) #[[ATTR18]] // DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR19]]
// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract float @__ocml_sincospi_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]] // DEFAULT-NEXT: [[CALL_I:%.*]] = call contract float @__ocml_sincospi_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]]
// DEFAULT-NEXT: store float [[CALL_I]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA13]] // DEFAULT-NEXT: store float [[CALL_I]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA13]]
// DEFAULT-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA13]] // DEFAULT-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA13]]
// DEFAULT-NEXT: store float [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA13]] // DEFAULT-NEXT: store float [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA13]]
// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]] // DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR19]]
// DEFAULT-NEXT: ret void // DEFAULT-NEXT: ret void
// //
// FINITEONLY-LABEL: @test_sincospif( // FINITEONLY-LABEL: @test_sincospif(
// FINITEONLY-NEXT: entry: // FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5) // FINITEONLY-NEXT: [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5)
// FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]] // FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract float @__ocml_sincospi_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]] // FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract float @__ocml_sincospi_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]]
// FINITEONLY-NEXT: store float [[CALL_I]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA13]] // FINITEONLY-NEXT: store float [[CALL_I]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA13]]
// FINITEONLY-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA13]] // FINITEONLY-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA13]]
// FINITEONLY-NEXT: store float [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA13]] // FINITEONLY-NEXT: store float [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA13]]
// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]] // FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
// FINITEONLY-NEXT: ret void // FINITEONLY-NEXT: ret void
// //
extern "C" __device__ void test_sincospif(float x, float *y, float *z) { extern "C" __device__ void test_sincospif(float x, float *y, float *z) {
sincospif(x, y, z); sincospif(x, y, z);
} }
// DEFAULT-LABEL: @test_sincospi( // DEFAULT-LABEL: @test_sincospi(
// DEFAULT-NEXT: entry: // DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca double, align 8, addrspace(5) // 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]]) #[[ATTR18]] // DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR19]]
// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract double @__ocml_sincospi_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]] // DEFAULT-NEXT: [[CALL_I:%.*]] = call contract double @__ocml_sincospi_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]]
// DEFAULT-NEXT: store double [[CALL_I]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA15]] // DEFAULT-NEXT: store double [[CALL_I]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA15]]
// DEFAULT-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA15]] // DEFAULT-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA15]]
// DEFAULT-NEXT: store double [[TMP0]], ptr [[Z:%.*]], align 8, !tbaa [[TBAA15]] // DEFAULT-NEXT: store double [[TMP0]], ptr [[Z:%.*]], align 8, !tbaa [[TBAA15]]
// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]] // DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR19]]
// DEFAULT-NEXT: ret void // DEFAULT-NEXT: ret void
// //
// FINITEONLY-LABEL: @test_sincospi( // FINITEONLY-LABEL: @test_sincospi(
// FINITEONLY-NEXT: entry: // FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[__TMP_I:%.*]] = alloca double, align 8, addrspace(5) // FINITEONLY-NEXT: [[__TMP_I:%.*]] = alloca double, align 8, addrspace(5)
// FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]] // FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract double @__ocml_sincospi_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]] // FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract double @__ocml_sincospi_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]]
// FINITEONLY-NEXT: store double [[CALL_I]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA15]] // FINITEONLY-NEXT: store double [[CALL_I]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA15]]
▲ Show 20 Lines • Show All 821 Lines • Show Last 20 Lines