[SYCLomatic] Refine the impl of nearbyint

clang/lib/DPCT/RuleInfra/CallExprRewriter.h

@@ -1141,6 +1141,14 @@ template <class ET> class ParenExprPrinter {
 public:
   ParenExprPrinter(ET &&E) : E(std::forward<ET>(E)) {}
   template <class StreamT> void print(StreamT &Stream) const {
+    if constexpr (std::is_base_of_v<
+                      clang::Stmt,
+                      std::remove_cv_t<std::remove_pointer_t<ET>>>) {
+      if (isa<DeclRefExpr>(E->IgnoreImpCasts())) {
+        dpct::print(Stream, E);
+        return;
+      }
+    }
     PairedPrinter PP(Stream, "(", ")");
     dpct::print(Stream, E);
   }

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")

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

@@ -455,9 +455,6 @@ std::optional<std::string> 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("

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",

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"))))};
 }

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",

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<sycl::access::address_space::generic_space, sycl::access::decorated::yes>(&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<sycl::access::address_space::generic_space, sycl::access::decorated::yes>(&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<sycl::access::address_space::generic_space, sycl::access::decorated::yes>(&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: /*

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: