diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index 4249852b80182..f61c83297c8ed 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -315,7 +315,8 @@ class SYCLIntegrationHeader {
     kind_std_layout,
     kind_sampler,
     kind_pointer,
-    kind_last = kind_pointer
+    kind_specialization_constants_buffer,
+    kind_last = kind_specialization_constants_buffer
   };
 
 public:
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 2939bf4316a06..9ffa0ce900698 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -58,6 +58,8 @@ enum KernelInvocationKind {
 
 static constexpr llvm::StringLiteral InitMethodName = "__init";
 static constexpr llvm::StringLiteral InitESIMDMethodName = "__init_esimd";
+static constexpr llvm::StringLiteral InitSpecConstantsBuffer =
+    "__init_specialization_constants_buffer";
 static constexpr llvm::StringLiteral FinalizeMethodName = "__finalize";
 constexpr unsigned MaxKernelArgsSize = 2048;
 
@@ -109,6 +111,10 @@ class Util {
   /// specialization constant class.
   static bool isSyclSpecConstantType(const QualType &Ty);
 
+  /// Checks whether given clang type is a full specialization of the SYCL
+  /// kernel_handler class.
+  static bool isSyclKernelHandlerType(const QualType &Ty);
+
   // Checks declaration context hierarchy.
   /// \param DC     the context of the item to be checked.
   /// \param Scopes the declaration scopes leading from the item context to the
@@ -616,11 +622,16 @@ class FindPFWGLambdaFnVisitor
     auto *M = dyn_cast<CXXMethodDecl>(Call->getDirectCallee());
     if (!M || (M->getOverloadedOperator() != OO_Call))
       return true;
-    const int NumPFWGLambdaArgs = 2; // group and lambda obj
+
+    unsigned int NumPFWGLambdaArgs =
+        M->getNumParams() + 1; // group, optional kernel_handler and lambda obj
     if (Call->getNumArgs() != NumPFWGLambdaArgs)
       return true;
     if (!Util::isSyclType(Call->getArg(1)->getType(), "group", true /*Tmpl*/))
       return true;
+    if ((Call->getNumArgs() > 2) &&
+        !Util::isSyclKernelHandlerType(Call->getArg(2)->getType()))
+      return true;
     if (Call->getArg(0)->getType()->getAsCXXRecordDecl() != LambdaObjTy)
       return true;
     LambdaFn = M; // call to PFWG lambda found - record the lambda
@@ -732,12 +743,7 @@ static ParamDesc makeParamDesc(const FieldDecl *Src, QualType Ty) {
                          Ctx.getTrivialTypeSourceInfo(Ty));
 }
 
-static ParamDesc makeParamDesc(ASTContext &Ctx, const CXXBaseSpecifier &Src,
-                               QualType Ty) {
-  // TODO: There is no name for the base available, but duplicate names are
-  // seemingly already possible, so we'll give them all the same name for now.
-  // This only happens with the accessor types.
-  std::string Name = "_arg__base";
+static ParamDesc makeParamDesc(ASTContext &Ctx, StringRef Name, QualType Ty) {
   return std::make_tuple(Ty, &Ctx.Idents.get(Name),
                          Ctx.getTrivialTypeSourceInfo(Ty));
 }
@@ -777,6 +783,28 @@ constructKernelName(Sema &S, FunctionDecl *KernelCallerFunc,
                                       KernelNameType)};
 }
 
+static bool isDefaultSPIRArch(ASTContext &Context) {
+  llvm::Triple T = Context.getTargetInfo().getTriple();
+  if (T.isSPIR() && T.getSubArch() == llvm::Triple::NoSubArch)
+    return true;
+  return false;
+}
+
+static ParmVarDecl *getSyclKernelHandlerArg(FunctionDecl *KernelCallerFunc) {
+  // Specialization constants in SYCL 2020 are not captured by lambda and
+  // accessed through new optional lambda argument kernel_handler
+  auto IsHandlerLambda = [](ParmVarDecl *PVD) {
+    return Util::isSyclKernelHandlerType(PVD->getType());
+  };
+
+  assert(llvm::count_if(KernelCallerFunc->parameters(), IsHandlerLambda) <= 1 &&
+         "Multiple kernel_handler parameters");
+
+  auto KHArg = llvm::find_if(KernelCallerFunc->parameters(), IsHandlerLambda);
+
+  return (KHArg != KernelCallerFunc->param_end()) ? *KHArg : nullptr;
+}
+
 // anonymous namespace so these don't get linkage.
 namespace {
 
@@ -1642,10 +1670,20 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
   }
 
   void addParam(const CXXBaseSpecifier &BS, QualType FieldTy) {
+    // TODO: There is no name for the base available, but duplicate names are
+    // seemingly already possible, so we'll give them all the same name for now.
+    // This only happens with the accessor types.
+    StringRef Name = "_arg__base";
     ParamDesc newParamDesc =
-        makeParamDesc(SemaRef.getASTContext(), BS, FieldTy);
+        makeParamDesc(SemaRef.getASTContext(), Name, FieldTy);
     addParam(newParamDesc, FieldTy);
   }
+  // Add a parameter with specified name and type
+  void addParam(StringRef Name, QualType ParamTy) {
+    ParamDesc newParamDesc =
+        makeParamDesc(SemaRef.getASTContext(), Name, ParamTy);
+    addParam(newParamDesc, ParamTy);
+  }
 
   void addParam(ParamDesc newParamDesc, QualType FieldTy) {
     // Create a new ParmVarDecl based on the new info.
@@ -1946,6 +1984,18 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
     return true;
   }
 
+  // Generate kernel argument to intialize specialization constants. This
+  // argument is only generated when the target has no native support for
+  // specialization constants
+  void handleSyclKernelHandlerType() {
+    ASTContext &Context = SemaRef.getASTContext();
+    if (isDefaultSPIRArch(Context))
+      return;
+
+    StringRef Name = "_arg__specialization_constants_buffer";
+    addParam(Name, Context.getPointerType(Context.CharTy));
+  }
+
   void setBody(CompoundStmt *KB) { KernelDecl->setBody(KB); }
 
   FunctionDecl *getKernelDecl() { return KernelDecl; }
@@ -2091,28 +2141,46 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
   // pointer-struct-wrapping code to ensure that we don't try to wrap
   // non-top-level pointers.
   uint64_t StructDepth = 0;
+  VarDecl *KernelHandlerClone = nullptr;
+
+  Stmt *replaceWithLocalClone(ParmVarDecl *OriginalParam, VarDecl *LocalClone,
+                              Stmt *FunctionBody) {
+    // DeclRefExpr with valid source location but with decl which is not marked
+    // as used is invalid.
+    LocalClone->setIsUsed();
+    std::pair<DeclaratorDecl *, DeclaratorDecl *> MappingPair =
+        std::make_pair(OriginalParam, LocalClone);
+    KernelBodyTransform KBT(MappingPair, SemaRef);
+    return KBT.TransformStmt(FunctionBody).get();
+  }
 
   // Using the statements/init expressions that we've created, this generates
   // the kernel body compound stmt. CompoundStmt needs to know its number of
   // statements in advance to allocate it, so we cannot do this as we go along.
   CompoundStmt *createKernelBody() {
+    // Push the Kernel function scope to ensure the scope isn't empty
+    SemaRef.PushFunctionScope();
+
+    // Initialize kernel object local clone
     assert(CollectionInitExprs.size() == 1 &&
            "Should have been popped down to just the first one");
     KernelObjClone->setInit(CollectionInitExprs.back());
-    Stmt *FunctionBody = KernelCallerFunc->getBody();
-
-    ParmVarDecl *KernelObjParam = *(KernelCallerFunc->param_begin());
 
-    // DeclRefExpr with valid source location but with decl which is not marked
-    // as used is invalid.
-    KernelObjClone->setIsUsed();
-    std::pair<DeclaratorDecl *, DeclaratorDecl *> MappingPair =
-        std::make_pair(KernelObjParam, KernelObjClone);
-
-    // Push the Kernel function scope to ensure the scope isn't empty
-    SemaRef.PushFunctionScope();
-    KernelBodyTransform KBT(MappingPair, SemaRef);
-    Stmt *NewBody = KBT.TransformStmt(FunctionBody).get();
+    // Replace references to the kernel object in kernel body, to use the
+    // compiler generated local clone
+    Stmt *NewBody =
+        replaceWithLocalClone(KernelCallerFunc->getParamDecl(0), KernelObjClone,
+                              KernelCallerFunc->getBody());
+
+    // If kernel_handler argument is passed by SYCL kernel, replace references
+    // to this argument in kernel body, to use the compiler generated local
+    // clone
+    if (ParmVarDecl *KernelHandlerParam =
+            getSyclKernelHandlerArg(KernelCallerFunc))
+      NewBody = replaceWithLocalClone(KernelHandlerParam, KernelHandlerClone,
+                                      NewBody);
+
+    // Use transformed body (with clones) as kernel body
     BodyStmts.push_back(NewBody);
 
     BodyStmts.insert(BodyStmts.end(), FinalizeStmts.begin(),
@@ -2412,6 +2480,39 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
     return true;
   }
 
+  // Generate __init call for kernel handler argument
+  void handleSpecialType(QualType KernelHandlerTy) {
+    DeclRefExpr *KernelHandlerCloneRef =
+        DeclRefExpr::Create(SemaRef.Context, NestedNameSpecifierLoc(),
+                            KernelCallerSrcLoc, KernelHandlerClone, false,
+                            DeclarationNameInfo(), KernelHandlerTy, VK_LValue);
+    const auto *RecordDecl =
+        KernelHandlerClone->getType()->getAsCXXRecordDecl();
+    MemberExprBases.push_back(KernelHandlerCloneRef);
+    createSpecialMethodCall(RecordDecl, InitSpecConstantsBuffer, BodyStmts);
+    MemberExprBases.pop_back();
+  }
+
+  void createKernelHandlerClone(ASTContext &Ctx, DeclContext *DC,
+                                ParmVarDecl *KernelHandlerArg) {
+    QualType Ty = KernelHandlerArg->getType();
+    TypeSourceInfo *TSInfo = Ctx.getTrivialTypeSourceInfo(Ty);
+    KernelHandlerClone =
+        VarDecl::Create(Ctx, DC, KernelCallerSrcLoc, KernelCallerSrcLoc,
+                        KernelHandlerArg->getIdentifier(), Ty, TSInfo, SC_None);
+
+    // Default initialize clone
+    InitializedEntity VarEntity =
+        InitializedEntity::InitializeVariable(KernelHandlerClone);
+    InitializationKind InitKind =
+        InitializationKind::CreateDefault(KernelCallerSrcLoc);
+    InitializationSequence InitSeq(SemaRef, VarEntity, InitKind, None);
+    ExprResult Init = InitSeq.Perform(SemaRef, VarEntity, InitKind, None);
+    KernelHandlerClone->setInit(
+        SemaRef.MaybeCreateExprWithCleanups(Init.get()));
+    KernelHandlerClone->setInitStyle(VarDecl::CallInit);
+  }
+
 public:
   static constexpr const bool VisitInsideSimpleContainers = false;
   SyclKernelBodyCreator(Sema &S, SyclKernelDeclCreator &DC,
@@ -2516,6 +2617,28 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
     return true;
   }
 
+  // Default inits the type, then calls the init-method in the body
+  void handleSyclKernelHandlerType(ParmVarDecl *KernelHandlerArg) {
+
+    // Create and default initialize local clone of kernel handler
+    createKernelHandlerClone(SemaRef.getASTContext(),
+                             DeclCreator.getKernelDecl(), KernelHandlerArg);
+
+    // Add declaration statement to openCL kernel body
+    Stmt *DS =
+        new (SemaRef.Context) DeclStmt(DeclGroupRef(KernelHandlerClone),
+                                       KernelCallerSrcLoc, KernelCallerSrcLoc);
+    BodyStmts.push_back(DS);
+
+    // Generate
+    // KernelHandlerClone.__init_specialization_constants_buffer(specialization_constants_buffer)
+    // call if target does not have native support for specialization constants.
+    // Here, specialization_constants_buffer is the compiler generated kernel
+    // argument of type char*.
+    if (!isDefaultSPIRArch(SemaRef.Context))
+      handleSpecialType(KernelHandlerArg->getType());
+  }
+
   bool enterStream(const CXXRecordDecl *RD, FieldDecl *FD, QualType Ty) final {
     ++StructDepth;
     // Add a dummy init expression to catch the accessor initializers.
@@ -2870,6 +2993,22 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
     return true;
   }
 
+  void handleSyclKernelHandlerType(QualType Ty) {
+    // The compiler generated kernel argument used to initialize SYCL 2020
+    // specialization constants, `specialization_constants_buffer`, should
+    // have corresponding entry in integration header. This argument is
+    // only generated when target has no native support for specialization
+    // constants.
+    ASTContext &Context = SemaRef.getASTContext();
+    if (isDefaultSPIRArch(Context))
+      return;
+
+    // Offset is zero since kernel_handler argument is not part of
+    // kernel object (i.e. it is not captured)
+    addParam(Context.getPointerType(Context.CharTy),
+             SYCLIntegrationHeader::kind_specialization_constants_buffer, 0);
+  }
+
   bool enterStream(const CXXRecordDecl *, FieldDecl *FD, QualType Ty) final {
     ++StructDepth;
     CurOffset += offsetOf(FD, Ty);
@@ -3257,6 +3396,13 @@ void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc,
   KernelObjVisitor Visitor{*this};
   Visitor.VisitRecordBases(KernelObj, kernel_decl, kernel_body, int_header);
   Visitor.VisitRecordFields(KernelObj, kernel_decl, kernel_body, int_header);
+
+  if (ParmVarDecl *KernelHandlerArg =
+          getSyclKernelHandlerArg(KernelCallerFunc)) {
+    kernel_decl.handleSyclKernelHandlerType();
+    kernel_body.handleSyclKernelHandlerType(KernelHandlerArg);
+    int_header.handleSyclKernelHandlerType(KernelHandlerArg->getType());
+  }
 }
 
 void Sema::MarkDevice(void) {
@@ -3504,6 +3650,7 @@ static const char *paramKind2Str(KernelParamKind K) {
     CASE(accessor);
     CASE(std_layout);
     CASE(sampler);
+    CASE(specialization_constants_buffer);
     CASE(pointer);
   }
   return "<ERROR>";
@@ -4089,6 +4236,15 @@ bool Util::isSyclSpecConstantType(const QualType &Ty) {
   return matchQualifiedTypeName(Ty, Scopes);
 }
 
+bool Util::isSyclKernelHandlerType(const QualType &Ty) {
+  const StringRef &Name = "kernel_handler";
+  std::array<DeclContextDesc, 3> Scopes = {
+      Util::DeclContextDesc{clang::Decl::Kind::Namespace, "cl"},
+      Util::DeclContextDesc{clang::Decl::Kind::Namespace, "sycl"},
+      Util::DeclContextDesc{Decl::Kind::CXXRecord, Name}};
+  return matchQualifiedTypeName(Ty, Scopes);
+}
+
 bool Util::isSyclBufferLocationType(const QualType &Ty) {
   const StringRef &PropertyName = "buffer_location";
   const StringRef &InstanceName = "instance";
diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp
index 5462a2d13d1f6..fc3ca2c146ad6 100644
--- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp
+++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp
@@ -291,14 +291,23 @@ class spec_constant {
 } // namespace experimental
 } // namespace ONEAPI
 
+class kernel_handler {
+  void __init_specialization_constants_buffer(char *specialization_constants_buffer) {}
+};
+
 #define ATTR_SYCL_KERNEL __attribute__((sycl_kernel))
 template <typename KernelName = auto_name, typename KernelType>
-ATTR_SYCL_KERNEL void kernel_single_task(const KernelType &kernelFunc) {
+ATTR_SYCL_KERNEL void kernel_single_task(const KernelType &kernelFunc) { // #KernelSingleTask
   kernelFunc();
 }
 
 template <typename KernelName = auto_name, typename KernelType>
-ATTR_SYCL_KERNEL void kernel_single_task_2017(KernelType kernelFunc) {
+ATTR_SYCL_KERNEL void kernel_single_task(const KernelType &kernelFunc, kernel_handler kh) {
+  kernelFunc(kh);
+}
+
+template <typename KernelName = auto_name, typename KernelType>
+ATTR_SYCL_KERNEL void kernel_single_task_2017(KernelType kernelFunc) { // #KernelSingleTask2017
   kernelFunc();
 }
 
@@ -347,6 +356,16 @@ class handler {
 #endif
   }
 
+  template <typename KernelName = auto_name, typename KernelType>
+  void single_task(const KernelType &kernelFunc, kernel_handler kh) {
+    using NameT = typename get_kernel_name_t<KernelName, KernelType>::name;
+#ifdef __SYCL_DEVICE_ONLY__
+    kernel_single_task<NameT>(kernelFunc, kh);
+#else
+    kernelFunc(kh);
+#endif
+  }
+
   template <typename KernelName = auto_name, typename KernelType>
   void single_task_2017(KernelType kernelFunc) {
     using NameT = typename get_kernel_name_t<KernelName, KernelType>::name;
diff --git a/clang/test/CodeGenSYCL/int_header_sycl2020_spec_const.cpp b/clang/test/CodeGenSYCL/int_header_sycl2020_spec_const.cpp
new file mode 100644
index 0000000000000..55733076e3efc
--- /dev/null
+++ b/clang/test/CodeGenSYCL/int_header_sycl2020_spec_const.cpp
@@ -0,0 +1,29 @@
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple nvptx64-unknown-unknown -fsycl-int-header=%t.h %s -o %t.out %s -o %t.out
+// RUN: FileCheck -input-file=%t.h %s --check-prefix=NONATIVESUPPORT --check-prefix=ALL
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -fsycl-int-header=%t.h %s -o %t.out %s -o %t.out
+// RUN: FileCheck -input-file=%t.h %s --check-prefix=NATIVESUPPORT --check-prefix=ALL
+
+// This test checks that the compiler generates required information
+// in integration header for kernel_handler type (SYCL 2020 specialization
+// constants).
+
+#include "sycl.hpp"
+
+using namespace cl::sycl;
+queue q;
+
+int main() {
+  q.submit([&](handler &h) {
+    int a;
+    kernel_handler kh;
+
+    h.single_task<class test_kernel_handler>(
+        [=](auto) {
+          int local = a;
+        },
+        kh);
+  });
+}
+// ALL: const kernel_param_desc_t kernel_signatures[] = {
+// NONATIVESUPPORT: { kernel_param_kind_t::kind_specialization_constants_buffer, 8, 0 }
+// NATIVESUPPORT-NOT: { kernel_param_kind_t::kind_specialization_constants_buffer, 8, 0 }
diff --git a/clang/test/CodeGenSYCL/kernel-by-reference.cpp b/clang/test/CodeGenSYCL/kernel-by-reference.cpp
index 056eea081cadf..720290c7b05cb 100644
--- a/clang/test/CodeGenSYCL/kernel-by-reference.cpp
+++ b/clang/test/CodeGenSYCL/kernel-by-reference.cpp
@@ -15,7 +15,7 @@ int simple_add(int i) {
 int main() {
   queue q;
 #if defined(SYCL2020)
-  // expected-warning@Inputs/sycl.hpp:301 {{Passing kernel functions by value is deprecated in SYCL 2020}}
+  // expected-warning@#KernelSingleTask2017 {{Passing kernel functions by value is deprecated in SYCL 2020}}
   // expected-note@+3 {{in instantiation of function template specialization}}
 #endif
   q.submit([&](handler &h) {
@@ -23,7 +23,7 @@ int main() {
   });
 
 #if defined(SYCL2017)
-  // expected-warning@Inputs/sycl.hpp:296 {{Passing of kernel functions by reference is a SYCL 2020 extension}}
+  // expected-warning@#KernelSingleTask {{Passing of kernel functions by reference is a SYCL 2020 extension}}
   // expected-note@+3 {{in instantiation of function template specialization}}
 #endif
   q.submit([&](handler &h) {
diff --git a/clang/test/CodeGenSYCL/kernel-handler.cpp b/clang/test/CodeGenSYCL/kernel-handler.cpp
new file mode 100644
index 0000000000000..9c24acc31a74c
--- /dev/null
+++ b/clang/test/CodeGenSYCL/kernel-handler.cpp
@@ -0,0 +1,38 @@
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple nvptx64-unknown-unknown -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s --check-prefix=NONATIVESUPPORT
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s --check-prefix=NATIVESUPPORT
+
+// This test checks IR generated when kernel_handler argument
+// (used to handle SYCL 2020 specialization constants) is passed
+// by kernel
+
+#include "sycl.hpp"
+
+using namespace cl::sycl;
+
+void test(int val) {
+  queue q;
+  q.submit([&](handler &h) {
+    int a;
+    kernel_handler kh;
+    h.single_task<class test_kernel_handler>(
+        [=](auto) {
+          int local = a;
+        },
+        kh);
+  });
+}
+
+// NONATIVESUPPORT: define dso_local void @"{{.*}}test_kernel_handler{{.*}}"
+// NONATIVESUPPORT-SAME: (i32 %_arg_, i8* %_arg__specialization_constants_buffer)
+// NONATIVESUPPORT: %kh = alloca %"class.[[MANGLEDCLASS:[a-zA-Z0-9_]+]].cl::sycl::kernel_handler", align 1
+// NONATIVESUPPORT: %[[KH:[0-9]+]] = load i8*, i8** %_arg__specialization_constants_buffer.addr, align 8
+// NONATIVESUPPORT: call void @{{.*}}__init_specialization_constants_buffer{{.*}}(%"class.[[MANGLEDCLASS]].cl::sycl::kernel_handler"* nonnull dereferenceable(1) %kh, i8* %[[KH]])
+// NONATIVESUPPORT: void @"[[MANGLEDKERNELCALL:[a-zA-Z0-9_$]+]]"
+// NONATIVESUPPORT-SAME: byval(%"class.[[MANGLEDCLASS]].cl::sycl::kernel_handler")
+
+// NATIVESUPPORT: define dso_local spir_kernel void @"{{.*}}test_kernel_handler{{.*}}"
+// NATIVESUPPORT-SAME: (i32 %_arg_)
+// NATIVESUPPORT: %kh = alloca %"class.[[MANGLEDCLASS:[a-zA-Z0-9_]+]].cl::sycl::kernel_handler"
+// NATIVESUPPORT-NOT: __init_specialization_constants_buffer
+// NATIVE-SUPPORT: call spir_func void @"[[MANGLEDKERNELCALL:[a-zA-Z0-9_$]+]]"
+// NATIVE-SUPPORT-SAME: byval(%"class.[[MANGLEDCLASS]].cl::sycl::kernel_handler")
diff --git a/clang/test/SemaSYCL/Inputs/sycl.hpp b/clang/test/SemaSYCL/Inputs/sycl.hpp
index a1a8d626b7c40..4782bafa74ce5 100644
--- a/clang/test/SemaSYCL/Inputs/sycl.hpp
+++ b/clang/test/SemaSYCL/Inputs/sycl.hpp
@@ -207,6 +207,16 @@ struct get_kernel_name_t<auto_name, Type> {
   using name = Type;
 };
 
+template <int dimensions = 1>
+class group {
+public:
+  group() = default; // fake constructor
+};
+
+class kernel_handler {
+  void __init_specialization_constants_buffer(char *specialization_constants_buffer) {}
+};
+
 // Used when parallel_for range is rounded-up.
 template <typename Type> class __pf_kernel_wrapper;
 
@@ -218,21 +228,39 @@ template <typename Type> struct get_kernel_wrapper_name_t {
 #define ATTR_SYCL_KERNEL __attribute__((sycl_kernel))
 template <typename KernelName = auto_name, typename KernelType>
 ATTR_SYCL_KERNEL void kernel_single_task(const KernelType &kernelFunc) {
-  kernelFunc();
+  kernelFunc(); // #KernelSingleTaskKernelFuncCall
+}
+template <typename KernelName = auto_name, typename KernelType>
+ATTR_SYCL_KERNEL void kernel_single_task(const KernelType &kernelFunc, kernel_handler kh) {
+  kernelFunc(kh);
 }
 template <typename KernelName = auto_name, typename KernelType>
 ATTR_SYCL_KERNEL void kernel_parallel_for(const KernelType &kernelFunc) {
   kernelFunc();
 }
+template <typename KernelName, typename KernelType>
+ATTR_SYCL_KERNEL void kernel_parallel_for_work_group(const KernelType &KernelFunc, kernel_handler kh) {
+  KernelFunc(group<1>(), kh);
+}
+
 class handler {
 public:
   template <typename KernelName = auto_name, typename KernelType>
   void single_task(const KernelType &kernelFunc) {
     using NameT = typename get_kernel_name_t<KernelName, KernelType>::name;
 #ifdef __SYCL_DEVICE_ONLY__
-    kernel_single_task<NameT>(kernelFunc);
+    kernel_single_task<NameT>(kernelFunc); // #KernelSingleTask
 #else
     kernelFunc();
+#endif
+  }
+  template <typename KernelName = auto_name, typename KernelType>
+  void single_task(const KernelType &kernelFunc, kernel_handler kh) {
+    using NameT = typename get_kernel_name_t<KernelName, KernelType>::name;
+#ifdef __SYCL_DEVICE_ONLY__
+    kernel_single_task<NameT>(kernelFunc, kh);
+#else
+    kernelFunc(kh);
 #endif
   }
   template <typename KernelName = auto_name, typename KernelType>
@@ -243,6 +271,16 @@ class handler {
     kernel_parallel_for<NameT>(kernelObj);
 #else
     kernelObj();
+#endif
+  }
+  template <typename KernelName = auto_name, typename KernelType>
+  void parallel_for_work_group(const KernelType &kernelFunc, kernel_handler kh) {
+    using NameT = typename get_kernel_name_t<KernelName, KernelType>::name;
+#ifdef __SYCL_DEVICE_ONLY__
+    kernel_parallel_for_work_group<NameT>(kernelFunc, kh);
+#else
+    group<1> G;
+    kernelFunc(G, kh);
 #endif
   }
 };
diff --git a/clang/test/SemaSYCL/args-size-overflow.cpp b/clang/test/SemaSYCL/args-size-overflow.cpp
index fe28acde50d7e..d3d88d39ca4e0 100644
--- a/clang/test/SemaSYCL/args-size-overflow.cpp
+++ b/clang/test/SemaSYCL/args-size-overflow.cpp
@@ -11,9 +11,9 @@ queue q;
 using Accessor =
     accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer>;
 #ifdef SPIR64
-// expected-warning@Inputs/sycl.hpp:233 {{size of kernel arguments (7994 bytes) may exceed the supported maximum of 2048 bytes on some devices}}
+// expected-warning@#KernelSingleTask {{size of kernel arguments (7994 bytes) may exceed the supported maximum of 2048 bytes on some devices}}
 #elif SPIR32
-// expected-warning@Inputs/sycl.hpp:233 {{size of kernel arguments (7986 bytes) may exceed the supported maximum of 2048 bytes on some devices}}
+// expected-warning@#KernelSingleTask {{size of kernel arguments (7986 bytes) may exceed the supported maximum of 2048 bytes on some devices}}
 #endif
 
 void use() {
diff --git a/clang/test/SemaSYCL/deferred-diagnostics-aux-builtin.cpp b/clang/test/SemaSYCL/deferred-diagnostics-aux-builtin.cpp
index 6e0f8d8659f38..3caee03756cad 100644
--- a/clang/test/SemaSYCL/deferred-diagnostics-aux-builtin.cpp
+++ b/clang/test/SemaSYCL/deferred-diagnostics-aux-builtin.cpp
@@ -12,7 +12,7 @@ int main(int argc, char **argv) {
   _mm_prefetch("test", 8); // expected-error {{argument value 8 is outside the valid range [0, 7]}}
 
   deviceQueue.submit([&](sycl::handler &h) {
-    // expected-note@Inputs/sycl.hpp:221 {{called by 'kernel_single_task<AName, (lambda}}
+    // expected-note@#KernelSingleTaskKernelFuncCall {{called by 'kernel_single_task<AName, (lambda}}
     h.single_task<class AName>([]() {
       _mm_prefetch("test", 4); // expected-error {{builtin is not supported on this target}}
       _mm_prefetch("test", 8); // expected-error {{argument value 8 is outside the valid range [0, 7]}} expected-error {{builtin is not supported on this target}}
@@ -20,4 +20,4 @@ int main(int argc, char **argv) {
   });
 
   return 0;
-}
\ No newline at end of file
+}
diff --git a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp
index 220dfde50e507..63bdc3b226e6e 100644
--- a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp
+++ b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp
@@ -64,7 +64,7 @@ template <typename T>
 void setup_sycl_operation(const T VA[]) {
 
   deviceQueue.submit([&](sycl::handler &h) {
-    // expected-note@Inputs/sycl.hpp:221 {{called by 'kernel_single_task<AName, (lambda}}
+    // expected-note@#KernelSingleTaskKernelFuncCall {{called by 'kernel_single_task<AName, (lambda}}
     h.single_task<class AName>([]() {
       // ======= Zero Length Arrays Not Allowed in Kernel ==========
       // expected-error@+1 {{zero-length arrays are not permitted in C++}}
@@ -156,7 +156,7 @@ int main(int argc, char **argv) {
 
   // --- direct lambda testing ---
   deviceQueue.submit([&](sycl::handler &h) {
-    // expected-note@Inputs/sycl.hpp:221 2 {{called by 'kernel_single_task<AName, (lambda}}
+    // expected-note@#KernelSingleTaskKernelFuncCall 2 {{called by 'kernel_single_task<AName, (lambda}}
     h.single_task<class AName>([]() {
       // expected-error@+1 {{zero-length arrays are not permitted in C++}}
       int BadArray[0];
diff --git a/clang/test/SemaSYCL/float128.cpp b/clang/test/SemaSYCL/float128.cpp
index ff6b2f71b311d..da7be1573151d 100644
--- a/clang/test/SemaSYCL/float128.cpp
+++ b/clang/test/SemaSYCL/float128.cpp
@@ -71,7 +71,7 @@ int main() {
   __float128 CapturedToDevice = 1;
   host_ok();
   deviceQueue.submit([&](sycl::handler &h) {
-    // expected-note@Inputs/sycl.hpp:221 {{called by 'kernel_single_task<variables, (lambda}}
+    // expected-note@#KernelSingleTaskKernelFuncCall {{called by 'kernel_single_task<variables, (lambda}}
     h.single_task<class variables>([=]() {
       // expected-error@+1 {{'__float128' is not supported on this target}}
       decltype(CapturedToDevice) D;
@@ -88,7 +88,7 @@ int main() {
   });
 
   deviceQueue.submit([&](sycl::handler &h) {
-    // expected-note@Inputs/sycl.hpp:221 4{{called by 'kernel_single_task<functions, (lambda}}
+    // expected-note@#KernelSingleTaskKernelFuncCall 4{{called by 'kernel_single_task<functions, (lambda}}
     h.single_task<class functions>([=]() {
       // expected-note@+1 2{{called by 'operator()'}}
       usage();
@@ -104,7 +104,7 @@ int main() {
   });
 
   deviceQueue.submit([&](sycl::handler &h) {
-    // expected-note@Inputs/sycl.hpp:221 {{called by 'kernel_single_task<ok, (lambda}}
+    // expected-note@#KernelSingleTaskKernelFuncCall {{called by 'kernel_single_task<ok, (lambda}}
     h.single_task<class ok>([=]() {
       // expected-note@+1 3{{used here}}
       Z<__float128> S;
diff --git a/clang/test/SemaSYCL/implicit_kernel_type.cpp b/clang/test/SemaSYCL/implicit_kernel_type.cpp
index d178a4c2d1256..9f28a950affbe 100644
--- a/clang/test/SemaSYCL/implicit_kernel_type.cpp
+++ b/clang/test/SemaSYCL/implicit_kernel_type.cpp
@@ -25,12 +25,12 @@ int main() {
   queue q;
 
 #if defined(WARN)
-  // expected-error@Inputs/sycl.hpp:233 {{'InvalidKernelName1' is an invalid kernel name type}}
-  // expected-note@Inputs/sycl.hpp:233 {{'InvalidKernelName1' should be globally-visible}}
+  // expected-error@#KernelSingleTask {{'InvalidKernelName1' is an invalid kernel name type}}
+  // expected-note@#KernelSingleTask {{'InvalidKernelName1' should be globally-visible}}
   // expected-note@+8 {{in instantiation of function template specialization}}
 #elif defined(ERROR)
-  // expected-error@Inputs/sycl.hpp:233 {{'InvalidKernelName1' is an invalid kernel name type}}
-  // expected-note@Inputs/sycl.hpp:233 {{'InvalidKernelName1' should be globally-visible}}
+  // expected-error@#KernelSingleTask {{'InvalidKernelName1' is an invalid kernel name type}}
+  // expected-note@#KernelSingleTask {{'InvalidKernelName1' should be globally-visible}}
   // expected-note@+4 {{in instantiation of function template specialization}}
 #endif
   class InvalidKernelName1 {};
@@ -39,9 +39,9 @@ int main() {
   });
 
 #if defined(WARN)
-  // expected-warning@Inputs/sycl.hpp:233 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}}
+  // expected-warning@#KernelSingleTask {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}}
 #elif defined(ERROR)
-  // expected-error@Inputs/sycl.hpp:233 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}}
+  // expected-error@#KernelSingleTask {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}}
 #endif
 
   q.submit([&](handler &h) {
@@ -53,9 +53,9 @@ int main() {
   });
 
 #if defined(WARN)
-  // expected-warning@Inputs/sycl.hpp:233 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}}
+  // expected-warning@#KernelSingleTask {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}}
 #elif defined(ERROR)
-  // expected-error@Inputs/sycl.hpp:233 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}}
+  // expected-error@#KernelSingleTask {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}}
 #endif
 
   q.submit([&](handler &h) {
diff --git a/clang/test/SemaSYCL/int128.cpp b/clang/test/SemaSYCL/int128.cpp
index 196f2dfad66d4..a833ccd6e448f 100644
--- a/clang/test/SemaSYCL/int128.cpp
+++ b/clang/test/SemaSYCL/int128.cpp
@@ -80,7 +80,7 @@ int main() {
   __int128 CapturedToDevice = 1;
   host_ok();
   deviceQueue.submit([&](sycl::handler &h) {
-    // expected-note@Inputs/sycl.hpp:221 {{called by 'kernel_single_task<variables, (lambda}}
+    // expected-note@#KernelSingleTaskKernelFuncCall {{called by 'kernel_single_task<variables, (lambda}}
     h.single_task<class variables>([=]() {
       // expected-error@+1 {{'__int128' is not supported on this target}}
       decltype(CapturedToDevice) D;
@@ -97,7 +97,7 @@ int main() {
   });
 
   deviceQueue.submit([&](sycl::handler &h) {
-    // expected-note@Inputs/sycl.hpp:221 5{{called by 'kernel_single_task<functions, (lambda}}
+    // expected-note@#KernelSingleTaskKernelFuncCall 5{{called by 'kernel_single_task<functions, (lambda}}
     h.single_task<class functions>([=]() {
       // expected-note@+1 2{{called by 'operator()'}}
       usage();
@@ -115,7 +115,7 @@ int main() {
   });
 
   deviceQueue.submit([&](sycl::handler &h) {
-    // expected-note@Inputs/sycl.hpp:221 {{called by 'kernel_single_task<ok, (lambda}}
+    // expected-note@#KernelSingleTaskKernelFuncCall {{called by 'kernel_single_task<ok, (lambda}}
     h.single_task<class ok>([=]() {
       // expected-note@+1 3{{used here}}
       Z<__int128> S;
diff --git a/clang/test/SemaSYCL/kernel-handler.cpp b/clang/test/SemaSYCL/kernel-handler.cpp
new file mode 100644
index 0000000000000..a5df0e186e10b
--- /dev/null
+++ b/clang/test/SemaSYCL/kernel-handler.cpp
@@ -0,0 +1,132 @@
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple nvptx64-unknown-unknown -ast-dump %s | FileCheck %s --check-prefix=NONATIVESUPPORT
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -ast-dump %s | FileCheck %s --check-prefix=NATIVESUPPORT
+
+// This test checks that the compiler handles kernel_handler type (for
+// SYCL 2020 specialization constants) correctly.
+
+#include "sycl.hpp"
+
+using namespace cl::sycl;
+queue q;
+
+int main() {
+  q.submit([&](handler &h) {
+    int a;
+    kernel_handler kh;
+
+    h.single_task<class test_kernel_handler>(
+        [=](auto) {
+          int local = a;
+        },
+        kh);
+    h.parallel_for_work_group<class test_pfwg_kernel_handler>(
+        [=](group<1> G, kernel_handler kh) {
+          int local = a;
+        },
+        kh);
+  });
+}
+
+// Check test_kernel_handler parameters
+// NONATIVESUPPORT: FunctionDecl {{.*}}test_kernel_handler{{.*}} 'void (int, char *)'
+// NONATIVESUPPORT-NEXT: ParmVarDecl {{.*}} used _arg_ 'int'
+// NONATIVESUPPORT-NEXT: ParmVarDecl {{.*}} used  _arg__specialization_constants_buffer 'char *'
+
+// Check declaration and initialization of kernel object local clone
+// NONATIVESUPPORT-NEXT: CompoundStmt
+// NONATIVESUPPORT-NEXT: DeclStmt
+// NONATIVESUPPORT-NEXT: VarDecl {{.*}} cinit
+// NONATIVESUPPORT-NEXT: InitListExpr
+// NONATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
+// NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_' 'int'
+
+// Check declaration and initialization of kernel handler local clone using default constructor
+// NONATIVESUPPORT-NEXT: DeclStmt
+// NONATIVESUPPORT-NEXT: VarDecl {{.*}} callinit
+// NONATIVESUPPORT-NEXT: CXXConstructExpr {{.*}} 'sycl::kernel_handler' 'void () noexcept'
+
+// Check call to __init_specialization_constants_buffer
+// NONATIVESUPPORT-NEXT: CXXMemberCallExpr {{.*}} 'void'
+// NONATIVESUPPORT-NEXT: MemberExpr {{.*}} 'void (char *)' lvalue .__init_specialization_constants_buffer
+// NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'sycl::kernel_handler' lvalue Var {{.*}} 'kh'
+// NONATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'char *' <LValueToRValue>
+// NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'char *' lvalue ParmVar {{.*}} '_arg__specialization_constants_buffer' 'char *'
+// NONATIVESUPPORT-NEXT: CompoundStmt
+// NONATIVESUPPORT-NEXT: CXXOperatorCallExpr
+// NONATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'void (*)(sycl::kernel_handler) const' <FunctionToPointerDecay>
+// NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'void (sycl::kernel_handler) const' lvalue CXXMethod {{.*}} 'operator()' 'void (sycl::kernel_handler) const'
+// Kernel body with clones
+// NONATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'const (lambda at {{.*}}kernel-handler.cpp{{.*}})' lvalue
+// NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}kernel-handler.cpp{{.*}})' lvalue Var {{.*}} '(lambda at {{.*}}kernel-handler.cpp{{.*}})'
+// NONATIVESUPPORT-NEXT: CXXConstructExpr {{.*}} 'sycl::kernel_handler':'sycl::kernel_handler' 'void (const sycl::kernel_handler &) noexcept'
+// NONATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'const sycl::kernel_handler' lvalue
+// NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'sycl::kernel_handler' lvalue Var {{.*}} 'kh' 'sycl::kernel_handler'
+
+// Check test_pfwg_kernel_handler parameters
+// NONATIVESUPPORT: FunctionDecl {{.*}}test_pfwg_kernel_handler{{.*}} 'void (int, char *)'
+// NONATIVESUPPORT-NEXT: ParmVarDecl {{.*}} used _arg_ 'int'
+// NONATIVESUPPORT-NEXT: ParmVarDecl {{.*}} used  _arg__specialization_constants_buffer 'char *'
+
+// Check declaration and initialization of kernel object local clone
+// NONATIVESUPPORT-NEXT: CompoundStmt
+// NONATIVESUPPORT-NEXT: DeclStmt
+// NONATIVESUPPORT-NEXT: VarDecl {{.*}} cinit
+// NONATIVESUPPORT-NEXT: InitListExpr
+// NONATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
+// NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_' 'int'
+
+// Check declaration and initialization of kernel handler local clone using default constructor
+// NONATIVESUPPORT-NEXT: DeclStmt
+// NONATIVESUPPORT-NEXT: VarDecl {{.*}} callinit
+// NONATIVESUPPORT-NEXT: CXXConstructExpr {{.*}} 'sycl::kernel_handler' 'void () noexcept'
+
+// Check call to __init_specialization_constants_buffer
+// NONATIVESUPPORT-NEXT: CXXMemberCallExpr {{.*}} 'void'
+// NONATIVESUPPORT-NEXT: MemberExpr {{.*}} 'void (char *)' lvalue .__init_specialization_constants_buffer
+// NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'sycl::kernel_handler' lvalue Var {{.*}} 'kh'
+// NONATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'char *' <LValueToRValue>
+// NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'char *' lvalue ParmVar {{.*}} '_arg__specialization_constants_buffer' 'char *'
+// NONATIVESUPPORT-NEXT: CompoundStmt
+// NONATIVESUPPORT-NEXT: ExprWithCleanups
+// NONATIVESUPPORT-NEXT: CXXOperatorCallExpr
+// NONATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'void (*)(group<1>, sycl::kernel_handler) const' <FunctionToPointerDecay>
+// NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'void (group<1>, sycl::kernel_handler) const' lvalue CXXMethod {{.*}} 'operator()' 'void (group<1>, sycl::kernel_handler) const'
+
+// Kernel body with clones
+// NONATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'const (lambda at {{.*}}kernel-handler.cpp{{.*}})' lvalue
+// NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}kernel-handler.cpp{{.*}})' lvalue Var {{.*}} '(lambda at {{.*}}kernel-handler.cpp{{.*}})'
+// NONATIVESUPPORT-NEXT: CXXConstructExpr {{.*}} 'group<1>':'sycl::group<1>' 'void (sycl::group<1> &&) noexcept'
+// NONATIVESUPPORT-NEXT: MaterializeTemporaryExpr
+// NONATIVESUPPORT-NEXT: CXXTemporaryObjectExpr
+// NONATIVESUPPORT-NEXT: CXXConstructExpr {{.*}} 'sycl::kernel_handler' 'void (const sycl::kernel_handler &) noexcept'
+// NONATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'const sycl::kernel_handler' lvalue
+// NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'sycl::kernel_handler' lvalue Var {{.*}} 'kh' 'sycl::kernel_handler'
+
+// Test AST for default SPIR architecture
+
+// Check test_kernel_handler parameters
+// NATIVESUPPORT: FunctionDecl {{.*}}test_kernel_handler{{.*}} 'void (int)'
+// NATIVESUPPORT-NEXT: ParmVarDecl {{.*}} used _arg_ 'int'
+
+// Check declaration and initialization of kernel object local clone
+// NATIVESUPPORT-NEXT: CompoundStmt
+// NATIVESUPPORT-NEXT: DeclStmt
+// NATIVESUPPORT-NEXT: VarDecl {{.*}} cinit
+// NATIVESUPPORT-NEXT: InitListExpr
+// NATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
+// NATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_' 'int'
+
+// Check declaration and initialization of kernel handler local clone using default constructor
+// NATIVESUPPORT-NEXT: DeclStmt
+// NATIVESUPPORT-NEXT: VarDecl {{.*}} callinit
+// NATIVESUPPORT-NEXT: CXXConstructExpr {{.*}} 'sycl::kernel_handler' 'void () noexcept'
+
+// Check no call to __init_specialization_constants_buffer
+// NATIVESUPPORT-NOT: MemberExpr {{.*}} 'void (char *)' lvalue .__init_specialization_constants_buffer
+
+// Kernel body with clones
+// NATIVESUPPORT: ImplicitCastExpr {{.*}} 'const (lambda at {{.*}}kernel-handler.cpp{{.*}})' lvalue
+// NATIVESUPPORT-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}kernel-handler.cpp{{.*}})' lvalue Var {{.*}} '(lambda at {{.*}}kernel-handler.cpp{{.*}})'
+// NATIVESUPPORT-NEXT: CXXConstructExpr {{.*}} 'sycl::kernel_handler':'sycl::kernel_handler' 'void (const sycl::kernel_handler &) noexcept'
+// NATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'const sycl::kernel_handler' lvalue
+// NATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'sycl::kernel_handler' lvalue Var {{.*}} 'kh' 'sycl::kernel_handler'
diff --git a/clang/test/SemaSYCL/kernelname-enum.cpp b/clang/test/SemaSYCL/kernelname-enum.cpp
index 62e54d61f51c8..b4f224729db8a 100644
--- a/clang/test/SemaSYCL/kernelname-enum.cpp
+++ b/clang/test/SemaSYCL/kernelname-enum.cpp
@@ -67,15 +67,15 @@ int main() {
   });
 
   q.submit([&](cl::sycl::handler &cgh) {
-    // expected-error@Inputs/sycl.hpp:233 {{'dummy_functor_2<val_3>' is an invalid kernel name type}}
-    // expected-note@Inputs/sycl.hpp:233 {{unscoped enum 'unscoped_enum_no_type_set' requires fixed underlying type}}
+    // expected-error@#KernelSingleTask {{'dummy_functor_2<val_3>' is an invalid kernel name type}}
+    // expected-note@#KernelSingleTask {{unscoped enum 'unscoped_enum_no_type_set' requires fixed underlying type}}
     // expected-note@+1{{in instantiation of function template specialization}}
     cgh.single_task(f2);
   });
 
   q.submit([&](cl::sycl::handler &cgh) {
-    // expected-error@Inputs/sycl.hpp:233 {{'templated_functor<dummy_functor_2>' is an invalid kernel name type}}
-    // expected-note@Inputs/sycl.hpp:233 {{unscoped enum 'unscoped_enum_no_type_set' requires fixed underlying type}}
+    // expected-error@#KernelSingleTask {{'templated_functor<dummy_functor_2>' is an invalid kernel name type}}
+    // expected-note@#KernelSingleTask {{unscoped enum 'unscoped_enum_no_type_set' requires fixed underlying type}}
     // expected-note@+1{{in instantiation of function template specialization}}
     cgh.single_task(f5);
   });
diff --git a/clang/test/SemaSYCL/stdtypes_kernel_type.cpp b/clang/test/SemaSYCL/stdtypes_kernel_type.cpp
index 1209c4083ace0..aef5333dbc176 100644
--- a/clang/test/SemaSYCL/stdtypes_kernel_type.cpp
+++ b/clang/test/SemaSYCL/stdtypes_kernel_type.cpp
@@ -25,38 +25,38 @@ queue q;
 int main() {
 #ifdef CHECK_ERROR
   q.submit([&](handler &h) {
-    // expected-error@Inputs/sycl.hpp:233 {{'nullptr_t' is an invalid kernel name type}}
-    // expected-note@Inputs/sycl.hpp:233 {{type 'nullptr_t' cannot be in the "std" namespace}}
+    // expected-error@#KernelSingleTask {{'nullptr_t' is an invalid kernel name type}}
+    // expected-note@#KernelSingleTask {{type 'nullptr_t' cannot be in the "std" namespace}}
     // expected-note@+1{{in instantiation of function template specialization}}
     h.single_task<std::nullptr_t>([=] {});
   });
   q.submit([&](handler &h) {
-    // expected-error@Inputs/sycl.hpp:233 {{'std::T' is an invalid kernel name type}}
-    // expected-note@Inputs/sycl.hpp:233 {{type 'std::T' cannot be in the "std" namespace}}
+    // expected-error@#KernelSingleTask {{'std::T' is an invalid kernel name type}}
+    // expected-note@#KernelSingleTask {{type 'std::T' cannot be in the "std" namespace}}
     // expected-note@+1{{in instantiation of function template specialization}}
     h.single_task<std::T>([=] {});
   });
   q.submit([&](handler &h) {
-    // expected-error@Inputs/sycl.hpp:233 {{'Templated_kernel_name<nullptr_t>' is an invalid kernel name type}}
-    // expected-note@Inputs/sycl.hpp:233 {{type 'nullptr_t' cannot be in the "std" namespace}}
+    // expected-error@#KernelSingleTask {{'Templated_kernel_name<nullptr_t>' is an invalid kernel name type}}
+    // expected-note@#KernelSingleTask {{type 'nullptr_t' cannot be in the "std" namespace}}
     // expected-note@+1{{in instantiation of function template specialization}}
     h.single_task<Templated_kernel_name<std::nullptr_t>>([=] {});
   });
   q.submit([&](handler &h) {
-    // expected-error@Inputs/sycl.hpp:233 {{'Templated_kernel_name<std::U>' is an invalid kernel name type}}
-    // expected-note@Inputs/sycl.hpp:233 {{type 'std::U' cannot be in the "std" namespace}}
+    // expected-error@#KernelSingleTask {{'Templated_kernel_name<std::U>' is an invalid kernel name type}}
+    // expected-note@#KernelSingleTask {{type 'std::U' cannot be in the "std" namespace}}
     // expected-note@+1{{in instantiation of function template specialization}}
     h.single_task<Templated_kernel_name<std::U>>([=] {});
   });
   q.submit([&](handler &cgh) {
-    // expected-error@Inputs/sycl.hpp:233 {{'Templated_kernel_name2<Templated_kernel_name<std::Foo>>' is an invalid kernel name type}}
-    // expected-note@Inputs/sycl.hpp:233{{type 'std::Foo' cannot be in the "std" namespace}}
+    // expected-error@#KernelSingleTask {{'Templated_kernel_name2<Templated_kernel_name<std::Foo>>' is an invalid kernel name type}}
+    // expected-note@#KernelSingleTask{{type 'std::Foo' cannot be in the "std" namespace}}
     // expected-note@+1{{in instantiation of function template specialization}}
     cgh.single_task<Templated_kernel_name2<Templated_kernel_name<std::Foo>>>([]() {});
   });
   q.submit([&](handler &cgh) {
-    // expected-error@Inputs/sycl.hpp:233 {{'TemplParamPack<int, float, nullptr_t, double>' is an invalid kernel name type}}
-    // expected-note@Inputs/sycl.hpp:233 {{type 'nullptr_t' cannot be in the "std" namespace}}
+    // expected-error@#KernelSingleTask {{'TemplParamPack<int, float, nullptr_t, double>' is an invalid kernel name type}}
+    // expected-note@#KernelSingleTask {{type 'nullptr_t' cannot be in the "std" namespace}}
     // expected-note@+1{{in instantiation of function template specialization}}
     cgh.single_task<TemplParamPack<int, float, std::nullptr_t, double>>([]() {});
   });
diff --git a/clang/test/SemaSYCL/unnamed-kernel.cpp b/clang/test/SemaSYCL/unnamed-kernel.cpp
index bb3dc914ab345..d728f60dbc72a 100644
--- a/clang/test/SemaSYCL/unnamed-kernel.cpp
+++ b/clang/test/SemaSYCL/unnamed-kernel.cpp
@@ -33,8 +33,8 @@ struct MyWrapper {
   void test() {
     cl::sycl::queue q;
 #ifndef __SYCL_UNNAMED_LAMBDA__
-    // expected-error@Inputs/sycl.hpp:233 {{'InvalidKernelName1' is an invalid kernel name type}}
-    // expected-note@Inputs/sycl.hpp:233 {{'InvalidKernelName1' should be globally-visible}}
+    // expected-error@#KernelSingleTask {{'InvalidKernelName1' is an invalid kernel name type}}
+    // expected-note@#KernelSingleTask {{'InvalidKernelName1' should be globally-visible}}
     // expected-note@+4{{in instantiation of function template specialization}}
 #endif
     class InvalidKernelName1 {};
@@ -43,8 +43,8 @@ struct MyWrapper {
     });
 
 #ifndef __SYCL_UNNAMED_LAMBDA__
-    // expected-error@Inputs/sycl.hpp:233 {{'namespace1::KernelName<InvalidKernelName2>' is an invalid kernel name type}}
-    // expected-note@Inputs/sycl.hpp:233 {{'InvalidKernelName2' should be globally-visible}}
+    // expected-error@#KernelSingleTask {{'namespace1::KernelName<InvalidKernelName2>' is an invalid kernel name type}}
+    // expected-note@#KernelSingleTask {{'InvalidKernelName2' should be globally-visible}}
     // expected-note@+4{{in instantiation of function template specialization}}
 #endif
     class InvalidKernelName2 {};
@@ -53,8 +53,8 @@ struct MyWrapper {
     });
 
 #ifndef __SYCL_UNNAMED_LAMBDA__
-    // expected-error@Inputs/sycl.hpp:233 {{'MyWrapper::InvalidKernelName0' is an invalid kernel name type}}
-    // expected-note@Inputs/sycl.hpp:233 {{'MyWrapper::InvalidKernelName0' should be globally-visible}}
+    // expected-error@#KernelSingleTask {{'MyWrapper::InvalidKernelName0' is an invalid kernel name type}}
+    // expected-note@#KernelSingleTask {{'MyWrapper::InvalidKernelName0' should be globally-visible}}
     // expected-note@+3{{in instantiation of function template specialization}}
 #endif
     q.submit([&](cl::sycl::handler &h) {
@@ -62,8 +62,8 @@ struct MyWrapper {
     });
 
 #ifndef __SYCL_UNNAMED_LAMBDA__
-    // expected-error@Inputs/sycl.hpp:233 {{'namespace1::KernelName<MyWrapper::InvalidKernelName3>' is an invalid kernel name type}}
-    // expected-note@Inputs/sycl.hpp:233 {{'MyWrapper::InvalidKernelName3' should be globally-visible}}
+    // expected-error@#KernelSingleTask {{'namespace1::KernelName<MyWrapper::InvalidKernelName3>' is an invalid kernel name type}}
+    // expected-note@#KernelSingleTask {{'MyWrapper::InvalidKernelName3' should be globally-visible}}
     // expected-note@+3{{in instantiation of function template specialization}}
 #endif
     q.submit([&](cl::sycl::handler &h) {
@@ -76,8 +76,8 @@ struct MyWrapper {
     });
 
 #ifndef __SYCL_UNNAMED_LAMBDA__
-    // expected-error@Inputs/sycl.hpp:233 {{'std::max_align_t' is an invalid kernel name type}}
-    // expected-note@Inputs/sycl.hpp:233 {{type 'std::max_align_t' cannot be in the "std" namespace}}
+    // expected-error@#KernelSingleTask {{'std::max_align_t' is an invalid kernel name type}}
+    // expected-note@#KernelSingleTask {{type 'std::max_align_t' cannot be in the "std" namespace}}
     // expected-note@+3{{in instantiation of function template specialization}}
 #endif
     q.submit([&](cl::sycl::handler &h) {
@@ -86,8 +86,8 @@ struct MyWrapper {
 
     using InvalidAlias = InvalidKernelName4;
 #ifndef __SYCL_UNNAMED_LAMBDA__
-    // expected-error@Inputs/sycl.hpp:233 {{'MyWrapper::InvalidKernelName4' is an invalid kernel name type}}
-    // expected-note@Inputs/sycl.hpp:233 {{'MyWrapper::InvalidKernelName4' should be globally-visible}}
+    // expected-error@#KernelSingleTask {{'MyWrapper::InvalidKernelName4' is an invalid kernel name type}}
+    // expected-note@#KernelSingleTask {{'MyWrapper::InvalidKernelName4' should be globally-visible}}
     // expected-note@+3{{in instantiation of function template specialization}}
 #endif
     q.submit([&](cl::sycl::handler &h) {
@@ -96,16 +96,16 @@ struct MyWrapper {
 
     using InvalidAlias1 = InvalidKernelName5;
 #ifndef __SYCL_UNNAMED_LAMBDA__
-    // expected-error@Inputs/sycl.hpp:233 {{'namespace1::KernelName<MyWrapper::InvalidKernelName5>' is an invalid kernel name type}}
-    // expected-note@Inputs/sycl.hpp:233 {{'MyWrapper::InvalidKernelName5' should be globally-visible}}
+    // expected-error@#KernelSingleTask {{'namespace1::KernelName<MyWrapper::InvalidKernelName5>' is an invalid kernel name type}}
+    // expected-note@#KernelSingleTask {{'MyWrapper::InvalidKernelName5' should be globally-visible}}
     // expected-note@+3{{in instantiation of function template specialization}}
 #endif
     q.submit([&](cl::sycl::handler &h) {
       h.single_task<namespace1::KernelName<InvalidAlias1>>([] {});
     });
 #ifndef __SYCL_UNNAMED_LAMBDA__
-    // expected-error@Inputs/sycl.hpp:233 {{'Templated_kernel_name2<Templated_kernel_name<InvalidKernelName1>>' is an invalid kernel name type}}
-    // expected-note@Inputs/sycl.hpp:233 {{'InvalidKernelName1' should be globally-visible}}
+    // expected-error@#KernelSingleTask {{'Templated_kernel_name2<Templated_kernel_name<InvalidKernelName1>>' is an invalid kernel name type}}
+    // expected-note@#KernelSingleTask {{'InvalidKernelName1' should be globally-visible}}
     // expected-note@+3{{in instantiation of function template specialization}}
 #endif
     q.submit([&](cl::sycl::handler &h) {
@@ -117,8 +117,8 @@ struct MyWrapper {
 int main() {
   cl::sycl::queue q;
 #ifndef __SYCL_UNNAMED_LAMBDA__
-  // expected-error-re@Inputs/sycl.hpp:233 {{'(lambda at {{.*}}unnamed-kernel.cpp{{.*}}' is an invalid kernel name type}}
-  // expected-note@Inputs/sycl.hpp:233 {{unnamed type used in a SYCL kernel name}}
+  // expected-error-re@#KernelSingleTask {{'(lambda at {{.*}}unnamed-kernel.cpp{{.*}}' is an invalid kernel name type}}
+  // expected-note@#KernelSingleTask {{unnamed type used in a SYCL kernel name}}
   // expected-note@+2{{in instantiation of function template specialization}}
 #endif
   q.submit([&](cl::sycl::handler &h) { h.single_task([] {}); });