Skip to content

Commit 99fcbd9

Browse files
zhiweij1zhiweij1
authored
[SYCLomatic] Refine the migration of std::nearbyint (#2845)
Signed-off-by: Jiang, Zhiwei <zhiwei.jiang@intel.com>
1 parent 91dfbb3 commit 99fcbd9

9 files changed

+100
-17
lines changed

‎clang/lib/DPCT/RuleInfra/CallExprRewriter.h

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1149,6 +1149,14 @@ template <class ET> class ParenExprPrinter {
11491149
public:
11501150
ParenExprPrinter(ET &&E) : E(std::forward<ET>(E)) {}
11511151
template <class StreamT> void print(StreamT &Stream) const {
1152+
if constexpr (std::is_base_of_v<
1153+
clang::Stmt,
1154+
std::remove_cv_t<std::remove_pointer_t<ET>>>) {
1155+
if (isa<DeclRefExpr>(E->IgnoreImpCasts())) {
1156+
dpct::print(Stream, E);
1157+
return;
1158+
}
1159+
}
11521160
PairedPrinter PP(Stream, "(", ")");
11531161
dpct::print(Stream, E);
11541162
}

‎clang/lib/DPCT/RulesLang/APINamesMath.inc

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -253,7 +253,7 @@ ENTRY_RENAMED_NO_REWRITE("__umul64hi", MapNames::getClNamespace(false, true) + "
253253
ENTRY_EMULATED("frexpf", MapNames::getClNamespace(false, true) + "frexp")
254254
ENTRY_EMULATED("modff", MapNames::getClNamespace(false, true) + "modf")
255255
ENTRY_EMULATED("nanf", MapNames::getClNamespace(false, true) + "nan")
256-
ENTRY_EMULATED("nearbyintf", MapNames::getClNamespace(false, true) + "floor")
256+
ENTRY_REWRITE("nearbyintf")
257257
ENTRY_EMULATED("remquof", MapNames::getClNamespace(false, true) + "remquo")
258258
ENTRY_EMULATED("rhypotf", MapNames::getClNamespace(false, true) + "hypot")
259259
ENTRY_REWRITE("sincosf")
@@ -266,7 +266,7 @@ ENTRY_EMULATED("__powf", MapNames::getDpctNamespace() + "pow")
266266
ENTRY_EMULATED("frexp", MapNames::getClNamespace(false, true) + "frexp")
267267
ENTRY_EMULATED("modf", MapNames::getClNamespace(false, true) + "modf")
268268
ENTRY_EMULATED("nan", MapNames::getClNamespace(false, true) + "nan")
269-
ENTRY_EMULATED("nearbyint", MapNames::getClNamespace(false, true) + "floor")
269+
ENTRY_REWRITE("nearbyint")
270270
ENTRY_EMULATED("remquo", MapNames::getClNamespace(false, true) + "remquo")
271271
ENTRY_EMULATED("rhypot", MapNames::getClNamespace(false, true) + "hypot")
272272
ENTRY_REWRITE("sincos")

‎clang/lib/DPCT/RulesLang/Math/CallExprRewriterMath.cpp

Lines changed: 0 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -455,9 +455,6 @@ std::optional<std::string> MathSimulatedRewriter::rewrite() {
455455
<< MapNames::getClNamespace() + "access::address_space::generic_space, "
456456
<< MapNames::getClNamespace() + "access::decorated::yes>("
457457
<< MigratedArg2 << "))";
458-
} else if (FuncName == "nearbyint" || FuncName == "nearbyintf") {
459-
OS << MapNames::getClNamespace(false, true) + "floor(" << MigratedArg0
460-
<< " + 0.5)";
461458
} else if (FuncName == "rhypot" || FuncName == "rhypotf") {
462459
auto MigratedArg1 = getMigratedArg(1);
463460
OS << "1 / " + MapNames::getClNamespace(false, true) + "hypot("

‎clang/lib/DPCT/RulesLang/Math/RewriterDoublePrecisionMathematicalFunctions.cpp

Lines changed: 16 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -349,6 +349,22 @@ RewriterMap dpct::createDoublePrecisionMathematicalFunctionsRewriterMap() {
349349
makeArgWithAddressSpaceCast(2)))),
350350
Diagnostics::MATH_EMULATION, std::string("sincospi"),
351351
MapNames::getClNamespace() + std::string("sincos")))
352+
// nearbyint
353+
MATH_API_REWRITERS_V2(
354+
"nearbyint",
355+
MATH_API_REWRITER_PAIR(
356+
math::Tag::device_std,
357+
CALL_FACTORY_ENTRY("nearbyint", CALL("std::nearbyint", ARG(0)))),
358+
MATH_API_REWRITER_PAIR(
359+
math::Tag::emulation,
360+
WARNING_FACTORY_ENTRY(
361+
"nearbyint",
362+
CALL_FACTORY_ENTRY("nearbyint",
363+
CALL(MapNames::getClNamespace() + "floor",
364+
BO(BinaryOperatorKind::BO_Add,
365+
PAREN(ARG(0)), ARG("0.5")))),
366+
Diagnostics::MATH_EMULATION, std::string("nearbyint"),
367+
MapNames::getClNamespace() + std::string("floor"))))
352368
// sinpi
353369
CALL_FACTORY_ENTRY(
354370
"sinpi",

‎clang/lib/DPCT/RulesLang/Math/RewriterSTDFunctions.cpp

Lines changed: 37 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -55,5 +55,41 @@ RewriterMap dpct::createSTDFunctionsRewriterMap() {
5555
ARG(0)))),
5656
MATH_API_REWRITER_PAIR(
5757
math::Tag::device_std,
58-
CALL_FACTORY_ENTRY("std::fabs", CALL("std::fabs", ARG(0)))))};
58+
CALL_FACTORY_ENTRY("std::fabs", CALL("std::fabs", ARG(0)))))
59+
MATH_API_REWRITERS_V2(
60+
"std::nearbyint",
61+
MATH_API_REWRITER_PAIR(
62+
math::Tag::device_std,
63+
CALL_FACTORY_ENTRY("std::nearbyint",
64+
CALL("std::nearbyint", ARG(0)))),
65+
MATH_API_REWRITER_PAIR(
66+
math::Tag::emulation,
67+
WARNING_FACTORY_ENTRY(
68+
"std::nearbyint",
69+
CALL_FACTORY_ENTRY(
70+
"std::nearbyint",
71+
CALL(MapNames::getClNamespace() + "floor",
72+
BO(BinaryOperatorKind::BO_Add, PAREN(ARG(0)),
73+
ARG("0.5")))),
74+
Diagnostics::MATH_EMULATION,
75+
std::string("std::nearbyint"),
76+
MapNames::getClNamespace() + std::string("floor"))))
77+
MATH_API_REWRITERS_V2(
78+
"std::nearbyintf",
79+
MATH_API_REWRITER_PAIR(
80+
math::Tag::device_std,
81+
CALL_FACTORY_ENTRY("std::nearbyintf",
82+
CALL("std::nearbyintf", ARG(0)))),
83+
MATH_API_REWRITER_PAIR(
84+
math::Tag::emulation,
85+
WARNING_FACTORY_ENTRY(
86+
"std::nearbyintf",
87+
CALL_FACTORY_ENTRY(
88+
"std::nearbyintf",
89+
CALL(MapNames::getClNamespace() + "floor",
90+
BO(BinaryOperatorKind::BO_Add, PAREN(ARG(0)),
91+
ARG("0.5f")))),
92+
Diagnostics::MATH_EMULATION,
93+
std::string("std::nearbyintf"),
94+
MapNames::getClNamespace() + std::string("floor"))))};
5995
}

‎clang/lib/DPCT/RulesLang/Math/RewriterSinglePrecisionMathematicalFunctions.cpp

Lines changed: 17 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -308,6 +308,23 @@ RewriterMap dpct::createSinglePrecisionMathematicalFunctionsRewriterMap() {
308308
makeArgWithAddressSpaceCast(2)))),
309309
Diagnostics::MATH_EMULATION, std::string("sincospif"),
310310
MapNames::getClNamespace() + std::string("sincos"))
311+
// nearbyintf
312+
MATH_API_REWRITERS_V2(
313+
"nearbyintf",
314+
MATH_API_REWRITER_PAIR(
315+
math::Tag::device_std,
316+
CALL_FACTORY_ENTRY("nearbyintf",
317+
CALL("std::nearbyintf", ARG(0)))),
318+
MATH_API_REWRITER_PAIR(
319+
math::Tag::emulation,
320+
WARNING_FACTORY_ENTRY(
321+
"nearbyintf",
322+
CALL_FACTORY_ENTRY("nearbyintf",
323+
CALL(MapNames::getClNamespace() + "floor",
324+
BO(BinaryOperatorKind::BO_Add,
325+
PAREN(ARG(0)), ARG("0.5f")))),
326+
Diagnostics::MATH_EMULATION, std::string("nearbyintf"),
327+
MapNames::getClNamespace() + std::string("floor"))))
311328
// sinpif
312329
MATH_API_REWRITERS_V2(
313330
"sinpif",

‎clang/test/dpct/math/cuda-math-intrinsics.cu

Lines changed: 10 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -2115,7 +2115,7 @@ __global__ void testSimulation() {
21152115
// CHECK: /*
21162116
// CHECK-NEXT: DPCT1017:{{[0-9]+}}: The sycl::floor call is used instead of the nearbyintf call. These two calls do not provide exactly the same functionality. Check the potential precision and/or performance issues for the generated code.
21172117
// CHECK-NEXT: */
2118-
// CHECK-NEXT: f = sycl::floor(f + 0.5);
2118+
// CHECK-NEXT: f = sycl::floor(f + 0.5f);
21192119
f = nearbyintf(f);
21202120

21212121
// CHECK: /*
@@ -2705,9 +2705,9 @@ __device__ void do_migration5() {
27052705
//CHECK-NEXT: std::min(i, i);
27062706
//CHECK-NEXT: sycl::fabs(f);
27072707
//CHECK-NEXT: /*
2708-
//CHECK-NEXT: DPCT1017:{{[0-9]+}}: The sycl::floor call is used instead of the nearbyintf call. These two calls do not provide exactly the same functionality. Check the potential precision and/or performance issues for the generated code.
2708+
//CHECK-NEXT: DPCT1017:{{[0-9]+}}: The sycl::floor call is used instead of the std::nearbyintf call. These two calls do not provide exactly the same functionality. Check the potential precision and/or performance issues for the generated code.
27092709
//CHECK-NEXT: */
2710-
//CHECK-NEXT: sycl::floor(f + 0.5);
2710+
//CHECK-NEXT: sycl::floor(f + 0.5f);
27112711
//CHECK-NEXT: /*
27122712
//CHECK-NEXT: DPCT1017:{{[0-9]+}}: The sycl::remquo call is used instead of the remquof call. These two calls do not provide exactly the same functionality. Check the potential precision and/or performance issues for the generated code.
27132713
//CHECK-NEXT: */
@@ -2724,7 +2724,7 @@ __device__ void do_migration5() {
27242724
//CHECK-NEXT: */
27252725
//CHECK-NEXT: sycl::modf(f, sycl::address_space_cast<sycl::access::address_space::generic_space, sycl::access::decorated::yes>(&f));
27262726
//CHECK-NEXT: /*
2727-
//CHECK-NEXT: DPCT1017:{{[0-9]+}}: The sycl::floor call is used instead of the nearbyint call. These two calls do not provide exactly the same functionality. Check the potential precision and/or performance issues for the generated code.
2727+
//CHECK-NEXT: DPCT1017:{{[0-9]+}}: The sycl::floor call is used instead of the std::nearbyint call. These two calls do not provide exactly the same functionality. Check the potential precision and/or performance issues for the generated code.
27282728
//CHECK-NEXT: */
27292729
//CHECK-NEXT: sycl::floor(f + 0.5);
27302730
//CHECK-NEXT: /*
@@ -2761,9 +2761,9 @@ __global__ void do_migration6() {
27612761
//CHECK-NEXT: std::min(i, i);
27622762
//CHECK-NEXT: sycl::fabs(f);
27632763
//CHECK-NEXT: /*
2764-
//CHECK-NEXT: DPCT1017:{{[0-9]+}}: The sycl::floor call is used instead of the nearbyintf call. These two calls do not provide exactly the same functionality. Check the potential precision and/or performance issues for the generated code.
2764+
//CHECK-NEXT: DPCT1017:{{[0-9]+}}: The sycl::floor call is used instead of the std::nearbyintf call. These two calls do not provide exactly the same functionality. Check the potential precision and/or performance issues for the generated code.
27652765
//CHECK-NEXT: */
2766-
//CHECK-NEXT: sycl::floor(f + 0.5);
2766+
//CHECK-NEXT: sycl::floor(f + 0.5f);
27672767
//CHECK-NEXT: /*
27682768
//CHECK-NEXT: DPCT1017:{{[0-9]+}}: The sycl::remquo call is used instead of the remquof call. These two calls do not provide exactly the same functionality. Check the potential precision and/or performance issues for the generated code.
27692769
//CHECK-NEXT: */
@@ -2780,7 +2780,7 @@ __global__ void do_migration6() {
27802780
//CHECK-NEXT: */
27812781
//CHECK-NEXT: sycl::modf(f, sycl::address_space_cast<sycl::access::address_space::generic_space, sycl::access::decorated::yes>(&f));
27822782
//CHECK-NEXT: /*
2783-
//CHECK-NEXT: DPCT1017:{{[0-9]+}}: The sycl::floor call is used instead of the nearbyint call. These two calls do not provide exactly the same functionality. Check the potential precision and/or performance issues for the generated code.
2783+
//CHECK-NEXT: DPCT1017:{{[0-9]+}}: The sycl::floor call is used instead of the std::nearbyint call. These two calls do not provide exactly the same functionality. Check the potential precision and/or performance issues for the generated code.
27842784
//CHECK-NEXT: */
27852785
//CHECK-NEXT: sycl::floor(f + 0.5);
27862786
//CHECK-NEXT: /*
@@ -2817,9 +2817,9 @@ __device__ __host__ void do_migration7() {
28172817
//CHECK-NEXT: std::min(i, i);
28182818
//CHECK-NEXT: sycl::fabs(f);
28192819
//CHECK-NEXT: /*
2820-
//CHECK-NEXT: DPCT1017:{{[0-9]+}}: The sycl::floor call is used instead of the nearbyintf call. These two calls do not provide exactly the same functionality. Check the potential precision and/or performance issues for the generated code.
2820+
//CHECK-NEXT: DPCT1017:{{[0-9]+}}: The sycl::floor call is used instead of the std::nearbyintf call. These two calls do not provide exactly the same functionality. Check the potential precision and/or performance issues for the generated code.
28212821
//CHECK-NEXT: */
2822-
//CHECK-NEXT: sycl::floor(f + 0.5);
2822+
//CHECK-NEXT: sycl::floor(f + 0.5f);
28232823
//CHECK-NEXT: /*
28242824
//CHECK-NEXT: DPCT1017:{{[0-9]+}}: The sycl::remquo call is used instead of the remquof call. These two calls do not provide exactly the same functionality. Check the potential precision and/or performance issues for the generated code.
28252825
//CHECK-NEXT: */
@@ -2835,7 +2835,7 @@ __device__ __host__ void do_migration7() {
28352835
//CHECK-NEXT: */
28362836
//CHECK-NEXT: sycl::modf(f, sycl::address_space_cast<sycl::access::address_space::generic_space, sycl::access::decorated::yes>(&f));
28372837
//CHECK-NEXT: /*
2838-
//CHECK-NEXT: DPCT1017:{{[0-9]+}}: The sycl::floor call is used instead of the nearbyint call. These two calls do not provide exactly the same functionality. Check the potential precision and/or performance issues for the generated code.
2838+
//CHECK-NEXT: DPCT1017:{{[0-9]+}}: The sycl::floor call is used instead of the std::nearbyint call. These two calls do not provide exactly the same functionality. Check the potential precision and/or performance issues for the generated code.
28392839
//CHECK-NEXT: */
28402840
//CHECK-NEXT: sycl::floor(f + 0.5);
28412841
//CHECK-NEXT: /*

‎clang/test/dpct/math_functions_std.cu

Lines changed: 9 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -31,6 +31,15 @@ __device__ void d() {
3131
i = std::abs(i2);
3232
l = std::abs(l2);
3333
ll = std::abs(ll2);
34+
35+
//CHECK: std::nearbyintf(f);
36+
//CHECK-NEXT: std::nearbyint(d);
37+
//CHECK-NEXT: std::nearbyintf(f);
38+
//CHECK-NEXT: std::nearbyint(d);
39+
nearbyintf(f);
40+
nearbyint(d);
41+
std::nearbyintf(f);
42+
std::nearbyint(d);
3443
}
3544

3645
void h() {

‎clang/test/dpct/query_api_mapping/Math/test_single_precision_mathematical_functions.cu

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -334,7 +334,7 @@
334334
// NEARBYINTF: CUDA API:
335335
// NEARBYINTF-NEXT: nearbyintf(f /*float*/);
336336
// NEARBYINTF-NEXT: Is migrated to:
337-
// NEARBYINTF-NEXT: sycl::floor(f + 0.5);
337+
// NEARBYINTF-NEXT: sycl::floor(f + 0.5f);
338338

339339
// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=nextafterf | FileCheck %s -check-prefix=NEXTAFTERF
340340
// NEXTAFTERF: CUDA API:

0 commit comments

Comments
 (0)