diff --git a/clang/include/clang/AST/RecursiveASTVisitor.h b/clang/include/clang/AST/RecursiveASTVisitor.h
index c8cc94fa1e86f..2e44c1371cdc0 100644
--- a/clang/include/clang/AST/RecursiveASTVisitor.h
+++ b/clang/include/clang/AST/RecursiveASTVisitor.h
@@ -2999,6 +2999,13 @@ DEF_TRAVERSE_STMT(ParenListExpr, {})
 DEF_TRAVERSE_STMT(SYCLUniqueStableNameExpr, {
   TRY_TO(TraverseTypeLoc(S->getTypeSourceInfo()->getTypeLoc()));
 })
+DEF_TRAVERSE_STMT(UnresolvedSYCLKernelCallStmt, {
+  if (getDerived().shouldVisitImplicitCode()) {
+    TRY_TO(TraverseStmt(S->getOriginalStmt()));
+    TRY_TO(TraverseStmt(S->getKernelLaunchIdExpr()));
+    ShouldVisitChildren = false;
+  }
+})
 DEF_TRAVERSE_STMT(OpenACCAsteriskSizeExpr, {})
 DEF_TRAVERSE_STMT(PredefinedExpr, {})
 DEF_TRAVERSE_STMT(ShuffleVectorExpr, {})
diff --git a/clang/include/clang/AST/StmtSYCL.h b/clang/include/clang/AST/StmtSYCL.h
index 70d8137992110..c8da4987321a9 100644
--- a/clang/include/clang/AST/StmtSYCL.h
+++ b/clang/include/clang/AST/StmtSYCL.h
@@ -99,6 +99,59 @@ class SYCLKernelCallStmt : public Stmt {
   }
 };
 
+// UnresolvedSYCLKernelCallStmt represents an invocation of a SYCL kernel in
+// a dependent context for which lookup of the sycl_enqueue_kernel_launch
+// identifier cannot be performed. These statements are transformed to
+// SYCLKernelCallStmt during template instantiation.
+class UnresolvedSYCLKernelCallStmt : public Stmt {
+  friend class ASTStmtReader;
+  Stmt *OriginalStmt = nullptr;
+  // KernelLaunchIdExpr stores an UnresolvedLookupExpr or UnresolvedMemberExpr
+  // corresponding to the SYCL kernel launch function for which a call
+  // will be synthesized during template instantiation.
+  Expr *KernelLaunchIdExpr = nullptr;
+  UnresolvedSYCLKernelCallStmt(CompoundStmt *CS, Expr *IdExpr)
+      : Stmt(UnresolvedSYCLKernelCallStmtClass), OriginalStmt(CS),
+        KernelLaunchIdExpr(IdExpr) {}
+
+  void setKernelLaunchIdExpr(Expr *IdExpr) { KernelLaunchIdExpr = IdExpr; }
+  void setOriginalStmt(CompoundStmt *CS) { OriginalStmt = CS; }
+
+public:
+  static UnresolvedSYCLKernelCallStmt *
+  Create(const ASTContext &C, CompoundStmt *CS, Expr *IdExpr) {
+    return new (C) UnresolvedSYCLKernelCallStmt(CS, IdExpr);
+  }
+
+  static UnresolvedSYCLKernelCallStmt *CreateEmpty(const ASTContext &C) {
+    return new (C) UnresolvedSYCLKernelCallStmt(nullptr, nullptr);
+  }
+
+  Expr *getKernelLaunchIdExpr() const { return KernelLaunchIdExpr; }
+  CompoundStmt *getOriginalStmt() { return cast<CompoundStmt>(OriginalStmt); }
+  const CompoundStmt *getOriginalStmt() const {
+    return cast<CompoundStmt>(OriginalStmt);
+  }
+
+  SourceLocation getBeginLoc() const LLVM_READONLY {
+    return getOriginalStmt()->getBeginLoc();
+  }
+
+  SourceLocation getEndLoc() const LLVM_READONLY {
+    return getOriginalStmt()->getEndLoc();
+  }
+  static bool classof(const Stmt *T) {
+    return T->getStmtClass() == UnresolvedSYCLKernelCallStmtClass;
+  }
+  child_range children() {
+    return child_range(&OriginalStmt, &OriginalStmt + 1);
+  }
+
+  const_child_range children() const {
+    return const_child_range(&OriginalStmt, &OriginalStmt + 1);
+  }
+};
+
 } // end namespace clang
 
 #endif // LLVM_CLANG_AST_STMTSYCL_H
diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td
index 4445a9094a9b7..7ebadfb25d9ca 100644
--- a/clang/include/clang/Basic/AttrDocs.td
+++ b/clang/include/clang/Basic/AttrDocs.td
@@ -554,7 +554,7 @@ follows.
   namespace sycl {
   class handler {
     template<typename KernelNameType, typename... Ts>
-    void sycl_enqueue_kernel_launch(const char *KernelName, Ts...) {
+    void sycl_kernel_launch(const char *KernelName, Ts...) {
       // Call functions appropriate for the desired offload backend
       // (OpenCL, CUDA, HIP, Level Zero, etc...) to enqueue kernel invocation.
     }
@@ -622,7 +622,7 @@ The offload kernel entry point for a SYCL kernel performs the following tasks:
 The ``sycl_kernel_entry_point`` attribute facilitates or automates these tasks
 by generating the offload kernel entry point, generating a unique symbol name
 for it, synthesizing code for kernel argument decomposition and reconstruction,
-and synthesizing a call to a ``sycl_enqueue_kernel_launch`` function template
+and synthesizing a call to a ``sycl_kernel_launch`` function template
 with the kernel name type, kernel symbol name, and (decomposed) kernel arguments
 passed as template or function arguments.
 
@@ -690,7 +690,7 @@ replaced with synthesized code that looks approximately as follows.
 
   sycl::stream sout = Kernel.sout;
   S s = Kernel.s;
-  sycl_enqueue_kernel_launch<KN>("kernel-symbol-name", sout, s);
+  sycl_kernel_launch<KN>("kernel-symbol-name", sout, s);
 
 There are a few items worthy of note:
 
@@ -701,16 +701,16 @@ There are a few items worthy of note:
 #. ``kernel-symbol-name`` is substituted for the actual symbol name that would
    be generated; these names are implementation details subject to change.
 
-#. Lookup for the ``sycl_enqueue_kernel_launch()`` function template is
+#. Lookup for the ``sycl_kernel_launch()`` function template is
    performed from the (possibly instantiated) location of the definition of
    ``kernel_entry_point()``. If overload resolution fails, the program is
    ill-formed. If the selected overload is a non-static member function, then
    ``this`` is passed for the implicit object parameter.
 
-#. Function arguments passed to ``sycl_enqueue_kernel_launch()`` are passed
+#. Function arguments passed to ``sycl_kernel_launch()`` are passed
    as if by ``std::forward<X>(x)``.
 
-#. The ``sycl_enqueue_kernel_launch()`` function is expected to be provided by
+#. The ``sycl_kernel_launch()`` function is expected to be provided by
    the SYCL library implementation. It is responsible for scheduling execution
    of the generated offload kernel entry point identified by
    ``kernel-symbol-name`` and copying the (decomposed) kernel arguments to
@@ -721,7 +721,7 @@ attribute to be called for the offload kernel entry point to be emitted. For
 inline functions and function templates, any ODR-use will suffice. For other
 functions, an ODR-use is not required; the offload kernel entry point will be
 emitted if the function is defined. In any case, a call to the function is
-required for the synthesized call to ``sycl_enqueue_kernel_launch()`` to occur.
+required for the synthesized call to ``sycl_kernel_launch()`` to occur.
 
 Functions declared with the ``sycl_kernel_entry_point`` attribute are not
 limited to the simple example shown above. They may have additional template
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index baf633b2f92cb..1cd9803ad8741 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -13015,6 +13015,15 @@ def err_sycl_entry_point_return_type : Error<
 def err_sycl_entry_point_deduced_return_type : Error<
   "the %0 attribute only applies to functions with a non-deduced 'void' return"
   " type">;
+def err_sycl_host_no_launch_function : Error<
+  "unable to find suitable 'sycl_kernel_launch' function for host code "
+  "synthesis">;
+def warn_sycl_device_no_host_launch_function : Warning<
+  "unable to find suitable 'sycl_kernel_launch' function for host code "
+  "synthesis">,
+  InGroup<DiagGroup<"sycl-host-launcher">>;
+def note_sycl_host_launch_function : Note<
+  "define 'sycl_kernel_launch' function template to fix this problem">;
 
 def warn_cuda_maxclusterrank_sm_90 : Warning<
   "maxclusterrank requires sm_90 or higher, CUDA arch provided: %0, ignoring "
diff --git a/clang/include/clang/Basic/StmtNodes.td b/clang/include/clang/Basic/StmtNodes.td
index c9c173f5c7469..1696f55b813ad 100644
--- a/clang/include/clang/Basic/StmtNodes.td
+++ b/clang/include/clang/Basic/StmtNodes.td
@@ -25,6 +25,7 @@ def CaseStmt : StmtNode<SwitchCase>;
 def DefaultStmt : StmtNode<SwitchCase>;
 def CapturedStmt : StmtNode<Stmt>;
 def SYCLKernelCallStmt : StmtNode<Stmt>;
+def UnresolvedSYCLKernelCallStmt : StmtNode<Stmt>;
 
 // Statements that might produce a value (for example, as the last non-null
 // statement in a GNU statement-expression).
diff --git a/clang/include/clang/Sema/ScopeInfo.h b/clang/include/clang/Sema/ScopeInfo.h
index 4f4d38c961140..f334f58ebd0a7 100644
--- a/clang/include/clang/Sema/ScopeInfo.h
+++ b/clang/include/clang/Sema/ScopeInfo.h
@@ -245,6 +245,10 @@ class FunctionScopeInfo {
   /// The set of GNU address of label extension "&&label".
   llvm::SmallVector<AddrLabelExpr *, 4> AddrLabels;
 
+  /// An unresolved identifier lookup expression for an implicit call
+  /// to a SYCL kernel launch function in a dependent context.
+  Expr *SYCLKernelLaunchIdExpr = nullptr;
+
 public:
   /// Represents a simple identification of a weak object.
   ///
diff --git a/clang/include/clang/Sema/SemaSYCL.h b/clang/include/clang/Sema/SemaSYCL.h
index 7ae556da2bec1..76046b765c0d6 100644
--- a/clang/include/clang/Sema/SemaSYCL.h
+++ b/clang/include/clang/Sema/SemaSYCL.h
@@ -66,7 +66,11 @@ class SemaSYCL : public SemaBase {
 
   void CheckSYCLExternalFunctionDecl(FunctionDecl *FD);
   void CheckSYCLEntryPointFunctionDecl(FunctionDecl *FD);
-  StmtResult BuildSYCLKernelCallStmt(FunctionDecl *FD, CompoundStmt *Body);
+  StmtResult BuildSYCLKernelCallStmt(FunctionDecl *FD, CompoundStmt *Body,
+                                     Expr *LaunchIdExpr);
+  ExprResult BuildSYCLKernelLaunchIdExpr(FunctionDecl *FD, QualType KNT);
+  StmtResult BuildUnresolvedSYCLKernelCallStmt(CompoundStmt *CS,
+                                                     Expr *IdExpr);
 };
 
 } // namespace clang
diff --git a/clang/include/clang/Serialization/ASTBitCodes.h b/clang/include/clang/Serialization/ASTBitCodes.h
index 441047d64f48c..0f170a40f05ca 100644
--- a/clang/include/clang/Serialization/ASTBitCodes.h
+++ b/clang/include/clang/Serialization/ASTBitCodes.h
@@ -1615,6 +1615,9 @@ enum StmtCode {
   /// A SYCLKernelCallStmt record.
   STMT_SYCLKERNELCALL,
 
+  /// A SYCLKernelCallStmt record.
+  STMT_UNRESOLVED_SYCL_KERNEL_CALL,
+
   /// A GCC-style AsmStmt record.
   STMT_GCCASM,
 
diff --git a/clang/lib/AST/ComputeDependence.cpp b/clang/lib/AST/ComputeDependence.cpp
index e0cf0deb12bd2..0192763c65f09 100644
--- a/clang/lib/AST/ComputeDependence.cpp
+++ b/clang/lib/AST/ComputeDependence.cpp
@@ -16,6 +16,7 @@
 #include "clang/AST/ExprConcepts.h"
 #include "clang/AST/ExprObjC.h"
 #include "clang/AST/ExprOpenMP.h"
+#include "clang/AST/StmtSYCL.h"
 #include "clang/Basic/ExceptionSpecificationType.h"
 #include "llvm/ADT/ArrayRef.h"
 
diff --git a/clang/lib/AST/StmtPrinter.cpp b/clang/lib/AST/StmtPrinter.cpp
index 5272ecba7fc93..9cc251a020096 100644
--- a/clang/lib/AST/StmtPrinter.cpp
+++ b/clang/lib/AST/StmtPrinter.cpp
@@ -1428,6 +1428,11 @@ void StmtPrinter::VisitSYCLUniqueStableNameExpr(
   OS << ")";
 }
 
+void StmtPrinter::VisitUnresolvedSYCLKernelCallStmt(
+    UnresolvedSYCLKernelCallStmt *Node) {
+  PrintStmt(Node->getOriginalStmt());
+}
+
 void StmtPrinter::VisitPredefinedExpr(PredefinedExpr *Node) {
   OS << PredefinedExpr::getIdentKindName(Node->getIdentKind());
 }
diff --git a/clang/lib/AST/StmtProfile.cpp b/clang/lib/AST/StmtProfile.cpp
index 2035fa7635f2a..f1fd1c869cc52 100644
--- a/clang/lib/AST/StmtProfile.cpp
+++ b/clang/lib/AST/StmtProfile.cpp
@@ -1374,6 +1374,11 @@ void StmtProfiler::VisitSYCLUniqueStableNameExpr(
   VisitType(S->getTypeSourceInfo()->getType());
 }
 
+void StmtProfiler::VisitUnresolvedSYCLKernelCallStmt(
+    const UnresolvedSYCLKernelCallStmt *S) {
+  VisitStmt(S);
+}
+
 void StmtProfiler::VisitPredefinedExpr(const PredefinedExpr *S) {
   VisitExpr(S);
   ID.AddInteger(llvm::to_underlying(S->getIdentKind()));
diff --git a/clang/lib/CodeGen/CGStmt.cpp b/clang/lib/CodeGen/CGStmt.cpp
index 9e26b145e0589..4b87f1622d1b5 100644
--- a/clang/lib/CodeGen/CGStmt.cpp
+++ b/clang/lib/CodeGen/CGStmt.cpp
@@ -117,6 +117,7 @@ void CodeGenFunction::EmitStmt(const Stmt *S, ArrayRef<const Attr *> Attrs) {
   case Stmt::CaseStmtClass:
   case Stmt::SEHLeaveStmtClass:
   case Stmt::SYCLKernelCallStmtClass:
+  case Stmt::UnresolvedSYCLKernelCallStmtClass:
     llvm_unreachable("should have emitted these statements as simple");
 
 #define STMT(Type, Base)
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index 67663ddec4cd2..2108e271bd430 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -15754,7 +15754,6 @@ Sema::ActOnStartOfFunctionDef(Scope *FnBodyScope, Declarator &D,
   if (!Bases.empty())
     OpenMP().ActOnFinishedFunctionDefinitionInOpenMPDeclareVariantScope(Dcl,
                                                                         Bases);
-
   return Dcl;
 }
 
@@ -16167,6 +16166,20 @@ Decl *Sema::ActOnStartOfFunctionDef(Scope *FnBodyScope, Decl *D,
 
   maybeAddDeclWithEffects(FD);
 
+  if (FD && !FD->isInvalidDecl() &&
+      FD->hasAttr<SYCLKernelEntryPointAttr>() && FnBodyScope) {
+    // Building KernelLaunchIdExpr requires performing an unqualified lookup
+    // which can only be done correctly while the stack of parsing scopes is
+    // alive, so we do it here when we start parsing function body even if it is
+    // a templated function.
+    const auto *SKEPAttr = FD->getAttr<SYCLKernelEntryPointAttr>();
+    if (!SKEPAttr->isInvalidAttr()) {
+      ExprResult LaunchIdExpr =
+          SYCL().BuildSYCLKernelLaunchIdExpr(FD, SKEPAttr->getKernelName());
+      getCurFunction()->SYCLKernelLaunchIdExpr = LaunchIdExpr.get();
+    }
+  }
+
   return D;
 }
 
@@ -16368,9 +16381,21 @@ Decl *Sema::ActOnFinishFunctionBody(Decl *dcl, Stmt *Body, bool IsInstantiation,
       SKEPAttr->setInvalidAttr();
     }
 
-    if (Body && !FD->isTemplated() && !SKEPAttr->isInvalidAttr()) {
-      StmtResult SR =
-          SYCL().BuildSYCLKernelCallStmt(FD, cast<CompoundStmt>(Body));
+    // We don't need to build SYCLKernelCallStmt for template instantiations
+    // since it was already created by template instantiator.
+    if (Body && !SKEPAttr->isInvalidAttr()) {
+      StmtResult SR;
+      if (FD->isTemplated()) {
+        SR = SYCL().BuildUnresolvedSYCLKernelCallStmt(
+            cast<CompoundStmt>(Body), getCurFunction()->SYCLKernelLaunchIdExpr);
+      } else if (FD->isTemplateInstantiation()) {
+        assert(isa<SYCLKernelCallStmt>(Body));
+        SR = Body;
+      } else {
+        SR = SYCL().BuildSYCLKernelCallStmt(
+            FD, cast<CompoundStmt>(Body),
+            getCurFunction()->SYCLKernelLaunchIdExpr);
+      }
       if (SR.isInvalid())
         return nullptr;
       Body = SR.get();
diff --git a/clang/lib/Sema/SemaExceptionSpec.cpp b/clang/lib/Sema/SemaExceptionSpec.cpp
index 15e3a39c86427..94eb3d0df3a4c 100644
--- a/clang/lib/Sema/SemaExceptionSpec.cpp
+++ b/clang/lib/Sema/SemaExceptionSpec.cpp
@@ -1261,6 +1261,9 @@ CanThrowResult Sema::canThrow(const Stmt *S) {
     return canSubStmtsThrow(*this, SKCS->getKernelLaunchStmt());
   }
 
+  case Stmt::UnresolvedSYCLKernelCallStmtClass:
+    return CT_Dependent;
+
     // ObjC message sends are like function calls, but never have exception
     // specs.
   case Expr::ObjCMessageExprClass:
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 56963fcb09853..dd332f06303ef 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -390,53 +390,23 @@ void SemaSYCL::CheckSYCLEntryPointFunctionDecl(FunctionDecl *FD) {
   }
 }
 
-namespace {
+ExprResult SemaSYCL::BuildSYCLKernelLaunchIdExpr(FunctionDecl *FD,
+                                                 QualType KNT) {
 
-CompoundStmt *BuildSYCLKernelLaunchStmt(Sema &SemaRef, FunctionDecl *FD,
-                                        const std::string &KernelName) {
   ASTContext &Ctx = SemaRef.getASTContext();
-  SmallVector<Stmt *> Stmts;
+  // Some routines need a valid source location to work correctly.
+  SourceLocation BodyLoc =
+      FD->getEndLoc().isValid() ? FD->getEndLoc() : FD->getLocation();
 
-  // Prepare a string literal that contains the kernel name.
-  QualType KernelNameCharTy = Ctx.CharTy.withConst();
-  llvm::APInt KernelNameSize(Ctx.getTypeSize(Ctx.getSizeType()),
-                             KernelName.size() + 1);
-  QualType KernelNameArrayTy = Ctx.getConstantArrayType(
-      KernelNameCharTy, KernelNameSize, nullptr, ArraySizeModifier::Normal, 0);
-  StringLiteral *KernelNameExpr = StringLiteral::Create(
-      Ctx, KernelName, StringLiteralKind::Ordinary,
-      /*Pascal*/ false, KernelNameArrayTy, SourceLocation());
-
-  // FIXME: An extern variable declaration with assignment to the kernel
-  // name expression is added to Stmts as a temporary measure to see results.
-  // reflected in tests. The kernel name expression will need to be passed as
-  // the first function argument in a call to sycl_enqueue_kernel_launch.
-  QualType ExternVarType = Ctx.getPointerType(Ctx.CharTy.withConst());
-  const IdentifierInfo *ExternVarName =
-      SemaRef.getPreprocessor().getIdentifierInfo("kernel_name");
-  VarDecl *ExternVarDecl = VarDecl::Create(
-      Ctx, FD, SourceLocation(), SourceLocation(), ExternVarName, ExternVarType,
-      /*TInfo*/ nullptr, SC_Extern);
-  DeclStmt *ExternVarDeclStmt = new (Ctx)
-      DeclStmt(DeclGroupRef(ExternVarDecl), SourceLocation(), SourceLocation());
-  Stmts.push_back(ExternVarDeclStmt);
-  DeclRefExpr *ExternVarDeclRef = new (Ctx) DeclRefExpr(
-      Ctx, ExternVarDecl, /*RefersToEnclosingVariableOrCapture*/ false,
-      ExternVarType, VK_LValue, SourceLocation());
-  ImplicitCastExpr *KernelNameArrayDecayExpr = new (Ctx) ImplicitCastExpr(
-      ImplicitCastExpr::OnStack, ExternVarType, CK_ArrayToPointerDecay,
-      KernelNameExpr, VK_PRValue, FPOptionsOverride());
-  BinaryOperator *AssignmentExpr = BinaryOperator::Create(
-      Ctx, ExternVarDeclRef, KernelNameArrayDecayExpr, BO_Assign, ExternVarType,
-      VK_LValue, OK_Ordinary, SourceLocation(), FPOptionsOverride());
-  Stmts.push_back(AssignmentExpr);
+  IdentifierInfo &LaunchFooName =
+      Ctx.Idents.get("sycl_kernel_launch", tok::TokenKind::identifier);
 
   // Perform overload resolution for a call to an accessible (member) function
-  // template named 'sycl_enqueue_kernel_launch' from within the definition of
-  // FD where:
+  // template named 'sycl_kernel_launch' "from within the definition of
+  // FD where":
   // - The kernel name type is passed as the first template argument.
-  // - Any remaining template parameters are deduced from the function arguments
-  //   or assigned by default template arguments.
+  // - Any remaining template parameters are deduced from the function
+  //   arguments or assigned by default template arguments.
   // - 'this' is passed as the implicit function argument if 'FD' is a
   //   non-static member function.
   // - The name of the kernel, expressed as a string literal, is passed as the
@@ -444,10 +414,108 @@ CompoundStmt *BuildSYCLKernelLaunchStmt(Sema &SemaRef, FunctionDecl *FD,
   // - The parameters of FD are forwarded as-if by 'std::forward()' as the
   //   remaining explicit function arguments.
   // - Any remaining function arguments are initialized by default arguments.
-  CompoundStmt *LaunchStmt = CompoundStmt::Create(
-      Ctx, Stmts, FPOptionsOverride(), SourceLocation(), SourceLocation());
+  LookupResult Result(SemaRef, &LaunchFooName, BodyLoc,
+                      Sema::LookupOrdinaryName);
+  CXXScopeSpec SS;
+  SemaRef.LookupTemplateName(Result, SemaRef.getCurScope(), SS,
+                             /*ObjectType=*/QualType(),
+                             /*EnteringContext=*/false, BodyLoc);
+
+  if (Result.empty() || Result.isAmbiguous()) {
+    SemaRef.Diag(BodyLoc, SemaRef.getLangOpts().SYCLIsHost
+                              ? diag::err_sycl_host_no_launch_function
+                              : diag::warn_sycl_device_no_host_launch_function);
+    SemaRef.Diag(BodyLoc, diag::note_sycl_host_launch_function);
+
+    return ExprError();
+  }
+
+  TemplateArgumentListInfo TALI{BodyLoc, BodyLoc};
+  TemplateArgument KNTA = TemplateArgument(KNT);
+  TemplateArgumentLoc TAL =
+      SemaRef.getTrivialTemplateArgumentLoc(KNTA, QualType(), BodyLoc);
+  TALI.addArgument(TAL);
+  ExprResult IdExpr;
+  if (SemaRef.isPotentialImplicitMemberAccess(SS, Result,
+                                              /*IsAddressOfOperand=*/false))
+    // BuildPossibleImplicitMemberExpr creates UnresolvedMemberExpr. Using it
+    // allows to pass implicit/explicit this argument automatically.
+    IdExpr = SemaRef.BuildPossibleImplicitMemberExpr(SS, BodyLoc, Result, &TALI,
+                                                     SemaRef.getCurScope());
+  else
+    IdExpr = SemaRef.BuildTemplateIdExpr(SS, BodyLoc, Result,
+                                         /*RequiresADL=*/true, &TALI);
+
+  // Can happen if SKEP attributed function is a static member, but the launcher
+  // is a regular member. Perhaps emit a note saying that we're in host code
+  // synthesis.
+  if (IdExpr.isInvalid())
+    return ExprError();
+
+  return IdExpr;
+}
 
-  return LaunchStmt;
+StmtResult SemaSYCL::BuildUnresolvedSYCLKernelCallStmt(CompoundStmt *CS,
+                                                             Expr *IdExpr) {
+  return UnresolvedSYCLKernelCallStmt::Create(SemaRef.getASTContext(), CS,
+                                                    IdExpr);
+}
+
+namespace {
+
+void PrepareKernelArgumentsForKernelLaunch(SmallVectorImpl<Expr *> &Args,
+                                                  const SYCLKernelInfo *SKI,
+                                                  Sema &SemaRef,
+                                                  SourceLocation Loc) {
+  assert(SKI && "Need a kernel!");
+  ASTContext &Ctx = SemaRef.getASTContext();
+
+  // Prepare a string literal that contains the kernel name.
+  const std::string KernelName = SKI->GetKernelName();
+  QualType KernelNameCharTy = Ctx.CharTy.withConst();
+  llvm::APInt KernelNameSize(Ctx.getTypeSize(Ctx.getSizeType()),
+                             KernelName.size() + 1);
+  QualType KernelNameArrayTy = Ctx.getConstantArrayType(
+      KernelNameCharTy, KernelNameSize, nullptr, ArraySizeModifier::Normal, 0);
+  Expr *KernelNameExpr =
+      StringLiteral::Create(Ctx, KernelName, StringLiteralKind::Ordinary,
+                            /*Pascal*/ false, KernelNameArrayTy, Loc);
+  Args.push_back(KernelNameExpr);
+
+  // Right now we simply forward the arguments of the skep-attributed function.
+  // With decomposition present there can be another logic.
+  // Make sure to use CurContext to avoid diagnostics that we're using a
+  // variable coming from another context. The function should be the same as in
+  // the kernel info though.
+  auto *FD = cast<FunctionDecl>(SemaRef.CurContext);
+  assert(declaresSameEntity(FD, SKI->getKernelEntryPointDecl()));
+  for (ParmVarDecl *PVD : FD->parameters()) {
+    QualType ParamType = PVD->getOriginalType().getNonReferenceType();
+    Expr *DRE = SemaRef.BuildDeclRefExpr(PVD, ParamType, VK_LValue, Loc);
+    assert(DRE);
+    Args.push_back(DRE);
+  }
+}
+
+StmtResult BuildSYCLKernelLaunchStmt(Sema &SemaRef,
+                                            const SYCLKernelInfo *SKI,
+                                            Expr *IdExpr, SourceLocation Loc) {
+  SmallVector<Stmt *> Stmts;
+  assert(SKI && "Need a Kernel!");
+
+  if (IdExpr) {
+    llvm::SmallVector<Expr *, 12> Args;
+    PrepareKernelArgumentsForKernelLaunch(Args, SKI, SemaRef, Loc);
+    ExprResult LaunchResult =
+        SemaRef.BuildCallExpr(SemaRef.getCurScope(), IdExpr, Loc, Args, Loc);
+    if (LaunchResult.isInvalid())
+      return StmtError();
+
+    Stmts.push_back(LaunchResult.get());
+  }
+
+  return CompoundStmt::Create(SemaRef.getASTContext(), Stmts,
+                              FPOptionsOverride(), Loc, Loc);
 }
 
 // The body of a function declared with the [[sycl_kernel_entry_point]]
@@ -535,11 +603,11 @@ OutlinedFunctionDecl *BuildSYCLKernelEntryPointOutline(Sema &SemaRef,
 } // unnamed namespace
 
 StmtResult SemaSYCL::BuildSYCLKernelCallStmt(FunctionDecl *FD,
-                                             CompoundStmt *Body) {
+                                             CompoundStmt *Body,
+                                             Expr *LaunchIdExpr) {
   assert(!FD->isInvalidDecl());
   assert(!FD->isTemplated());
   assert(FD->hasPrototype());
-
   // The current context must be the function definition context to ensure
   // that name lookup and parameter and local variable creation are performed
   // within the correct scope.
@@ -557,18 +625,19 @@ StmtResult SemaSYCL::BuildSYCLKernelCallStmt(FunctionDecl *FD,
   assert(declaresSameEntity(SKI.getKernelEntryPointDecl(), FD) &&
          "SYCL kernel name conflict");
 
-  // Build the kernel launch statement.
-  Stmt *LaunchStmt =
-      BuildSYCLKernelLaunchStmt(SemaRef, FD, SKI.GetKernelName());
-  assert(LaunchStmt);
-
   // Build the outline of the synthesized device entry point function.
   OutlinedFunctionDecl *OFD =
       BuildSYCLKernelEntryPointOutline(SemaRef, FD, Body);
   assert(OFD);
 
+  // Build host kernel launch stmt.
+  SourceLocation BodyLoc =
+      FD->getEndLoc().isValid() ? FD->getEndLoc() : FD->getLocation();
+  StmtResult LaunchRes =
+      BuildSYCLKernelLaunchStmt(SemaRef, &SKI, LaunchIdExpr, BodyLoc);
+
   Stmt *NewBody =
-      new (getASTContext()) SYCLKernelCallStmt(Body, LaunchStmt, OFD);
+      new (getASTContext()) SYCLKernelCallStmt(Body, LaunchRes.get(), OFD);
 
   return NewBody;
 }
diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h
index 1d14ead778446..9932d46824d2d 100644
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@@ -12933,6 +12933,25 @@ ExprResult TreeTransform<Derived>::TransformSYCLUniqueStableNameExpr(
       E->getLocation(), E->getLParenLocation(), E->getRParenLocation(), NewT);
 }
 
+template <typename Derived>
+StmtResult TreeTransform<Derived>::TransformUnresolvedSYCLKernelCallStmt(
+    UnresolvedSYCLKernelCallStmt *S) {
+  ExprResult IdExpr = getDerived().TransformExpr(S->getKernelLaunchIdExpr());
+
+  if (IdExpr.isInvalid())
+     return StmtError();
+
+  StmtResult Body = getDerived().TransformStmt(S->getOriginalStmt());
+  if (Body.isInvalid())
+     return StmtError();
+
+  StmtResult SR = SemaRef.SYCL().BuildSYCLKernelCallStmt(
+      cast<FunctionDecl>(SemaRef.CurContext), cast<CompoundStmt>(Body.get()),
+      IdExpr.get());
+
+  return SR;
+}
+
 template<typename Derived>
 ExprResult
 TreeTransform<Derived>::TransformPredefinedExpr(PredefinedExpr *E) {
diff --git a/clang/lib/Serialization/ASTReaderStmt.cpp b/clang/lib/Serialization/ASTReaderStmt.cpp
index 93276ce0b12ae..7bf0353797064 100644
--- a/clang/lib/Serialization/ASTReaderStmt.cpp
+++ b/clang/lib/Serialization/ASTReaderStmt.cpp
@@ -593,6 +593,14 @@ void ASTStmtReader::VisitSYCLUniqueStableNameExpr(SYCLUniqueStableNameExpr *E) {
   E->setTypeSourceInfo(Record.readTypeSourceInfo());
 }
 
+void ASTStmtReader::VisitUnresolvedSYCLKernelCallStmt(
+    UnresolvedSYCLKernelCallStmt *S) {
+  VisitStmt(S);
+
+  S->setOriginalStmt(cast<CompoundStmt>(Record.readSubStmt()));
+  S->setKernelLaunchIdExpr(Record.readExpr());
+}
+
 void ASTStmtReader::VisitPredefinedExpr(PredefinedExpr *E) {
   VisitExpr(E);
   bool HasFunctionName = Record.readInt();
@@ -3163,6 +3171,10 @@ Stmt *ASTReader::ReadStmtFromStream(ModuleFile &F) {
       S = SYCLUniqueStableNameExpr::CreateEmpty(Context);
       break;
 
+    case STMT_UNRESOLVED_SYCL_KERNEL_CALL:
+      S = UnresolvedSYCLKernelCallStmt::CreateEmpty(Context);
+      break;
+
     case EXPR_OPENACC_ASTERISK_SIZE:
       S = OpenACCAsteriskSizeExpr::CreateEmpty(Context);
       break;
diff --git a/clang/lib/Serialization/ASTWriterStmt.cpp b/clang/lib/Serialization/ASTWriterStmt.cpp
index 49b8ef178f93e..31d0e9796acf3 100644
--- a/clang/lib/Serialization/ASTWriterStmt.cpp
+++ b/clang/lib/Serialization/ASTWriterStmt.cpp
@@ -670,6 +670,16 @@ void ASTStmtWriter::VisitSYCLUniqueStableNameExpr(SYCLUniqueStableNameExpr *E) {
   Code = serialization::EXPR_SYCL_UNIQUE_STABLE_NAME;
 }
 
+void ASTStmtWriter::VisitUnresolvedSYCLKernelCallStmt(
+    UnresolvedSYCLKernelCallStmt *S) {
+  VisitStmt(S);
+
+  Record.AddStmt(S->getOriginalStmt());
+  Record.AddStmt(S->getKernelLaunchIdExpr());
+
+  Code = serialization::STMT_UNRESOLVED_SYCL_KERNEL_CALL;
+}
+
 void ASTStmtWriter::VisitPredefinedExpr(PredefinedExpr *E) {
   VisitExpr(E);
 
diff --git a/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp b/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp
index c853c00019c10..ecaee9ee0094c 100644
--- a/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp
+++ b/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp
@@ -1824,6 +1824,7 @@ void ExprEngine::Visit(const Stmt *S, ExplodedNode *Pred,
     case Stmt::OMPTargetParallelGenericLoopDirectiveClass:
     case Stmt::CapturedStmtClass:
     case Stmt::SYCLKernelCallStmtClass:
+    case Stmt::UnresolvedSYCLKernelCallStmtClass:
     case Stmt::OpenACCComputeConstructClass:
     case Stmt::OpenACCLoopConstructClass:
     case Stmt::OpenACCCombinedConstructClass:
diff --git a/clang/test/AST/ast-print-sycl-kernel-call.cpp b/clang/test/AST/ast-print-sycl-kernel-call.cpp
index 2243ee024be1a..64c6624b768c9 100644
--- a/clang/test/AST/ast-print-sycl-kernel-call.cpp
+++ b/clang/test/AST/ast-print-sycl-kernel-call.cpp
@@ -3,7 +3,7 @@
 
 struct sycl_kernel_launcher {
   template<typename KernelName, typename... Ts>
-  void sycl_enqueue_kernel_launch(const char *, Ts...) {}
+  void sycl_kernel_launch(const char *, Ts...) {}
 
   template<typename KernelName, typename KernelType>
   void kernel_entry_point(KernelType kernel) {
diff --git a/clang/test/ASTSYCL/ast-dump-sycl-kernel-call-stmt.cpp b/clang/test/ASTSYCL/ast-dump-sycl-kernel-call-stmt.cpp
index cdff3fc8821a1..8e9a6fe84961a 100644
--- a/clang/test/ASTSYCL/ast-dump-sycl-kernel-call-stmt.cpp
+++ b/clang/test/ASTSYCL/ast-dump-sycl-kernel-call-stmt.cpp
@@ -34,6 +34,8 @@ template<int> struct K {
   void operator()(Ts...) const {}
 };
 
+template <typename KernelName, typename... Ts>
+void sycl_kernel_launch(const char *, Ts...) {}
 
 [[clang::sycl_kernel_entry_point(KN<1>)]]
 void skep1() {
@@ -42,12 +44,11 @@ void skep1() {
 // CHECK-NEXT: | |-SYCLKernelCallStmt {{.*}}
 // CHECK-NEXT: | | |-CompoundStmt {{.*}}
 // CHECK-NEXT: | | |-CompoundStmt {{.*}}
-// CHECK-NEXT: | | | |-DeclStmt {{.*}}
-// CHECK-NEXT: | | | | `-VarDecl {{.*}} kernel_name 'const char *' extern
-// CHECK-NEXT: | | | `-BinaryOperator {{.*}} 'const char *' lvalue '='
-// CHECK-NEXT: | | |   |-DeclRefExpr {{.*}} 'const char *' lvalue Var {{.*}} 'kernel_name' 'const char *'
+// CHECK-NEXT: | | | `-CallExpr {{.*}}
+// CHECK-NEXT: | | |   |-ImplicitCastExpr {{.*}} 'void (*)(const char *)' <FunctionToPointerDecay>
+// CHECK-NEXT: | | |   | `-DeclRefExpr {{.*}} 'void (const char *)' lvalue Function {{.*}} 'sycl_kernel_launch' {{.*}}
 // CHECK-NEXT: | | |   `-ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
-// CHECK-NEXT: | | |     `-StringLiteral {{.*}} 'const char[14]' lvalue "_ZTS2KNILi1EE"
+// CHECK-NEXT: | | |       `-StringLiteral {{.*}} 'const char[14]' lvalue "_ZTS2KNILi1EE"
 // CHECK-NEXT: | | `-OutlinedFunctionDecl {{.*}}
 // CHECK-NEXT: | |   `-CompoundStmt {{.*}}
 // CHECK-NEXT: | `-SYCLKernelEntryPointAttr {{.*}} KN<1>
@@ -64,9 +65,10 @@ void skep2<KN<2>>(K<2>);
 // CHECK-NEXT: | |-TemplateTypeParmDecl {{.*}} KT
 // CHECK-NEXT: | |-FunctionDecl {{.*}} skep2 'void (KT)'
 // CHECK-NEXT: | | |-ParmVarDecl {{.*}} k 'KT'
-// CHECK-NEXT: | | |-CompoundStmt {{.*}}
-// CHECK-NEXT: | | | `-CallExpr {{.*}} '<dependent type>'
-// CHECK-NEXT: | | |   `-DeclRefExpr {{.*}} 'KT' lvalue ParmVar {{.*}} 'k' 'KT'
+// CHECK-NEXT: | | |-UnresolvedSYCLKernelCallStmt {{.*}}
+// CHECK-NEXT: | | | `-CompoundStmt {{.*}}
+// CHECK-NEXT: | | |   `-CallExpr {{.*}} '<dependent type>'
+// CHECK-NEXT: | | |     `-DeclRefExpr {{.*}} 'KT' lvalue ParmVar {{.*}} 'k' 'KT'
 // CHECK-NEXT: | | `-SYCLKernelEntryPointAttr {{.*}} KNT
 
 // CHECK-NEXT: | `-FunctionDecl {{.*}} skep2 'void (K<2>)' explicit_instantiation_definition
@@ -85,12 +87,14 @@ void skep2<KN<2>>(K<2>);
 // CHECK-NEXT: |   | |   `-ImplicitCastExpr {{.*}} 'const K<2>' lvalue <NoOp>
 // CHECK-NEXT: |   | |     `-DeclRefExpr {{.*}} 'K<2>' lvalue ParmVar {{.*}} 'k' 'K<2>'
 // CHECK-NEXT: |   | |-CompoundStmt {{.*}}
-// CHECK-NEXT: |   | | |-DeclStmt {{.*}}
-// CHECK-NEXT: |   | | | `-VarDecl {{.*}} kernel_name 'const char *' extern
-// CHECK-NEXT: |   | | `-BinaryOperator {{.*}} 'const char *' lvalue '='
-// CHECK-NEXT: |   | |   |-DeclRefExpr {{.*}} 'const char *' lvalue Var {{.*}} 'kernel_name' 'const char *'
-// CHECK-NEXT: |   | |   `-ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
-// CHECK-NEXT: |   | |     `-StringLiteral {{.*}} 'const char[14]' lvalue "_ZTS2KNILi2EE"
+// CHECK-NEXT: |   | | `-CallExpr {{.*}} 'void'
+// CHECK-NEXT: |   | |   |-ImplicitCastExpr {{.*}} <FunctionToPointerDecay>
+// CHECK-NEXT: |   | |   | `-DeclRefExpr {{.*}} 'void (const char *, K<2>)' lvalue Function {{.*}} 'sycl_kernel_launch' {{.*}}
+// CHECK-NEXT: |   | |   |-ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+// CHECK-NEXT: |   | |   | `-StringLiteral {{.*}} 'const char[14]' lvalue "_ZTS2KNILi2EE"
+// CHECK-NEXT: |   | |   `-CXXConstructExpr {{.*}} 'K<2>' 'void (const K<2> &) noexcept'
+// CHECK-NEXT: |   | |     `-ImplicitCastExpr {{.*}} <NoOp>
+// CHECK-NEXT: |   | |       `-DeclRefExpr {{.*}} 'K<2>' lvalue ParmVar {{.*}} 'k' 'K<2>'
 // CHECK-NEXT: |   | `-OutlinedFunctionDecl {{.*}}
 // CHECK-NEXT: |   |   |-ImplicitParamDecl {{.*}} implicit used k 'K<2>'
 // CHECK-NEXT: |   |   `-CompoundStmt {{.*}}
@@ -116,9 +120,10 @@ void skep3<KN<3>>(K<3> k) {
 // CHECK-NEXT: | |-TemplateTypeParmDecl {{.*}} KT
 // CHECK-NEXT: | |-FunctionDecl {{.*}} skep3 'void (KT)'
 // CHECK-NEXT: | | |-ParmVarDecl {{.*}} k 'KT'
-// CHECK-NEXT: | | |-CompoundStmt {{.*}}
-// CHECK-NEXT: | | | `-CallExpr {{.*}} '<dependent type>'
-// CHECK-NEXT: | | |   `-DeclRefExpr {{.*}} 'KT' lvalue ParmVar {{.*}} 'k' 'KT'
+// CHECK-NEXT: | | |-UnresolvedSYCLKernelCallStmt {{.*}}
+// CHECK-NEXT: | | | `-CompoundStmt {{.*}}
+// CHECK-NEXT: | | |   `-CallExpr {{.*}} '<dependent type>'
+// CHECK-NEXT: | | |     `-DeclRefExpr {{.*}} 'KT' lvalue ParmVar {{.*}} 'k' 'KT'
 // CHECK-NEXT: | | `-SYCLKernelEntryPointAttr {{.*}} KNT
 
 // CHECK-NEXT: | `-Function {{.*}} 'skep3' 'void (K<3>)'
@@ -138,12 +143,14 @@ void skep3<KN<3>>(K<3> k) {
 // CHECK-NEXT: | | |   `-ImplicitCastExpr {{.*}} 'const K<3>' lvalue <NoOp>
 // CHECK-NEXT: | | |     `-DeclRefExpr {{.*}} 'K<3>' lvalue ParmVar {{.*}} 'k' 'K<3>'
 // CHECK-NEXT: | | |-CompoundStmt {{.*}}
-// CHECK-NEXT: | | | |-DeclStmt {{.*}}
-// CHECK-NEXT: | | | | `-VarDecl {{.*}} kernel_name 'const char *' extern
-// CHECK-NEXT: | | | `-BinaryOperator {{.*}} 'const char *' lvalue '='
-// CHECK-NEXT: | | |   |-DeclRefExpr {{.*}} 'const char *' lvalue Var {{.*}} 'kernel_name' 'const char *'
-// CHECK-NEXT: | | |   `-ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
-// CHECK-NEXT: | | |     `-StringLiteral {{.*}} 'const char[14]' lvalue "_ZTS2KNILi3EE"
+// CHECK-NEXT: | | | `-CallExpr {{.*}}
+// CHECK-NEXT: | | |   |-ImplicitCastExpr {{.*}} 'void (*)(const char *, K<3>)' <FunctionToPointerDecay>
+// CHECK-NEXT: | | |   | `-DeclRefExpr {{.*}} 'void (const char *, K<3>)' lvalue Function {{.*}} 'sycl_kernel_launch' 'void (const char *, K<3>)' {{.*}}
+// CHECK-NEXT: | | |   |-ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+// CHECK-NEXT: | | |   | `-StringLiteral {{.*}} 'const char[14]' lvalue "_ZTS2KNILi3EE"
+// CHECK-NEXT: | | |   `-CXXConstructExpr {{.*}} 'K<3>' 'void (const K<3> &) noexcept'
+// CHECK-NEXT: | | |     `-ImplicitCastExpr {{.*}} 'const K<3>' lvalue <NoOp>
+// CHECK-NEXT: | | |       `-DeclRefExpr {{.*}} 'K<3>' lvalue ParmVar {{.*}} 'k' 'K<3>'
 // CHECK-NEXT: | | `-OutlinedFunctionDecl {{.*}}
 // CHECK-NEXT: | |   |-ImplicitParamDecl {{.*}} implicit used k 'K<3>'
 // CHECK-NEXT: | |   `-CompoundStmt {{.*}}
@@ -174,12 +181,18 @@ void skep4(K<4> k, int p1, int p2) {
 // CHECK-NEXT: | | |   `-ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
 // CHECK-NEXT: | | |     `-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'p2' 'int'
 // CHECK-NEXT: | | |-CompoundStmt {{.*}}
-// CHECK-NEXT: | | | |-DeclStmt {{.*}}
-// CHECK-NEXT: | | | | `-VarDecl {{.*}} kernel_name 'const char *' extern
-// CHECK-NEXT: | | | `-BinaryOperator {{.*}} 'const char *' lvalue '='
-// CHECK-NEXT: | | |   |-DeclRefExpr {{.*}} 'const char *' lvalue Var {{.*}} 'kernel_name' 'const char *'
-// CHECK-NEXT: | | |   `-ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
-// CHECK-NEXT: | | |     `-StringLiteral {{.*}} 'const char[14]' lvalue "_ZTS2KNILi4EE"
+// CHECK-NEXT: | | | `-CallExpr {{.*}} 'void'
+// CHECK-NEXT: | | |   |-ImplicitCastExpr {{.*}} 'void (*)(const char *, K<4>, int, int)' <FunctionToPointerDecay>
+// CHECK-NEXT: | | |   | `-DeclRefExpr {{.*}} 'void (const char *, K<4>, int, int)' lvalue Function {{.*}} 'sycl_kernel_launch' 'void (const char *, K<4>, int, int)' {{.*}}
+// CHECK-NEXT: | | |   |-ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+// CHECK-NEXT: | | |   | `-StringLiteral {{.*}} 'const char[14]' lvalue "_ZTS2KNILi4EE"
+// CHECK-NEXT: | | |   |-CXXConstructExpr {{.*}} 'K<4>' 'void (const K<4> &) noexcept'
+// CHECK-NEXT: | | |   | `-ImplicitCastExpr {{.*}} 'const K<4>' lvalue <NoOp>
+// CHECK-NEXT: | | |   |   `-DeclRefExpr {{.*}} 'K<4>' lvalue ParmVar {{.*}} 'k' 'K<4>'
+// CHECK-NEXT: | | |   |-ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
+// CHECK-NEXT: | | |   | `-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'p1' 'int'
+// CHECK-NEXT: | | |   `-ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
+// CHECK-NEXT: | | |     `-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'p2' 'int'
 // CHECK-NEXT: | | `-OutlinedFunctionDecl {{.*}}
 // CHECK-NEXT: | |   |-ImplicitParamDecl {{.*}} implicit used k 'K<4>'
 // CHECK-NEXT: | |   |-ImplicitParamDecl {{.*}} implicit used p1 'int'
@@ -211,12 +224,22 @@ void skep5(int unused1, K<5> k, int unused2, int p, int unused3) {
 // CHECK-NEXT: | |-SYCLKernelCallStmt {{.*}}
 // CHECK-NEXT: | | |-CompoundStmt {{.*}}
 // CHECK:      | | |-CompoundStmt {{.*}}
-// CHECK-NEXT: | | | |-DeclStmt {{.*}}
-// CHECK-NEXT: | | | | `-VarDecl {{.*}} kernel_name 'const char *' extern
-// CHECK-NEXT: | | | `-BinaryOperator {{.*}} 'const char *' lvalue '='
-// CHECK-NEXT: | | |   |-DeclRefExpr {{.*}} 'const char *' lvalue Var {{.*}} 'kernel_name' 'const char *'
-// CHECK-NEXT: | | |   `-ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
-// CHECK-NEXT: | | |     `-StringLiteral {{.*}} 'const char[14]' lvalue "_ZTS2KNILi5EE"
+// CHECK-NEXT: | | | `-CallExpr {{.*}} 'void'
+// CHECK-NEXT: | | |   |-ImplicitCastExpr {{.*}} 'void (*)(const char *, int, K<5>, int, int, int)' <FunctionToPointerDecay>
+// CHECK-NEXT: | | |   | `-DeclRefExpr {{.*}} 'void (const char *, int, K<5>, int, int, int)' lvalue Function {{.*}} 'sycl_kernel_launch' 'void (const char *, int, K<5>, int, int, int)' {{.*}}
+// CHECK-NEXT: | | |   |-ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+// CHECK-NEXT: | | |   | `-StringLiteral {{.*}} 'const char[14]' lvalue "_ZTS2KNILi5EE"
+// CHECK-NEXT: | | |   |-ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
+// CHECK-NEXT: | | |   | `-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'unused1' 'int'
+// CHECK-NEXT: | | |   |-CXXConstructExpr {{.*}} 'K<5>' 'void (const K<5> &) noexcept'
+// CHECK-NEXT: | | |   | `-ImplicitCastExpr {{.*}} 'const K<5>' lvalue <NoOp>
+// CHECK-NEXT: | | |   |   `-DeclRefExpr {{.*}} 'K<5>' lvalue ParmVar {{.*}} 'k' 'K<5>'
+// CHECK-NEXT: | | |   |-ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
+// CHECK-NEXT: | | |   | `-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'unused2' 'int'
+// CHECK-NEXT: | | |   |-ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
+// CHECK-NEXT: | | |   | `-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'p' 'int'
+// CHECK-NEXT: | | |   `-ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
+// CHECK-NEXT: | | |     `-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'unused3' 'int'
 // CHECK-NEXT: | | `-OutlinedFunctionDecl {{.*}}
 // CHECK-NEXT: | |   |-ImplicitParamDecl {{.*}} implicit unused1 'int'
 // CHECK-NEXT: | |   |-ImplicitParamDecl {{.*}} implicit used k 'K<5>'
@@ -263,12 +286,13 @@ void skep6(const S6 &k) {
 // CHECK-NEXT: | | |   | `-DeclRefExpr {{.*}} 'void () const' lvalue CXXMethod {{.*}} 'operator()' 'void () const'
 // CHECK-NEXT: | | |   `-DeclRefExpr {{.*}} 'const S6' lvalue ParmVar {{.*}} 'k' 'const S6 &'
 // CHECK-NEXT: | | |-CompoundStmt {{.*}}
-// CHECK-NEXT: | | | |-DeclStmt {{.*}}
-// CHECK-NEXT: | | | | `-VarDecl {{.*}} kernel_name 'const char *' extern
-// CHECK-NEXT: | | | `-BinaryOperator {{.*}} 'const char *' lvalue '='
-// CHECK-NEXT: | | |   |-DeclRefExpr {{.*}} 'const char *' lvalue Var {{.*}} 'kernel_name' 'const char *'
-// CHECK-NEXT: | | |   `-ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
-// CHECK-NEXT: | | |     `-StringLiteral {{.*}} 'const char[14]' lvalue "_ZTS2KNILi6EE"
+// CHECK-NEXT: | | | `-CallExpr {{.*}} 'void'
+// CHECK-NEXT: | | |   |-ImplicitCastExpr {{.*}} 'void (*)(const char *, S6)' <FunctionToPointerDecay>
+// CHECK-NEXT: | | |   | `-DeclRefExpr {{.*}} 'void (const char *, S6)' lvalue Function {{.*}} 'sycl_kernel_launch' 'void (const char *, S6)' {{.*}}
+// CHECK-NEXT: | | |   |-ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+// CHECK-NEXT: | | |   | `-StringLiteral {{.*}} 'const char[14]' lvalue "_ZTS2KNILi6EE"
+// CHECK-NEXT: | | |   `-CXXConstructExpr {{.*}} 'S6' 'void (const S6 &) noexcept'
+// CHECK-NEXT: | | |     `-DeclRefExpr {{.*}} 'const S6' lvalue ParmVar {{.*}} 'k' 'const S6 &'
 // CHECK-NEXT: | | `-OutlinedFunctionDecl {{.*}}
 // CHECK-NEXT: | |   |-ImplicitParamDecl {{.*}} implicit used k 'const S6 &'
 // CHECK-NEXT: | |   `-CompoundStmt {{.*}}
@@ -303,12 +327,14 @@ void skep7(S7 k) {
 // CHECK-NEXT: | | |   `-ImplicitCastExpr {{.*}} 'const S7' lvalue <NoOp>
 // CHECK-NEXT: | | |     `-DeclRefExpr {{.*}} 'S7' lvalue ParmVar {{.*}} 'k' 'S7'
 // CHECK-NEXT: | | |-CompoundStmt {{.*}}
-// CHECK-NEXT: | | | |-DeclStmt {{.*}}
-// CHECK-NEXT: | | | | `-VarDecl {{.*}} kernel_name 'const char *' extern
-// CHECK-NEXT: | | | `-BinaryOperator {{.*}} 'const char *' lvalue '='
-// CHECK-NEXT: | | |   |-DeclRefExpr {{.*}} 'const char *' lvalue Var {{.*}} 'kernel_name' 'const char *'
-// CHECK-NEXT: | | |   `-ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
-// CHECK-NEXT: | | |     `-StringLiteral {{.*}} 'const char[14]' lvalue "_ZTS2KNILi7EE"
+// CHECK-NEXT: | | | `-CallExpr {{.*}} 'void'
+// CHECK-NEXT: | | |   |-ImplicitCastExpr {{.*}} 'void (*)(const char *, S7)' <FunctionToPointerDecay>
+// CHECK-NEXT: | | |   | `-DeclRefExpr {{.*}} 'void (const char *, S7)' lvalue Function {{.*}} 'sycl_kernel_launch' 'void (const char *, S7)' {{.*}}
+// CHECK-NEXT: | | |   |-ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+// CHECK-NEXT: | | |   | `-StringLiteral {{.*}} 'const char[14]' lvalue "_ZTS2KNILi7EE"
+// CHECK-NEXT: | | |   `-CXXConstructExpr {{.*}} 'S7' 'void (const S7 &) noexcept'
+// CHECK-NEXT: | | |     `-ImplicitCastExpr {{.*}} 'const S7' lvalue <NoOp>
+// CHECK-NEXT: | | |       `-DeclRefExpr {{.*}} 'S7' lvalue ParmVar {{.*}} 'k' 'S7'
 // CHECK-NEXT: | | `-OutlinedFunctionDecl {{.*}}
 // CHECK-NEXT: | |   |-ImplicitParamDecl {{.*}} implicit used k 'S7'
 // CHECK-NEXT: | |   `-CompoundStmt {{.*}}
@@ -335,11 +361,96 @@ void skep8(S8 k) {
 // CHECK-NEXT: | |-SYCLKernelCallStmt {{.*}}
 // CHECK-NEXT: | | |-CompoundStmt {{.*}}
 // CHECK:      | | |-CompoundStmt {{.*}}
-// CHECK:      | | |   `-ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
-// CHECK-NEXT: | | |     `-StringLiteral {{.*}} 'const char[12]' lvalue "_ZTS6\316\264\317\204\317\207"
+// CHECK-NEXT: | | | `-CallExpr {{.*}} 'void'
+// CHECK-NEXT: | | |   |-ImplicitCastExpr {{.*}} 'void (*)(const char *, S8)' <FunctionToPointerDecay>
+// CHECK-NEXT: | | |   | `-DeclRefExpr {{.*}} 'void (const char *, S8)' lvalue Function {{.*}} 'sycl_kernel_launch' 'void (const char *, S8)' {{.*}}
+// CHECK-NEXT: | | |   |-ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+// CHECK-NEXT: | | |   | `-StringLiteral {{.*}} 'const char[12]' lvalue "_ZTS6\316\264\317\204\317\207"
+// CHECK-NEXT: | | |   `-CXXConstructExpr {{.*}} 'S8' 'void (const S8 &) noexcept'
+// CHECK-NEXT: | | |     `-ImplicitCastExpr {{.*}} 'const S8' lvalue <NoOp>
+// CHECK-NEXT: | | |       `-DeclRefExpr {{.*}} 'S8' lvalue ParmVar {{.*}} 'k' 'S8'
 // CHECK:      | | `-OutlinedFunctionDecl {{.*}}
 // CHECK:      | `-SYCLKernelEntryPointAttr {{.*}}
 
+class Handler {
+template <typename KernelName, typename... Ts>
+void sycl_kernel_launch(const char *, Ts...) {}
+public:
+template<typename KNT, typename KT>
+[[clang::sycl_kernel_entry_point(KNT)]]
+void skep9(KT k, int a, int b) {
+  k(a, b);
+}
+};
+void foo() {
+  Handler H;
+  H.skep9<KN<9>>([=](int a, int b){return a+b;}, 1, 2);
+}
+
+// CHECK: | |-FunctionTemplateDecl {{.*}} skep9
+// CHECK-NEXT: | | |-TemplateTypeParmDecl {{.*}} referenced typename depth 0 index 0 KNT
+// CHECK-NEXT: | | |-TemplateTypeParmDecl {{.*}} referenced typename depth 0 index 1 KT
+// CHECK-NEXT: | | |-CXXMethodDecl {{.*}} skep9 'void (KT, int, int)' implicit-inline
+// CHECK-NEXT: | | | |-ParmVarDecl {{.*}} referenced k 'KT'
+// CHECK-NEXT: | | | |-ParmVarDecl {{.*}} referenced a 'int'
+// CHECK-NEXT: | | | |-ParmVarDecl {{.*}} referenced b 'int'
+// CHECK-NEXT: | | | |-UnresolvedSYCLKernelCallStmt {{.*}}
+// CHECK-NEXT: | | | | `-CompoundStmt {{.*}}
+// CHECK-NEXT: | | | |   `-CallExpr {{.*}} '<dependent type>'
+// CHECK-NEXT: | | | |     |-DeclRefExpr {{.*}} 'KT' lvalue ParmVar {{.*}} 'k' 'KT'
+// CHECK-NEXT: | | | |     |-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'a' 'int'
+// CHECK-NEXT: | | | |     `-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'b' 'int'
+// CHECK-NEXT: | | | `-SYCLKernelEntryPointAttr {{.*}} KNT
+// CHECK-NEXT: | | `-CXXMethodDecl {{.*}} used skep9 {{.*}} implicit_instantiation implicit-inline
+// CHECK-NEXT: | |   |-TemplateArgument type 'KN<9>'
+// CHECK-NEXT: | |   | `-RecordType {{.*}} 'KN<9>' canonical
+// CHECK-NEXT: | |   |   `-ClassTemplateSpecialization {{.*}}'KN'
+// CHECK-NEXT: | |   |-TemplateArgument type {{.*}}
+// CHECK-NEXT: | |   | `-RecordType {{.*}}
+// CHECK-NEXT: | |   |   `-CXXRecord {{.*}}
+// CHECK-NEXT: | |   |-ParmVarDecl {{.*}} used k {{.*}}
+// CHECK-NEXT: | |   |-ParmVarDecl {{.*}} used a 'int'
+// CHECK-NEXT: | |   |-ParmVarDecl {{.*}} used b 'int'
+// CHECK-NEXT: | |   |-SYCLKernelCallStmt {{.*}}
+// CHECK-NEXT: | |   | |-CompoundStmt {{.*}}
+// CHECK-NEXT: | |   | | `-CXXOperatorCallExpr {{.*}} 'int' '()'
+// CHECK-NEXT: | |   | |   |-ImplicitCastExpr {{.*}} 'int (*)(int, int) const' <FunctionToPointerDecay>
+// CHECK-NEXT: | |   | |   | `-DeclRefExpr {{.*}} 'int (int, int) const' lvalue CXXMethod {{.*}} 'operator()' 'int (int, int) const'
+// CHECK-NEXT: | |   | |   |-ImplicitCastExpr {{.*}} lvalue <NoOp>
+// CHECK-NEXT: | |   | |   | `-DeclRefExpr {{.*}} lvalue ParmVar {{.*}} 'k' {{.*}}
+// CHECK-NEXT: | |   | |   |-ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
+// CHECK-NEXT: | |   | |   | `-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'a' 'int'
+// CHECK-NEXT: | |   | |   `-ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
+// CHECK-NEXT: | |   | |     `-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'b' 'int'
+// CHECK-NEXT: | |   | |-CompoundStmt {{.*}}
+// CHECK-NEXT: | |   | | `-CXXMemberCallExpr {{.*}} 'void'
+// CHECK-NEXT: | |   | |   |-MemberExpr {{.*}} '<bound member function type>' ->sycl_kernel_launch {{.*}}
+// CHECK-NEXT: | |   | |   | `-CXXThisExpr {{.*}} 'Handler *' implicit this
+// CHECK-NEXT: | |   | |   |-ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+// CHECK-NEXT: | |   | |   | `-StringLiteral {{.*}} 'const char[14]' lvalue "_ZTS2KNILi9EE"
+// CHECK-NEXT: | |   | |   |-CXXConstructExpr {{.*}}
+// CHECK-NEXT: | |   | |   | `-ImplicitCastExpr {{.*}} lvalue <NoOp>
+// CHECK-NEXT: | |   | |   |   `-DeclRefExpr {{.*}} lvalue ParmVar {{.*}} 'k' {{.*}}
+// CHECK-NEXT: | |   | |   |-ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
+// CHECK-NEXT: | |   | |   | `-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'a' 'int'
+// CHECK-NEXT: | |   | |   `-ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
+// CHECK-NEXT: | |   | |     `-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'b' 'int'
+// CHECK-NEXT: | |   | `-OutlinedFunctionDecl {{.*}}
+// CHECK-NEXT: | |   |   |-ImplicitParamDecl {{.*}} implicit used k {{.*}}
+// CHECK-NEXT: | |   |   |-ImplicitParamDecl {{.*}} implicit used a 'int'
+// CHECK-NEXT: | |   |   |-ImplicitParamDecl {{.*}} implicit used b 'int'
+// CHECK-NEXT: | |   |   `-CompoundStmt {{.*}}
+// CHECK-NEXT: | |   |     `-CXXOperatorCallExpr {{.*}} 'int' '()'
+// CHECK-NEXT: | |   |       |-ImplicitCastExpr {{.*}} 'int (*)(int, int) const' <FunctionToPointerDecay>
+// CHECK-NEXT: | |   |       | `-DeclRefExpr {{.*}} 'int (int, int) const' lvalue CXXMethod {{.*}} 'operator()' 'int (int, int) const'
+// CHECK-NEXT: | |   |       |-ImplicitCastExpr {{.*}} lvalue <NoOp>
+// CHECK-NEXT: | |   |       | `-DeclRefExpr {{.*}} lvalue ImplicitParam {{.*}} 'k' {{.*}}
+// CHECK-NEXT: | |   |       |-ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
+// CHECK-NEXT: | |   |       | `-DeclRefExpr {{.*}} 'int' lvalue ImplicitParam {{.*}} 'a' 'int'
+// CHECK-NEXT: | |   |       `-ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
+// CHECK-NEXT: | |   |         `-DeclRefExpr {{.*}} 'int' lvalue ImplicitParam {{.*}} 'b' 'int'
+// CHECK-NEXT: | |   `-SYCLKernelEntryPointAttr {{.*}} struct KN<9>
+
 
 void the_end() {}
 // CHECK:      `-FunctionDecl {{.*}} the_end 'void ()'
diff --git a/clang/test/ASTSYCL/ast-dump-sycl-kernel-entry-point.cpp b/clang/test/ASTSYCL/ast-dump-sycl-kernel-entry-point.cpp
index 1a82bdc1f5698..bfa3f764ceb0f 100644
--- a/clang/test/ASTSYCL/ast-dump-sycl-kernel-entry-point.cpp
+++ b/clang/test/ASTSYCL/ast-dump-sycl-kernel-entry-point.cpp
@@ -28,6 +28,9 @@
 // A unique kernel name type is required for each declared kernel entry point.
 template<int, int=0> struct KN;
 
+template <typename KernelName, typename... Tys>
+void sycl_kernel_launch(const char *, Tys &&...Args) {}
+
 [[clang::sycl_kernel_entry_point(KN<1>)]]
 void skep1() {
 }
diff --git a/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp b/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp
index e88e4b7cf3149..48c13240b1ffc 100644
--- a/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp
+++ b/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp
@@ -25,8 +25,13 @@
 
 // Test the generation of SYCL kernel caller functions. These functions are
 // generated from functions declared with the sycl_kernel_entry_point attribute
-// and emited during device compilation. They are not emitted during device
-// compilation.
+// and emited during device compilation.
+// Test the generation of SYCL kernel launch statements during host compilation.
+// These statements are calls to sycl_enqueus_kernel_launch functions or class
+// members in case skep-attributed functions are also members of the same class.
+
+template <typename KernelName, typename KernelObj>
+void sycl_kernel_launch(const char *, KernelObj) {}
 
 struct single_purpose_kernel_name;
 struct single_purpose_kernel {
@@ -47,6 +52,17 @@ void kernel_single_task(KernelType kernelFunc) {
 // Exercise code gen with kernel name types named with esoteric characters.
 struct \u03b4\u03c4\u03c7; // Delta Tau Chi (δτχ)
 
+class Handler {
+template <typename KernelName, typename... Ts>
+void sycl_kernel_launch(const char *, Ts...) {}
+public:
+template<typename KNT, typename KT>
+[[clang::sycl_kernel_entry_point(KNT)]]
+void skep(KT k, int a, int b) {
+  k(a, b);
+}
+};
+
 int main() {
   single_purpose_kernel obj;
   single_purpose_kernel_task(obj);
@@ -54,6 +70,8 @@ int main() {
   auto lambda = [=](auto) { (void) capture; };
   kernel_single_task<decltype(lambda)>(lambda);
   kernel_single_task<\u03b4\u03c4\u03c7>([](int){});
+  Handler H;
+  H.skep<class notaverygoodkernelname>([=](int a, int b){return a+b;}, 1, 2);
 }
 
 // Verify that SYCL kernel caller functions are not emitted during host
@@ -79,50 +97,105 @@ int main() {
 // CHECK-HOST-LINUX:      define dso_local void @_Z26single_purpose_kernel_task21single_purpose_kernel() #{{[0-9]+}} {
 // CHECK-HOST-LINUX-NEXT: entry:
 // CHECK-HOST-LINUX-NEXT:   %kernelFunc = alloca %struct.single_purpose_kernel, align 1
-// CHECK-HOST-LINUX-NEXT:   store ptr @.str, ptr @kernel_name, align 8
+// CHECK-HOST-LINUX-NEXT:   %agg.tmp = alloca %struct.single_purpose_kernel, align 1
+// CHECK-HOST-LINUX-NEXT:   call void @_Z18sycl_kernel_launchI26single_purpose_kernel_name21single_purpose_kernelEvPKcT0_(ptr noundef @.str)
 // CHECK-HOST-LINUX-NEXT:   ret void
 // CHECK-HOST-LINUX-NEXT: }
 //
 // CHECK-HOST-LINUX:      define internal void @_Z18kernel_single_taskIZ4mainEUlT_E_S1_EvT0_(i32 %kernelFunc.coerce) #{{[0-9]+}} {
 // CHECK-HOST-LINUX-NEXT: entry:
 // CHECK-HOST-LINUX-NEXT:   %kernelFunc = alloca %class.anon, align 4
+// CHECK-HOST-LINUX-NEXT:   %agg.tmp = alloca %class.anon, align 4
 // CHECK-HOST-LINUX-NEXT:   %coerce.dive = getelementptr inbounds nuw %class.anon, ptr %kernelFunc, i32 0, i32 0
 // CHECK-HOST-LINUX-NEXT:   store i32 %kernelFunc.coerce, ptr %coerce.dive, align 4
-// CHECK-HOST-LINUX-NEXT:   store ptr @.str.1, ptr @kernel_name, align 8
+// CHECK-HOST-LINUX-NEXT:   call void @llvm.memcpy.p0.p0.i64(ptr align 4 %agg.tmp, ptr align 4 %kernelFunc, i64 4, i1 false)
+// CHECK-HOST-LINUX-NEXT:   %coerce.dive1 = getelementptr inbounds nuw %class.anon, ptr %agg.tmp, i32 0, i32 0
+// CHECK-HOST-LINUX-NEXT:   %0 = load i32, ptr %coerce.dive1, align 4
+// CHECK-HOST-LINUX-NEXT:   call void @_Z18sycl_kernel_launchIZ4mainEUlT_E_S1_EvPKcT0_(ptr noundef @.str.1, i32 %0)
 // CHECK-HOST-LINUX-NEXT:   ret void
 // CHECK-HOST-LINUX-NEXT: }
 //
 // CHECK-HOST-LINUX:      define internal void @"_Z18kernel_single_taskI6\CE\B4\CF\84\CF\87Z4mainEUliE_EvT0_"() #{{[0-9]+}} {
 // CHECK-HOST-LINUX-NEXT: entry:
 // CHECK-HOST-LINUX-NEXT:   %kernelFunc = alloca %class.anon.0, align 1
-// CHECK-HOST-LINUX-NEXT:   store ptr @.str.2, ptr @kernel_name, align 8
+// CHECK-HOST-LINUX-NEXT:   %agg.tmp = alloca %class.anon.0, align 1
+// CHECK-HOST-LINUX-NEXT:   call void @"_Z18sycl_kernel_launchI6\CE\B4\CF\84\CF\87Z4mainEUliE_EvPKcT0_"(ptr noundef @.str.2)
 // CHECK-HOST-LINUX-NEXT:   ret void
 // CHECK-HOST-LINUX-NEXT: }
-//
+
+// CHECK-HOST-LINUX: define internal void @_ZN7Handler4skepIZ4mainE22notaverygoodkernelnameZ4mainEUliiE_EEvT0_ii(ptr noundef nonnull align 1 dereferenceable(1) %this, i32 noundef %a, i32 noundef %b) #0 align 2 {
+// CHECK-HOST-LINUX-NEXT: entry:
+// CHECK-HOST-LINUX-NEXT:   %k = alloca %class.anon.1, align 1
+// CHECK-HOST-LINUX-NEXT:   %this.addr = alloca ptr, align 8
+// CHECK-HOST-LINUX-NEXT:   %a.addr = alloca i32, align 4
+// CHECK-HOST-LINUX-NEXT:   %b.addr = alloca i32, align 4
+// CHECK-HOST-LINUX-NEXT:   %agg.tmp = alloca %class.anon.1, align 1
+// CHECK-HOST-LINUX-NEXT:   store ptr %this, ptr %this.addr, align 8
+// CHECK-HOST-LINUX-NEXT:   store i32 %a, ptr %a.addr, align 4
+// CHECK-HOST-LINUX-NEXT:   store i32 %b, ptr %b.addr, align 4
+// CHECK-HOST-LINUX-NEXT:   %this1 = load ptr, ptr %this.addr, align 8
+// CHECK-HOST-LINUX-NEXT:   %0 = load i32, ptr %a.addr, align 4
+// CHECK-HOST-LINUX-NEXT:   %1 = load i32, ptr %b.addr, align 4
+// CHECK-HOST-LINUX-NEXT:   call void @_ZN7Handler18sycl_kernel_launchIZ4mainE22notaverygoodkernelnameJZ4mainEUliiE_iiEEEvPKcDpT0_(ptr noundef nonnull align 1 dereferenceable(1) %this1, ptr noundef @.str.3, i32 noundef %0, i32 noundef %1)
+// CHECK-HOST-LINUX-NEXT:   ret void
+// CHECK-HOST-LINUX-NEXT: }
+
 // CHECK-HOST-WINDOWS:      define dso_local void @"?single_purpose_kernel_task@@YAXUsingle_purpose_kernel@@@Z"(i8 %kernelFunc.coerce) #{{[0-9]+}} {
 // CHECK-HOST-WINDOWS-NEXT: entry:
 // CHECK-HOST-WINDOWS-NEXT:   %kernelFunc = alloca %struct.single_purpose_kernel, align 1
+// CHECK-HOST-WINDOWS-NEXT:   %agg.tmp = alloca %struct.single_purpose_kernel, align 1
 // CHECK-HOST-WINDOWS-NEXT:   %coerce.dive = getelementptr inbounds nuw %struct.single_purpose_kernel, ptr %kernelFunc, i32 0, i32 0
 // CHECK-HOST-WINDOWS-NEXT:   store i8 %kernelFunc.coerce, ptr %coerce.dive, align 1
-// CHECK-HOST-WINDOWS-NEXT:   store ptr @"??_C@_0CB@KFIJOMLB@_ZTS26single_purpose_kernel_name@", ptr @"?kernel_name@?0??single_purpose_kernel_task@@YAXUsingle_purpose_kernel@@@Z@3PEBDEB", align 8
+// CHECK-HOST-WINDOWS-NEXT:   %coerce.dive1 = getelementptr inbounds nuw %struct.single_purpose_kernel, ptr %agg.tmp, i32 0, i32 0
+// CHECK-HOST-WINDOWS-NEXT:   %0 = load i8, ptr %coerce.dive1, align 1
+// CHECK-HOST-WINDOWS-NEXT:   call void @"??$sycl_kernel_launch@Usingle_purpose_kernel_name@@Usingle_purpose_kernel@@@@YAXPEBDUsingle_purpose_kernel@@@Z"(ptr noundef @"??_C@_0CB@KFIJOMLB@_ZTS26single_purpose_kernel_name@", i8 %0)
 // CHECK-HOST-WINDOWS-NEXT:   ret void
 // CHECK-HOST-WINDOWS-NEXT: }
 //
 // CHECK-HOST-WINDOWS:      define internal void @"??$kernel_single_task@V<lambda_1>@?0??main@@9@V1?0??2@9@@@YAXV<lambda_1>@?0??main@@9@@Z"(i32 %kernelFunc.coerce) #{{[0-9]+}} {
 // CHECK-HOST-WINDOWS-NEXT: entry:
 // CHECK-HOST-WINDOWS-NEXT:   %kernelFunc = alloca %class.anon, align 4
+// CHECK-HOST-WINDOWS-NEXT:   %agg.tmp = alloca %class.anon, align 4
 // CHECK-HOST-WINDOWS-NEXT:   %coerce.dive = getelementptr inbounds nuw %class.anon, ptr %kernelFunc, i32 0, i32 0
 // CHECK-HOST-WINDOWS-NEXT:   store i32 %kernelFunc.coerce, ptr %coerce.dive, align 4
-// CHECK-HOST-WINDOWS-NEXT:   store ptr @"??_C@_0BC@NHCDOLAA@_ZTSZ4mainEUlT_E_?$AA@", ptr @"?kernel_name@?0???$kernel_single_task@V<lambda_1>@?0??main@@9@V1?0??2@9@@@YAXV<lambda_1>@?0??main@@9@@Z@3PEBDEB", align 8
+// CHECK-HOST-WINDOWS-NEXT:   call void @llvm.memcpy.p0.p0.i64(ptr align 4 %agg.tmp, ptr align 4 %kernelFunc, i64 4, i1 false)
+// CHECK-HOST-WINDOWS-NEXT:   %coerce.dive1 = getelementptr inbounds nuw %class.anon, ptr %agg.tmp, i32 0, i32 0
+// CHECK-HOST-WINDOWS-NEXT:   %0 = load i32, ptr %coerce.dive1, align 4
+// CHECK-HOST-WINDOWS-NEXT:   call void @"??$sycl_kernel_launch@V<lambda_1>@?0??main@@9@V1?0??2@9@@@YAXPEBDV<lambda_1>@?0??main@@9@@Z"(ptr noundef @"??_C@_0BC@NHCDOLAA@_ZTSZ4mainEUlT_E_?$AA@", i32 %0)
+//
 // CHECK-HOST-WINDOWS-NEXT:   ret void
 // CHECK-HOST-WINDOWS-NEXT: }
 //
 // CHECK-HOST-WINDOWS:      define internal void @"??$kernel_single_task@U\CE\B4\CF\84\CF\87@@V<lambda_2>@?0??main@@9@@@YAXV<lambda_2>@?0??main@@9@@Z"(i8 %kernelFunc.coerce) #{{[0-9]+}} {
 // CHECK-HOST-WINDOWS-NEXT: entry:
 // CHECK-HOST-WINDOWS-NEXT:   %kernelFunc = alloca %class.anon.0, align 1
+// CHECK-HOST-WINDOWS-NEXT:   %agg.tmp = alloca %class.anon.0, align 1
 // CHECK-HOST-WINDOWS-NEXT:   %coerce.dive = getelementptr inbounds nuw %class.anon.0, ptr %kernelFunc, i32 0, i32 0
 // CHECK-HOST-WINDOWS-NEXT:   store i8 %kernelFunc.coerce, ptr %coerce.dive, align 1
-// CHECK-HOST-WINDOWS-NEXT:   store ptr @"??_C@_0M@BCGAEMBE@_ZTS6?N?$LE?O?$IE?O?$IH?$AA@", ptr @"?kernel_name@?0???$kernel_single_task@U\CE\B4\CF\84\CF\87@@V<lambda_2>@?0??main@@9@@@YAXV<lambda_2>@?0??main@@9@@Z@3PEBDEB", align 8
+// CHECK-HOST-WINDOWS-NEXT:   %coerce.dive1 = getelementptr inbounds nuw %class.anon.0, ptr %agg.tmp, i32 0, i32 0
+// CHECK-HOST-WINDOWS-NEXT:   %0 = load i8, ptr %coerce.dive1, align 1
+// CHECK-HOST-WINDOWS-NEXT:   call void @"??$sycl_kernel_launch@U\CE\B4\CF\84\CF\87@@V<lambda_2>@?0??main@@9@@@YAXPEBDV<lambda_2>@?0??main@@9@@Z"(ptr noundef @"??_C@_0M@BCGAEMBE@_ZTS6?N?$LE?O?$IE?O?$IH?$AA@", i8 %0)
+// CHECK-HOST-WINDOWS-NEXT:   ret void
+// CHECK-HOST-WINDOWS-NEXT: }
+
+// CHECK-HOST-WINDOWS: define internal void @"??$skep@Vnotaverygoodkernelname@?1??main@@9@V<lambda_3>@?0??2@9@@Handler@@QEAAXV<lambda_3>@?0??main@@9@HH@Z"(ptr noundef nonnull align 1 dereferenceable(1) %this, i8 %k.coerce, i32 noundef %a, i32 noundef %b) #0 align 2 {
+// CHECK-HOST-WINDOWS-NEXT: entry:
+// CHECK-HOST-WINDOWS-NEXT:   %k = alloca %class.anon.1, align 1
+// CHECK-HOST-WINDOWS-NEXT:   %b.addr = alloca i32, align 4
+// CHECK-HOST-WINDOWS-NEXT:   %a.addr = alloca i32, align 4
+// CHECK-HOST-WINDOWS-NEXT:   %this.addr = alloca ptr, align 8
+// CHECK-HOST-WINDOWS-NEXT:   %agg.tmp = alloca %class.anon.1, align 1
+// CHECK-HOST-WINDOWS-NEXT:   %coerce.dive = getelementptr inbounds nuw %class.anon.1, ptr %k, i32 0, i32 0
+// CHECK-HOST-WINDOWS-NEXT:   store i8 %k.coerce, ptr %coerce.dive, align 1
+// CHECK-HOST-WINDOWS-NEXT:   store i32 %b, ptr %b.addr, align 4
+// CHECK-HOST-WINDOWS-NEXT:   store i32 %a, ptr %a.addr, align 4
+// CHECK-HOST-WINDOWS-NEXT:   store ptr %this, ptr %this.addr, align 8
+// CHECK-HOST-WINDOWS-NEXT:   %this1 = load ptr, ptr %this.addr, align 8
+// CHECK-HOST-WINDOWS-NEXT:   %0 = load i32, ptr %b.addr, align 4
+// CHECK-HOST-WINDOWS-NEXT:   %1 = load i32, ptr %a.addr, align 4
+// CHECK-HOST-WINDOWS-NEXT:   %coerce.dive2 = getelementptr inbounds nuw %class.anon.1, ptr %agg.tmp, i32 0, i32 0
+// CHECK-HOST-WINDOWS-NEXT:   %2 = load i8, ptr %coerce.dive2, align 1
+// CHECK-HOST-WINDOWS-NEXT:   call void @"??$sycl_kernel_launch@Vnotaverygoodkernelname@?1??main@@9@V<lambda_3>@?0??2@9@HH@Handler@@AEAAXPEBDV<lambda_3>@?0??main@@9@HH@Z"(ptr noundef nonnull align 1 dereferenceable(1) %this1, ptr noundef @"??_C@_0CE@NJIGCEIA@_ZTSZ4mainE22notaverygoodkerneln@", i8 %2, i32 noundef %1, i32 noundef %0)
 // CHECK-HOST-WINDOWS-NEXT:   ret void
 // CHECK-HOST-WINDOWS-NEXT: }
 
diff --git a/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp b/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp
index 14366a092a1fe..63db83c02bbef 100644
--- a/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp
+++ b/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp
@@ -1,6 +1,8 @@
 // RUN: %clang_cc1 -triple spir64-unknown-unknown -aux-triple x86_64-pc-windows-msvc -fsycl-is-device -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s '-D$ADDRSPACE=addrspace(1) '
 // RUN: %clang_cc1 -triple x86_64-pc-windows-msvc -fsycl-is-host -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s '-D$ADDRSPACE='
 
+template <typename KernelName, typename KernelObj>
+void sycl_kernel_launch(const char *, KernelObj) {}
 
 template<typename KN, typename Func>
 [[clang::sycl_kernel_entry_point(KN)]] void kernel(Func F){
diff --git a/clang/test/SemaSYCL/sycl-host-kernel-launch.cpp b/clang/test/SemaSYCL/sycl-host-kernel-launch.cpp
new file mode 100644
index 0000000000000..eda0c4da489a9
--- /dev/null
+++ b/clang/test/SemaSYCL/sycl-host-kernel-launch.cpp
@@ -0,0 +1,199 @@
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++17 -fsyntax-only -fsycl-is-host -fcxx-exceptions -verify=host,expected %s
+// RUN: %clang_cc1 -triple x86_64-windows-msvc -std=c++17 -fsyntax-only -fsycl-is-host -fcxx-exceptions -verify=host,expected %s
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++17 -fsyntax-only -fsycl-is-device -fcxx-exceptions -verify=device,expected %s
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++20 -fsyntax-only -fsycl-is-host -fcxx-exceptions -verify=host,expected %s
+
+// A unique kernel name type is required for each declared kernel entry point.
+template<int, int = 0> struct KN;
+
+[[clang::sycl_kernel_entry_point(KN<1>)]]
+void nolauncher() {} 
+// host-error@-1 {{unable to find suitable 'sycl_kernel_launch' function for host code synthesis}}
+// device-warning@-2 {{unable to find suitable 'sycl_kernel_launch' function for host code synthesis}}
+// expected-note@-3 {{define 'sycl_kernel_launch' function template to fix}}
+
+void sycl_kernel_launch(const char *, int arg);
+// expected-note@-1 {{declared as a non-template here}}
+
+[[clang::sycl_kernel_entry_point(KN<2>)]]
+void nontemplatel() {}
+// host-error@-1 {{unable to find suitable 'sycl_kernel_launch' function for host code synthesis}}
+// device-warning@-2 {{unable to find suitable 'sycl_kernel_launch' function for host code synthesis}}
+// expected-note@-3 {{define 'sycl_kernel_launch' function template to fix}}
+// expected-error@-4 {{'sycl_kernel_launch' following the 'template' keyword does not refer to a template}}
+
+template <typename KernName>
+void sycl_kernel_launch(const char *, int arg);
+// expected-note@-1 {{candidate function template not viable: requires 2 arguments, but 1 was provided}}
+// expected-note@-2 2{{candidate function template not viable: no known conversion from 'Kern' to 'int' for 2nd argument}}
+
+[[clang::sycl_kernel_entry_point(KN<3>)]]
+void notenoughargs() {}
+// expected-error@-1 {{no matching function for call to 'sycl_kernel_launch'}}
+// FIXME: Should this also say "no suitable function for host code synthesis"?
+
+
+template <typename KernName>
+void sycl_kernel_launch(const char *, bool arg = 1);
+// expected-note@-1 2{{candidate function template not viable: no known conversion from 'Kern' to 'bool' for 2nd argument}}
+
+[[clang::sycl_kernel_entry_point(KN<4>)]]
+void enoughargs() {}
+
+namespace boop {
+template <typename KernName, typename KernelObj>
+void sycl_kernel_launch(const char *, KernelObj);
+
+template <typename KernName, typename KernelObj>
+[[clang::sycl_kernel_entry_point(KernName)]]
+void iboop(KernelObj Kernel) {
+  Kernel();
+}
+}
+
+template <typename KernName, typename KernelObj>
+[[clang::sycl_kernel_entry_point(KernName)]]
+void idontboop(KernelObj Kernel) {
+  Kernel();
+}
+// expected-error@-3 {{no matching function for call to 'sycl_kernel_launch'}}
+
+struct Kern {
+  int a;
+  int *b;
+  Kern(int _a, int* _b) : a(_a), b(_b) {}
+  void operator()(){ *b = a;}
+};
+
+void foo() {
+  int *a;
+  Kern b(1, a);
+  idontboop<KN<6>>(b);
+  // expected-note@-1 {{in instantiation of function template specialization 'idontboop<KN<6>, Kern>' requested here}}
+  boop::iboop<KN<7>>(b);
+}
+
+class MaybeHandler {
+
+template <typename KernName>
+void sycl_kernel_launch(const char *);
+
+template <typename KernName, typename... Tys>
+void sycl_kernel_launch(const char *, Tys ...Args);
+
+public:
+
+template <typename KernName, typename KernelObj>
+[[clang::sycl_kernel_entry_point(KernName)]]
+void entry(KernelObj Kernel) {
+  Kernel();
+}
+};
+
+class MaybeHandler2 {
+
+template <typename KernName, typename... Tys>
+static void sycl_kernel_launch(const char *, Tys ...Args);
+
+public:
+
+template <typename KernName, typename KernelObj>
+[[clang::sycl_kernel_entry_point(KernName)]]
+void entry(KernelObj Kernel) {
+  Kernel();
+}
+};
+
+class MaybeHandler3 {
+
+template <typename KernName, typename... Tys>
+static void sycl_kernel_launch(const char *, Tys ...Args);
+
+public:
+
+template <typename KernName, typename KernelObj>
+[[clang::sycl_kernel_entry_point(KernName)]]
+static void entry(KernelObj Kernel) {
+  Kernel();
+}
+};
+
+class MaybeHandler4 {
+
+template <typename KernName, typename... Tys>
+void sycl_kernel_launch(const char *, Tys ...Args);
+
+public:
+
+template <typename KernName, typename KernelObj>
+[[clang::sycl_kernel_entry_point(KernName)]]
+static void entry(KernelObj Kernel) { 
+  // expected-error@-1 {{call to non-static member function without an object argument}}
+  // FIXME: Should that be clearer?
+  Kernel();
+}
+};
+
+template<typename>
+struct base_handler {
+  template<typename KNT, typename... Ts>
+  void sycl_kernel_launch(const char*, Ts...) {}
+};
+struct derived_handler : base_handler<derived_handler> {
+  template<typename KNT, typename KT>
+  [[clang::sycl_kernel_entry_point(KNT)]]
+  void entry(KT k) { k(); }
+};
+
+template<int N>
+struct derived_handler_t : base_handler<derived_handler_t<N>> {
+  template<typename KNT, typename KT>
+// FIXME this fails because accessing members of dependent bases requires
+// explicit qualification.
+  [[clang::sycl_kernel_entry_point(KNT)]]
+  void entry(KT k) { k(); }
+  // expected-error@-1 {{no matching function for call to 'sycl_kernel_launch'}}
+};
+
+template<typename KNT>
+struct kernel_launcher {
+  template<typename... Ts>
+  void operator()(const char*, Ts...) const {}
+};
+
+namespace var {
+template<typename KNT>
+kernel_launcher<KNT> sycl_kernel_launch;
+
+struct handler {
+  template<typename KNT, typename KT>
+  [[clang::sycl_kernel_entry_point(KNT)]]
+  void entry(KT k) { k(); }
+};
+}
+
+
+void bar() {
+  int *a;
+  Kern b(1, a);
+  MaybeHandler H;
+  MaybeHandler2 H1;
+  MaybeHandler3 H2;
+  MaybeHandler4 H3;
+  H.entry<KN<8>>(b);
+  H1.entry<KN<9>>(b);
+  H2.entry<KN<10>>(b);
+  H3.entry<KN<11>>(b);
+
+  derived_handler H5;
+  H5.entry<KN<12>>(b);
+
+  derived_handler_t<13> H6;
+  H6.entry<KN<13>>(b); //expected-note {{in instantiation of function template specialization}}
+
+  var::handler h;
+  h.entry<KN<14>>(b);
+}
+
+
+
diff --git a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-appertainment.cpp b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-appertainment.cpp
index 3f07feb87c9a1..c181f76321a26 100644
--- a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-appertainment.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-appertainment.cpp
@@ -40,6 +40,10 @@ struct coroutine_traits {
 // A unique kernel name type is required for each declared kernel entry point.
 template<int, int = 0> struct KN;
 
+// A launcher function definition required for host code synthesis to silence
+// complains.
+template <typename KernelName, typename... Tys>
+void sycl_kernel_launch(const char *, Tys &&...Args) {}
 
 ////////////////////////////////////////////////////////////////////////////////
 // Valid declarations.
diff --git a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-grammar.cpp b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-grammar.cpp
index fd1f00ae05d7a..1cdd48f1e5840 100644
--- a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-grammar.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-grammar.cpp
@@ -10,6 +10,10 @@
 template<int> struct ST; // #ST-decl
 template<int N> using TTA = ST<N>; // #TTA-decl
 
+// A launcher function definition required for host code synthesis to silence
+// complains.
+template <typename KernelName, typename... Tys>
+void sycl_kernel_launch(const char *, Tys &&...Args) {}
 
 ////////////////////////////////////////////////////////////////////////////////
 // Valid declarations.
diff --git a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-kernel-name-module.cpp b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-kernel-name-module.cpp
index 8788e147a2ae4..44a3ce6f3640a 100644
--- a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-kernel-name-module.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-kernel-name-module.cpp
@@ -17,6 +17,11 @@ module M2 { header "m2.h" }
 #--- common.h
 template<int> struct KN;
 
+// A launcher function definition required for host code synthesis to silence
+// complains.
+template <typename KernelName, typename... Tys>
+void sycl_kernel_launch(const char *, Tys &&...Args) {}
+
 [[clang::sycl_kernel_entry_point(KN<1>)]]
 void common_test1() {}
 
@@ -25,7 +30,6 @@ template<typename T>
 void common_test2() {}
 template void common_test2<KN<2>>();
 
-
 #--- m1.h
 #include "common.h"
 
diff --git a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-kernel-name-pch.cpp b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-kernel-name-pch.cpp
index 0575a7a5a67eb..0e6d1a6c57e39 100644
--- a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-kernel-name-pch.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-kernel-name-pch.cpp
@@ -15,6 +15,11 @@
 #--- pch.h
 template<int> struct KN;
 
+// A launcher function definition required for host code synthesis to silence
+// complains.
+template <typename KernelName, typename... Tys>
+void sycl_kernel_launch(const char *, Tys &&...Args) {}
+
 [[clang::sycl_kernel_entry_point(KN<1>)]]
 void pch_test1() {} // << expected previous declaration note here.
 
@@ -26,11 +31,11 @@ template void pch_test2<KN<2>>();
 
 #--- test.cpp
 // expected-error@+3 {{the 'clang::sycl_kernel_entry_point' kernel name argument conflicts with a previous declaration}}
-// expected-note@pch.h:4 {{previous declaration is here}}
+// expected-note@pch.h:9 {{previous declaration is here}}
 [[clang::sycl_kernel_entry_point(KN<1>)]]
 void test1() {}
 
 // expected-error@+3 {{the 'clang::sycl_kernel_entry_point' kernel name argument conflicts with a previous declaration}}
-// expected-note@pch.h:8 {{previous declaration is here}}
+// expected-note@pch.h:13 {{previous declaration is here}}
 [[clang::sycl_kernel_entry_point(KN<2>)]]
 void test2() {}
diff --git a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-kernel-name.cpp b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-kernel-name.cpp
index 5a3b43be66daf..7b525abadd2c3 100644
--- a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-kernel-name.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-kernel-name.cpp
@@ -9,6 +9,12 @@
 // specification.
 
 struct S1;
+
+// A launcher function definition required for host code synthesis to silence
+// complains.
+template <typename KernelName, typename... Tys>
+void sycl_kernel_launch(const char *, Tys &&...Args) {}
+
 // expected-warning@+3 {{redundant 'clang::sycl_kernel_entry_point' attribute}}
 // expected-note@+1  {{previous attribute is here}}
 [[clang::sycl_kernel_entry_point(S1),
diff --git a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-sfinae.cpp b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-sfinae.cpp
index 3689adaab9b5b..9674dac456f9f 100644
--- a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-sfinae.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-sfinae.cpp
@@ -10,6 +10,11 @@
 // attribute during instantiation of a specialization unless that specialization
 // is selected by overload resolution.
 
+// A launcher function definition required for host code synthesis to silence
+// complains.
+template <typename KernelName, typename... Tys>
+void sycl_kernel_launch(const char *, Tys &&...Args) {}
+
 // FIXME: C++23 [temp.expl.spec]p12 states:
 // FIXME:   ... Similarly, attributes appearing in the declaration of a template
 // FIXME:   have no effect on an explicit specialization of that template.
diff --git a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-this.cpp b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-this.cpp
index fc0640e1900cb..7b8fc6c9a7630 100644
--- a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-this.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-this.cpp
@@ -22,6 +22,11 @@ struct type_info {
 };
 } // namespace std
 
+// A launcher function definition required for host code synthesis to silence
+// complains.
+template <typename KernelName, typename... Tys>
+void sycl_kernel_launch(const char *, Tys &&...Args) {}
+
 ////////////////////////////////////////////////////////////////////////////////
 // Valid declarations.
 ////////////////////////////////////////////////////////////////////////////////
diff --git a/clang/tools/libclang/CXCursor.cpp b/clang/tools/libclang/CXCursor.cpp
index 3c4062410eac1..ad44aea6e3f5e 100644
--- a/clang/tools/libclang/CXCursor.cpp
+++ b/clang/tools/libclang/CXCursor.cpp
@@ -376,6 +376,7 @@ CXCursor cxcursor::MakeCXCursor(const Stmt *S, const Decl *Parent,
     break;
 
   case Stmt::SYCLKernelCallStmtClass:
+  case Stmt::UnresolvedSYCLKernelCallStmtClass:
     K = CXCursor_UnexposedStmt;
     break;