[SYCLomaitc] Use the rewrite to impl cospif (#2715)

Signed-off-by: Chen, Sheng S <sheng.s.chen@intel.com>

Author: ShengchenJ

clang/lib/DPCT/RulesLang/APINamesMath.inc

@@ -62,7 +62,7 @@ ENTRY_REWRITE("ceilf")
 ENTRY_RENAMED_SINGLE("copysignf", MapNames::getClNamespace(false, true) + "copysign")
 ENTRY_RENAMED_SINGLE("cosf", MapNames::getClNamespace(false, true) + "cos")
 ENTRY_RENAMED_SINGLE("coshf", MapNames::getClNamespace(false, true) + "cosh")
-ENTRY_RENAMED_SINGLE("cospif", MapNames::getClNamespace(false, true) + "cospi")
+ENTRY_REWRITE("cospif")
 ENTRY_RENAMED_SINGLE("erfcf", MapNames::getClNamespace(false, true) + "erfc")
 ENTRY_RENAMED_SINGLE("erff", MapNames::getClNamespace(false, true) + "erf")
 ENTRY_RENAMED_SINGLE("exp10f", MapNames::getClNamespace(false, true) + "exp10")

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

@@ -251,6 +251,7 @@ enum class Tag : size_t {
   ext_experimental, // device API using experimental feature
   host_perf,        // host API for performance
   host_normal,      // host API
+  host_device,      // host deivce API
   unsupported_warning,
   no_rewrite,
   tag_size
@@ -280,6 +281,8 @@ class MathRewriterFactory final : public CallExprRewriterFactoryBase {
       MathAPIRewriters[static_cast<size_t>(math::Tag::math_libdevice)];
   element_t &DeviceStdRewriter =
       MathAPIRewriters[static_cast<size_t>(math::Tag::device_std)];
+  element_t &HostDeviceRewriter =
+      MathAPIRewriters[static_cast<size_t>(math::Tag::host_device)];
   element_t &EmulationRewriter =
       MathAPIRewriters[static_cast<size_t>(math::Tag::emulation)];
   element_t &ExtExperimentalRewriter =
@@ -378,6 +381,9 @@ class MathRewriterFactory final : public CallExprRewriterFactoryBase {
     if (EmulationRewriter && EmulationRewriter.value().first(C))
       return EmulationRewriter.value().second.second->create(C);
 
+    if (HostDeviceRewriter && HostDeviceRewriter.value().first(C))
+      return HostDeviceRewriter.value().second.second->create(C);
+
     if (UnsupportedWarningRewriter &&
         UnsupportedWarningRewriter.value().first(C))
       return UnsupportedWarningRewriter.value().second.second->create(C);

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

@@ -295,6 +295,14 @@ RewriterMap dpct::createSinglePrecisionMathematicalFunctionsRewriterMap() {
                   "sinpif",
                   CALL(MapNames::getClNamespace() + "sinpi",
                        CAST_IF_NOT_SAME(makeLiteral("float"), ARG(0))))))
+      // cospif
+      MATH_API_REWRITERS_V2(
+          "cospif", MATH_API_REWRITER_PAIR(
+                        math::Tag::host_device,
+                        CALL_FACTORY_ENTRY(
+                            "cospif", CALL(MapNames::getClNamespace() + "cospi",
+                                           CAST_IF_NOT_SAME(
+                                               makeLiteral("float"), ARG(0))))))
       // tanhf
       MATH_API_REWRITERS_V2(
           "tanhf",

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

@@ -3539,3 +3539,14 @@ __global__ void foo7(float aa) {
   // CHECK: sycl::sinpi(aa);
   ::sinpif(aa);
 }
+
+void foo8(float aa) {
+  // CHECK: sycl::cospi(aa);
+  ::cospif(aa);
+}
+
+__global__ void foo9(float aa) {
+  // CHECK: sycl::cospi(aa);
+  ::cospif(aa);
+}
+