[SYCLomatic] Add an empty rewriter in case no migration required for math API (#2846)

Signed-off-by: Jiang, Zhiwei <zhiwei.jiang@intel.com>

Author: zhiweij1

clang/lib/DPCT/RuleInfra/CallExprRewriter.h

@@ -586,6 +586,14 @@ class NoRewriteFuncNameRewriter : public CallExprRewriter {
   std::optional<std::string> rewrite() override { return NewFuncName; }
 };
 
+// No replacement generated
+class EmptyRewriter : public CallExprRewriter {
+public:
+  EmptyRewriter(const CallExpr *, StringRef, StringRef)
+      : CallExprRewriter(Call, SourceCalleeName) {}
+  std::optional<std::string> rewrite() override { return std::nullopt; }
+};
+
 struct ThrustFunctor {
   ThrustFunctor(const clang::Expr *E) : E(E) {}
   const clang::Expr *E;

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

@@ -31,6 +31,8 @@ using WarpFunctionRewriterFactory =
     CallExprRewriterFactory<WarpFunctionRewriter, std::string>;
 using NoRewriteFuncNameRewriterFactory =
     CallExprRewriterFactory<NoRewriteFuncNameRewriter, std::string>;
+using EmptyRewriterFactory =
+    CallExprRewriterFactory<EmptyRewriter, std::string>;
 
 /// Base class for rewriting math function calls
 class MathCallExprRewriter : public FuncCallExprRewriter {
@@ -349,10 +351,9 @@ class MathRewriterFactory final : public CallExprRewriterFactoryBase {
       : Name(Name), MathAPIRewriters(MathAPIRewritersInput) {
     NoRewriteRewriter = std::make_pair(
         TrueFunctor,
-        std::make_pair(Name,
-                       std::dynamic_pointer_cast<CallExprRewriterFactoryBase>(
-                           std::make_shared<NoRewriteFuncNameRewriterFactory>(
-                               Name, Name))));
+        std::make_pair(
+            Name, std::dynamic_pointer_cast<CallExprRewriterFactoryBase>(
+                      std::make_shared<EmptyRewriterFactory>(Name, Name))));
   }
   // a. Host API priority:
   //   1. host_perf
@@ -561,22 +562,19 @@ createMathAPIRewriterDevice(
           math::IsDefinedInCUDA(),
           std::move(createMathAPIRewriterDeviceImpl(Name, PerfPred, DevicePerf,
                                                     DeviceNodes)),
-          {Name,
-           std::make_shared<NoRewriteFuncNameRewriterFactory>(Name, Name)}),
+          {Name, std::make_shared<EmptyRewriterFactory>(Name, Name)}),
       createConditionalFactory(
           math::IsUnresolvedLookupExpr,
           createConditionalFactory(
               math::IsDirectCallerPureDevice,
               std::move(createMathAPIRewriterDeviceImpl(
                   Name, PerfPred, DevicePerf, DeviceNodes)),
-              {Name,
-               std::make_shared<NoRewriteFuncNameRewriterFactory>(Name, Name)}),
+              {Name, std::make_shared<EmptyRewriterFactory>(Name, Name)}),
           createConditionalFactory(
               math::IsDefinedInCUDA(),
               std::move(createMathAPIRewriterDeviceImpl(
                   Name, PerfPred, DevicePerf, DeviceNodes)),
-              {Name, std::make_shared<NoRewriteFuncNameRewriterFactory>(
-                         Name, Name)})));
+              {Name, std::make_shared<EmptyRewriterFactory>(Name, Name)})));
 }
 
 inline std::pair<std::string, std::shared_ptr<CallExprRewriterFactoryBase>>
@@ -590,21 +588,17 @@ createMathAPIRewriterDevice(
       createConditionalFactory(
           math::IsDefinedInCUDA(),
           std::move(createMathAPIRewriterDeviceImpl(Name, DeviceNodes)),
-          {Name,
-           std::make_shared<NoRewriteFuncNameRewriterFactory>(Name, Name)}),
+          {Name, std::make_shared<EmptyRewriterFactory>(Name, Name)}),
       createConditionalFactory(
           math::IsUnresolvedLookupExpr,
           createConditionalFactory(
               math::IsDirectCallerPureDevice,
               std::move(createMathAPIRewriterDeviceImpl(Name, DeviceNodes)),
-              {Name,
-               std::make_shared<NoRewriteFuncNameRewriterFactory>(Name, Name)}),
+              {Name, std::make_shared<EmptyRewriterFactory>(Name, Name)}),
           createConditionalFactory(
               math::IsDefinedInCUDA(),
-              std::move(
-                  createMathAPIRewriterDeviceImpl(Name, DeviceNodes)),
-              {Name, std::make_shared<NoRewriteFuncNameRewriterFactory>(
-                         Name, Name)})));
+              std::move(createMathAPIRewriterDeviceImpl(Name, DeviceNodes)),
+              {Name, std::make_shared<EmptyRewriterFactory>(Name, Name)})));
 }
 
 template <class T>
@@ -620,13 +614,11 @@ createMathAPIRewriterExperimentalBfloat16(
     if (math::useExtBFloat16Math() && Rewriter1.second)
       return createConditionalFactory(
           math::IsDefinedInCUDA(), std::move(Rewriter1),
-          {Name,
-           std::make_shared<NoRewriteFuncNameRewriterFactory>(Name, Name)});
+          {Name, std::make_shared<EmptyRewriterFactory>(Name, Name)});
     if (Rewriter2.second)
       return createConditionalFactory(
           math::IsDefinedInCUDA(), std::move(Rewriter2),
-          {Name,
-           std::make_shared<NoRewriteFuncNameRewriterFactory>(Name, Name)});
+          {Name, std::make_shared<EmptyRewriterFactory>(Name, Name)});
   }
   // report unsupport
   return std::pair<std::string, std::shared_ptr<CallExprRewriterFactoryBase>>(
@@ -643,7 +635,7 @@ createMathAPIRewriterHost(
     T) {
   return createConditionalFactory(
       math::IsDefinedInCUDA(), std::move(HostNormal),
-      {Name, std::make_shared<NoRewriteFuncNameRewriterFactory>(Name, Name)});
+      {Name, std::make_shared<EmptyRewriterFactory>(Name, Name)});
 }
 
 template <class T>
@@ -660,7 +652,7 @@ createMathAPIRewriterHost(
       math::IsDefinedInCUDA(),
       createConditionalFactory(makeCheckAnd(math::IsPerf, PerfPred),
                                std::move(HostPerf), std::move(HostNormal)),
-      {Name, std::make_shared<NoRewriteFuncNameRewriterFactory>(Name, Name)});
+      {Name, std::make_shared<EmptyRewriterFactory>(Name, Name)});
 }
 
 template <bool IsDouble> std::string getPiString() {

clang/test/dpct/complex.cu

@@ -17,9 +17,6 @@
 // CHECK-NEXT: #define COMPLEX_D_DIV(a, b) dpct::cdiv<double>(a, b)
 // CHECK-NEXT: #define COMPLEX_D_FMA(a, b, c) dpct::cmul<double>(a, b) + c
 // CHECK-NEXT: #define COMPLEX_D_ABS(a) dpct::cabs<double>(a)
-// CHECK-NEXT: /*
-// CHECK-NEXT: DPCT1064:{{[0-9]+}}: Migrated fabs call is used in a macro/template definition and may not be valid for all macro/template uses. Adjust the code.
-// CHECK-NEXT: */
 // CHECK-NEXT: #define COMPLEX_D_ABS1(a) (sycl::fabs((a).x()) + sycl::fabs((a).y()))
 // CHECK-NEXT: #define COMPLEX_D_CONJ(a) dpct::conj<double>(a)
 #define COMPLEX_D_MAKE(r,i) make_cuDoubleComplex(r, i)
@@ -47,9 +44,6 @@
 // CHECK-NEXT: #define COMPLEX_F_DIV(a, b) dpct::cdiv<float>(a, b)
 // CHECK-NEXT: #define COMPLEX_F_FMA(a, b, c) dpct::cmul<float>(a, b) + c
 // CHECK-NEXT: #define COMPLEX_F_ABS(a) dpct::cabs<float>(a)
-// CHECK-NEXT: /*
-// CHECK-NEXT: DPCT1064:{{[0-9]+}}: Migrated fabsf call is used in a macro/template definition and may not be valid for all macro/template uses. Adjust the code.
-// CHECK-NEXT: */
 // CHECK-NEXT: #define COMPLEX_F_ABS1(a) (sycl::fabs((a).x()) + sycl::fabs((a).y()))
 // CHECK-NEXT: #define COMPLEX_F_CONJ(a) dpct::conj<float>(a)
 #define COMPLEX_F_MAKE(r,i) make_cuFloatComplex(r, i)

clang/test/dpct/cuda_arch_test/math_in_cuda_arch.cu

@@ -4,7 +4,7 @@
 
 // CHECK: float f(float x) {
 // CHECK-EMPTY: 
-// CHECK-NEXT:   return expf(x);
+// CHECK-NEXT:   return ::expf(x);
 // CHECK-EMPTY: 
 // CHECK-NEXT: }
 // CHECK-NEXT: float f_host_ct0(float x) {

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

@@ -3584,3 +3584,15 @@ template <typename T> class AAA_st {
     return 0;
   }
 };
+
+template <typename t1> void saturate() { printf("default saturate\n"); }
+template <> void saturate<int>() { printf("int saturate\n"); }
+
+// CHECK: void foo11() {
+// CHECK-NEXT:   saturate<int>();
+// CHECK-NEXT:   saturate<float>();
+// CHECK-NEXT: }
+void foo11() {
+  saturate<int>();
+  saturate<float>();
+}