llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-backend-x86 Author: Matt Arsenault (arsenm) <details> <summary>Changes</summary> We have special case handling for the logb builtins, so use them. --- Full diff: https://github.com/llvm/llvm-project/pull/171186.diff 2 Files Affected: - (modified) clang/lib/Headers/__clang_hip_math.h (+2-2) - (modified) clang/test/Headers/__clang_hip_math.hip (+116-46) ``````````diff diff --git a/clang/lib/Headers/__clang_hip_math.h b/clang/lib/Headers/__clang_hip_math.h index 759e742c9d012..03c2721b4ad3c 100644 --- a/clang/lib/Headers/__clang_hip_math.h +++ b/clang/lib/Headers/__clang_hip_math.h @@ -498,7 +498,7 @@ __DEVICE__ float log2f(float __x) { return __FAST_OR_SLOW(__log2f, __builtin_log2f)(__x); } __DEVICE__ -float logbf(float __x) { return __ocml_logb_f32(__x); } +float logbf(float __x) { return __builtin_logbf(__x); } __DEVICE__ float logf(float __x) { return __FAST_OR_SLOW(__logf, __builtin_logf)(__x); } @@ -901,7 +901,7 @@ __DEVICE__ double log2(double __x) { return __ocml_log2_f64(__x); } __DEVICE__ -double logb(double __x) { return __ocml_logb_f64(__x); } +double logb(double __x) { return __builtin_logb(__x); } __DEVICE__ long int lrint(double __x) { return __builtin_rint(__x); } diff --git a/clang/test/Headers/__clang_hip_math.hip b/clang/test/Headers/__clang_hip_math.hip index 4163666811c91..426e5af319cbf 100644 --- a/clang/test/Headers/__clang_hip_math.hip +++ b/clang/test/Headers/__clang_hip_math.hip @@ -3871,69 +3871,139 @@ extern "C" __device__ double test_log2(double x) { return log2(x); } -// DEFAULT-LABEL: define dso_local noundef float @test_logbf( -// DEFAULT-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR4]] { +// DEFAULT-LABEL: define dso_local float @test_logbf( +// DEFAULT-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR3]] { // DEFAULT-NEXT: [[ENTRY:.*:]] -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_logb_f32(float noundef [[X]]) #[[ATTR13]] -// DEFAULT-NEXT: ret float [[CALL_I]] -// -// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) float @test_logbf( -// FINITEONLY-SAME: float noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR4]] { +// DEFAULT-NEXT: [[TMP0:%.*]] = tail call { float, i32 } @llvm.frexp.f32.i32(float [[X]]) +// DEFAULT-NEXT: [[TMP1:%.*]] = extractvalue { float, i32 } [[TMP0]], 1 +// DEFAULT-NEXT: [[TMP2:%.*]] = add nsw i32 [[TMP1]], -1 +// DEFAULT-NEXT: [[TMP3:%.*]] = sitofp i32 [[TMP2]] to float +// DEFAULT-NEXT: [[TMP4:%.*]] = tail call contract float @llvm.fabs.f32(float [[X]]) +// DEFAULT-NEXT: [[TMP5:%.*]] = fcmp contract one float [[TMP4]], 0x7FF0000000000000 +// DEFAULT-NEXT: [[TMP6:%.*]] = select contract i1 [[TMP5]], float [[TMP3]], float [[TMP4]] +// DEFAULT-NEXT: [[TMP7:%.*]] = fcmp contract oeq float [[X]], 0.000000e+00 +// DEFAULT-NEXT: [[TMP8:%.*]] = select contract i1 [[TMP7]], float 0xFFF0000000000000, float [[TMP6]] +// DEFAULT-NEXT: ret float [[TMP8]] +// +// FINITEONLY-LABEL: define dso_local nofpclass(nan inf) float @test_logbf( +// FINITEONLY-SAME: float noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR3]] { // FINITEONLY-NEXT: [[ENTRY:.*:]] -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_logb_f32(float noundef nofpclass(nan inf) [[X]]) #[[ATTR13]] -// FINITEONLY-NEXT: ret float [[CALL_I]] +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call { float, i32 } @llvm.frexp.f32.i32(float nofpclass(nan inf) [[X]]) +// FINITEONLY-NEXT: [[TMP1:%.*]] = extractvalue { float, i32 } [[TMP0]], 1 +// FINITEONLY-NEXT: [[TMP2:%.*]] = add nsw i32 [[TMP1]], -1 +// FINITEONLY-NEXT: [[TMP3:%.*]] = sitofp i32 [[TMP2]] to float +// FINITEONLY-NEXT: ret float [[TMP3]] // -// APPROX-LABEL: define dso_local noundef float @test_logbf( -// APPROX-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR4]] { +// APPROX-LABEL: define dso_local float @test_logbf( +// APPROX-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR3]] { // APPROX-NEXT: [[ENTRY:.*:]] -// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_logb_f32(float noundef [[X]]) #[[ATTR13]] -// APPROX-NEXT: ret float [[CALL_I]] -// -// NCRDIV-LABEL: define dso_local noundef float @test_logbf( -// NCRDIV-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR4]] { +// APPROX-NEXT: [[TMP0:%.*]] = tail call { float, i32 } @llvm.frexp.f32.i32(float [[X]]) +// APPROX-NEXT: [[TMP1:%.*]] = extractvalue { float, i32 } [[TMP0]], 1 +// APPROX-NEXT: [[TMP2:%.*]] = add nsw i32 [[TMP1]], -1 +// APPROX-NEXT: [[TMP3:%.*]] = sitofp i32 [[TMP2]] to float +// APPROX-NEXT: [[TMP4:%.*]] = tail call contract float @llvm.fabs.f32(float [[X]]) +// APPROX-NEXT: [[TMP5:%.*]] = fcmp contract one float [[TMP4]], 0x7FF0000000000000 +// APPROX-NEXT: [[TMP6:%.*]] = select contract i1 [[TMP5]], float [[TMP3]], float [[TMP4]] +// APPROX-NEXT: [[TMP7:%.*]] = fcmp contract oeq float [[X]], 0.000000e+00 +// APPROX-NEXT: [[TMP8:%.*]] = select contract i1 [[TMP7]], float 0xFFF0000000000000, float [[TMP6]] +// APPROX-NEXT: ret float [[TMP8]] +// +// NCRDIV-LABEL: define dso_local float @test_logbf( +// NCRDIV-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR3]] { // NCRDIV-NEXT: [[ENTRY:.*:]] -// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_logb_f32(float noundef [[X]]) #[[ATTR13]] -// NCRDIV-NEXT: ret float [[CALL_I]] -// -// AMDGCNSPIRV-LABEL: define spir_func noundef float @test_logbf( -// AMDGCNSPIRV-SAME: float noundef [[X:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR4]] { +// NCRDIV-NEXT: [[TMP0:%.*]] = tail call { float, i32 } @llvm.frexp.f32.i32(float [[X]]) +// NCRDIV-NEXT: [[TMP1:%.*]] = extractvalue { float, i32 } [[TMP0]], 1 +// NCRDIV-NEXT: [[TMP2:%.*]] = add nsw i32 [[TMP1]], -1 +// NCRDIV-NEXT: [[TMP3:%.*]] = sitofp i32 [[TMP2]] to float +// NCRDIV-NEXT: [[TMP4:%.*]] = tail call contract float @llvm.fabs.f32(float [[X]]) +// NCRDIV-NEXT: [[TMP5:%.*]] = fcmp contract one float [[TMP4]], 0x7FF0000000000000 +// NCRDIV-NEXT: [[TMP6:%.*]] = select contract i1 [[TMP5]], float [[TMP3]], float [[TMP4]] +// NCRDIV-NEXT: [[TMP7:%.*]] = fcmp contract oeq float [[X]], 0.000000e+00 +// NCRDIV-NEXT: [[TMP8:%.*]] = select contract i1 [[TMP7]], float 0xFFF0000000000000, float [[TMP6]] +// NCRDIV-NEXT: ret float [[TMP8]] +// +// AMDGCNSPIRV-LABEL: define spir_func float @test_logbf( +// AMDGCNSPIRV-SAME: float noundef [[X:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR3]] { // AMDGCNSPIRV-NEXT: [[ENTRY:.*:]] -// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_logb_f32(float noundef [[X]]) #[[ATTR13]] -// AMDGCNSPIRV-NEXT: ret float [[CALL_I]] +// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call addrspace(4) { float, i32 } @llvm.frexp.f32.i32(float [[X]]) +// AMDGCNSPIRV-NEXT: [[TMP1:%.*]] = extractvalue { float, i32 } [[TMP0]], 1 +// AMDGCNSPIRV-NEXT: [[TMP2:%.*]] = add nsw i32 [[TMP1]], -1 +// AMDGCNSPIRV-NEXT: [[TMP3:%.*]] = sitofp i32 [[TMP2]] to float +// AMDGCNSPIRV-NEXT: [[TMP4:%.*]] = tail call contract addrspace(4) float @llvm.fabs.f32(float [[X]]) +// AMDGCNSPIRV-NEXT: [[TMP5:%.*]] = fcmp contract one float [[TMP4]], 0x7FF0000000000000 +// AMDGCNSPIRV-NEXT: [[TMP6:%.*]] = select contract i1 [[TMP5]], float [[TMP3]], float [[TMP4]] +// AMDGCNSPIRV-NEXT: [[TMP7:%.*]] = fcmp contract oeq float [[X]], 0.000000e+00 +// AMDGCNSPIRV-NEXT: [[TMP8:%.*]] = select contract i1 [[TMP7]], float 0xFFF0000000000000, float [[TMP6]] +// AMDGCNSPIRV-NEXT: ret float [[TMP8]] // extern "C" __device__ float test_logbf(float x) { return logbf(x); } -// DEFAULT-LABEL: define dso_local noundef double @test_logb( -// DEFAULT-SAME: double noundef [[X:%.*]]) local_unnamed_addr #[[ATTR4]] { +// DEFAULT-LABEL: define dso_local double @test_logb( +// DEFAULT-SAME: double noundef [[X:%.*]]) local_unnamed_addr #[[ATTR3]] { // DEFAULT-NEXT: [[ENTRY:.*:]] -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_logb_f64(double noundef [[X]]) #[[ATTR13]] -// DEFAULT-NEXT: ret double [[CALL_I]] -// -// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) double @test_logb( -// FINITEONLY-SAME: double noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR4]] { +// DEFAULT-NEXT: [[TMP0:%.*]] = tail call { double, i32 } @llvm.frexp.f64.i32(double [[X]]) +// DEFAULT-NEXT: [[TMP1:%.*]] = extractvalue { double, i32 } [[TMP0]], 1 +// DEFAULT-NEXT: [[TMP2:%.*]] = add nsw i32 [[TMP1]], -1 +// DEFAULT-NEXT: [[TMP3:%.*]] = sitofp i32 [[TMP2]] to double +// DEFAULT-NEXT: [[TMP4:%.*]] = tail call contract double @llvm.fabs.f64(double [[X]]) +// DEFAULT-NEXT: [[TMP5:%.*]] = fcmp contract one double [[TMP4]], 0x7FF0000000000000 +// DEFAULT-NEXT: [[TMP6:%.*]] = select contract i1 [[TMP5]], double [[TMP3]], double [[TMP4]] +// DEFAULT-NEXT: [[TMP7:%.*]] = fcmp contract oeq double [[X]], 0.000000e+00 +// DEFAULT-NEXT: [[TMP8:%.*]] = select contract i1 [[TMP7]], double 0xFFF0000000000000, double [[TMP6]] +// DEFAULT-NEXT: ret double [[TMP8]] +// +// FINITEONLY-LABEL: define dso_local nofpclass(nan inf) double @test_logb( +// FINITEONLY-SAME: double noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR3]] { // FINITEONLY-NEXT: [[ENTRY:.*:]] -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @__ocml_logb_f64(double noundef nofpclass(nan inf) [[X]]) #[[ATTR13]] -// FINITEONLY-NEXT: ret double [[CALL_I]] +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call { double, i32 } @llvm.frexp.f64.i32(double nofpclass(nan inf) [[X]]) +// FINITEONLY-NEXT: [[TMP1:%.*]] = extractvalue { double, i32 } [[TMP0]], 1 +// FINITEONLY-NEXT: [[TMP2:%.*]] = add nsw i32 [[TMP1]], -1 +// FINITEONLY-NEXT: [[TMP3:%.*]] = sitofp i32 [[TMP2]] to double +// FINITEONLY-NEXT: ret double [[TMP3]] // -// APPROX-LABEL: define dso_local noundef double @test_logb( -// APPROX-SAME: double noundef [[X:%.*]]) local_unnamed_addr #[[ATTR4]] { +// APPROX-LABEL: define dso_local double @test_logb( +// APPROX-SAME: double noundef [[X:%.*]]) local_unnamed_addr #[[ATTR3]] { // APPROX-NEXT: [[ENTRY:.*:]] -// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_logb_f64(double noundef [[X]]) #[[ATTR13]] -// APPROX-NEXT: ret double [[CALL_I]] -// -// NCRDIV-LABEL: define dso_local noundef double @test_logb( -// NCRDIV-SAME: double noundef [[X:%.*]]) local_unnamed_addr #[[ATTR4]] { +// APPROX-NEXT: [[TMP0:%.*]] = tail call { double, i32 } @llvm.frexp.f64.i32(double [[X]]) +// APPROX-NEXT: [[TMP1:%.*]] = extractvalue { double, i32 } [[TMP0]], 1 +// APPROX-NEXT: [[TMP2:%.*]] = add nsw i32 [[TMP1]], -1 +// APPROX-NEXT: [[TMP3:%.*]] = sitofp i32 [[TMP2]] to double +// APPROX-NEXT: [[TMP4:%.*]] = tail call contract double @llvm.fabs.f64(double [[X]]) +// APPROX-NEXT: [[TMP5:%.*]] = fcmp contract one double [[TMP4]], 0x7FF0000000000000 +// APPROX-NEXT: [[TMP6:%.*]] = select contract i1 [[TMP5]], double [[TMP3]], double [[TMP4]] +// APPROX-NEXT: [[TMP7:%.*]] = fcmp contract oeq double [[X]], 0.000000e+00 +// APPROX-NEXT: [[TMP8:%.*]] = select contract i1 [[TMP7]], double 0xFFF0000000000000, double [[TMP6]] +// APPROX-NEXT: ret double [[TMP8]] +// +// NCRDIV-LABEL: define dso_local double @test_logb( +// NCRDIV-SAME: double noundef [[X:%.*]]) local_unnamed_addr #[[ATTR3]] { // NCRDIV-NEXT: [[ENTRY:.*:]] -// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_logb_f64(double noundef [[X]]) #[[ATTR13]] -// NCRDIV-NEXT: ret double [[CALL_I]] -// -// AMDGCNSPIRV-LABEL: define spir_func noundef double @test_logb( -// AMDGCNSPIRV-SAME: double noundef [[X:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR4]] { +// NCRDIV-NEXT: [[TMP0:%.*]] = tail call { double, i32 } @llvm.frexp.f64.i32(double [[X]]) +// NCRDIV-NEXT: [[TMP1:%.*]] = extractvalue { double, i32 } [[TMP0]], 1 +// NCRDIV-NEXT: [[TMP2:%.*]] = add nsw i32 [[TMP1]], -1 +// NCRDIV-NEXT: [[TMP3:%.*]] = sitofp i32 [[TMP2]] to double +// NCRDIV-NEXT: [[TMP4:%.*]] = tail call contract double @llvm.fabs.f64(double [[X]]) +// NCRDIV-NEXT: [[TMP5:%.*]] = fcmp contract one double [[TMP4]], 0x7FF0000000000000 +// NCRDIV-NEXT: [[TMP6:%.*]] = select contract i1 [[TMP5]], double [[TMP3]], double [[TMP4]] +// NCRDIV-NEXT: [[TMP7:%.*]] = fcmp contract oeq double [[X]], 0.000000e+00 +// NCRDIV-NEXT: [[TMP8:%.*]] = select contract i1 [[TMP7]], double 0xFFF0000000000000, double [[TMP6]] +// NCRDIV-NEXT: ret double [[TMP8]] +// +// AMDGCNSPIRV-LABEL: define spir_func double @test_logb( +// AMDGCNSPIRV-SAME: double noundef [[X:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR3]] { // AMDGCNSPIRV-NEXT: [[ENTRY:.*:]] -// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_logb_f64(double noundef [[X]]) #[[ATTR13]] -// AMDGCNSPIRV-NEXT: ret double [[CALL_I]] +// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call addrspace(4) { double, i32 } @llvm.frexp.f64.i32(double [[X]]) +// AMDGCNSPIRV-NEXT: [[TMP1:%.*]] = extractvalue { double, i32 } [[TMP0]], 1 +// AMDGCNSPIRV-NEXT: [[TMP2:%.*]] = add nsw i32 [[TMP1]], -1 +// AMDGCNSPIRV-NEXT: [[TMP3:%.*]] = sitofp i32 [[TMP2]] to double +// AMDGCNSPIRV-NEXT: [[TMP4:%.*]] = tail call contract addrspace(4) double @llvm.fabs.f64(double [[X]]) +// AMDGCNSPIRV-NEXT: [[TMP5:%.*]] = fcmp contract one double [[TMP4]], 0x7FF0000000000000 +// AMDGCNSPIRV-NEXT: [[TMP6:%.*]] = select contract i1 [[TMP5]], double [[TMP3]], double [[TMP4]] +// AMDGCNSPIRV-NEXT: [[TMP7:%.*]] = fcmp contract oeq double [[X]], 0.000000e+00 +// AMDGCNSPIRV-NEXT: [[TMP8:%.*]] = select contract i1 [[TMP7]], double 0xFFF0000000000000, double [[TMP6]] +// AMDGCNSPIRV-NEXT: ret double [[TMP8]] // extern "C" __device__ double test_logb(double x) { return logb(x); `````````` </details> https://github.com/llvm/llvm-project/pull/171186 _______________________________________________ llvm-branch-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
