Skip to content

Commit 71be91e

Browse files
arsenmarsenm
committed
HIP: Directly call rint builtins
1 parent d031ff3 commit 71be91e

File tree

2 files changed

+30
-30
lines changed

2 files changed

+30
-30
lines changed

‎clang/lib/Headers/__clang_hip_math.h‎

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -309,7 +309,7 @@ __DEVICE__
309309
float lgammaf(float __x) { return __ocml_lgamma_f32(__x); }
310310

311311
__DEVICE__
312-
long long int llrintf(float __x) { return __ocml_rint_f32(__x); }
312+
long long int llrintf(float __x) { return __builtin_rintf(__x); }
313313

314314
__DEVICE__
315315
long long int llroundf(float __x) { return __builtin_roundf(__x); }
@@ -330,7 +330,7 @@ __DEVICE__
330330
float logf(float __x) { return __builtin_logf(__x); }
331331

332332
__DEVICE__
333-
long int lrintf(float __x) { return __ocml_rint_f32(__x); }
333+
long int lrintf(float __x) { return __builtin_rintf(__x); }
334334

335335
__DEVICE__
336336
long int lroundf(float __x) { return __builtin_roundf(__x); }
@@ -435,7 +435,7 @@ __DEVICE__
435435
float rhypotf(float __x, float __y) { return __ocml_rhypot_f32(__x, __y); }
436436

437437
__DEVICE__
438-
float rintf(float __x) { return __ocml_rint_f32(__x); }
438+
float rintf(float __x) { return __builtin_rintf(__x); }
439439

440440
__DEVICE__
441441
float rnorm3df(float __x, float __y, float __z) {
@@ -857,7 +857,7 @@ __DEVICE__
857857
double lgamma(double __x) { return __ocml_lgamma_f64(__x); }
858858

859859
__DEVICE__
860-
long long int llrint(double __x) { return __ocml_rint_f64(__x); }
860+
long long int llrint(double __x) { return __builtin_rint(__x); }
861861

862862
__DEVICE__
863863
long long int llround(double __x) { return __builtin_round(__x); }
@@ -878,7 +878,7 @@ __DEVICE__
878878
double logb(double __x) { return __ocml_logb_f64(__x); }
879879

880880
__DEVICE__
881-
long int lrint(double __x) { return __ocml_rint_f64(__x); }
881+
long int lrint(double __x) { return __builtin_rint(__x); }
882882

883883
__DEVICE__
884884
long int lround(double __x) { return __builtin_round(__x); }
@@ -991,7 +991,7 @@ __DEVICE__
991991
double rhypot(double __x, double __y) { return __ocml_rhypot_f64(__x, __y); }
992992

993993
__DEVICE__
994-
double rint(double __x) { return __ocml_rint_f64(__x); }
994+
double rint(double __x) { return __builtin_rint(__x); }
995995

996996
__DEVICE__
997997
double rnorm(int __dim,

‎clang/test/Headers/__clang_hip_math.hip‎

Lines changed: 24 additions & 24 deletions
Original file line numberDiff line numberDiff line change
@@ -1479,14 +1479,14 @@ extern "C" __device__ double test_lgamma(double x) {
14791479

14801480
// DEFAULT-LABEL: @test_llrintf(
14811481
// DEFAULT-NEXT: entry:
1482-
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rint_f32(float noundef [[X:%.*]]) #[[ATTR14]]
1483-
// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64
1482+
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.rint.f32(float [[X:%.*]])
1483+
// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi float [[TMP0]] to i64
14841484
// DEFAULT-NEXT: ret i64 [[CONV_I]]
14851485
//
14861486
// FINITEONLY-LABEL: @test_llrintf(
14871487
// FINITEONLY-NEXT: entry:
1488-
// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_rint_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]]
1489-
// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64
1488+
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.rint.f32(float [[X:%.*]])
1489+
// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi float [[TMP0]] to i64
14901490
// FINITEONLY-NEXT: ret i64 [[CONV_I]]
14911491
//
14921492
extern "C" __device__ long long int test_llrintf(float x) {
@@ -1495,14 +1495,14 @@ extern "C" __device__ long long int test_llrintf(float x) {
14951495

14961496
// DEFAULT-LABEL: @test_llrint(
14971497
// DEFAULT-NEXT: entry:
1498-
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rint_f64(double noundef [[X:%.*]]) #[[ATTR14]]
1499-
// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64
1498+
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.rint.f64(double [[X:%.*]])
1499+
// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi double [[TMP0]] to i64
15001500
// DEFAULT-NEXT: ret i64 [[CONV_I]]
15011501
//
15021502
// FINITEONLY-LABEL: @test_llrint(
15031503
// FINITEONLY-NEXT: entry:
1504-
// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_rint_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]]
1505-
// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64
1504+
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.rint.f64(double [[X:%.*]])
1505+
// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi double [[TMP0]] to i64
15061506
// FINITEONLY-NEXT: ret i64 [[CONV_I]]
15071507
//
15081508
extern "C" __device__ long long int test_llrint(double x) {
@@ -1655,14 +1655,14 @@ extern "C" __device__ double test_logb(double x) {
16551655

16561656
// DEFAULT-LABEL: @test_lrintf(
16571657
// DEFAULT-NEXT: entry:
1658-
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rint_f32(float noundef [[X:%.*]]) #[[ATTR14]]
1659-
// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64
1658+
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.rint.f32(float [[X:%.*]])
1659+
// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi float [[TMP0]] to i64
16601660
// DEFAULT-NEXT: ret i64 [[CONV_I]]
16611661
//
16621662
// FINITEONLY-LABEL: @test_lrintf(
16631663
// FINITEONLY-NEXT: entry:
1664-
// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_rint_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]]
1665-
// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64
1664+
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.rint.f32(float [[X:%.*]])
1665+
// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi float [[TMP0]] to i64
16661666
// FINITEONLY-NEXT: ret i64 [[CONV_I]]
16671667
//
16681668
extern "C" __device__ long int test_lrintf(float x) {
@@ -1671,14 +1671,14 @@ extern "C" __device__ long int test_lrintf(float x) {
16711671

16721672
// DEFAULT-LABEL: @test_lrint(
16731673
// DEFAULT-NEXT: entry:
1674-
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rint_f64(double noundef [[X:%.*]]) #[[ATTR14]]
1675-
// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64
1674+
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.rint.f64(double [[X:%.*]])
1675+
// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi double [[TMP0]] to i64
16761676
// DEFAULT-NEXT: ret i64 [[CONV_I]]
16771677
//
16781678
// FINITEONLY-LABEL: @test_lrint(
16791679
// FINITEONLY-NEXT: entry:
1680-
// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_rint_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]]
1681-
// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64
1680+
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.rint.f64(double [[X:%.*]])
1681+
// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi double [[TMP0]] to i64
16821682
// FINITEONLY-NEXT: ret i64 [[CONV_I]]
16831683
//
16841684
extern "C" __device__ long int test_lrint(double x) {
@@ -2440,27 +2440,27 @@ extern "C" __device__ double test_rhypot(double x, double y) {
24402440

24412441
// DEFAULT-LABEL: @test_rintf(
24422442
// DEFAULT-NEXT: entry:
2443-
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rint_f32(float noundef [[X:%.*]]) #[[ATTR14]]
2444-
// DEFAULT-NEXT: ret float [[CALL_I]]
2443+
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.rint.f32(float [[X:%.*]])
2444+
// DEFAULT-NEXT: ret float [[TMP0]]
24452445
//
24462446
// FINITEONLY-LABEL: @test_rintf(
24472447
// FINITEONLY-NEXT: entry:
2448-
// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_rint_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]]
2449-
// FINITEONLY-NEXT: ret float [[CALL_I]]
2448+
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.rint.f32(float [[X:%.*]])
2449+
// FINITEONLY-NEXT: ret float [[TMP0]]
24502450
//
24512451
extern "C" __device__ float test_rintf(float x) {
24522452
return rintf(x);
24532453
}
24542454

24552455
// DEFAULT-LABEL: @test_rint(
24562456
// DEFAULT-NEXT: entry:
2457-
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rint_f64(double noundef [[X:%.*]]) #[[ATTR14]]
2458-
// DEFAULT-NEXT: ret double [[CALL_I]]
2457+
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.rint.f64(double [[X:%.*]])
2458+
// DEFAULT-NEXT: ret double [[TMP0]]
24592459
//
24602460
// FINITEONLY-LABEL: @test_rint(
24612461
// FINITEONLY-NEXT: entry:
2462-
// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_rint_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]]
2463-
// FINITEONLY-NEXT: ret double [[CALL_I]]
2462+
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.rint.f64(double [[X:%.*]])
2463+
// FINITEONLY-NEXT: ret double [[TMP0]]
24642464
//
24652465
extern "C" __device__ double test_rint(double x) {
24662466
return rint(x);

0 commit comments

Comments
 (0)