Author: Matt Arsenault Date: 2023-07-07T09:26:40-04:00 New Revision: 9a317516a515f3c5b15f9060329a503e8f261c7f
URL: https://github.com/llvm/llvm-project/commit/9a317516a515f3c5b15f9060329a503e8f261c7f DIFF: https://github.com/llvm/llvm-project/commit/9a317516a515f3c5b15f9060329a503e8f261c7f.diff LOG: HIP: Directly call round 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 c67f20873253b0..8a90c4acd94105 100644 --- a/clang/lib/Headers/__clang_hip_math.h +++ b/clang/lib/Headers/__clang_hip_math.h @@ -312,7 +312,7 @@ __DEVICE__ long long int llrintf(float __x) { return __ocml_rint_f32(__x); } __DEVICE__ -long long int llroundf(float __x) { return __ocml_round_f32(__x); } +long long int llroundf(float __x) { return __builtin_roundf(__x); } __DEVICE__ float log10f(float __x) { return __ocml_log10_f32(__x); } @@ -333,7 +333,7 @@ __DEVICE__ long int lrintf(float __x) { return __ocml_rint_f32(__x); } __DEVICE__ -long int lroundf(float __x) { return __ocml_round_f32(__x); } +long int lroundf(float __x) { return __builtin_roundf(__x); } __DEVICE__ float modff(float __x, float *__iptr) { @@ -460,7 +460,7 @@ float rnormf(int __dim, } __DEVICE__ -float roundf(float __x) { return __ocml_round_f32(__x); } +float roundf(float __x) { return __builtin_roundf(__x); } __DEVICE__ float rsqrtf(float __x) { return __ocml_rsqrt_f32(__x); } @@ -860,7 +860,7 @@ __DEVICE__ long long int llrint(double __x) { return __ocml_rint_f64(__x); } __DEVICE__ -long long int llround(double __x) { return __ocml_round_f64(__x); } +long long int llround(double __x) { return __builtin_round(__x); } __DEVICE__ double log(double __x) { return __ocml_log_f64(__x); } @@ -881,7 +881,7 @@ __DEVICE__ long int lrint(double __x) { return __ocml_rint_f64(__x); } __DEVICE__ -long int lround(double __x) { return __ocml_round_f64(__x); } +long int lround(double __x) { return __builtin_round(__x); } __DEVICE__ double modf(double __x, double *__iptr) { @@ -1016,7 +1016,7 @@ double rnorm4d(double __x, double __y, double __z, double __w) { } __DEVICE__ -double round(double __x) { return __ocml_round_f64(__x); } +double round(double __x) { return __builtin_round(__x); } __DEVICE__ double rsqrt(double __x) { return __ocml_rsqrt_f64(__x); } diff --git a/clang/test/Headers/__clang_hip_math.hip b/clang/test/Headers/__clang_hip_math.hip index a6797b01f25b27..35f637b912f439 100644 --- a/clang/test/Headers/__clang_hip_math.hip +++ b/clang/test/Headers/__clang_hip_math.hip @@ -1515,14 +1515,14 @@ extern "C" __device__ long long int test_llrint(double x) { // DEFAULT-LABEL: @test_llroundf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_round_f32(float noundef [[X:%.*]]) #[[ATTR14]] -// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64 +// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.round.f32(float [[X:%.*]]) +// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi float [[TMP0]] to i64 // DEFAULT-NEXT: ret i64 [[CONV_I]] // // FINITEONLY-LABEL: @test_llroundf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_round_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] -// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64 +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.round.f32(float [[X:%.*]]) +// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi float [[TMP0]] to i64 // FINITEONLY-NEXT: ret i64 [[CONV_I]] // extern "C" __device__ long long int test_llroundf(float x) { @@ -1531,14 +1531,14 @@ extern "C" __device__ long long int test_llroundf(float x) { // DEFAULT-LABEL: @test_llround( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_round_f64(double noundef [[X:%.*]]) #[[ATTR14]] -// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64 +// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.round.f64(double [[X:%.*]]) +// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi double [[TMP0]] to i64 // DEFAULT-NEXT: ret i64 [[CONV_I]] // // FINITEONLY-LABEL: @test_llround( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_round_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] -// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64 +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.round.f64(double [[X:%.*]]) +// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi double [[TMP0]] to i64 // FINITEONLY-NEXT: ret i64 [[CONV_I]] // extern "C" __device__ long long int test_llround(double x) { @@ -1691,14 +1691,14 @@ extern "C" __device__ long int test_lrint(double x) { // DEFAULT-LABEL: @test_lroundf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_round_f32(float noundef [[X:%.*]]) #[[ATTR14]] -// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64 +// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.round.f32(float [[X:%.*]]) +// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi float [[TMP0]] to i64 // DEFAULT-NEXT: ret i64 [[CONV_I]] // // FINITEONLY-LABEL: @test_lroundf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_round_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] -// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64 +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.round.f32(float [[X:%.*]]) +// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi float [[TMP0]] to i64 // FINITEONLY-NEXT: ret i64 [[CONV_I]] // extern "C" __device__ long int test_lroundf(float x) { @@ -1707,14 +1707,14 @@ extern "C" __device__ long int test_lroundf(float x) { // DEFAULT-LABEL: @test_lround( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_round_f64(double noundef [[X:%.*]]) #[[ATTR14]] -// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64 +// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.round.f64(double [[X:%.*]]) +// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi double [[TMP0]] to i64 // DEFAULT-NEXT: ret i64 [[CONV_I]] // // FINITEONLY-LABEL: @test_lround( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_round_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] -// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64 +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.round.f64(double [[X:%.*]]) +// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi double [[TMP0]] to i64 // FINITEONLY-NEXT: ret i64 [[CONV_I]] // extern "C" __device__ long int test_lround(double x) { @@ -2616,13 +2616,13 @@ extern "C" __device__ double test_rnorm4d(double x, double y, double z, double w // DEFAULT-LABEL: @test_roundf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_round_f32(float noundef [[X:%.*]]) #[[ATTR14]] -// DEFAULT-NEXT: ret float [[CALL_I]] +// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.round.f32(float [[X:%.*]]) +// DEFAULT-NEXT: ret float [[TMP0]] // // FINITEONLY-LABEL: @test_roundf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_round_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] -// FINITEONLY-NEXT: ret float [[CALL_I]] +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.round.f32(float [[X:%.*]]) +// FINITEONLY-NEXT: ret float [[TMP0]] // extern "C" __device__ float test_roundf(float x) { return roundf(x); @@ -2630,13 +2630,13 @@ extern "C" __device__ float test_roundf(float x) { // DEFAULT-LABEL: @test_round( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_round_f64(double noundef [[X:%.*]]) #[[ATTR14]] -// DEFAULT-NEXT: ret double [[CALL_I]] +// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.round.f64(double [[X:%.*]]) +// DEFAULT-NEXT: ret double [[TMP0]] // // FINITEONLY-LABEL: @test_round( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_round_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] -// FINITEONLY-NEXT: ret double [[CALL_I]] +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.round.f64(double [[X:%.*]]) +// FINITEONLY-NEXT: ret double [[TMP0]] // extern "C" __device__ double test_round(double x) { return round(x); _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
