Skip to content

Commit

Permalink
HIP: Directly call rint builtins
Browse files Browse the repository at this point in the history
  • Loading branch information
@arsenm
arsenm committed Jul 25, 2023
1 parent d031ff3 commit 71be91e
Show file tree
Hide file tree
Showing 2 changed files with 30 additions and 30 deletions.
12 changes: 6 additions & 6 deletions clang/lib/Headers/__clang_hip_math.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -309,7 +309,7 @@ __DEVICE__
float lgammaf(float __x) { return __ocml_lgamma_f32(__x); }

__DEVICE__
long long int llrintf(float __x) { return __ocml_rint_f32(__x); }
long long int llrintf(float __x) { return __builtin_rintf(__x); }

__DEVICE__
long long int llroundf(float __x) { return __builtin_roundf(__x); }
Expand All @@ -330,7 +330,7 @@ __DEVICE__
float logf(float __x) { return __builtin_logf(__x); }

__DEVICE__
long int lrintf(float __x) { return __ocml_rint_f32(__x); }
long int lrintf(float __x) { return __builtin_rintf(__x); }

__DEVICE__
long int lroundf(float __x) { return __builtin_roundf(__x); }
Expand Down Expand Up @@ -435,7 +435,7 @@ __DEVICE__
float rhypotf(float __x, float __y) { return __ocml_rhypot_f32(__x, __y); }

__DEVICE__
float rintf(float __x) { return __ocml_rint_f32(__x); }
float rintf(float __x) { return __builtin_rintf(__x); }

__DEVICE__
float rnorm3df(float __x, float __y, float __z) {
Expand Down Expand Up @@ -857,7 +857,7 @@ __DEVICE__
double lgamma(double __x) { return __ocml_lgamma_f64(__x); }

__DEVICE__
long long int llrint(double __x) { return __ocml_rint_f64(__x); }
long long int llrint(double __x) { return __builtin_rint(__x); }

__DEVICE__
long long int llround(double __x) { return __builtin_round(__x); }
Expand All @@ -878,7 +878,7 @@ __DEVICE__
double logb(double __x) { return __ocml_logb_f64(__x); }

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

__DEVICE__
long int lround(double __x) { return __builtin_round(__x); }
Expand Down Expand Up @@ -991,7 +991,7 @@ __DEVICE__
double rhypot(double __x, double __y) { return __ocml_rhypot_f64(__x, __y); }

__DEVICE__
double rint(double __x) { return __ocml_rint_f64(__x); }
double rint(double __x) { return __builtin_rint(__x); }

__DEVICE__
double rnorm(int __dim,
Expand Down
48 changes: 24 additions & 24 deletions clang/test/Headers/__clang_hip_math.hip
View file
Original file line number Diff line number Diff line change
Expand Up @@ -1479,14 +1479,14 @@ extern "C" __device__ double test_lgamma(double x) {

// DEFAULT-LABEL: @test_llrintf(
// DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rint_f32(float noundef [[X:%.*]]) #[[ATTR14]]
// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.rint.f32(float [[X:%.*]])
// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi float [[TMP0]] to i64
// DEFAULT-NEXT: ret i64 [[CONV_I]]
//
// FINITEONLY-LABEL: @test_llrintf(
// FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_rint_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.rint.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_llrintf(float x) {
Expand All @@ -1495,14 +1495,14 @@ extern "C" __device__ long long int test_llrintf(float x) {

// DEFAULT-LABEL: @test_llrint(
// DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rint_f64(double noundef [[X:%.*]]) #[[ATTR14]]
// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.rint.f64(double [[X:%.*]])
// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi double [[TMP0]] to i64
// DEFAULT-NEXT: ret i64 [[CONV_I]]
//
// FINITEONLY-LABEL: @test_llrint(
// FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_rint_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.rint.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_llrint(double x) {
Expand Down Expand Up @@ -1655,14 +1655,14 @@ extern "C" __device__ double test_logb(double x) {

// DEFAULT-LABEL: @test_lrintf(
// DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rint_f32(float noundef [[X:%.*]]) #[[ATTR14]]
// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.rint.f32(float [[X:%.*]])
// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi float [[TMP0]] to i64
// DEFAULT-NEXT: ret i64 [[CONV_I]]
//
// FINITEONLY-LABEL: @test_lrintf(
// FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_rint_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.rint.f32(float [[X:%.*]])
// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi float [[TMP0]] to i64
// FINITEONLY-NEXT: ret i64 [[CONV_I]]
//
extern "C" __device__ long int test_lrintf(float x) {
Expand All @@ -1671,14 +1671,14 @@ extern "C" __device__ long int test_lrintf(float x) {

// DEFAULT-LABEL: @test_lrint(
// DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rint_f64(double noundef [[X:%.*]]) #[[ATTR14]]
// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.rint.f64(double [[X:%.*]])
// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi double [[TMP0]] to i64
// DEFAULT-NEXT: ret i64 [[CONV_I]]
//
// FINITEONLY-LABEL: @test_lrint(
// FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_rint_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.rint.f64(double [[X:%.*]])
// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi double [[TMP0]] to i64
// FINITEONLY-NEXT: ret i64 [[CONV_I]]
//
extern "C" __device__ long int test_lrint(double x) {
Expand Down Expand Up @@ -2440,27 +2440,27 @@ extern "C" __device__ double test_rhypot(double x, double y) {

// DEFAULT-LABEL: @test_rintf(
// DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rint_f32(float noundef [[X:%.*]]) #[[ATTR14]]
// DEFAULT-NEXT: ret float [[CALL_I]]
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.rint.f32(float [[X:%.*]])
// DEFAULT-NEXT: ret float [[TMP0]]
//
// FINITEONLY-LABEL: @test_rintf(
// FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_rint_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.rint.f32(float [[X:%.*]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
extern "C" __device__ float test_rintf(float x) {
return rintf(x);
}

// DEFAULT-LABEL: @test_rint(
// DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rint_f64(double noundef [[X:%.*]]) #[[ATTR14]]
// DEFAULT-NEXT: ret double [[CALL_I]]
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.rint.f64(double [[X:%.*]])
// DEFAULT-NEXT: ret double [[TMP0]]
//
// FINITEONLY-LABEL: @test_rint(
// FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_rint_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.rint.f64(double [[X:%.*]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
extern "C" __device__ double test_rint(double x) {
return rint(x);
Expand Down

0 comments on commit 71be91e

Please sign in to comment.