[SYCLomatic] Fix the issue that kernel function reference not replaced by wrapper and wrapper_register when used as function pointer (#2765)

Signed-off-by: intwanghao <hao3.wang@intel.com>

Author: intwanghao

clang/lib/DPCT/RulesLang/RulesLang.cpp

@@ -4512,10 +4512,14 @@ void StreamAPICallRule::runRule(const MatchFinder::MatchResult &Result) {
 }
 
 void KernelCallRefRule::registerMatcher(ast_matchers::MatchFinder &MF) {
-  MF.addMatcher(declRefExpr(allOf(to(functionDecl(hasAttr(attr::CUDAGlobal))),
-                                  unless(hasAncestor(cudaKernelCallExpr()))))
-                    .bind("kernelRef"),
-                this);
+  MF.addMatcher(
+      functionDecl(
+          forEachDescendant(
+              declRefExpr(allOf(to(functionDecl(hasAttr(attr::CUDAGlobal))),
+                                unless(hasAncestor(cudaKernelCallExpr()))))
+                  .bind("kernelRef")))
+          .bind("outerFunc"),
+      this);
   MF.addMatcher(unresolvedLookupExpr(unless(hasAncestor(cudaKernelCallExpr())))
                     .bind("unresolvedRef"),
                 this);
@@ -4591,6 +4595,11 @@ void KernelCallRefRule::insertWrapperPostfix(const T *Node,
 void KernelCallRefRule::runRule(
     const ast_matchers::MatchFinder::MatchResult &Result) {
   if (auto DRE = getAssistNodeAsType<DeclRefExpr>(Result, "kernelRef")) {
+    const FunctionDecl *OuterFD =
+        getAssistNodeAsType<FunctionDecl>(Result, "outerFunc");
+    if (!OuterFD) {
+      return;
+    }
     if (auto ParentCE = DpctGlobalInfo::findAncestor<CallExpr>(DRE)) {
       if (auto Callee = ParentCE->getDirectCallee()) {
         if (dpct::DpctGlobalInfo::isInCudaPath(Callee->getBeginLoc())) {
@@ -4614,31 +4623,25 @@ void KernelCallRefRule::runRule(
         DFI->collectInfoForWrapper(FD);
       }
     }
-    if (auto *OuterFD = DpctGlobalInfo::findAncestor<FunctionDecl>(DRE)) {
-      if ((OuterFD->getTemplatedKind() ==
-           FunctionDecl::TemplatedKind::TK_NonTemplate) ||
-          (OuterFD->getTemplatedKind() ==
-           FunctionDecl::TemplatedKind::TK_FunctionTemplate)) {
-        std::string TypeRepl;
-        if (DpctGlobalInfo::isCVersionCUDALaunchUsed()) {
-          if ((IsTemplateRelated &&
-               (!DRE->hasExplicitTemplateArgs() ||
-                (DRE->getNumTemplateArgs() <= TemplateParamNum))) ||
-              DRE->hadMultipleCandidates()) {
-            TypeRepl = getTypeRepl(DRE);
-          }
+    if ((OuterFD->getTemplatedKind() ==
+         FunctionDecl::TemplatedKind::TK_NonTemplate) ||
+        (OuterFD->getTemplatedKind() ==
+         FunctionDecl::TemplatedKind::TK_FunctionTemplate)) {
+      std::string TypeRepl;
+      if (DpctGlobalInfo::isCVersionCUDALaunchUsed()) {
+        if ((IsTemplateRelated &&
+             (!DRE->hasExplicitTemplateArgs() ||
+              (DRE->getNumTemplateArgs() <= TemplateParamNum))) ||
+            DRE->hadMultipleCandidates()) {
+          TypeRepl = getTypeRepl(DRE);
         }
-        insertWrapperPostfix<DeclRefExpr>(
-            DRE, std::move(TypeRepl),
-            DpctGlobalInfo::isCVersionCUDALaunchUsed());
       }
+      insertWrapperPostfix<DeclRefExpr>(
+          DRE, std::move(TypeRepl), DpctGlobalInfo::isCVersionCUDALaunchUsed());
     }
   }
   if (auto ULE =
           getAssistNodeAsType<UnresolvedLookupExpr>(Result, "unresolvedRef")) {
-    if (!DpctGlobalInfo::isCVersionCUDALaunchUsed()) {
-      return;
-    }
     bool KernelRefFound = false;
     for (auto *D : ULE->decls()) {
       const FunctionDecl *FD = dyn_cast<FunctionDecl>(D);
@@ -4670,7 +4673,8 @@ void KernelCallRefRule::runRule(
         }
       }
     }
-    insertWrapperPostfix<UnresolvedLookupExpr>(ULE, getTypeRepl(ULE), true);
+    insertWrapperPostfix<UnresolvedLookupExpr>(
+        ULE, getTypeRepl(ULE), DpctGlobalInfo::isCVersionCUDALaunchUsed());
   }
 }
 

clang/test/dpct/function_pointer2.cu

@@ -0,0 +1,45 @@
+// RUN: dpct --format-range=none -out-root %T/function_pointer2 %s --cuda-include-path="%cuda-path/include" -- -x cuda --cuda-host-only
+// RUN: FileCheck --input-file %T/function_pointer2/function_pointer2.dp.cpp --match-full-lines %s
+// RUN: %if build_lit %{icpx -c -fsycl %T/function_pointer2/function_pointer2.dp.cpp -o %T/function_pointer2/function_pointer2.dp.o %}
+
+#include <cuda_runtime.h>
+#include <iostream>
+
+template<typename T>
+__global__ static inline void vectorTemplateAdd(const T *A, T *B, T *C, int N) {
+    int i = blockIdx.x * blockDim.x + threadIdx.x;
+    if (i < N) {
+        C[i] = A[i] + B[i];
+    }
+}
+
+template <typename T>
+using fpt = void(*)(const T *, T*, T*, int);
+
+template<typename T>
+void foo() {
+    int *d_A, *d_B, *d_C;
+    // CHECK:  fpt<T> fp = &vectorTemplateAdd_wrapper<T>;
+    fpt<T> fp = &vectorTemplateAdd<T>;
+    // CHECK:  dpct::kernel_launcher::launch(fp, 1, 10, 0, 0, d_A, d_B, d_C, 10);
+    fp<<<1, 10>>>(d_A, d_B, d_C, 10);
+}
+
+static __global__ void setup_kernel(int p){}
+
+template<typename T>
+void goo();
+
+template<typename T>
+void goo() {
+  // CHECK: auto a = (void *)setup_kernel_wrapper;
+  auto a = (void *)setup_kernel;
+}
+
+template void goo<int>();
+
+int main() {
+  foo<int>();
+  std::cout << "test success" << std::endl;
+  return 0;
+}

clang/test/dpct/kernel_without_name.cu

@@ -293,7 +293,7 @@ template <typename T> struct kernel_type_t {
 // CHECK-NEXT:                  int b,
 // CHECK-NEXT:                  Tk *mem) {
 // CHECK-NEXT:   using Tk = typename kernel_type_t<T>::Type;
-template <typename T> __global__ 
+template <typename T> __device__ 
 void foo_device7(int a,
                  int b) {
   using Tk = typename kernel_type_t<T>::Type;