From 2c50caafe1b1cc7e7518bb35befa01a69aa50705 Mon Sep 17 00:00:00 2001 From: "Jiang, Zhiwei" Date: Fri, 16 May 2025 09:47:11 +0800 Subject: [PATCH] WIP Signed-off-by: Jiang, Zhiwei --- clang/lib/DPCT/RuleInfra/CallExprRewriter.h | 8 ++++ clang/lib/DPCT/RulesLang/APINamesMath.inc | 4 +- .../RulesLang/Math/CallExprRewriterMath.cpp | 3 -- ...erDoublePrecisionMathematicalFunctions.cpp | 16 ++++++++ .../RulesLang/Math/RewriterSTDFunctions.cpp | 38 ++++++++++++++++++- ...erSinglePrecisionMathematicalFunctions.cpp | 17 +++++++++ clang/test/dpct/math/cuda-math-intrinsics.cu | 20 +++++----- ...single_precision_mathematical_functions.cu | 2 +- 8 files changed, 91 insertions(+), 17 deletions(-) diff --git a/clang/lib/DPCT/RuleInfra/CallExprRewriter.h b/clang/lib/DPCT/RuleInfra/CallExprRewriter.h index b8d859a0e357..3db3d8923243 100644 --- a/clang/lib/DPCT/RuleInfra/CallExprRewriter.h +++ b/clang/lib/DPCT/RuleInfra/CallExprRewriter.h @@ -1141,6 +1141,14 @@ template class ParenExprPrinter { public: ParenExprPrinter(ET &&E) : E(std::forward(E)) {} template void print(StreamT &Stream) const { + if constexpr (std::is_base_of_v< + clang::Stmt, + std::remove_cv_t>>) { + if (isa(E->IgnoreImpCasts())) { + dpct::print(Stream, E); + return; + } + } PairedPrinter PP(Stream, "(", ")"); dpct::print(Stream, E); } diff --git a/clang/lib/DPCT/RulesLang/APINamesMath.inc b/clang/lib/DPCT/RulesLang/APINamesMath.inc index 052947aab571..a8707d92c1db 100644 --- a/clang/lib/DPCT/RulesLang/APINamesMath.inc +++ b/clang/lib/DPCT/RulesLang/APINamesMath.inc @@ -253,7 +253,7 @@ ENTRY_RENAMED_NO_REWRITE("__umul64hi", MapNames::getClNamespace(false, true) + " ENTRY_EMULATED("frexpf", MapNames::getClNamespace(false, true) + "frexp") ENTRY_EMULATED("modff", MapNames::getClNamespace(false, true) + "modf") ENTRY_EMULATED("nanf", MapNames::getClNamespace(false, true) + "nan") -ENTRY_EMULATED("nearbyintf", MapNames::getClNamespace(false, true) + "floor") +ENTRY_REWRITE("nearbyintf") ENTRY_EMULATED("remquof", MapNames::getClNamespace(false, true) + "remquo") ENTRY_EMULATED("rhypotf", MapNames::getClNamespace(false, true) + "hypot") ENTRY_REWRITE("sincosf") @@ -266,7 +266,7 @@ ENTRY_EMULATED("__powf", MapNames::getDpctNamespace() + "pow") ENTRY_EMULATED("frexp", MapNames::getClNamespace(false, true) + "frexp") ENTRY_EMULATED("modf", MapNames::getClNamespace(false, true) + "modf") ENTRY_EMULATED("nan", MapNames::getClNamespace(false, true) + "nan") -ENTRY_EMULATED("nearbyint", MapNames::getClNamespace(false, true) + "floor") +ENTRY_REWRITE("nearbyint") ENTRY_EMULATED("remquo", MapNames::getClNamespace(false, true) + "remquo") ENTRY_EMULATED("rhypot", MapNames::getClNamespace(false, true) + "hypot") ENTRY_REWRITE("sincos") diff --git a/clang/lib/DPCT/RulesLang/Math/CallExprRewriterMath.cpp b/clang/lib/DPCT/RulesLang/Math/CallExprRewriterMath.cpp index 6e33118d45c2..0444ecc0b5b0 100644 --- a/clang/lib/DPCT/RulesLang/Math/CallExprRewriterMath.cpp +++ b/clang/lib/DPCT/RulesLang/Math/CallExprRewriterMath.cpp @@ -455,9 +455,6 @@ std::optional MathSimulatedRewriter::rewrite() { << MapNames::getClNamespace() + "access::address_space::generic_space, " << MapNames::getClNamespace() + "access::decorated::yes>(" << MigratedArg2 << "))"; - } else if (FuncName == "nearbyint" || FuncName == "nearbyintf") { - OS << MapNames::getClNamespace(false, true) + "floor(" << MigratedArg0 - << " + 0.5)"; } else if (FuncName == "rhypot" || FuncName == "rhypotf") { auto MigratedArg1 = getMigratedArg(1); OS << "1 / " + MapNames::getClNamespace(false, true) + "hypot(" diff --git a/clang/lib/DPCT/RulesLang/Math/RewriterDoublePrecisionMathematicalFunctions.cpp b/clang/lib/DPCT/RulesLang/Math/RewriterDoublePrecisionMathematicalFunctions.cpp index fdc5f2257fd3..245d95437b58 100644 --- a/clang/lib/DPCT/RulesLang/Math/RewriterDoublePrecisionMathematicalFunctions.cpp +++ b/clang/lib/DPCT/RulesLang/Math/RewriterDoublePrecisionMathematicalFunctions.cpp @@ -349,6 +349,22 @@ RewriterMap dpct::createDoublePrecisionMathematicalFunctionsRewriterMap() { makeArgWithAddressSpaceCast(2)))), Diagnostics::MATH_EMULATION, std::string("sincospi"), MapNames::getClNamespace() + std::string("sincos"))) + // nearbyint + MATH_API_REWRITERS_V2( + "nearbyint", + MATH_API_REWRITER_PAIR( + math::Tag::device_std, + CALL_FACTORY_ENTRY("nearbyint", CALL("std::nearbyint", ARG(0)))), + MATH_API_REWRITER_PAIR( + math::Tag::emulation, + WARNING_FACTORY_ENTRY( + "nearbyint", + CALL_FACTORY_ENTRY("nearbyint", + CALL(MapNames::getClNamespace() + "floor", + BO(BinaryOperatorKind::BO_Add, + PAREN(ARG(0)), ARG("0.5")))), + Diagnostics::MATH_EMULATION, std::string("nearbyint"), + MapNames::getClNamespace() + std::string("floor")))) // sinpi CALL_FACTORY_ENTRY( "sinpi", diff --git a/clang/lib/DPCT/RulesLang/Math/RewriterSTDFunctions.cpp b/clang/lib/DPCT/RulesLang/Math/RewriterSTDFunctions.cpp index 780bb693dd14..fdb487d62085 100644 --- a/clang/lib/DPCT/RulesLang/Math/RewriterSTDFunctions.cpp +++ b/clang/lib/DPCT/RulesLang/Math/RewriterSTDFunctions.cpp @@ -55,5 +55,41 @@ RewriterMap dpct::createSTDFunctionsRewriterMap() { ARG(0)))), MATH_API_REWRITER_PAIR( math::Tag::device_std, - CALL_FACTORY_ENTRY("std::fabs", CALL("std::fabs", ARG(0)))))}; + CALL_FACTORY_ENTRY("std::fabs", CALL("std::fabs", ARG(0))))) + MATH_API_REWRITERS_V2( + "std::nearbyint", + MATH_API_REWRITER_PAIR( + math::Tag::device_std, + CALL_FACTORY_ENTRY("std::nearbyint", + CALL("std::nearbyint", ARG(0)))), + MATH_API_REWRITER_PAIR( + math::Tag::emulation, + WARNING_FACTORY_ENTRY( + "std::nearbyint", + CALL_FACTORY_ENTRY( + "std::nearbyint", + CALL(MapNames::getClNamespace() + "floor", + BO(BinaryOperatorKind::BO_Add, PAREN(ARG(0)), + ARG("0.5")))), + Diagnostics::MATH_EMULATION, + std::string("std::nearbyint"), + MapNames::getClNamespace() + std::string("floor")))) + MATH_API_REWRITERS_V2( + "std::nearbyintf", + MATH_API_REWRITER_PAIR( + math::Tag::device_std, + CALL_FACTORY_ENTRY("std::nearbyintf", + CALL("std::nearbyintf", ARG(0)))), + MATH_API_REWRITER_PAIR( + math::Tag::emulation, + WARNING_FACTORY_ENTRY( + "std::nearbyintf", + CALL_FACTORY_ENTRY( + "std::nearbyintf", + CALL(MapNames::getClNamespace() + "floor", + BO(BinaryOperatorKind::BO_Add, PAREN(ARG(0)), + ARG("0.5f")))), + Diagnostics::MATH_EMULATION, + std::string("std::nearbyintf"), + MapNames::getClNamespace() + std::string("floor"))))}; } diff --git a/clang/lib/DPCT/RulesLang/Math/RewriterSinglePrecisionMathematicalFunctions.cpp b/clang/lib/DPCT/RulesLang/Math/RewriterSinglePrecisionMathematicalFunctions.cpp index bd9a2727c627..082f690ace17 100644 --- a/clang/lib/DPCT/RulesLang/Math/RewriterSinglePrecisionMathematicalFunctions.cpp +++ b/clang/lib/DPCT/RulesLang/Math/RewriterSinglePrecisionMathematicalFunctions.cpp @@ -308,6 +308,23 @@ RewriterMap dpct::createSinglePrecisionMathematicalFunctionsRewriterMap() { makeArgWithAddressSpaceCast(2)))), Diagnostics::MATH_EMULATION, std::string("sincospif"), MapNames::getClNamespace() + std::string("sincos")) + // nearbyintf + MATH_API_REWRITERS_V2( + "nearbyintf", + MATH_API_REWRITER_PAIR( + math::Tag::device_std, + CALL_FACTORY_ENTRY("nearbyintf", + CALL("std::nearbyintf", ARG(0)))), + MATH_API_REWRITER_PAIR( + math::Tag::emulation, + WARNING_FACTORY_ENTRY( + "nearbyintf", + CALL_FACTORY_ENTRY("nearbyintf", + CALL(MapNames::getClNamespace() + "floor", + BO(BinaryOperatorKind::BO_Add, + PAREN(ARG(0)), ARG("0.5f")))), + Diagnostics::MATH_EMULATION, std::string("nearbyintf"), + MapNames::getClNamespace() + std::string("floor")))) // sinpif MATH_API_REWRITERS_V2( "sinpif", diff --git a/clang/test/dpct/math/cuda-math-intrinsics.cu b/clang/test/dpct/math/cuda-math-intrinsics.cu index ca94098a8895..7a87d66e3322 100644 --- a/clang/test/dpct/math/cuda-math-intrinsics.cu +++ b/clang/test/dpct/math/cuda-math-intrinsics.cu @@ -2115,7 +2115,7 @@ __global__ void testSimulation() { // CHECK: /* // 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. // CHECK-NEXT: */ - // CHECK-NEXT: f = sycl::floor(f + 0.5); + // CHECK-NEXT: f = sycl::floor(f + 0.5f); f = nearbyintf(f); // CHECK: /* @@ -2705,9 +2705,9 @@ __device__ void do_migration5() { //CHECK-NEXT: std::min(i, i); //CHECK-NEXT: sycl::fabs(f); //CHECK-NEXT: /* - //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. + //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. //CHECK-NEXT: */ - //CHECK-NEXT: sycl::floor(f + 0.5); + //CHECK-NEXT: sycl::floor(f + 0.5f); //CHECK-NEXT: /* //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. //CHECK-NEXT: */ @@ -2724,7 +2724,7 @@ __device__ void do_migration5() { //CHECK-NEXT: */ //CHECK-NEXT: sycl::modf(f, sycl::address_space_cast(&f)); //CHECK-NEXT: /* - //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. + //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. //CHECK-NEXT: */ //CHECK-NEXT: sycl::floor(f + 0.5); //CHECK-NEXT: /* @@ -2761,9 +2761,9 @@ __global__ void do_migration6() { //CHECK-NEXT: std::min(i, i); //CHECK-NEXT: sycl::fabs(f); //CHECK-NEXT: /* - //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. + //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. //CHECK-NEXT: */ - //CHECK-NEXT: sycl::floor(f + 0.5); + //CHECK-NEXT: sycl::floor(f + 0.5f); //CHECK-NEXT: /* //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. //CHECK-NEXT: */ @@ -2780,7 +2780,7 @@ __global__ void do_migration6() { //CHECK-NEXT: */ //CHECK-NEXT: sycl::modf(f, sycl::address_space_cast(&f)); //CHECK-NEXT: /* - //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. + //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. //CHECK-NEXT: */ //CHECK-NEXT: sycl::floor(f + 0.5); //CHECK-NEXT: /* @@ -2817,9 +2817,9 @@ __device__ __host__ void do_migration7() { //CHECK-NEXT: std::min(i, i); //CHECK-NEXT: sycl::fabs(f); //CHECK-NEXT: /* - //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. + //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. //CHECK-NEXT: */ - //CHECK-NEXT: sycl::floor(f + 0.5); + //CHECK-NEXT: sycl::floor(f + 0.5f); //CHECK-NEXT: /* //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. //CHECK-NEXT: */ @@ -2835,7 +2835,7 @@ __device__ __host__ void do_migration7() { //CHECK-NEXT: */ //CHECK-NEXT: sycl::modf(f, sycl::address_space_cast(&f)); //CHECK-NEXT: /* - //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. + //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. //CHECK-NEXT: */ //CHECK-NEXT: sycl::floor(f + 0.5); //CHECK-NEXT: /* diff --git a/clang/test/dpct/query_api_mapping/Math/test_single_precision_mathematical_functions.cu b/clang/test/dpct/query_api_mapping/Math/test_single_precision_mathematical_functions.cu index a150ee042042..16d4be32392b 100644 --- a/clang/test/dpct/query_api_mapping/Math/test_single_precision_mathematical_functions.cu +++ b/clang/test/dpct/query_api_mapping/Math/test_single_precision_mathematical_functions.cu @@ -334,7 +334,7 @@ // NEARBYINTF: CUDA API: // NEARBYINTF-NEXT: nearbyintf(f /*float*/); // NEARBYINTF-NEXT: Is migrated to: -// NEARBYINTF-NEXT: sycl::floor(f + 0.5); +// NEARBYINTF-NEXT: sycl::floor(f + 0.5f); // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=nextafterf | FileCheck %s -check-prefix=NEXTAFTERF // NEXTAFTERF: CUDA API: