diff --git a/clang/lib/DPCT/ASTTraversal.cpp b/clang/lib/DPCT/ASTTraversal.cpp
index a2df28f96c14..22f6c722c8a0 100644
--- a/clang/lib/DPCT/ASTTraversal.cpp
+++ b/clang/lib/DPCT/ASTTraversal.cpp
@@ -109,6 +109,7 @@ REGISTER_RULE(EventAPICallRule, PassKind::PK_Migration)
 REGISTER_RULE(ProfilingEnableOnDemandRule, PassKind::PK_Analysis)
 REGISTER_RULE(StreamAPICallRule, PassKind::PK_Migration)
 REGISTER_RULE(KernelCallRule, PassKind::PK_Analysis)
+REGISTER_RULE(KernelCallRefRule, PassKind::PK_Migration)
 REGISTER_RULE(DeviceFunctionDeclRule, PassKind::PK_Analysis)
 REGISTER_RULE(MemVarRefMigrationRule, PassKind::PK_Migration)
 REGISTER_RULE(ConstantMemVarMigrationRule, PassKind::PK_Migration)
diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp
index 458724d38b1a..ecf76b2bc2a2 100644
--- a/clang/lib/DPCT/AnalysisInfo.cpp
+++ b/clang/lib/DPCT/AnalysisInfo.cpp
@@ -1573,13 +1573,7 @@ std::shared_ptr<DeviceFunctionDecl> DpctGlobalInfo::insertDeviceFunctionDecl(
       ->insertNode<ExplicitInstantiationDecl, DeviceFunctionDecl>(
           LocInfo.second, FTL, Attrs, Specialization, TAList);
 }
-std::shared_ptr<DeviceFunctionDecl>
-DpctGlobalInfo::insertDeviceFunctionDeclInModule(const FunctionDecl *FD) {
-  auto LocInfo = getLocInfo(FD);
-  return insertFile(LocInfo.first)
-      ->insertNode<DeviceFunctionDeclInModule, DeviceFunctionDecl>(
-          LocInfo.second, FD);
-}
+
 void DpctGlobalInfo::buildKernelInfo() {
   for (auto &File : FileMap)
     File.second->buildKernelInfo();
@@ -2494,6 +2488,7 @@ unsigned DpctGlobalInfo::ExperimentalFlag = 0;
 unsigned DpctGlobalInfo::HelperFuncPreferenceFlag = 0;
 bool DpctGlobalInfo::AnalysisModeFlag = false;
 bool DpctGlobalInfo::UseSYCLCompatFlag = false;
+bool DpctGlobalInfo::CVersionCUDALaunchUsedFlag = false;
 unsigned int DpctGlobalInfo::ColorOption = 1;
 std::unordered_map<int, std::shared_ptr<DeviceFunctionInfo>>
     DpctGlobalInfo::CubPlaceholderIndexMap;
@@ -4826,6 +4821,9 @@ DeviceFunctionDecl::DeviceFunctionDecl(
       FD->getTypeSourceInfo()->getTypeLoc().getAs<FunctionTypeLoc>(),
       FD->hasAttrs() ? FD->getAttrs() : NullAttrs);
   buildTextureObjectParamsInfo(FD->parameters());
+  if (FD->hasAttr<CUDAGlobalAttr>()) {
+    collectInfoForWrapper(FD);
+  }
 }
 DeviceFunctionDecl::DeviceFunctionDecl(
     unsigned Offset, const clang::tooling::UnifiedPath &FilePathIn,
@@ -4841,6 +4839,9 @@ DeviceFunctionDecl::DeviceFunctionDecl(
 
   buildReplaceLocInfo(FTL, Attrs);
   buildTextureObjectParamsInfo(FTL.getParams());
+  if (Specialization->hasAttr<CUDAGlobalAttr>()) {
+    collectInfoForWrapper(Specialization);
+  }
 }
 std::shared_ptr<DeviceFunctionInfo>
 DeviceFunctionDecl::LinkUnresolved(const UnresolvedLookupExpr *ULE,
@@ -4940,6 +4941,9 @@ void DeviceFunctionDecl::emplaceReplacement() {
       Obj->addParamDeclReplacement();
     }
   }
+  if (FuncInfo->getDeviceFunctionInfoForWrapper()) {
+    insertWrapper();
+  }
 }
 void DeviceFunctionDecl::LinkDecl(const FunctionDecl *FD, DeclList &List,
                                   std::shared_ptr<DeviceFunctionInfo> &Info) {
@@ -4973,11 +4977,9 @@ void DeviceFunctionDecl::LinkDecl(const FunctionDecl *FD, DeclList &List,
     return;
   }
   std::shared_ptr<DeviceFunctionDecl> D;
-  if (isModuleFunction(FD)) {
-    D = DpctGlobalInfo::getInstance().insertDeviceFunctionDeclInModule(FD);
-  } else {
-    D = DpctGlobalInfo::getInstance().insertDeviceFunctionDecl(FD);
-  }
+
+  D = DpctGlobalInfo::getInstance().insertDeviceFunctionDecl(FD);
+
   if (Info) {
     if (auto FuncInfo = D->getFuncInfo())
       Info->merge(FuncInfo);
@@ -4986,6 +4988,11 @@ void DeviceFunctionDecl::LinkDecl(const FunctionDecl *FD, DeclList &List,
     Info = FuncInfo;
   else
     List.push_back(D);
+
+  if (Info && isModuleFunction(FD) && FD->hasAttr<CUDAGlobalAttr>()) {
+    Info->collectInfoForWrapper(FD);
+    Info->setModuleUsed();
+  }
 }
 void DeviceFunctionDecl::LinkDecl(const NamedDecl *ND, DeclList &List,
                                   std::shared_ptr<DeviceFunctionInfo> &Info) {
@@ -5180,84 +5187,173 @@ class KernelPrinter {
                          Result.length() - Indent.length() - NL.length());
   }
 };
-///// class DeviceFunctionDeclInModule /////
-void DeviceFunctionDeclInModule::insertWrapper() {
+
+void DeviceFunctionDecl::insertWrapper() {
   auto NL = std::string(getNL());
   std::string WrapperStr = "";
   llvm::raw_string_ostream OS(WrapperStr);
   KernelPrinter Printer(NL, "", OS);
   Printer.newLine();
   Printer.newLine();
-  Printer.line("extern \"C\" {");
+  auto InfoForWrapper = FuncInfo->getDeviceFunctionInfoForWrapper();
+  bool ModuleUsed = FuncInfo->isModuleUsed();
+  auto &TParamsInfo = InfoForWrapper->TemplateParametersInfo;
+  auto &ParamsInfo = InfoForWrapper->ParametersInfo;
+  std::string FuncName = FuncInfo->getFunctionName();
+  if (ModuleUsed) {
+    Printer.line("extern \"C\" {");
+  }
   {
     auto FunctionBlock = Printer.block();
     Printer.indent();
     requestFeature(HelperFeatureEnum::device_ext);
-    Printer << "DPCT_EXPORT void " << FuncName << "_wrapper("
-            << MapNames::getClNamespace() << "queue &queue, const "
-            << MapNames::getClNamespace()
-            << "nd_range<3> &nr, unsigned int localMemSize, void "
-               "**kernelParams, void **extra)";
-    if (HasBody) {
-      auto for_each_parameter = [&](auto F) {
-        auto it = getParametersInfo().begin();
-        for (int i = 0; it != getParametersInfo().end(); ++it, ++i) {
-          F(i, it->second);
+    // 1.Generate wrapper signature
+    if (ModuleUsed) {
+      Printer << "DPCT_EXPORT void " << FuncName << "_wrapper("
+              << MapNames::getClNamespace() << "queue &queue, const "
+              << MapNames::getClNamespace()
+              << "nd_range<3> &nr, unsigned int localMemSize, void "
+                 "**kernelParams, void **extra)";
+    } else {
+      Printer.line("// Auto generated SYCL kernel wrapper used to migration "
+                   "kernel function pointer.");
+      if (!TParamsInfo.empty()) {
+        Printer << "template<";
+        for (size_t i = 0; i < TParamsInfo.size(); i++) {
+          Printer << (i == 0 ? "" : " ,") << TParamsInfo[i].first << " "
+                  << TParamsInfo[i].second
+                  << TemplateParameterDefaultValueMap[i];
         }
-      };
-
-      Printer << " {";
-      {
-        auto BodyBlock = Printer.block();
+        Printer << ">";
         Printer.newLine();
-        auto DefaultParamNum = ParamsNum - NonDefaultParamNum;
-        Printer.line(llvm::formatv(
-            "// {0} non-default parameters, {1} default parameters",
-            NonDefaultParamNum, DefaultParamNum));
-        Printer.line(llvm::formatv("{0}args_selector<{1}, {2}, decltype({3})> "
-                                   "selector(kernelParams, extra);",
-                                   MapNames::getDpctNamespace(),
-                                   NonDefaultParamNum, DefaultParamNum,
-                                   FuncName));
-        for_each_parameter([&](auto &&i, auto &&p) {
-          Printer.line("auto& " + p + " = selector.get<" + std::to_string(i) +
-                       ">();");
-        });
-
-        Kernel->buildInfo();
-        Printer.line(Kernel->getReplacement());
       }
-      Printer.line("}");
+      Printer << "void " << FuncName << "_wrapper(";
+      for (size_t i = 0; i < ParamsInfo.size(); i++) {
+        Printer << (i == 0 ? "" : " ,") << ParamsInfo[i].first << " "
+                << ParamsInfo[i].second << ParameterDefaultValueMap[i];
+      }
+      Printer << ")";
+    }
+    // 2.Generate wrapper body
+    if (HasBody) {
+      if (ModuleUsed) {
+        auto for_each_parameter = [&](auto F) {
+          auto it = ParamsInfo.begin();
+          for (int i = 0; it != ParamsInfo.end(); ++it, ++i) {
+            F(i, it->second);
+          }
+        };
+        Printer << " {";
+        {
+          auto BodyBlock = Printer.block();
+          Printer.newLine();
+          auto DefaultParamNum = ParamsNum - NonDefaultParamNum;
+          Printer.line(llvm::formatv(
+              "// {0} non-default parameters, {1} default parameters",
+              NonDefaultParamNum, DefaultParamNum));
+          Printer.line(
+              llvm::formatv("{0}args_selector<{1}, {2}, decltype({3})> "
+                            "selector(kernelParams, extra);",
+                            MapNames::getDpctNamespace(), NonDefaultParamNum,
+                            DefaultParamNum, FuncName));
+          for_each_parameter([&](auto &&i, auto &&p) {
+            Printer.line("auto& " + p + " = selector.get<" + std::to_string(i) +
+                         ">();");
+          });
+          (InfoForWrapper->KernelForWrapper)->buildInfo();
+          Printer.line((InfoForWrapper->KernelForWrapper)->getReplacement());
+        }
+        Printer.line("}");
+      } else {
+        Printer << " {";
+        {
+          auto BodyBlock = Printer.block();
+          Printer.newLine();
+          Printer.line(MapNames::getClNamespace() + "queue queue = *" +
+                       MapNames::getDpctNamespace() + "kernel_launcher::_que;");
+          Printer.line(
+              "unsigned int localMemSize = " + MapNames::getDpctNamespace() +
+              "kernel_launcher::_local_mem_size;");
+          Printer.line(MapNames::getClNamespace() + "nd_range<3> nr = " +
+                       MapNames::getDpctNamespace() + "kernel_launcher::_nr;");
+          Printer.newLine();
+          (InfoForWrapper->KernelForWrapper)->buildInfo();
+          Printer.line((InfoForWrapper->KernelForWrapper)->getReplacement());
+        }
+        Printer.line("}");
+      }
     } else {
       Printer << ";";
       Printer.newLine();
     }
   }
-
-  Printer << "}";
+  if (ModuleUsed) {
+    Printer << "}";
+  }
 
   auto Repl = std::make_shared<ExtReplacement>(FilePath, DeclEnd, 0, WrapperStr,
                                                nullptr);
   Repl->setBlockLevelFormatFlag();
   DpctGlobalInfo::getInstance().addReplacement(Repl);
 }
-void DeviceFunctionDeclInModule::buildParameterInfo(const FunctionDecl *FD) {
-  for (auto It = FD->param_begin(); It != FD->param_end(); It++) {
-    ParametersInfo.push_back(std::pair<std::string, std::string>(
-        (*It)->getOriginalType().getAsString(), (*It)->getNameAsString()));
+void DeviceFunctionDecl::collectInfoForWrapper(const FunctionDecl *FD) {
+  if ((FD->getTemplatedKind() != FunctionDecl::TemplatedKind::TK_NonTemplate) &&
+      (FD->getTemplatedKind() !=
+       FunctionDecl::TemplatedKind::TK_FunctionTemplate)) {
+    return;
   }
-}
-void DeviceFunctionDeclInModule::buildWrapperInfo(const FunctionDecl *FD) {
-  auto &SM = DpctGlobalInfo::getSourceManager();
   const FunctionDecl *Def;
   HasBody = FD->hasBody(Def);
   if (HasBody && FD != Def) {
     HasBody = false;
   }
 
-  FuncName = FD->getNameAsString();
+  auto analyzeTypeLoc = [](const TypeLoc &TL) {
+    ExprAnalysis EA;
+    EA.analyze(TL);
+    return EA.getReplacedString();
+  };
+
+  if (auto FTD = FD->getDescribedFunctionTemplate()) {
+    if (auto TemplateParmsList = FTD->getTemplateParameters()) {
+      for (size_t i = 0; i < TemplateParmsList->size(); ++i) {
+        auto TemplateParm = TemplateParmsList->getParam(i);
+        if (auto TTPD = dyn_cast<TemplateTypeParmDecl>(TemplateParm)) {
+          if (TTPD->hasDefaultArgument() &&
+              !TTPD->defaultArgumentWasInherited()) {
+            TemplateParameterDefaultValueMap[i] =
+                " = " + analyzeTypeLoc(TTPD->getDefaultArgument()
+                                           .getTypeSourceInfo()
+                                           ->getTypeLoc());
+          }
+        } else if (auto NTTPD =
+                       dyn_cast<NonTypeTemplateParmDecl>(TemplateParm)) {
+          if (NTTPD->hasDefaultArgument() &&
+              !NTTPD->defaultArgumentWasInherited()) {
+            TemplateParameterDefaultValueMap[i] =
+                " = " + ExprAnalysis::ref(
+                            NTTPD->getDefaultArgument().getSourceExpression());
+          }
+        }
+      }
+    }
+  }
+  for (size_t i = 0; i < FD->param_size(); i++) {
+    auto PDecl = FD->getParamDecl(i);
+    if (!PDecl->hasInheritedDefaultArg()) {
+      if (PDecl->hasUninstantiatedDefaultArg()) {
+        ParameterDefaultValueMap[i] =
+            " = " + ExprAnalysis::ref(PDecl->getUninstantiatedDefaultArg());
+      } else if (PDecl->hasDefaultArg()) {
+        ParameterDefaultValueMap[i] =
+            " = " + ExprAnalysis::ref(PDecl->getDefaultArg());
+      }
+    }
+  }
+
   // FD has relatively large range, which is likely to be straddle,
   // getDefinitionRange may not work as good as getExpansionRange
+  auto &SM = DpctGlobalInfo::getSourceManager();
   auto EndLoc =
       SM.getSpellingLoc(SM.getExpansionRange(FD->getEndLoc()).getEnd());
   auto LastTokenLen = Lexer::MeasureTokenLength(
@@ -5270,30 +5366,7 @@ void DeviceFunctionDeclInModule::buildWrapperInfo(const FunctionDecl *FD) {
   }
   DeclEnd = SM.getFileOffset(EndLoc);
 }
-void DeviceFunctionDeclInModule::buildCallInfo(const FunctionDecl *FD) {
-  Kernel = KernelCallExpr::buildForWrapper(FilePath, FD, getFuncInfo(FD));
-}
-DeviceFunctionDeclInModule::DeviceFunctionDeclInModule(
-    unsigned Offset, const clang::tooling::UnifiedPath &FilePathIn,
-    const FunctionTypeLoc &FTL, const ParsedAttributes &Attrs,
-    const FunctionDecl *FD)
-    : DeviceFunctionDecl(Offset, FilePathIn, FTL, Attrs, FD) {
-  buildParameterInfo(FD);
-  buildWrapperInfo(FD);
-  buildCallInfo(FD);
-}
-DeviceFunctionDeclInModule::DeviceFunctionDeclInModule(
-    unsigned Offset, const clang::tooling::UnifiedPath &FilePathIn,
-    const FunctionDecl *FD)
-    : DeviceFunctionDecl(Offset, FilePathIn, FD) {
-  buildParameterInfo(FD);
-  buildWrapperInfo(FD);
-  buildCallInfo(FD);
-}
-void DeviceFunctionDeclInModule::emplaceReplacement() {
-  DeviceFunctionDecl::emplaceReplacement();
-  insertWrapper();
-}
+
 ///// class DeviceFunctionInfo /////
 DeviceFunctionInfo::DeviceFunctionInfo(size_t ParamsNum,
                                        size_t NonDefaultParamNum,
@@ -5304,6 +5377,63 @@ DeviceFunctionInfo::DeviceFunctionInfo(size_t ParamsNum,
       FunctionName(FunctionName), IsLambda(false) {
   ParametersProps.resize(ParamsNum);
 }
+
+void DeviceFunctionInfo::collectInfoForWrapper(const FunctionDecl *FD) {
+  if (!DFInfoForWrapper) {
+    DFInfoForWrapper = std::make_shared<DeviceFunctionInfoForWrapper>();
+    auto LocInfo = DpctGlobalInfo::getLocInfo(FD->getBeginLoc());
+    auto &TemplateParametersInfo = DFInfoForWrapper->TemplateParametersInfo;
+    auto &ParametersInfo = DFInfoForWrapper->ParametersInfo;
+    auto analyzeTypeLoc = [](const TypeLoc &TL) {
+      ExprAnalysis EA;
+      EA.analyze(TL);
+      return EA.getReplacedString();
+    };
+
+    auto &Context = dpct::DpctGlobalInfo::getContext();
+    auto Parents = Context.getParents(*FD);
+    if (Parents.size()) {
+      if (auto FTD = Parents[0].get<FunctionTemplateDecl>()) {
+        FD = FTD->getTemplatedDecl();
+        if (auto TemplateParmsList = FTD->getTemplateParameters()) {
+          for (size_t i = 0; i < TemplateParmsList->size(); ++i) {
+            auto TemplateParm = TemplateParmsList->getParam(i);
+            if (auto TTPD = dyn_cast<TemplateTypeParmDecl>(TemplateParm)) {
+              TemplateParametersInfo.push_back(
+                  {std::string(TTPD->wasDeclaredWithTypename() ? "typename"
+                                                               : "class") +
+                       std::string(TTPD->isParameterPack() ? "... " : ""),
+                   TTPD->getNameAsString()});
+            } else if (auto NTTPD =
+                           dyn_cast<NonTypeTemplateParmDecl>(TemplateParm)) {
+              std::string DefVal;
+              TemplateParametersInfo.push_back(
+                  {analyzeTypeLoc(NTTPD->getTypeSourceInfo()->getTypeLoc()),
+                   NTTPD->getNameAsString()});
+            }
+          }
+        }
+      }
+    }
+    DFInfoForWrapper->KernelForWrapper =
+        KernelCallExpr::buildForWrapper(LocInfo.first, FD);
+    std::string TemplateArgsStr;
+    for (size_t i = 0; i < TemplateParametersInfo.size(); i++) {
+      TemplateArgsStr +=
+          (i == 0 ? "" : ", ") + TemplateParametersInfo[i].second;
+    }
+    if (!TemplateArgsStr.empty()) {
+      DFInfoForWrapper->KernelForWrapper->setTemplateArgsStrForWrapper(
+          TemplateArgsStr);
+    }
+    for (auto It = FD->param_begin(); It != FD->param_end(); It++) {
+      ParametersInfo.push_back(
+          {DpctGlobalInfo::getReplacedTypeName((*It)->getType()),
+           (*It)->getNameAsString()});
+    }
+  }
+}
+
 std::shared_ptr<CallFunctionExpr>
 DeviceFunctionInfo::findCallee(const CallExpr *C) {
   auto CallLocInfo = DpctGlobalInfo::getLocInfo(C);
@@ -5744,11 +5874,15 @@ void KernelCallExpr::printSubmitLambda(KernelPrinter &Printer) {
 void KernelCallExpr::printParallelFor(KernelPrinter &Printer, bool IsInSubmit) {
   std::string TemplateArgsStr;
   if (DpctGlobalInfo::isSyclNamedLambda() && hasTemplateArgs()) {
-    bool IsNeedWarning = false;
-    TemplateArgsStr = getTemplateArguments(IsNeedWarning, false, true);
-    if (!TemplateArgsStr.empty() && IsNeedWarning) {
-      printWarningMessage(Printer, Diagnostics::UNDEDUCED_TYPE,
-                          "dpct_kernel_name");
+    if (IsForWrapper) {
+      TemplateArgsStr = TemplateArgsStrForWrapper;
+    } else {
+      bool IsNeedWarning = false;
+      TemplateArgsStr = getTemplateArguments(IsNeedWarning, false, true);
+      if (!TemplateArgsStr.empty() && IsNeedWarning) {
+        printWarningMessage(Printer, Diagnostics::UNDEDUCED_TYPE,
+                            "dpct_kernel_name");
+      }
     }
   }
   if (IsInSubmit) {
@@ -5862,7 +5996,11 @@ void KernelCallExpr::printKernel(KernelPrinter &Printer) {
     Printer.line(S.StmtStr);
   }
   std::string TemplateArgsStr;
-  if (hasWrittenTemplateArgs()) {
+  if (IsForWrapper) {
+    if (!TemplateArgsStrForWrapper.empty()) {
+      TemplateArgsStr = "<" + TemplateArgsStrForWrapper + ">";
+    }
+  } else if (hasWrittenTemplateArgs()) {
     bool IsNeedWarning = false;
     TemplateArgsStr =
         buildString("<", getTemplateArguments(IsNeedWarning), ">");
@@ -6030,10 +6168,9 @@ std::shared_ptr<KernelCallExpr> KernelCallExpr::buildFromCudaLaunchKernel(
                              CE->getArg(5)},
       CE);
   Kernel->buildNeedBracesInfo(CE);
-  if (auto Callee = getAddressedRef(CE->getArg(0))) {
+  const FunctionDecl *FD = nullptr;
+  if (auto Callee = getAddressedRef(CE->getArg(0), true, &FD)) {
     Kernel->buildCalleeInfo(Callee, std::nullopt);
-    auto FD =
-        dyn_cast_or_null<FunctionDecl>(Callee->getReferencedDeclOfCallee());
     auto FuncInfo = Kernel->getFuncInfo();
     if (FD && FuncInfo) {
       auto ArgsArray = ExprAnalysis::ref(CE->getArg(3));
@@ -6061,13 +6198,13 @@ std::shared_ptr<KernelCallExpr> KernelCallExpr::buildFromCudaLaunchKernel(
 }
 std::shared_ptr<KernelCallExpr>
 KernelCallExpr::buildForWrapper(clang::tooling::UnifiedPath FilePath,
-                                const FunctionDecl *FD,
-                                std::shared_ptr<DeviceFunctionInfo> FuncInfo) {
+                                const FunctionDecl *FD) {
   auto &SM = DpctGlobalInfo::getSourceManager();
   auto Kernel =
       std::shared_ptr<KernelCallExpr>(new KernelCallExpr(0, FilePath));
+  Kernel->IsForWrapper = true;
   Kernel->Name = FD->getNameAsString();
-  Kernel->setFuncInfo(FuncInfo);
+  Kernel->setFuncInfo(DeviceFunctionDecl::LinkRedecls(FD));
   Kernel->ExecutionConfig.Config[0] = "";
   Kernel->ExecutionConfig.Config[1] = "";
   Kernel->ExecutionConfig.Config[2] = "localMemSize";
@@ -6077,6 +6214,7 @@ KernelCallExpr::buildForWrapper(clang::tooling::UnifiedPath FilePath,
   Kernel->ExecutionConfig.IsQueuePtr = false;
   Kernel->NeedBraces = false;
   Kernel->getFuncInfo()->getVarMap().Dim = 3;
+  Kernel->resizeTextureObjectList(FD->getNumParams());
   for (auto &Parm : FD->parameters()) {
     Kernel->ArgsInfo.emplace_back(Parm, Kernel.get());
   }
diff --git a/clang/lib/DPCT/AnalysisInfo.h b/clang/lib/DPCT/AnalysisInfo.h
index 9d3170c90533..fd3b2842571f 100644
--- a/clang/lib/DPCT/AnalysisInfo.h
+++ b/clang/lib/DPCT/AnalysisInfo.h
@@ -88,7 +88,6 @@ class KernelCallExpr;
 class DeviceFunctionInfo;
 class CallFunctionExpr;
 class DeviceFunctionDecl;
-class DeviceFunctionDeclInModule;
 class MemVarInfo;
 class VarInfo;
 class ExplicitInstantiationDecl;
@@ -239,6 +238,12 @@ struct RnnBackwardFuncInfo {
   std::vector<std::string> FuncArgs;
 };
 
+struct DeviceFunctionInfoForWrapper {
+  std::vector<std::pair<std::string, std::string>> ParametersInfo;
+  std::vector<std::pair<std::string, std::string>> TemplateParametersInfo;
+  std::shared_ptr<KernelCallExpr> KernelForWrapper;
+};
+
 // <function name, Info>
 using HDFuncInfoMap = std::unordered_map<std::string, HostDeviceFuncInfo>;
 // <file path, <Offset, Info>>
@@ -1000,10 +1005,14 @@ class DpctGlobalInfo {
       return Cur.get<TargetTy>();
     });
   }
-  template <class TargetTy, class NodeTy>
+  template <class TargetTy, class NodeTy, class... SkipNodeTy>
   static auto findParent(const NodeTy *Node) {
-    return findAncestor<TargetTy>(
-        Node, [](const DynTypedNode &Cur) -> bool { return true; });
+    return findAncestor<TargetTy>(Node, [](const DynTypedNode &Cur) -> bool {
+      if ((... || Cur.get<SkipNodeTy>())) {
+        return false;
+      }
+      return true;
+    });
   }
 
   template <typename TargetTy, typename NodeTy>
@@ -1136,8 +1145,6 @@ class DpctGlobalInfo {
   std::shared_ptr<DeviceFunctionDecl> insertDeviceFunctionDecl(
       const FunctionDecl *Specialization, const FunctionTypeLoc &FTL,
       const ParsedAttributes &Attrs, const TemplateArgumentListInfo &TAList);
-  std::shared_ptr<DeviceFunctionDecl>
-  insertDeviceFunctionDeclInModule(const FunctionDecl *FD);
 
   // Build kernel and device function declaration replacements and store
   // them.
@@ -1340,6 +1347,8 @@ class DpctGlobalInfo {
   static bool useNoQueueDevice() {
     return getHelperFuncPreference(HelperFuncPreference::NoQueueDevice);
   }
+  static void setCVersionCUDALaunchUsed() { CVersionCUDALaunchUsedFlag = true; }
+  static bool isCVersionCUDALaunchUsed() { return CVersionCUDALaunchUsedFlag; }
   static void setUseSYCLCompat(bool Flag = true) { UseSYCLCompatFlag = Flag; }
   static bool useSYCLCompat() { return UseSYCLCompatFlag; }
   static bool useEnqueueBarrier() {
@@ -1662,6 +1671,7 @@ class DpctGlobalInfo {
   static unsigned HelperFuncPreferenceFlag;
   static bool AnalysisModeFlag;
   static bool UseSYCLCompatFlag;
+  static bool CVersionCUDALaunchUsedFlag;
   static unsigned int ColorOption;
   static std::unordered_map<int, std::shared_ptr<DeviceFunctionInfo>>
       CubPlaceholderIndexMap;
@@ -2565,7 +2575,8 @@ class DeviceFunctionDecl {
       LinkDecl(D, List, Info);
   }
   void setFuncInfo(std::shared_ptr<DeviceFunctionInfo> Info);
-
+  void insertWrapper();
+  void collectInfoForWrapper(const FunctionDecl *FD);
   virtual ~DeviceFunctionDecl() = default;
 
 protected:
@@ -2588,7 +2599,10 @@ class DeviceFunctionDecl {
   bool IsDefFilePathNeeded = false;
   std::vector<std::shared_ptr<TextureObjectInfo>> TextureObjectList;
   FormatInfo FormatInformation;
-
+  bool HasBody = false;
+  size_t DeclEnd;
+  std::map<int, std::string> TemplateParameterDefaultValueMap;
+  std::map<int, std::string> ParameterDefaultValueMap;
   static std::shared_ptr<DeviceFunctionInfo> &getFuncInfo(const FunctionDecl *);
   static std::unordered_map<std::string, std::shared_ptr<DeviceFunctionInfo>>
       FuncInfoMap;
@@ -2619,32 +2633,6 @@ class ExplicitInstantiationDecl : public DeviceFunctionDecl {
   std::string getExtraParameters(LocInfo LI) override;
 };
 
-class DeviceFunctionDeclInModule : public DeviceFunctionDecl {
-  void insertWrapper();
-  bool HasBody = false;
-  size_t DeclEnd;
-  std::string FuncName;
-  std::vector<std::pair<std::string, std::string>> ParametersInfo;
-  std::shared_ptr<KernelCallExpr> Kernel;
-  void buildParameterInfo(const FunctionDecl *FD);
-  void buildWrapperInfo(const FunctionDecl *FD);
-  void buildCallInfo(const FunctionDecl *FD);
-  std::vector<std::pair<std::string, std::string>> &getParametersInfo() {
-    return ParametersInfo;
-  }
-
-public:
-  DeviceFunctionDeclInModule(unsigned Offset,
-                             const clang::tooling::UnifiedPath &FilePathIn,
-                             const FunctionTypeLoc &FTL,
-                             const ParsedAttributes &Attrs,
-                             const FunctionDecl *FD);
-  DeviceFunctionDeclInModule(unsigned Offset,
-                             const clang::tooling::UnifiedPath &FilePathIn,
-                             const FunctionDecl *FD);
-  void emplaceReplacement() override;
-};
-
 // device function info includes parameters num, memory variable and call
 // expression in the function.
 class DeviceFunctionInfo {
@@ -2744,6 +2732,13 @@ class DeviceFunctionInfo {
   bool isParameterReferenced(unsigned int Index);
   void setParameterReferencedStatus(unsigned int Index, bool IsReferenced);
   std::string getFunctionName() { return FunctionName; }
+  void collectInfoForWrapper(const FunctionDecl *FD);
+  void setModuleUsed() { ModuleUsed = true; }
+  bool isModuleUsed() { return ModuleUsed; }
+  std::shared_ptr<DeviceFunctionInfoForWrapper>
+  getDeviceFunctionInfoForWrapper() {
+    return DFInfoForWrapper;
+  }
 
 private:
   void mergeCalledTexObj(
@@ -2776,12 +2771,15 @@ class DeviceFunctionInfo {
   bool CallGroupFunctionInControlFlow = false;
   bool HasCheckedCallGroupFunctionInControlFlow = false;
   OverloadedOperatorKind OO_Kind = OverloadedOperatorKind::OO_None;
+  bool ModuleUsed = false;
+  std::shared_ptr<DeviceFunctionInfoForWrapper> DFInfoForWrapper = nullptr;
 };
 
 class KernelCallExpr : public CallFunctionExpr {
 public:
   bool IsInMacroDefine = false;
   bool NeedLambda = false;
+  bool IsForWrapper = false;
   bool NeedDefaultRetValue = false;
 
 private:
@@ -2854,8 +2852,10 @@ class KernelCallExpr : public CallFunctionExpr {
       const std::pair<clang::tooling::UnifiedPath, unsigned> &LocInfo,
       const CallExpr *, bool IsAssigned = false);
   static std::shared_ptr<KernelCallExpr>
-  buildForWrapper(clang::tooling::UnifiedPath, const FunctionDecl *,
-                  std::shared_ptr<DeviceFunctionInfo>);
+  buildForWrapper(clang::tooling::UnifiedPath, const FunctionDecl *);
+  void setTemplateArgsStrForWrapper(std::string Str) {
+    TemplateArgsStrForWrapper = std::move(Str);
+  }
   unsigned int GridDim = 3;
   unsigned int BlockDim = 3;
   void setEmitSizeofWarningFlag(bool Flag) { EmitSizeofWarning = Flag; }
@@ -2959,6 +2959,7 @@ class KernelCallExpr : public CallFunctionExpr {
   OuterStmtsList OuterStmts;
   StmtList KernelStmts;
   std::string KernelArgs;
+  std::string TemplateArgsStrForWrapper;
   int TotalArgsSize = 0;
   bool EmitSizeofWarning = false;
   unsigned int SizeOfHighestDimension = 0;
diff --git a/clang/lib/DPCT/RulesLang/RulesLang.cpp b/clang/lib/DPCT/RulesLang/RulesLang.cpp
index 68ffdaf523b7..fc76e1eaceca 100644
--- a/clang/lib/DPCT/RulesLang/RulesLang.cpp
+++ b/clang/lib/DPCT/RulesLang/RulesLang.cpp
@@ -4439,6 +4439,168 @@ 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(unresolvedLookupExpr(unless(hasAncestor(cudaKernelCallExpr())))
+                    .bind("unresolvedRef"),
+                this);
+}
+
+std::string KernelCallRefRule::getTypeRepl(const Expr *E) {
+  std::string TypeRef;
+  if (auto BO = DpctGlobalInfo::findParent<BinaryOperator, Expr,
+                                           ImplicitCastExpr, ParenExpr>(E)) {
+    TypeRef = "decltype(" + ExprAnalysis::ref(BO->getLHS()) + ")";
+  } else if (auto VD =
+                 DpctGlobalInfo::findParent<VarDecl, Expr, ImplicitCastExpr,
+                                            ParenExpr>(E)) {
+    TypeRef = "decltype(" + VD->getNameAsString() + ")";
+  } else if (auto RS =
+                 DpctGlobalInfo::findParent<ReturnStmt, Expr, ImplicitCastExpr,
+                                            ParenExpr>(E)) {
+    auto FD = DpctGlobalInfo::findAncestor<FunctionDecl>(RS);
+    TypeRef =
+        getStmtSpelling(FD->getReturnTypeSourceRange(), FD->getSourceRange());
+  } else if (auto CE =
+                 DpctGlobalInfo::findParent<CallExpr, Expr, ImplicitCastExpr,
+                                            ParenExpr>(E)) {
+    size_t N = 0;
+    for (auto Arg : CE->arguments()) {
+      if (Arg == E) {
+        break;
+      }
+      N++;
+    }
+    TypeRef = "typename " + MapNames::getDpctNamespace() +
+              "nth_argument_type<decltype(" +
+              ExprAnalysis::ref(CE->getCallee()) + "), " + std::to_string(N) +
+              ">::type";
+  }
+  if (!TypeRef.empty()) {
+    return "<" + TypeRef + ">";
+  }
+  return TypeRef;
+}
+
+template <typename T>
+void KernelCallRefRule::insertWrapperPostfix(const T *Node,
+                                             std::string &&TypeRepl,
+                                             bool isInsertWrapperRegister) {
+  auto NLoc = DpctGlobalInfo::getSourceManager().getSpellingLoc(
+      Node->getNameInfo().getBeginLoc());
+  emplaceTransformation(new InsertText(
+      NLoc.getLocWithOffset(Node->getNameInfo().getAsString().length()),
+      "_wrapper"));
+
+  if (!isInsertWrapperRegister) {
+    return;
+  }
+  const Expr *E = Node;
+  if (auto UO = DpctGlobalInfo::findParent<UnaryOperator, T, ImplicitCastExpr,
+                                           ParenExpr>(Node)) {
+    if (UO->getOpcode() == UO_AddrOf) {
+      E = UO;
+    }
+  } else if (auto COC = DpctGlobalInfo::findParent<CXXOperatorCallExpr, T,
+                                                   ImplicitCastExpr, ParenExpr>(
+                 Node)) {
+    if (COC->getOperator() == clang::OO_Amp) {
+      E = COC;
+    }
+  }
+  emplaceTransformation(new InsertBeforeStmt(
+      E, MapNames::getDpctNamespace() + "wrapper_register" + TypeRepl + "("));
+  emplaceTransformation(new InsertAfterStmt(E, ").get()"));
+}
+
+void KernelCallRefRule::runRule(
+    const ast_matchers::MatchFinder::MatchResult &Result) {
+  if (auto DRE = getAssistNodeAsType<DeclRefExpr>(Result, "kernelRef")) {
+    if (auto ParentCE = DpctGlobalInfo::findAncestor<CallExpr>(DRE)) {
+      if (auto Callee = ParentCE->getDirectCallee()) {
+        if (dpct::DpctGlobalInfo::isInCudaPath(Callee->getBeginLoc())) {
+          return;
+        }
+      }
+    }
+    const FunctionDecl *FD = dyn_cast<FunctionDecl>(DRE->getDecl());
+    bool IsTemplateRelated = false;
+    size_t TemplateParamNum = 0;
+    if (FD) {
+      if (FD->getTemplatedKind() !=
+          FunctionDecl::TemplatedKind::TK_NonTemplate) {
+        IsTemplateRelated = true;
+      }
+      if (auto FTD = FD->getPrimaryTemplate()) {
+        if (auto TP = FTD->getTemplateParameters())
+          TemplateParamNum = TP->size();
+      }
+      if (auto DFI = DeviceFunctionDecl::LinkRedecls(FD)) {
+        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);
+          }
+        }
+        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);
+      if (!FD) {
+        if (const FunctionTemplateDecl *FTD =
+                dyn_cast<FunctionTemplateDecl>(D)) {
+          FD = FTD->getTemplatedDecl();
+        }
+      }
+      if (FD && FD->hasAttr<CUDAGlobalAttr>()) {
+        KernelRefFound = true;
+        break;
+      }
+    }
+    if (!KernelRefFound) {
+      return;
+    }
+    if (auto ParentCE = dpct::DpctGlobalInfo::findAncestor<CallExpr>(ULE)) {
+      if (auto Callee = ParentCE->getDirectCallee()) {
+        if (dpct::DpctGlobalInfo::isInCudaPath(Callee->getBeginLoc())) {
+          return;
+        }
+      } else if (auto PULE =
+                     dyn_cast<UnresolvedLookupExpr>(ParentCE->getCallee())) {
+        for (auto *D : PULE->decls()) {
+          if (dpct::DpctGlobalInfo::isInCudaPath(D->getBeginLoc())) {
+            return;
+          }
+        }
+      }
+    }
+    insertWrapperPostfix<UnresolvedLookupExpr>(ULE, getTypeRepl(ULE), true);
+  }
+}
 
 // kernel call information collection
 void KernelCallRule::registerMatcher(ast_matchers::MatchFinder &MF) {
@@ -4542,6 +4704,43 @@ void KernelCallRule::runRule(
              false);
       return;
     }
+    bool IsDirectCall = true;
+    if (!KCall->getDirectCallee() &&
+        !dyn_cast<UnresolvedLookupExpr>(KCall->getCallee())) {
+      IsDirectCall = false;
+      std::string ReplStr;
+      llvm::raw_string_ostream OS(ReplStr);
+      OS << MapNames::getDpctNamespace() + "kernel_launcher::launch("
+         << ExprAnalysis::ref(KCall->getCallee());
+      if (const CallExpr *Configs = KCall->getConfig()) {
+        size_t ConfigArgsNum = Configs->getNumArgs();
+        for (size_t i = 0; i < ConfigArgsNum; i++) {
+          if (auto Config = Configs->getArg(i)) {
+            if (i == 0 || i == 1) {
+              OS << ", " << ExprAnalysis::ref(Config);
+            } else if (i == 2 || i == 3) {
+              if (Config->isDefaultArgument()) {
+                OS << ", 0";
+              } else {
+                OS << ", " << ExprAnalysis::ref(Config);
+              }
+            }
+          }
+        }
+        size_t ArgsNum = KCall->getNumArgs();
+        for (size_t i = 0; i < ArgsNum; i++) {
+          if (auto Arg = KCall->getArg(i)) {
+            if (!Arg->isDefaultArgument()) {
+              OS << ", " << ExprAnalysis::ref(Arg);
+            } else {
+              break;
+            }
+          }
+        }
+        OS << ")";
+        emplaceTransformation(new ReplaceStmt(KCall, OS.str()));
+      }
+    }
 
     const auto &SM = (*Result.Context).getSourceManager();
 
@@ -4549,17 +4748,19 @@ void KernelCallRule::runRule(
       // Report warning message
       report(KCall->getBeginLoc(), Diagnostics::KERNEL_CALLEE_MACRO_ARG, false);
     }
-
-    // Remove KCall in the original location
-    auto KCallSpellingRange = getTheLastCompleteImmediateRange(
-        KCall->getBeginLoc(), KCall->getEndLoc());
-    auto KCallLen = SM.getCharacterData(KCallSpellingRange.second) -
-                    SM.getCharacterData(KCallSpellingRange.first) +
-                    Lexer::MeasureTokenLength(KCallSpellingRange.second, SM,
-                                              Result.Context->getLangOpts());
-    emplaceTransformation(
-        new ReplaceText(KCallSpellingRange.first, KCallLen, ""));
-    auto EpilogLocation = removeTrailingSemicolon(KCall, Result);
+    if (IsDirectCall) {
+      // Remove KCall in the original location
+      auto KCallSpellingRange = getTheLastCompleteImmediateRange(
+          KCall->getBeginLoc(), KCall->getEndLoc());
+      auto KCallLen = SM.getCharacterData(KCallSpellingRange.second) -
+                      SM.getCharacterData(KCallSpellingRange.first) +
+                      Lexer::MeasureTokenLength(KCallSpellingRange.second, SM,
+                                                Result.Context->getLangOpts());
+      emplaceTransformation(
+          new ReplaceText(KCallSpellingRange.first, KCallLen, ""));
+    }
+    auto EpilogLocation = findAndRemoveTrailingSemicolon(
+        KCall, Result, IsDirectCall ? true : false);
     if (DpctGlobalInfo::isCodePinEnabled()) {
       instrumentKernelLogsForCodePin(KCall, EpilogLocation);
     }
@@ -4579,11 +4780,11 @@ void KernelCallRule::runRule(
 
     // Add kernel call to map,
     // will do code generation in Global.buildReplacements();
-    if (!FD->isTemplateInstantiation()){
+    if (IsDirectCall && !FD->isTemplateInstantiation()) {
       DpctGlobalInfo::getInstance().insertKernelCallExpr(KCall);
     }
     const CallExpr *Config = KCall->getConfig();
-    if (Config) {
+    if (IsDirectCall && Config) {
       if (Config->getNumArgs() > 2) {
         const Expr *SharedMemSize = Config->getArg(2);
         if (containSizeOfType(SharedMemSize)) {
@@ -4626,8 +4827,51 @@ void KernelCallRule::runRule(
     }
     if (!LaunchKernelCall)
       return;
+    const Expr *CalleeDRE = LaunchKernelCall->getArg(0);
+    bool IsFuncTypeErased = true;
+    auto QT = CalleeDRE->getType();
+
+    if (QT->isPointerType()) {
+      QT = QT->getPointeeType();
+    }
+    if (QT->isFunctionType()) {
+      IsFuncTypeErased = false;
+    }
+
+    if (!getAddressedRef(CalleeDRE)) {
+      if (IsFuncTypeErased) {
+        DpctGlobalInfo::setCVersionCUDALaunchUsed();
+      }
+      std::string ReplStr;
+      llvm::raw_string_ostream OS(ReplStr);
+      if (IsAssigned) {
+        OS << MapNames::getCheckErrorMacroName() << "(";
+      }
+      OS << MapNames::getDpctNamespace() << "kernel_launcher::launch(";
+      size_t ArgsNum = LaunchKernelCall->getNumArgs();
+      for (size_t i = 0; i < ArgsNum; i++) {
+        if (auto Arg = LaunchKernelCall->getArg(i)) {
+          if (i == 0) {
+            if (auto E = getAddressedRef(CalleeDRE, false, nullptr)) {
+              OS << ExprAnalysis::ref(E);
+            } else {
+              OS << ExprAnalysis::ref(Arg);
+            }
+          } else {
+            OS << ", " << ExprAnalysis::ref(Arg);
+          }
+        }
+      }
+      OS << ")";
+      if (IsAssigned) {
+        OS << ")";
+      }
+      emplaceTransformation(new ReplaceStmt(LaunchKernelCall, OS.str()));
+      return;
+    }
+
     if (!IsAssigned)
-      removeTrailingSemicolon(LaunchKernelCall, Result);
+      findAndRemoveTrailingSemicolon(LaunchKernelCall, Result);
     if (DpctGlobalInfo::getInstance().buildLaunchKernelInfo(LaunchKernelCall,
                                                             IsAssigned)) {
       emplaceTransformation(new ReplaceStmt(LaunchKernelCall, true, false, ""));
@@ -4636,16 +4880,18 @@ void KernelCallRule::runRule(
 }
 
 // Find and remove the semicolon after the kernel call
-SourceLocation KernelCallRule::removeTrailingSemicolon(
-    const CallExpr *KCall,
-    const ast_matchers::MatchFinder::MatchResult &Result) {
+SourceLocation KernelCallRule::findAndRemoveTrailingSemicolon(
+    const CallExpr *KCall, const ast_matchers::MatchFinder::MatchResult &Result,
+    bool Remove) {
   const auto &SM = (*Result.Context).getSourceManager();
   auto KELoc =
       getTheLastCompleteImmediateRange(KCall->getBeginLoc(), KCall->getEndLoc())
           .second;
   auto Tok = Lexer::findNextToken(KELoc, SM, LangOptions()).value();
   if (Tok.is(tok::TokenKind::semi)) {
-    emplaceTransformation(new ReplaceToken(Tok.getLocation(), ""));
+    if (Remove) {
+      emplaceTransformation(new ReplaceToken(Tok.getLocation(), ""));
+    }
     return Lexer::findNextToken(Tok.getLocation(), SM, LangOptions())
         .value()
         .getLocation();
@@ -4653,8 +4899,6 @@ SourceLocation KernelCallRule::removeTrailingSemicolon(
   return Tok.getLocation();
 }
 
-
-
 bool isRecursiveDeviceFuncDecl(const FunctionDecl* FD) {
   // Build call graph for FunctionDecl and look for cycles in call graph.
   // Emit the warning message when the recursive call exists in kernel function.
diff --git a/clang/lib/DPCT/RulesLang/RulesLang.h b/clang/lib/DPCT/RulesLang/RulesLang.h
index 685ecae99ed2..afda6edd1bf3 100644
--- a/clang/lib/DPCT/RulesLang/RulesLang.h
+++ b/clang/lib/DPCT/RulesLang/RulesLang.h
@@ -442,13 +442,32 @@ class KernelCallRule : public NamedMigrationRule<KernelCallRule> {
 public:
   void registerMatcher(ast_matchers::MatchFinder &MF) override;
   void runRule(const ast_matchers::MatchFinder::MatchResult &Result);
-  SourceLocation
-  removeTrailingSemicolon(const CallExpr *KCall,
-                          const ast_matchers::MatchFinder::MatchResult &Result);
+  SourceLocation findAndRemoveTrailingSemicolon(
+      const CallExpr *KCall,
+      const ast_matchers::MatchFinder::MatchResult &Result, bool Remove = true);
   void instrumentKernelLogsForCodePin(const CUDAKernelCallExpr *KCall,
                                       SourceLocation &EpilogLocation);
 };
 
+/// Migration rule for kernel function references.
+/// This rule handles kernel functions that are used as function pointers.
+/// For such kernel functions, a wrapper is generated with a `_wrapper`
+/// postfix added to the kernel function name. Additionally, if the kernel
+/// function pointer is used in a context where its original type information is
+/// erased (e.g., raw pointer usage), an extra wrapper registration is required.
+/// This ensures that the raw pointer is associated with the appropriate wrapper
+/// and retains the necessary type information.
+class KernelCallRefRule : public NamedMigrationRule<KernelCallRefRule> {
+  std::string getTypeRepl(const Expr *E);
+  template <typename T>
+  void insertWrapperPostfix(const T *Node, std::string &&TypeRepl,
+                            bool isInsertWrapperRegister);
+
+public:
+  void registerMatcher(ast_matchers::MatchFinder &MF) override;
+  void runRule(const ast_matchers::MatchFinder::MatchResult &Result);
+};
+
 /// Migration rule for device function calls
 class DeviceFunctionDeclRule
     : public NamedMigrationRule<DeviceFunctionDeclRule> {
diff --git a/clang/lib/DPCT/Utility.cpp b/clang/lib/DPCT/Utility.cpp
index 8630a2cf317e..26c9bd3eab77 100644
--- a/clang/lib/DPCT/Utility.cpp
+++ b/clang/lib/DPCT/Utility.cpp
@@ -4998,21 +4998,57 @@ int isArgumentInitialized(
 
   return DeclsRequireInit.empty();
 }
-const DeclRefExpr *getAddressedRef(const Expr *E) {
+const Expr *getAddressedRef(const Expr *E, bool IsCheckFunctionDecl,
+                            const FunctionDecl **FuncDecl) {
   E = E->IgnoreImplicitAsWritten();
   if (auto DRE = dyn_cast<DeclRefExpr>(E)) {
-    if (DRE->getDecl()->getKind() == Decl::Function) {
+    if (IsCheckFunctionDecl) {
+      if (auto FD = dyn_cast<FunctionDecl>(DRE->getDecl())) {
+        if (FuncDecl) {
+          *FuncDecl = FD;
+        }
+        return DRE;
+      }
+    } else {
       return DRE;
     }
+  } else if (auto ULE = dyn_cast<UnresolvedLookupExpr>(E)) {
+    if (IsCheckFunctionDecl) {
+      for (auto *D : ULE->decls()) {
+        const FunctionDecl *FD = dyn_cast<FunctionDecl>(D);
+        if (!FD) {
+          if (const FunctionTemplateDecl *FTD =
+                  dyn_cast<FunctionTemplateDecl>(D)) {
+            FD = FTD->getTemplatedDecl();
+          }
+        }
+        if (FD) {
+          if (FuncDecl) {
+            *FuncDecl = FD;
+          }
+          return ULE;
+        }
+      }
+    } else {
+      return ULE;
+    }
   } else if (auto Paren = dyn_cast<ParenExpr>(E)) {
-    return getAddressedRef(Paren->getSubExpr());
+    return getAddressedRef(Paren->getSubExpr(), IsCheckFunctionDecl, FuncDecl);
   } else if (auto Cast = dyn_cast<CastExpr>(E)) {
-    return getAddressedRef(Cast->getSubExprAsWritten());
+    return getAddressedRef(Cast->getSubExprAsWritten(), IsCheckFunctionDecl,
+                           FuncDecl);
   } else if (auto UO = dyn_cast<UnaryOperator>(E)) {
     if (UO->getOpcode() == UO_AddrOf) {
-      return getAddressedRef(UO->getSubExpr());
+      return getAddressedRef(UO->getSubExpr(), IsCheckFunctionDecl, FuncDecl);
+    }
+  } else if (auto COC = dyn_cast<CXXOperatorCallExpr>(E)) {
+    if (COC->getOperator() == clang::OO_Amp) {
+      return getAddressedRef(COC->getArg(0), IsCheckFunctionDecl, FuncDecl);
     }
   }
+  if (FuncDecl) {
+    *FuncDecl = nullptr;
+  }
   return nullptr;
 }
 
diff --git a/clang/lib/DPCT/Utility.h b/clang/lib/DPCT/Utility.h
index 62b30f92747d..8df4fc43241a 100644
--- a/clang/lib/DPCT/Utility.h
+++ b/clang/lib/DPCT/Utility.h
@@ -375,7 +375,8 @@ findTheOuterMostCompoundStmtUntilMeetControlFlowNodes(
 const clang::NamedDecl *getNamedDecl(const clang::Type *TypePtr);
 const clang::LambdaExpr *
 getImmediateOuterLambdaExpr(const clang::FunctionDecl *FuncDecl);
-const DeclRefExpr *getAddressedRef(const Expr *E);
+const Expr *getAddressedRef(const Expr *E, bool IsCheckFunctionDecl = true,
+                            const FunctionDecl **FuncDecl = nullptr);
 const clang::FunctionDecl *findTheOuterMostFunctionDecl(const clang::Decl *D);
 
 // Source Range & location, offset.
diff --git a/clang/runtime/dpct-rt/include/dpct/kernel.hpp b/clang/runtime/dpct-rt/include/dpct/kernel.hpp
index 4081f1f0ec57..ec42bd63e0ff 100644
--- a/clang/runtime/dpct-rt/include/dpct/kernel.hpp
+++ b/clang/runtime/dpct-rt/include/dpct/kernel.hpp
@@ -444,6 +444,157 @@ static inline void invoke_kernel_function(dpct::kernel_function &function,
            localMemSize, kernelParams, extra);
 }
 
+/// Utility class for launching SYCL kernels through kernel
+/// function wrapper.
+/// For example:
+/// A SYCL kernel function:
+///   void kernel_func(int *ptr, sycl::nd_item<3> item);
+/// Kernel function wrapper:
+///   void kernel_func_wrapper(int *ptr) {
+///     sycl::queue queue = *dpct::kernel_launcher::_que;
+///     unsigned int localMemSize = dpct::kernel_launcher::_local_mem_size;
+///     sycl::nd_range<3> nr = dpct::kernel_launcher::_nr;
+///     queue.parallel_for(
+///       nr,
+///       [=](sycl::nd_item<3> item_ct1) {
+///         kernel_func(ptr, item_ct1);
+///       });
+///   }
+/// Then launch the kernel through wrapper like:
+///   typedef void(*fpt)(int *);
+///   fpt fp = kernel_func_wrapper;
+///   dpct::kernel_launcher::launch(fp, dpct::dim3(1), dpct::dim3(1), 0, 0,
+///   device_ptr);
+/// If the origin function type is erased, then need to register it first:
+///   void *fp = (void *)wrapper_register(&kernel_func_wrapper).get();
+///   dpct::kernel_launcher::launch(fp, dpct::dim3(1), dpct::dim3(1), args, 0,
+///   0);
+class kernel_launcher {
+  template <typename FuncT, typename ArgSelector, std::size_t... Index>
+  static void launch_helper(FuncT &&func, ArgSelector &selector,
+                            std::index_sequence<Index...>) {
+    func(selector.template get<Index>()...);
+  }
+  static void set_execution_config(dim3 group_range, dim3 local_range,
+                                   unsigned int local_mem_size, queue_ptr que) {
+    if (que) {
+      _que = que;
+    } else {
+      _que = &get_default_queue();
+    }
+    _nr = sycl::nd_range<3>(
+        static_cast<sycl::range<3>>(group_range * local_range),
+        static_cast<sycl::range<3>>(local_range));
+    _local_mem_size = local_mem_size;
+  };
+
+public:
+  /// Variables for storing execution configuration.
+  static inline thread_local sycl::queue *_que = nullptr;
+  static inline thread_local sycl::nd_range<3> _nr = sycl::nd_range<3>();
+  static inline thread_local unsigned int _local_mem_size = 0;
+  /// Map for retrieving launchable functor from a raw pointer.
+  static inline std::map<
+      const void *,
+      std::function<void(dim3, dim3, void **, unsigned int, queue_ptr)>>
+      kernel_function_ptr_map = {};
+
+  /// Registers a kernel function pointer with a corresponding launchable
+  /// functor.
+  /// \param [in] func Pointer to the kernel function.
+  /// \param [in] launcher Functor to handle kernel invocation.
+  static void register_kernel_ptr(
+      const void *func,
+      std::function<void(dim3, dim3, void **, unsigned int, queue_ptr)>
+          launcher) {
+    kernel_function_ptr_map[func] = std::move(launcher);
+  }
+  /// Launches a kernel function with arguments provided directly through
+  /// kernel function wrapper.
+  /// \tparam FuncT Type of the kernel function wrapper.
+  /// \tparam ArgsT Types of kernel arguments.
+  /// \param [in] func Pointer to the kernel function wrapper.
+  /// \param [in] group_range SYCL group range.
+  /// \param [in] local_range SYCL local range.
+  /// \param [in] local_mem_size The size of local memory required by the kernel
+  /// function.
+  /// \param [in] que SYCL queue used to execute kernel.
+  /// \param [in] args Kernel arguments.
+  template <typename FuncT, typename... ArgsT>
+  static void launch(FuncT *func, dim3 group_range, dim3 local_range,
+                     unsigned int local_mem_size, queue_ptr que,
+                     ArgsT... args) {
+    set_execution_config(group_range, local_range, local_mem_size, que);
+    func(args...);
+  }
+  /// Launches a kernel function through registered kernel function wrapper.
+  /// \param [in] func Pointer to the registered kernel function wrapper.
+  /// \param [in] group_range SYCL group range.
+  /// \param [in] local_range SYCL local range.
+  /// \param [in] args Array of pointers to kernel arguments.
+  /// \param [in] local_mem_size The size of local memory required by the kernel
+  /// function.
+  /// \param [in] que SYCL queue used to execute kernel.
+  static void launch(const void *func, dim3 group_range, dim3 local_range,
+                     void **args, unsigned int local_mem_size, queue_ptr que) {
+    kernel_function_ptr_map[func](group_range, local_range, args,
+                                  local_mem_size, que);
+  }
+  /// Launches a kernel function with packed arguments through kernel
+  /// function wrapper.
+  /// \tparam FuncT Type of the kernel function wrapper.
+  /// \param [in] func Pointer to the kernel function wrapper.
+  /// \param [in] group_range SYCL group range.
+  /// \param [in] local_range SYCL local range.
+  /// \param [in] args Array of pointers to kernel arguments.
+  /// \param [in] local_mem_size The size of local memory required by the kernel
+  /// function.
+  /// \param [in] que SYCL queue used to execute kernel.
+  template <typename FuncT>
+  static void launch(FuncT *func, dim3 group_range, dim3 local_range,
+                     void **args, unsigned int local_mem_size, queue_ptr que) {
+    constexpr size_t p_num = args_selector<0, 0, FuncT>::params_num;
+    set_execution_config(group_range, local_range, local_mem_size, que);
+    args_selector<p_num, p_num, FuncT> selector(args, nullptr);
+    launch_helper(func, selector, std::make_index_sequence<p_num>{});
+  }
+};
+
+/// Helper class to register and invoke kernel functions through a wrapper.
+template <typename F> class wrapper_register;
+template <typename Ret, typename... Args>
+class wrapper_register<Ret (*)(Args...)> {
+public:
+  typedef Ret (*FT)(Args...);
+  FT func;
+  /// Constructor to register a kernel function pointer.
+  /// \param [in] fp Pointer to the kernel function.
+  wrapper_register(FT fp) : func(fp) {
+    kernel_launcher::register_kernel_ptr((void *)func, *this);
+  }
+  /// Invokes the kernel function through the stored kernel function wrapper.
+  /// \param [in] group_range SYCL group range.
+  /// \param [in] local_range SYCL local range.
+  /// \param [in] args Array of pointers to kernel arguments.
+  /// \param [in] local_mem_size The size of local memory required by the kernel
+  /// function.
+  /// \param [in] que SYCL queue used to execute kernel.
+  void operator()(dim3 group_range, dim3 local_range, void **args,
+                  unsigned int local_mem_size, queue_ptr que) {
+    kernel_launcher::launch(func, group_range, local_range, args,
+                            local_mem_size, que);
+  }
+  /// Retrieves the original kernel function pointer.
+  /// \return The original kernel function pointer.
+  const FT &get() const noexcept { return func; }
+  /// Implicit conversion to the original kernel function pointer.
+  /// \return The original kernel function pointer.
+  operator FT() const noexcept { return func; }
+};
+/// Deduction guide for wrapper_register.
+template <typename Ret, typename... Args>
+wrapper_register(Ret (*)(Args...)) -> wrapper_register<Ret (*)(Args...)>;
+
 /// Find image wrapper in a kernel library and return its address.
 /// \param [in] library Handle to the kernel library.
 /// \param [in] name Name of the target image wrapper.
diff --git a/clang/runtime/dpct-rt/include/dpct/util.hpp b/clang/runtime/dpct-rt/include/dpct/util.hpp
index bf42003d1a25..91f40c479976 100644
--- a/clang/runtime/dpct-rt/include/dpct/util.hpp
+++ b/clang/runtime/dpct-rt/include/dpct/util.hpp
@@ -1125,6 +1125,8 @@ class args_selector<n_nondefault_params, n_default_params, R(Ts...)> {
   template <int i>
   using arg_type = std::tuple_element_t<account_for_default_params<i>(),
 					  std::tuple<Ts...>>;
+  static constexpr int params_num = sizeof...(Ts);
+
 private:
   template <int i>
   static constexpr int get_offset() {
@@ -1185,6 +1187,24 @@ class args_selector<n_nondefault_params, n_default_params, R(Ts...)> {
   }
 };
 
+/// \brief This struct template used to get the type of the N-th argument of a
+/// callable type `Func`. It supports both function types (e.g., `void(int,
+/// double)`) and callable objects such as lambdas or functors.
+///
+/// \tparam Func The callable type from which to extract the argument type.
+/// \tparam N The index of the argument to retrieve.
+///
+/// Example:
+/// using Func = void(int, double, const char*);
+/// static_assert(std::is_same<nth_argument_type<Func, 0>::type, int>::value,
+/// "Unexpected type");
+template <typename Func, std::size_t N> struct nth_argument_type {
+  template <typename R, typename... Args>
+  static auto
+      helper(R(Args...)) -> std::tuple_element_t<N, std::tuple<Args...>>;
+  using type = decltype(helper(std::declval<Func>()));
+};
+
 #ifdef _WIN32
 #define DPCT_EXPORT __declspec(dllexport)
 #else
diff --git a/clang/test/dpct/function_pointer.cu b/clang/test/dpct/function_pointer.cu
new file mode 100644
index 000000000000..b6af360240f7
--- /dev/null
+++ b/clang/test/dpct/function_pointer.cu
@@ -0,0 +1,187 @@
+// RUN: dpct --format-range=none -out-root %T/function_pointer %s --cuda-include-path="%cuda-path/include" -- -x cuda --cuda-host-only
+// RUN: FileCheck --input-file %T/function_pointer/function_pointer.dp.cpp --match-full-lines %s
+// RUN: %if build_lit %{icpx -c -fsycl %T/function_pointer/function_pointer.dp.cpp -o %T/function_pointer/function_pointer.dp.o %}
+
+#include <cuda_runtime.h>
+#include <iostream>
+
+__global__ void vectorAdd(const int *A, int *B, int *C, int N) {
+    int i = blockIdx.x * blockDim.x + threadIdx.x;
+    if (i < N) {
+        C[i] = A[i] + B[i];
+    }
+}
+
+// CHECK:  void vectorAdd_wrapper(const int * A ,int * B ,int * C ,int N) {
+// CHECK:        sycl::queue queue = *dpct::kernel_launcher::_que;
+// CHECK:        unsigned int localMemSize = dpct::kernel_launcher::_local_mem_size;
+// CHECK:        sycl::nd_range<3> nr = dpct::kernel_launcher::_nr;
+// CHECK:        queue.parallel_for(
+// CHECK:          nr,
+// CHECK:          [=](sycl::nd_item<3> item_ct1) {
+// CHECK:            vectorAdd(A, B, C, N, item_ct1);
+// CHECK:          });
+// CHECK:  }
+
+template<typename T>
+__global__ 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];
+    }
+}
+
+// CHECK:  template<typename T>
+// CHECK:  void vectorTemplateAdd_wrapper(const T * A ,T * B ,T * C ,int N) {
+// CHECK:      sycl::queue queue = *dpct::kernel_launcher::_que;
+// CHECK:      unsigned int localMemSize = dpct::kernel_launcher::_local_mem_size;
+// CHECK:      sycl::nd_range<3> nr = dpct::kernel_launcher::_nr;
+// CHECK:      queue.parallel_for(
+// CHECK:        nr,
+// CHECK:        [=](sycl::nd_item<3> item_ct1) {
+// CHECK:          vectorTemplateAdd<T>(A, B, C, N, item_ct1);
+// CHECK:        });
+// CHECK:  }
+
+template <typename T>
+using fpt = void(*)(const T *, T*, T*, int);
+
+void foo() {
+    int N = 10;
+    size_t size = N * sizeof(int);
+
+    int *h_A = new int[N];
+    int *h_B = new int[N];
+    int *h_C = new int[N];
+
+    for (int i = 0; i < N; ++i) {
+        h_A[i] = static_cast<int>(i);
+        h_B[i] = static_cast<int>(i * 2);
+    }
+
+    int *d_A, *d_B, *d_C;
+    cudaMalloc(&d_A, size);
+    cudaMalloc(&d_B, size);
+    cudaMalloc(&d_C, size);
+
+    cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
+    cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);
+
+// CHECK:  fpt<int> fp = dpct::wrapper_register(vectorAdd_wrapper).get();
+// CHECK:  dpct::kernel_launcher::launch(fp, 1, 10, 0, 0, d_A, d_B, d_C, N);
+    fpt<int> fp = vectorAdd;
+    fp<<<1, 10>>>(d_A, d_B, d_C, N);
+
+    cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);
+
+    std::cout << "Result: " << std::endl;
+    for (int i = 0; i < N; ++i) {
+        if(h_A[i] + h_B[i] != h_C[i]) {
+            std::cout << "test failed" << std::endl;
+            exit(-1);
+        }
+        std::cout << h_A[i] << " + " << h_B[i] << " = " << h_C[i] << std::endl;
+    }
+    
+    void *args[4];
+    args[0] = &d_A;
+    args[1] = &d_B;
+    args[2] = &d_C;
+    args[3] = &N;
+    // CHECK:  dpct::kernel_launcher::launch(fp, 1, 10, args, 0, 0);
+    cudaLaunchKernel((void *)fp, 1, 10, args, 0, 0);
+
+    cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);
+
+    std::cout << "Result: " << std::endl;
+    for (int i = 0; i < N; ++i) {
+        if(h_A[i] + h_B[i] != h_C[i]) {
+            std::cout << "test failed" << std::endl;
+            exit(-1);
+        }
+        std::cout << h_A[i] << " + " << h_B[i] << " = " << h_C[i] << std::endl;
+    }
+
+    // CHECK:  dpct::kernel_launcher::launch(fp, 1, 10, args, 0, 0);
+    cudaLaunchKernel<void(const int*, int*, int*, int)>(fp, 1, 10, args, 0, 0);
+
+    cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);
+
+    std::cout << "Result: " << std::endl;
+    for (int i = 0; i < N; ++i) {
+        if(h_A[i] + h_B[i] != h_C[i]) {
+            std::cout << "test failed" << std::endl;
+            exit(-1);
+        }
+        std::cout << h_A[i] << " + " << h_B[i] << " = " << h_C[i] << std::endl;
+    }
+
+
+    cudaFree(d_A);
+    cudaFree(d_B);
+    cudaFree(d_C);
+
+    delete[] h_A;
+    delete[] h_B;
+    delete[] h_C;
+}
+
+template<typename T>
+void goo(fpt<T> p) {
+    int N = 10;
+    size_t size = N * sizeof(int);
+
+    int *h_A = new int[N];
+    int *h_B = new int[N];
+    int *h_C = new int[N];
+
+    for (int i = 0; i < N; ++i) {
+        h_A[i] = static_cast<int>(i);
+        h_B[i] = static_cast<int>(i * 2);
+    }
+
+    int *d_A, *d_B, *d_C;
+    cudaMalloc(&d_A, size);
+    cudaMalloc(&d_B, size);
+    cudaMalloc(&d_C, size);
+
+    cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
+    cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);
+    // CHECK:  dpct::kernel_launcher::launch(p, 1, 10, 0, 0, d_A, d_B, d_C, N);
+    p<<<1, 10>>>(d_A, d_B, d_C, N);
+
+    cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);
+
+    std::cout << "Result: " << std::endl;
+    for (int i = 0; i < N; ++i) {
+        if(h_A[i] + h_B[i] != h_C[i]) {
+            std::cout << "test failed" << std::endl;
+            exit(-1);
+        }
+        std::cout << h_A[i] << " + " << h_B[i] << " = " << h_C[i] << std::endl;
+    }
+
+    cudaFree(d_A);
+    cudaFree(d_B);
+    cudaFree(d_C);
+
+    delete[] h_A;
+    delete[] h_B;
+    delete[] h_C;
+}
+
+template <typename T>
+void hoo() {
+    // CHECK: fpt<int> a = dpct::wrapper_register<decltype(a)>(vectorTemplateAdd_wrapper).get();
+  fpt<int> a = vectorTemplateAdd;
+  goo<int>(a);
+  // CHECK:  goo<T>(dpct::wrapper_register<typename dpct::nth_argument_type<decltype(goo<T>), 0>::type>(vectorTemplateAdd_wrapper).get());
+  goo<T>(vectorTemplateAdd);
+}
+
+int main() {
+  hoo<int>();
+  foo();
+  std::cout << "test success" << std::endl;
+  return 0;
+}
diff --git a/clang/test/dpct/launch-kernel-cooperative-usm.cu b/clang/test/dpct/launch-kernel-cooperative-usm.cu
index 8261186a36a8..2f5a89b68451 100644
--- a/clang/test/dpct/launch-kernel-cooperative-usm.cu
+++ b/clang/test/dpct/launch-kernel-cooperative-usm.cu
@@ -23,6 +23,23 @@ __global__ void kernel(int *d, cudaTextureObject_t tex) {
   int gtid = blockIdx.x * blockDim.x + threadIdx.x;
   tex1D(d + gtid, tex, gtid);
 }
+// CHECK: // Auto generated SYCL kernel wrapper used to migration kernel function pointer.
+// CHECK: void kernel_wrapper(int * d ,dpct::image_wrapper_base_p tex) {
+// CHECK:   sycl::queue queue = *dpct::kernel_launcher::_que;
+// CHECK:   unsigned int localMemSize = dpct::kernel_launcher::_local_mem_size;
+// CHECK:   sycl::nd_range<3> nr = dpct::kernel_launcher::_nr;
+// CHECK:     static_cast<dpct::image_wrapper<int, 1> *>(tex)->create_image(queue);
+// CHECK:     queue.submit(
+// CHECK:       [&](sycl::handler &cgh) {
+// CHECK:         auto tex_acc = static_cast<dpct::image_wrapper<int, 1> *>(tex)->get_access(cgh, queue);
+// CHECK:         auto tex_smpl = (tex)->get_sampler();
+// CHECK:         cgh.parallel_for(
+// CHECK:           nr,
+// CHECK:           [=](sycl::nd_item<3> item_ct1) {
+// CHECK:             kernel(d, dpct::image_accessor_ext<int, 1>(tex_smpl, tex_acc), item_ct1);
+// CHECK:           });
+// CHECK:       });
+// CHECK: }
 
 int main() {
   int *d_data;
@@ -80,16 +97,10 @@ int main() {
   // CHECK-NEXT:  });
   cudaLaunchCooperativeKernel((const void *)&template_kernel<int>, dim3(16), dim3(16), args, 32, stream);
 
+  // CHECK: void *kernel_func = (void *)dpct::wrapper_register(&kernel_wrapper).get();
   void *kernel_func = (void *)&kernel;
 
-  // CHECK: /*
-  // CHECK-NEXT: DPCT1123:{{[0-9]+}}: The kernel function pointer cannot be used in the device code. You need to call the kernel function with the correct argument(s) directly. According to the kernel function definition, adjusting the dimension of the sycl::nd_item may also be required.
-  // CHECK-NEXT: */
-  // CHECK-NEXT: q_ct1.parallel_for(
-  // CHECK-NEXT:   sycl::nd_range<3>(sycl::range<3>(1, 1, 16) * sycl::range<3>(1, 1, 16), sycl::range<3>(1, 1, 16)), 
-  // CHECK-NEXT:   [=](sycl::nd_item<3> item_ct1) {
-  // CHECK-NEXT:     kernel_func();
-  // CHECK-NEXT:   });
+  // CHECK: dpct::kernel_launcher::launch(kernel_func, dpct::dim3(16), dpct::dim3(16), args, 0, 0);
   cudaLaunchCooperativeKernel(kernel_func, dim3(16), dim3(16), args, 0, 0);
 
   cudaStreamDestroy(stream);
diff --git a/clang/test/dpct/launch-kernel-cooperative.cu b/clang/test/dpct/launch-kernel-cooperative.cu
index 4954f87307fe..11e398e0265e 100644
--- a/clang/test/dpct/launch-kernel-cooperative.cu
+++ b/clang/test/dpct/launch-kernel-cooperative.cu
@@ -78,17 +78,10 @@ int main() {
   // CHECK-NEXT:       });
   // CHECK-NEXT:   });
   cudaLaunchCooperativeKernel((const void *)&template_kernel<int>, dim3(16), dim3(16), args, 32, stream);
-
+  // CHECK: void *kernel_func = (void *)dpct::wrapper_register(&kernel_wrapper).get();
   void *kernel_func = (void *)&kernel;
 
-  // CHECK: /*
-  // CHECK-NEXT: DPCT1123:{{[0-9]+}}: The kernel function pointer cannot be used in the device code. You need to call the kernel function with the correct argument(s) directly. According to the kernel function definition, adjusting the dimension of the sycl::nd_item may also be required.
-  // CHECK-NEXT: */
-  // CHECK-NEXT: q_ct1.parallel_for(
-  // CHECK-NEXT:   sycl::nd_range<3>(sycl::range<3>(1, 1, 16) * sycl::range<3>(1, 1, 16), sycl::range<3>(1, 1, 16)), 
-  // CHECK-NEXT:   [=](sycl::nd_item<3> item_ct1) {
-  // CHECK-NEXT:     kernel_func();
-  // CHECK-NEXT:   });
+  // CHECK: dpct::kernel_launcher::launch(kernel_func, dpct::dim3(16), dpct::dim3(16), args, 0, 0);
   cudaLaunchCooperativeKernel(kernel_func, dim3(16), dim3(16), args, 0, 0);
 
   cudaStreamDestroy(stream);
diff --git a/clang/test/dpct/launch-kernel-usm.cu b/clang/test/dpct/launch-kernel-usm.cu
index 13c209edb037..17cdc4fe6923 100644
--- a/clang/test/dpct/launch-kernel-usm.cu
+++ b/clang/test/dpct/launch-kernel-usm.cu
@@ -76,28 +76,15 @@ int main() {
   // CHECK-NEXT:      });
   // CHECK-NEXT:  });
   cudaLaunchKernel((const void *)&template_kernel<int>, dim3(16), dim3(16), args, 32, stream);
-
+  // CHECK: void *kernel_func = (void *)dpct::wrapper_register(&kernel_wrapper).get();
   void *kernel_func = (void *)&kernel;
-  // CHECK: /*
-  // CHECK-NEXT: DPCT1123:{{[0-9]+}}: The kernel function pointer cannot be used in the device code. You need to call the kernel function with the correct argument(s) directly. According to the kernel function definition, adjusting the dimension of the sycl::nd_item may also be required.
-  // CHECK-NEXT: */
-  // CHECK-NEXT: q_ct1.parallel_for(
-  // CHECK-NEXT:   sycl::nd_range<3>(sycl::range<3>(1, 1, 16) * sycl::range<3>(1, 1, 16), sycl::range<3>(1, 1, 16)), 
-  // CHECK-NEXT:   [=](sycl::nd_item<3> item_ct1) {
-  // CHECK-NEXT:     kernel_func();
-  // CHECK-NEXT:   });
+  // CHECK: dpct::kernel_launcher::launch(kernel_func, dpct::dim3(16), dpct::dim3(16), args, 0, 0);
   cudaLaunchKernel(kernel_func, dim3(16), dim3(16), args, 0, 0);
 
   void *kernel_array[100];
+  // CHECK: kernel_array[10] = (void *)dpct::wrapper_register(&kernel_wrapper).get();
   kernel_array[10] = (void *)&kernel;
-  // CHECK: /*
-  // CHECK-NEXT: DPCT1123:{{[0-9]+}}: The kernel function pointer cannot be used in the device code. You need to call the kernel function with the correct argument(s) directly. According to the kernel function definition, adjusting the dimension of the sycl::nd_item may also be required.
-  // CHECK-NEXT: */
-  // CHECK-NEXT: q_ct1.parallel_for(
-  // CHECK-NEXT:   sycl::nd_range<3>(sycl::range<3>(1, 1, 16) * sycl::range<3>(1, 1, 16), sycl::range<3>(1, 1, 16)), 
-  // CHECK-NEXT:   [=](sycl::nd_item<3> item_ct1) {
-  // CHECK-NEXT:     (kernel_array[10])();
-  // CHECK-NEXT:   });
+  // CHECK: dpct::kernel_launcher::launch(kernel_array[10], dpct::dim3(16), dpct::dim3(16), args, 0, 0);
   cudaLaunchKernel(kernel_array[10], dim3(16), dim3(16), args, 0, 0);
 
   cudaStreamDestroy(stream);
diff --git a/clang/test/dpct/launch-kernel.cu b/clang/test/dpct/launch-kernel.cu
index 416ee0d19fb6..5523a7d20803 100644
--- a/clang/test/dpct/launch-kernel.cu
+++ b/clang/test/dpct/launch-kernel.cu
@@ -76,16 +76,9 @@ int main() {
   // CHECK-NEXT:       });
   // CHECK-NEXT:   });
   cudaLaunchKernel((const void *)&template_kernel<int>, dim3(16), dim3(16), args, 32, stream);
-
+  // CHECK:  void *kernel_func = (void *)dpct::wrapper_register(&kernel_wrapper).get();
   void *kernel_func = (void *)&kernel;
-  // CHECK: /*
-  // CHECK-NEXT: DPCT1123:{{[0-9]+}}: The kernel function pointer cannot be used in the device code. You need to call the kernel function with the correct argument(s) directly. According to the kernel function definition, adjusting the dimension of the sycl::nd_item may also be required.
-  // CHECK-NEXT: */
-  // CHECK-NEXT: q_ct1.parallel_for(
-  // CHECK-NEXT:   sycl::nd_range<3>(sycl::range<3>(1, 1, 16) * sycl::range<3>(1, 1, 16), sycl::range<3>(1, 1, 16)), 
-  // CHECK-NEXT:   [=](sycl::nd_item<3> item_ct1) {
-  // CHECK-NEXT:     kernel_func();
-  // CHECK-NEXT:   });
+  // CHECK:  dpct::kernel_launcher::launch(kernel_func, dpct::dim3(16), dpct::dim3(16), args, 0, 0);
   cudaLaunchKernel(kernel_func, dim3(16), dim3(16), args, 0, 0);
 
   cudaStreamDestroy(stream);
diff --git a/clang/test/dpct/launch_kernel/k.inc b/clang/test/dpct/launch_kernel/k.inc
index 6817c9b9bbdf..acc90a7f8f55 100644
--- a/clang/test/dpct/launch_kernel/k.inc
+++ b/clang/test/dpct/launch_kernel/k.inc
@@ -26,14 +26,11 @@ void foo38() {
   //CHECK-NEXT:/*
   //CHECK-NEXT:DPCT1049:{{[0-9]+}}: The work-group size passed to the SYCL kernel may exceed the limit. To get the device limit, query info::device::max_work_group_size. Adjust the work-group size if needed.
   //CHECK-NEXT:*/
-  //CHECK-NEXT:/*
-  //CHECK-NEXT:DPCT1123:{{[0-9]+}}: The kernel function pointer cannot be used in the device code. You need to call the kernel function with the correct argument(s) directly. According to the kernel function definition, adjusting the dimension of the sycl::nd_item may also be required.
-  //CHECK-NEXT:*/
   //CHECK-NEXT:CHECK_1([&](){
   //CHECK-NEXT:stream->parallel_for(
   //CHECK-NEXT:sycl::nd_range<3>(sycl::range<3>(z, y, x) * sycl::range<3>(1, 1, block), sycl::range<3>(1, 1, block)), 
   //CHECK-NEXT:[=](sycl::nd_item<3> item_ct1) {
-  //CHECK-NEXT: ((void*)&kernel38<T>)();
+  //CHECK-NEXT: kernel38<T>();
   //CHECK-NEXT:});
   //CHECK-NEXT:  return 0;
   //CHECK-NEXT:}());
diff --git a/clang/test/dpct/macro_test.cu b/clang/test/dpct/macro_test.cu
index cbedd59037f3..645a78d8254a 100644
--- a/clang/test/dpct/macro_test.cu
+++ b/clang/test/dpct/macro_test.cu
@@ -1332,7 +1332,7 @@ void foo38() {
   //CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(z, y, x) * sycl::range<3>(1, 1, block),
   //CHECK-NEXT:                   sycl::range<3>(1, 1, block)),
   //CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) {
-  //CHECK-NEXT:   ((void *)&kernel38<T>)();
+  //CHECK-NEXT:   kernel38<T>();
   //CHECK-NEXT: });
   //CHECK-NEXT: return 0;
   //CHECK-NEXT: }());
@@ -1342,7 +1342,7 @@ void foo38() {
   //CHECK-NEXT:       sycl::nd_range<3>(sycl::range<3>(z, y, x) * sycl::range<3>(1, 1, block),
   //CHECK-NEXT:                         sycl::range<3>(1, 1, block)),
   //CHECK-NEXT:       [=](sycl::nd_item<3> item_ct1) {
-  //CHECK-NEXT:         ((void *)&kernel38<T>)();
+  //CHECK-NEXT:         kernel38<T>();
   //CHECK-NEXT:       });
   //CHECK-NEXT:   return 0;
   //CHECK-NEXT: }();
diff --git a/clang/test/dpct/query_api_mapping/Runtime/test-after9.cu b/clang/test/dpct/query_api_mapping/Runtime/test-after9.cu
index 0f80158e8e2f..263ee5dc3c06 100644
--- a/clang/test/dpct/query_api_mapping/Runtime/test-after9.cu
+++ b/clang/test/dpct/query_api_mapping/Runtime/test-after9.cu
@@ -16,8 +16,4 @@
 // CUDALAUNCHCOOPERATIVEKERNEL-NEXT:                               blockDim /*dim3*/, args /*void ***/,
 // CUDALAUNCHCOOPERATIVEKERNEL-NEXT:                               sharedMem /*size_t*/, s /*cudaStream_t*/);
 // CUDALAUNCHCOOPERATIVEKERNEL-NEXT: Is migrated to:
-// CUDALAUNCHCOOPERATIVEKERNEL-NEXT:   s->parallel_for(
-// CUDALAUNCHCOOPERATIVEKERNEL-NEXT:     sycl::nd_range<3>(gridDim * blockDim, blockDim),
-// CUDALAUNCHCOOPERATIVEKERNEL-NEXT:     [=](sycl::nd_item<3> item_ct1) {
-// CUDALAUNCHCOOPERATIVEKERNEL-NEXT:       f();
-// CUDALAUNCHCOOPERATIVEKERNEL-NEXT:     });
+// CUDALAUNCHCOOPERATIVEKERNEL-NEXT:   dpct::kernel_launcher::launch(f, gridDim, blockDim, args, sharedMem, s);
diff --git a/clang/test/dpct/query_api_mapping/Runtime/test.cu b/clang/test/dpct/query_api_mapping/Runtime/test.cu
index 37dabec57b9e..35eaf29d89b8 100644
--- a/clang/test/dpct/query_api_mapping/Runtime/test.cu
+++ b/clang/test/dpct/query_api_mapping/Runtime/test.cu
@@ -351,11 +351,7 @@
 // CUDALAUNCHKERNEL-NEXT:   cudaLaunchKernel(f /*cudaError_t*/, gridDim /*dim3*/, blockDim /*dim3*/,
 // CUDALAUNCHKERNEL-NEXT:                    args /*void ***/, sharedMem /*size_t*/, s /*cudaStream_t*/);
 // CUDALAUNCHKERNEL-NEXT: Is migrated to:
-// CUDALAUNCHKERNEL-NEXT:   s->parallel_for(
-// CUDALAUNCHKERNEL-NEXT:     sycl::nd_range<3>(gridDim * blockDim, blockDim),
-// CUDALAUNCHKERNEL-NEXT:     [=](sycl::nd_item<3> item_ct1) {
-// CUDALAUNCHKERNEL-NEXT:       f();
-// CUDALAUNCHKERNEL-NEXT:     });
+// CUDALAUNCHKERNEL-NEXT:   dpct::kernel_launcher::launch(f, gridDim, blockDim, args, sharedMem, s);
 
 /// Occupancy