[SYCLomatic] Don't migrate API defined by user and refine unresolved decl API migration (#2781)

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

Author: zhiweij1

clang/lib/DPCT/RuleInfra/CallExprRewriterCommon.h

@@ -2132,6 +2132,16 @@ class IsDefinedInCUDA {
     return isFromCUDA(FD);
   }
 };
+class IsDefinedByUser {
+public:
+  IsDefinedByUser() {}
+  bool operator()(const CallExpr *C) {
+    auto FD = C->getDirectCallee();
+    if (!FD)
+      return false;
+    return DpctGlobalInfo::isInRoot(FD->getLocation());
+  }
+};
 } // namespace math
 } // namespace dpct
 

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

@@ -388,29 +388,29 @@ class MathRewriterFactory final : public CallExprRewriterFactoryBase {
             return Rewriter.value();
         }
       }
-      if (math::IsUnresolvedLookupExpr(C)) {
-        if (math::IsDirectCallerPureDevice(C)) {
-          if (Rewriter = getDeviceRewriter(C))
-            return Rewriter.value();
-        }
-      }
       if (math::IsDefinedInCUDA()(C)) {
         if (Rewriter = getDeviceRewriter(C))
           return Rewriter.value();
       }
     }
-
-    // Host and device
-    if (HostDeviceRewriter && HostDeviceRewriter.value().first(C))
-      return HostDeviceRewriter.value().second.second->create(C);
-
-    if (EmulationRewriter && EmulationRewriter.value().first(C))
-      return EmulationRewriter.value().second.second->create(C);
-
-    if (UnsupportedWarningRewriter &&
-        UnsupportedWarningRewriter.value().first(C))
-      return UnsupportedWarningRewriter.value().second.second->create(C);
-
+    if (math::IsUnresolvedLookupExpr(C)) {
+      if (math::IsDirectCallerPureDevice(C)) {
+        if (auto Rewriter = getDeviceRewriter(C))
+          return Rewriter.value();
+      } else if (math::IsDirectCallerPureHost(C)) {
+        return NoRewriteRewriter.value().second.second->create(C);
+      }
+    }
+    if (!math::IsDefinedByUser()(C)) {
+      // Host and device
+      if (HostDeviceRewriter && HostDeviceRewriter.value().first(C))
+        return HostDeviceRewriter.value().second.second->create(C);
+      if (EmulationRewriter && EmulationRewriter.value().first(C))
+        return EmulationRewriter.value().second.second->create(C);
+      if (UnsupportedWarningRewriter &&
+          UnsupportedWarningRewriter.value().first(C))
+        return UnsupportedWarningRewriter.value().second.second->create(C);
+    }
     return NoRewriteRewriter.value().second.second->create(C);
   }
 };

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

@@ -3550,3 +3550,36 @@ __global__ void foo9(float aa) {
   ::cospif(aa);
 }
 
+inline __host__ __device__ float4 operator-(float b, float4 a) {
+  return make_float4(b - a.x, b - a.y, b - a.z, b - a.w);
+}
+
+inline __host__ __device__ float4 fabs(float4 v) {
+  return make_float4(fabs(v.x), fabs(v.y), fabs(v.z), fabs(v.w));
+}
+
+__device__ void foo10(float4 a, float b) {
+  // CHECK: sycl::float4 c = fabs(dpct_operator_overloading::operator-(b , a));
+  float4 c = fabs(b - a);
+}
+
+// CHECK: template <typename T> class AAA_st {
+// CHECK-NEXT:   struct BBB_st {
+// CHECK-NEXT:     float f;
+// CHECK-NEXT:   };
+// CHECK-NEXT:   int foo() {
+// CHECK-NEXT:     BBB_st *ptr = new BBB_st;
+// CHECK-NEXT:     fabs(ptr->f);
+// CHECK-NEXT:     return 0;
+// CHECK-NEXT:   }
+// CHECK-NEXT: };
+template <typename T> class AAA_st {
+  struct BBB_st {
+    float f;
+  };
+  int foo() {
+    BBB_st *ptr = new BBB_st;
+    fabs(ptr->f);
+    return 0;
+  }
+};