Skip to content

Commit 6303b69

Browse files
arsenmHoney Goyal
arsenm
authored and
Honey Goyal
committed
clang/HIP: Avoid using ocml logb (llvm#171186)
We have special case handling for the logb builtins, so use them.
1 parent 62b860c commit 6303b69

File tree

2 files changed

+118
-48
lines changed

2 files changed

+118
-48
lines changed

‎clang/lib/Headers/__clang_hip_math.h‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -498,7 +498,7 @@ __DEVICE__
498498
float log2f(float __x) { return __FAST_OR_SLOW(__log2f, __builtin_log2f)(__x); }
499499

500500
__DEVICE__
501-
float logbf(float __x) { return __ocml_logb_f32(__x); }
501+
float logbf(float __x) { return __builtin_logbf(__x); }
502502

503503
__DEVICE__
504504
float logf(float __x) { return __FAST_OR_SLOW(__logf, __builtin_logf)(__x); }
@@ -901,7 +901,7 @@ __DEVICE__
901901
double log2(double __x) { return __ocml_log2_f64(__x); }
902902

903903
__DEVICE__
904-
double logb(double __x) { return __ocml_logb_f64(__x); }
904+
double logb(double __x) { return __builtin_logb(__x); }
905905

906906
__DEVICE__
907907
long int lrint(double __x) { return __builtin_rint(__x); }

‎clang/test/Headers/__clang_hip_math.hip‎

Lines changed: 116 additions & 46 deletions
Original file line numberDiff line numberDiff line change
@@ -3871,69 +3871,139 @@ extern "C" __device__ double test_log2(double x) {
38713871
return log2(x);
38723872
}
38733873

3874-
// DEFAULT-LABEL: define dso_local noundef float @test_logbf(
3875-
// DEFAULT-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR4]] {
3874+
// DEFAULT-LABEL: define dso_local float @test_logbf(
3875+
// DEFAULT-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
38763876
// DEFAULT-NEXT: [[ENTRY:.*:]]
3877-
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_logb_f32(float noundef [[X]]) #[[ATTR13]]
3878-
// DEFAULT-NEXT: ret float [[CALL_I]]
3879-
//
3880-
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) float @test_logbf(
3881-
// FINITEONLY-SAME: float noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR4]] {
3877+
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call { float, i32 } @llvm.frexp.f32.i32(float [[X]])
3878+
// DEFAULT-NEXT: [[TMP1:%.*]] = extractvalue { float, i32 } [[TMP0]], 1
3879+
// DEFAULT-NEXT: [[TMP2:%.*]] = add nsw i32 [[TMP1]], -1
3880+
// DEFAULT-NEXT: [[TMP3:%.*]] = sitofp i32 [[TMP2]] to float
3881+
// DEFAULT-NEXT: [[TMP4:%.*]] = tail call contract float @llvm.fabs.f32(float [[X]])
3882+
// DEFAULT-NEXT: [[TMP5:%.*]] = fcmp contract one float [[TMP4]], 0x7FF0000000000000
3883+
// DEFAULT-NEXT: [[TMP6:%.*]] = select contract i1 [[TMP5]], float [[TMP3]], float [[TMP4]]
3884+
// DEFAULT-NEXT: [[TMP7:%.*]] = fcmp contract oeq float [[X]], 0.000000e+00
3885+
// DEFAULT-NEXT: [[TMP8:%.*]] = select contract i1 [[TMP7]], float 0xFFF0000000000000, float [[TMP6]]
3886+
// DEFAULT-NEXT: ret float [[TMP8]]
3887+
//
3888+
// FINITEONLY-LABEL: define dso_local nofpclass(nan inf) float @test_logbf(
3889+
// FINITEONLY-SAME: float noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
38823890
// FINITEONLY-NEXT: [[ENTRY:.*:]]
3883-
// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_logb_f32(float noundef nofpclass(nan inf) [[X]]) #[[ATTR13]]
3884-
// FINITEONLY-NEXT: ret float [[CALL_I]]
3891+
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call { float, i32 } @llvm.frexp.f32.i32(float nofpclass(nan inf) [[X]])
3892+
// FINITEONLY-NEXT: [[TMP1:%.*]] = extractvalue { float, i32 } [[TMP0]], 1
3893+
// FINITEONLY-NEXT: [[TMP2:%.*]] = add nsw i32 [[TMP1]], -1
3894+
// FINITEONLY-NEXT: [[TMP3:%.*]] = sitofp i32 [[TMP2]] to float
3895+
// FINITEONLY-NEXT: ret float [[TMP3]]
38853896
//
3886-
// APPROX-LABEL: define dso_local noundef float @test_logbf(
3887-
// APPROX-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR4]] {
3897+
// APPROX-LABEL: define dso_local float @test_logbf(
3898+
// APPROX-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
38883899
// APPROX-NEXT: [[ENTRY:.*:]]
3889-
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_logb_f32(float noundef [[X]]) #[[ATTR13]]
3890-
// APPROX-NEXT: ret float [[CALL_I]]
3891-
//
3892-
// NCRDIV-LABEL: define dso_local noundef float @test_logbf(
3893-
// NCRDIV-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR4]] {
3900+
// APPROX-NEXT: [[TMP0:%.*]] = tail call { float, i32 } @llvm.frexp.f32.i32(float [[X]])
3901+
// APPROX-NEXT: [[TMP1:%.*]] = extractvalue { float, i32 } [[TMP0]], 1
3902+
// APPROX-NEXT: [[TMP2:%.*]] = add nsw i32 [[TMP1]], -1
3903+
// APPROX-NEXT: [[TMP3:%.*]] = sitofp i32 [[TMP2]] to float
3904+
// APPROX-NEXT: [[TMP4:%.*]] = tail call contract float @llvm.fabs.f32(float [[X]])
3905+
// APPROX-NEXT: [[TMP5:%.*]] = fcmp contract one float [[TMP4]], 0x7FF0000000000000
3906+
// APPROX-NEXT: [[TMP6:%.*]] = select contract i1 [[TMP5]], float [[TMP3]], float [[TMP4]]
3907+
// APPROX-NEXT: [[TMP7:%.*]] = fcmp contract oeq float [[X]], 0.000000e+00
3908+
// APPROX-NEXT: [[TMP8:%.*]] = select contract i1 [[TMP7]], float 0xFFF0000000000000, float [[TMP6]]
3909+
// APPROX-NEXT: ret float [[TMP8]]
3910+
//
3911+
// NCRDIV-LABEL: define dso_local float @test_logbf(
3912+
// NCRDIV-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
38943913
// NCRDIV-NEXT: [[ENTRY:.*:]]
3895-
// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_logb_f32(float noundef [[X]]) #[[ATTR13]]
3896-
// NCRDIV-NEXT: ret float [[CALL_I]]
3897-
//
3898-
// AMDGCNSPIRV-LABEL: define spir_func noundef float @test_logbf(
3899-
// AMDGCNSPIRV-SAME: float noundef [[X:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR4]] {
3914+
// NCRDIV-NEXT: [[TMP0:%.*]] = tail call { float, i32 } @llvm.frexp.f32.i32(float [[X]])
3915+
// NCRDIV-NEXT: [[TMP1:%.*]] = extractvalue { float, i32 } [[TMP0]], 1
3916+
// NCRDIV-NEXT: [[TMP2:%.*]] = add nsw i32 [[TMP1]], -1
3917+
// NCRDIV-NEXT: [[TMP3:%.*]] = sitofp i32 [[TMP2]] to float
3918+
// NCRDIV-NEXT: [[TMP4:%.*]] = tail call contract float @llvm.fabs.f32(float [[X]])
3919+
// NCRDIV-NEXT: [[TMP5:%.*]] = fcmp contract one float [[TMP4]], 0x7FF0000000000000
3920+
// NCRDIV-NEXT: [[TMP6:%.*]] = select contract i1 [[TMP5]], float [[TMP3]], float [[TMP4]]
3921+
// NCRDIV-NEXT: [[TMP7:%.*]] = fcmp contract oeq float [[X]], 0.000000e+00
3922+
// NCRDIV-NEXT: [[TMP8:%.*]] = select contract i1 [[TMP7]], float 0xFFF0000000000000, float [[TMP6]]
3923+
// NCRDIV-NEXT: ret float [[TMP8]]
3924+
//
3925+
// AMDGCNSPIRV-LABEL: define spir_func float @test_logbf(
3926+
// AMDGCNSPIRV-SAME: float noundef [[X:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR3]] {
39003927
// AMDGCNSPIRV-NEXT: [[ENTRY:.*:]]
3901-
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_logb_f32(float noundef [[X]]) #[[ATTR13]]
3902-
// AMDGCNSPIRV-NEXT: ret float [[CALL_I]]
3928+
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call addrspace(4) { float, i32 } @llvm.frexp.f32.i32(float [[X]])
3929+
// AMDGCNSPIRV-NEXT: [[TMP1:%.*]] = extractvalue { float, i32 } [[TMP0]], 1
3930+
// AMDGCNSPIRV-NEXT: [[TMP2:%.*]] = add nsw i32 [[TMP1]], -1
3931+
// AMDGCNSPIRV-NEXT: [[TMP3:%.*]] = sitofp i32 [[TMP2]] to float
3932+
// AMDGCNSPIRV-NEXT: [[TMP4:%.*]] = tail call contract addrspace(4) float @llvm.fabs.f32(float [[X]])
3933+
// AMDGCNSPIRV-NEXT: [[TMP5:%.*]] = fcmp contract one float [[TMP4]], 0x7FF0000000000000
3934+
// AMDGCNSPIRV-NEXT: [[TMP6:%.*]] = select contract i1 [[TMP5]], float [[TMP3]], float [[TMP4]]
3935+
// AMDGCNSPIRV-NEXT: [[TMP7:%.*]] = fcmp contract oeq float [[X]], 0.000000e+00
3936+
// AMDGCNSPIRV-NEXT: [[TMP8:%.*]] = select contract i1 [[TMP7]], float 0xFFF0000000000000, float [[TMP6]]
3937+
// AMDGCNSPIRV-NEXT: ret float [[TMP8]]
39033938
//
39043939
extern "C" __device__ float test_logbf(float x) {
39053940
return logbf(x);
39063941
}
39073942

3908-
// DEFAULT-LABEL: define dso_local noundef double @test_logb(
3909-
// DEFAULT-SAME: double noundef [[X:%.*]]) local_unnamed_addr #[[ATTR4]] {
3943+
// DEFAULT-LABEL: define dso_local double @test_logb(
3944+
// DEFAULT-SAME: double noundef [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
39103945
// DEFAULT-NEXT: [[ENTRY:.*:]]
3911-
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_logb_f64(double noundef [[X]]) #[[ATTR13]]
3912-
// DEFAULT-NEXT: ret double [[CALL_I]]
3913-
//
3914-
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) double @test_logb(
3915-
// FINITEONLY-SAME: double noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR4]] {
3946+
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call { double, i32 } @llvm.frexp.f64.i32(double [[X]])
3947+
// DEFAULT-NEXT: [[TMP1:%.*]] = extractvalue { double, i32 } [[TMP0]], 1
3948+
// DEFAULT-NEXT: [[TMP2:%.*]] = add nsw i32 [[TMP1]], -1
3949+
// DEFAULT-NEXT: [[TMP3:%.*]] = sitofp i32 [[TMP2]] to double
3950+
// DEFAULT-NEXT: [[TMP4:%.*]] = tail call contract double @llvm.fabs.f64(double [[X]])
3951+
// DEFAULT-NEXT: [[TMP5:%.*]] = fcmp contract one double [[TMP4]], 0x7FF0000000000000
3952+
// DEFAULT-NEXT: [[TMP6:%.*]] = select contract i1 [[TMP5]], double [[TMP3]], double [[TMP4]]
3953+
// DEFAULT-NEXT: [[TMP7:%.*]] = fcmp contract oeq double [[X]], 0.000000e+00
3954+
// DEFAULT-NEXT: [[TMP8:%.*]] = select contract i1 [[TMP7]], double 0xFFF0000000000000, double [[TMP6]]
3955+
// DEFAULT-NEXT: ret double [[TMP8]]
3956+
//
3957+
// FINITEONLY-LABEL: define dso_local nofpclass(nan inf) double @test_logb(
3958+
// FINITEONLY-SAME: double noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
39163959
// FINITEONLY-NEXT: [[ENTRY:.*:]]
3917-
// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @__ocml_logb_f64(double noundef nofpclass(nan inf) [[X]]) #[[ATTR13]]
3918-
// FINITEONLY-NEXT: ret double [[CALL_I]]
3960+
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call { double, i32 } @llvm.frexp.f64.i32(double nofpclass(nan inf) [[X]])
3961+
// FINITEONLY-NEXT: [[TMP1:%.*]] = extractvalue { double, i32 } [[TMP0]], 1
3962+
// FINITEONLY-NEXT: [[TMP2:%.*]] = add nsw i32 [[TMP1]], -1
3963+
// FINITEONLY-NEXT: [[TMP3:%.*]] = sitofp i32 [[TMP2]] to double
3964+
// FINITEONLY-NEXT: ret double [[TMP3]]
39193965
//
3920-
// APPROX-LABEL: define dso_local noundef double @test_logb(
3921-
// APPROX-SAME: double noundef [[X:%.*]]) local_unnamed_addr #[[ATTR4]] {
3966+
// APPROX-LABEL: define dso_local double @test_logb(
3967+
// APPROX-SAME: double noundef [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
39223968
// APPROX-NEXT: [[ENTRY:.*:]]
3923-
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_logb_f64(double noundef [[X]]) #[[ATTR13]]
3924-
// APPROX-NEXT: ret double [[CALL_I]]
3925-
//
3926-
// NCRDIV-LABEL: define dso_local noundef double @test_logb(
3927-
// NCRDIV-SAME: double noundef [[X:%.*]]) local_unnamed_addr #[[ATTR4]] {
3969+
// APPROX-NEXT: [[TMP0:%.*]] = tail call { double, i32 } @llvm.frexp.f64.i32(double [[X]])
3970+
// APPROX-NEXT: [[TMP1:%.*]] = extractvalue { double, i32 } [[TMP0]], 1
3971+
// APPROX-NEXT: [[TMP2:%.*]] = add nsw i32 [[TMP1]], -1
3972+
// APPROX-NEXT: [[TMP3:%.*]] = sitofp i32 [[TMP2]] to double
3973+
// APPROX-NEXT: [[TMP4:%.*]] = tail call contract double @llvm.fabs.f64(double [[X]])
3974+
// APPROX-NEXT: [[TMP5:%.*]] = fcmp contract one double [[TMP4]], 0x7FF0000000000000
3975+
// APPROX-NEXT: [[TMP6:%.*]] = select contract i1 [[TMP5]], double [[TMP3]], double [[TMP4]]
3976+
// APPROX-NEXT: [[TMP7:%.*]] = fcmp contract oeq double [[X]], 0.000000e+00
3977+
// APPROX-NEXT: [[TMP8:%.*]] = select contract i1 [[TMP7]], double 0xFFF0000000000000, double [[TMP6]]
3978+
// APPROX-NEXT: ret double [[TMP8]]
3979+
//
3980+
// NCRDIV-LABEL: define dso_local double @test_logb(
3981+
// NCRDIV-SAME: double noundef [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
39283982
// NCRDIV-NEXT: [[ENTRY:.*:]]
3929-
// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_logb_f64(double noundef [[X]]) #[[ATTR13]]
3930-
// NCRDIV-NEXT: ret double [[CALL_I]]
3931-
//
3932-
// AMDGCNSPIRV-LABEL: define spir_func noundef double @test_logb(
3933-
// AMDGCNSPIRV-SAME: double noundef [[X:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR4]] {
3983+
// NCRDIV-NEXT: [[TMP0:%.*]] = tail call { double, i32 } @llvm.frexp.f64.i32(double [[X]])
3984+
// NCRDIV-NEXT: [[TMP1:%.*]] = extractvalue { double, i32 } [[TMP0]], 1
3985+
// NCRDIV-NEXT: [[TMP2:%.*]] = add nsw i32 [[TMP1]], -1
3986+
// NCRDIV-NEXT: [[TMP3:%.*]] = sitofp i32 [[TMP2]] to double
3987+
// NCRDIV-NEXT: [[TMP4:%.*]] = tail call contract double @llvm.fabs.f64(double [[X]])
3988+
// NCRDIV-NEXT: [[TMP5:%.*]] = fcmp contract one double [[TMP4]], 0x7FF0000000000000
3989+
// NCRDIV-NEXT: [[TMP6:%.*]] = select contract i1 [[TMP5]], double [[TMP3]], double [[TMP4]]
3990+
// NCRDIV-NEXT: [[TMP7:%.*]] = fcmp contract oeq double [[X]], 0.000000e+00
3991+
// NCRDIV-NEXT: [[TMP8:%.*]] = select contract i1 [[TMP7]], double 0xFFF0000000000000, double [[TMP6]]
3992+
// NCRDIV-NEXT: ret double [[TMP8]]
3993+
//
3994+
// AMDGCNSPIRV-LABEL: define spir_func double @test_logb(
3995+
// AMDGCNSPIRV-SAME: double noundef [[X:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR3]] {
39343996
// AMDGCNSPIRV-NEXT: [[ENTRY:.*:]]
3935-
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_logb_f64(double noundef [[X]]) #[[ATTR13]]
3936-
// AMDGCNSPIRV-NEXT: ret double [[CALL_I]]
3997+
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call addrspace(4) { double, i32 } @llvm.frexp.f64.i32(double [[X]])
3998+
// AMDGCNSPIRV-NEXT: [[TMP1:%.*]] = extractvalue { double, i32 } [[TMP0]], 1
3999+
// AMDGCNSPIRV-NEXT: [[TMP2:%.*]] = add nsw i32 [[TMP1]], -1
4000+
// AMDGCNSPIRV-NEXT: [[TMP3:%.*]] = sitofp i32 [[TMP2]] to double
4001+
// AMDGCNSPIRV-NEXT: [[TMP4:%.*]] = tail call contract addrspace(4) double @llvm.fabs.f64(double [[X]])
4002+
// AMDGCNSPIRV-NEXT: [[TMP5:%.*]] = fcmp contract one double [[TMP4]], 0x7FF0000000000000
4003+
// AMDGCNSPIRV-NEXT: [[TMP6:%.*]] = select contract i1 [[TMP5]], double [[TMP3]], double [[TMP4]]
4004+
// AMDGCNSPIRV-NEXT: [[TMP7:%.*]] = fcmp contract oeq double [[X]], 0.000000e+00
4005+
// AMDGCNSPIRV-NEXT: [[TMP8:%.*]] = select contract i1 [[TMP7]], double 0xFFF0000000000000, double [[TMP6]]
4006+
// AMDGCNSPIRV-NEXT: ret double [[TMP8]]
39374007
//
39384008
extern "C" __device__ double test_logb(double x) {
39394009
return logb(x);

0 commit comments

Comments
 (0)