[SYCLomatic] Do not migrate user-defined API which has the same name as CUDA texture API (#2802)

Signed-off-by: Jiang, Zhiwei <[email protected]>

Author: zhiweij1

clang/lib/DPCT/RuleInfra/CallExprRewriterCommon.h

@@ -2141,7 +2141,7 @@ class IsDefinedByUser {
     auto FD = C->getDirectCallee();
     if (!FD)
       return false;
-    return DpctGlobalInfo::isInRoot(FD->getLocation());
+    return isUserDefinedDecl(FD);
   }
 };
 } // namespace math

clang/lib/DPCT/RulesLang/RulesLangTexture.cpp

@@ -780,10 +780,12 @@ bool TextureRule::processTexVarDeclInDevice(const VarDecl *VD) {
 }
 
 void TextureRule::runRule(const MatchFinder::MatchResult &Result) {
-
   if (getAssistNodeAsType<UnresolvedLookupExpr>(Result,
                                                 "unresolvedLookupExpr")) {
     const CallExpr *CE = getAssistNodeAsType<CallExpr>(Result, "callExpr");
+    if (const auto *FD = CE->getDirectCallee())
+      if (isUserDefinedDecl(FD))
+        return;
     ExprAnalysis A;
     A.analyze(CE);
     emplaceTransformation(A.getReplacement());
@@ -932,7 +934,10 @@ void TextureRule::runRule(const MatchFinder::MatchResult &Result) {
     if (!ReplType.empty())
       emplaceTransformation(new ReplaceToken(TL->getBeginLoc(), TL->getEndLoc(),
                                              std::string(ReplType)));
-  } else if (auto CE = getNodeAsType<CallExpr>(Result, "call")) {
+  } else if (const auto *CE = getNodeAsType<CallExpr>(Result, "call")) {
+    if (const auto *FD = CE->getDirectCallee())
+      if (isUserDefinedDecl(FD))
+        return;
     auto Name = CE->getDirectCallee()->getNameAsString();
     if (DpctGlobalInfo::useSYCLCompat()) {
       report(CE->getBeginLoc(), Diagnostics::UNSUPPORT_SYCLCOMPAT, false, Name);

clang/test/dpct/user_api/user_api.cu

@@ -1,16 +1,31 @@
-// RUN: dpct --process-all -in-root %S --format-range=none -out-root %T/user_api %s --cuda-include-path="%cuda-path/include" --sycl-named-lambda -- -x cuda --cuda-host-only
-// RUN: FileCheck --match-full-lines --input-file %T/user_api/user_api.dp.cpp %s
-// RUN: %if build_lit %{icpx -c -fsycl %T/user_api/user_api.dp.cpp -o %T/user_api/user_api.dp.o %}
+// RUN: dpct --in-root %S --format-range=none --out-root %T %s --cuda-include-path="%cuda-path/include"
+// RUN: FileCheck --match-full-lines --input-file %T/user_api.dp.cpp %s
+// RUN: %if build_lit %{icpx -c -fsycl %T/user_api.dp.cpp -o %T/user_api.dp.o %}
 
 // CHECK: #include <sycl/sycl.hpp>
 // CHECK-NEXT: #include <dpct/dpct.hpp>
-// CHECK-NEXT: void create_maxpool_cudnn_tensors(){}
-// CHECK-NEXT: void foo(){
+#include <cuda.h>
+
+// CHECK: void create_maxpool_cudnn_tensors() {}
+// CHECK-NEXT: void foo1() {
 // CHECK-NEXT:   create_maxpool_cudnn_tensors();
 // CHECK-NEXT: }
-#include <cuda.h>
-void create_maxpool_cudnn_tensors(){}
-void foo(){
+void create_maxpool_cudnn_tensors() {}
+void foo1() {
   create_maxpool_cudnn_tensors();
 }
 
+void tex2D(int a, int b) {}
+void tex2D(int a, int b, int c) {}
+void tex3D(int a, int b, int c) {}
+
+// CHECK: void foo2() {
+// CHECK-NEXT:   tex2D(0, 0);
+// CHECK-NEXT:   tex2D(0, 0, 0);
+// CHECK-NEXT:   tex3D(0, 0, 0);
+// CHECK-NEXT: }
+void foo2() {
+  tex2D(0, 0);
+  tex2D(0, 0, 0);
+  tex3D(0, 0, 0);
+}