diff --git a/clang/include/clang/AST/Decl.h b/clang/include/clang/AST/Decl.h
index 7db74e0803cec..9940636192f1d 100644
--- a/clang/include/clang/AST/Decl.h
+++ b/clang/include/clang/AST/Decl.h
@@ -283,15 +283,16 @@ class NamedDecl : public Decl {
   /// Creating this name is expensive, so it should be called only when
   /// performance doesn't matter.
   void printQualifiedName(raw_ostream &OS) const;
-  void printQualifiedName(raw_ostream &OS, const PrintingPolicy &Policy) const;
+  void printQualifiedName(raw_ostream &OS, const PrintingPolicy &Policy,
+                          bool WithGlobalNsPrefix = false) const;
 
   /// Print only the nested name specifier part of a fully-qualified name,
   /// including the '::' at the end. E.g.
   ///    when `printQualifiedName(D)` prints "A::B::i",
   ///    this function prints "A::B::".
   void printNestedNameSpecifier(raw_ostream &OS) const;
-  void printNestedNameSpecifier(raw_ostream &OS,
-                                const PrintingPolicy &Policy) const;
+  void printNestedNameSpecifier(raw_ostream &OS, const PrintingPolicy &Policy,
+                                bool WithGlobalNsPrefix = false) const;
 
   // FIXME: Remove string version.
   std::string getQualifiedNameAsString() const;
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index f4f20dfe13966..9afff5381820d 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -10761,7 +10761,9 @@ def err_builtin_launder_invalid_arg : Error<
 // SYCL-specific diagnostics
 def err_sycl_kernel_incorrectly_named : Error<
   "kernel %select{name is missing"
-  "|needs to have a globally-visible name}0">;
+  "|needs to have a globally-visible name"
+  "|name is invalid. Unscoped enum requires fixed underlying type"
+  "}0">;
 def err_sycl_restrict : Error<
   "SYCL kernel cannot "
 "%select{use a non-const global variable"
diff --git a/clang/lib/AST/Decl.cpp b/clang/lib/AST/Decl.cpp
index 9c1b99d30e788..e4961c964ff47 100644
--- a/clang/lib/AST/Decl.cpp
+++ b/clang/lib/AST/Decl.cpp
@@ -1538,14 +1538,14 @@ void NamedDecl::printQualifiedName(raw_ostream &OS) const {
   printQualifiedName(OS, getASTContext().getPrintingPolicy());
 }
 
-void NamedDecl::printQualifiedName(raw_ostream &OS,
-                                   const PrintingPolicy &P) const {
+void NamedDecl::printQualifiedName(raw_ostream &OS, const PrintingPolicy &P,
+                                   bool WithGlobalNsPrefix) const {
   if (getDeclContext()->isFunctionOrMethod()) {
     // We do not print '(anonymous)' for function parameters without name.
     printName(OS);
     return;
   }
-  printNestedNameSpecifier(OS, P);
+  printNestedNameSpecifier(OS, P, WithGlobalNsPrefix);
   if (getDeclName())
     OS << *this;
   else {
@@ -1566,7 +1566,8 @@ void NamedDecl::printNestedNameSpecifier(raw_ostream &OS) const {
 }
 
 void NamedDecl::printNestedNameSpecifier(raw_ostream &OS,
-                                         const PrintingPolicy &P) const {
+                                         const PrintingPolicy &P,
+                                         bool WithGlobalNsPrefix) const {
   const DeclContext *Ctx = getDeclContext();
 
   // For ObjC methods and properties, look through categories and use the
@@ -1593,6 +1594,9 @@ void NamedDecl::printNestedNameSpecifier(raw_ostream &OS,
     Ctx = Ctx->getParent();
   }
 
+  if (WithGlobalNsPrefix)
+    OS << "::";
+
   for (const DeclContext *DC : llvm::reverse(Contexts)) {
     if (const auto *Spec = dyn_cast<ClassTemplateSpecializationDecl>(DC)) {
       OS << Spec->getName();
@@ -1605,8 +1609,7 @@ void NamedDecl::printNestedNameSpecifier(raw_ostream &OS,
       if (ND->isAnonymousNamespace()) {
         OS << (P.MSVCFormatting ? "`anonymous namespace\'"
                                 : "(anonymous namespace)");
-      }
-      else
+      } else
         OS << *ND;
     } else if (const auto *RD = dyn_cast<RecordDecl>(DC)) {
       if (!RD->getIdentifier())
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 30537b847b015..3258278234ddc 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -1637,6 +1637,18 @@ static std::string eraseAnonNamespace(std::string S) {
   return S;
 }
 
+static bool checkEnumTemplateParameter(const EnumDecl *ED,
+                                       DiagnosticsEngine &Diag,
+                                       SourceLocation KernelLocation) {
+  if (!ED->isScoped() && !ED->isFixed()) {
+    Diag.Report(KernelLocation, diag::err_sycl_kernel_incorrectly_named) << 2;
+    Diag.Report(ED->getSourceRange().getBegin(), diag::note_entity_declared_at)
+        << ED;
+    return true;
+  }
+  return false;
+}
+
 // Emits a forward declaration
 void SYCLIntegrationHeader::emitFwdDecl(raw_ostream &O, const Decl *D,
                                         SourceLocation KernelLocation) {
@@ -1690,10 +1702,22 @@ void SYCLIntegrationHeader::emitFwdDecl(raw_ostream &O, const Decl *D,
   PrintingPolicy P(D->getASTContext().getLangOpts());
   P.adjustForCPlusPlusFwdDecl();
   P.SuppressTypedefs = true;
+  P.SuppressUnwrittenScope = true;
   std::string S;
   llvm::raw_string_ostream SO(S);
   D->print(SO, P);
-  O << SO.str() << ";\n";
+  O << SO.str();
+
+  if (const auto *ED = dyn_cast<EnumDecl>(D)) {
+    QualType T = ED->getIntegerType();
+    // Backup since getIntegerType() returns null for enum forward
+    // declaration with no fixed underlying type
+    if (T.isNull())
+      T = ED->getPromotionType();
+    O << " : " << T.getAsString();
+  }
+
+  O << ";\n";
 
   // print closing braces for namespaces if needed
   for (unsigned I = 0; I < NamespaceCnt; ++I)
@@ -1762,8 +1786,20 @@ void SYCLIntegrationHeader::emitForwardClassDecls(
 
       switch (Arg.getKind()) {
       case TemplateArgument::ArgKind::Type:
-        emitForwardClassDecls(O, Arg.getAsType(), KernelLocation, Printed);
+      case TemplateArgument::ArgKind::Integral: {
+        QualType T = (Arg.getKind() == TemplateArgument::ArgKind::Type)
+                         ? Arg.getAsType()
+                         : Arg.getIntegralType();
+
+        // Handle Kernel Name Type templated using enum type and value.
+        if (const auto *ET = T->getAs<EnumType>()) {
+          const EnumDecl *ED = ET->getDecl();
+          if (!checkEnumTemplateParameter(ED, Diag, KernelLocation))
+            emitFwdDecl(O, ED, KernelLocation);
+        } else if (Arg.getKind() == TemplateArgument::ArgKind::Type)
+          emitForwardClassDecls(O, T, KernelLocation, Printed);
         break;
+      }
       case TemplateArgument::ArgKind::Pack: {
         ArrayRef<TemplateArgument> Pack = Arg.getPackAsArray();
 
@@ -1822,6 +1858,97 @@ static std::string getCPPTypeString(QualType Ty) {
   return eraseAnonNamespace(Ty.getAsString(P));
 }
 
+static void printArguments(ASTContext &Ctx, raw_ostream &ArgOS,
+                           ArrayRef<TemplateArgument> Args,
+                           const PrintingPolicy &P);
+
+static void printArgument(ASTContext &Ctx, raw_ostream &ArgOS,
+                          TemplateArgument Arg, const PrintingPolicy &P) {
+  switch (Arg.getKind()) {
+  case TemplateArgument::ArgKind::Pack: {
+    printArguments(Ctx, ArgOS, Arg.getPackAsArray(), P);
+    break;
+  }
+  case TemplateArgument::ArgKind::Integral: {
+    QualType T = Arg.getIntegralType();
+    const EnumType *ET = T->getAs<EnumType>();
+
+    if (ET) {
+      const llvm::APSInt &Val = Arg.getAsIntegral();
+      ArgOS << "(" << ET->getDecl()->getQualifiedNameAsString() << ")" << Val;
+    } else {
+      Arg.print(P, ArgOS);
+    }
+    break;
+  }
+  case TemplateArgument::ArgKind::Type: {
+    LangOptions LO;
+    PrintingPolicy TypePolicy(LO);
+    TypePolicy.SuppressTypedefs = true;
+    TypePolicy.SuppressTagKeyword = true;
+    QualType T = Arg.getAsType();
+    QualType FullyQualifiedType = TypeName::getFullyQualifiedType(T, Ctx, true);
+    ArgOS << FullyQualifiedType.getAsString(TypePolicy);
+    break;
+  }
+  default:
+    Arg.print(P, ArgOS);
+  }
+}
+
+static void printArguments(ASTContext &Ctx, raw_ostream &ArgOS,
+                           ArrayRef<TemplateArgument> Args,
+                           const PrintingPolicy &P) {
+  for (unsigned I = 0; I < Args.size(); I++) {
+    const TemplateArgument &Arg = Args[I];
+
+    if (I != 0)
+      ArgOS << ", ";
+
+    printArgument(Ctx, ArgOS, Arg, P);
+  }
+}
+
+static void printTemplateArguments(ASTContext &Ctx, raw_ostream &ArgOS,
+                                   ArrayRef<TemplateArgument> Args,
+                                   const PrintingPolicy &P) {
+  ArgOS << "<";
+  printArguments(Ctx, ArgOS, Args, P);
+  ArgOS << ">";
+}
+
+static std::string getKernelNameTypeString(QualType T) {
+
+  const CXXRecordDecl *RD = T->getAsCXXRecordDecl();
+
+  if (!RD)
+    return getCPPTypeString(T);
+
+  // If kernel name type is a template specialization with enum type
+  // template parameters, enumerators in name type string should be
+  // replaced  with their underlying value since the enum definition
+  // is not visible in integration header.
+  if (const auto *TSD = dyn_cast<ClassTemplateSpecializationDecl>(RD)) {
+    LangOptions LO;
+    PrintingPolicy P(LO);
+    P.SuppressTypedefs = true;
+    SmallString<64> Buf;
+    llvm::raw_svector_ostream ArgOS(Buf);
+
+    // Print template class name
+    TSD->printQualifiedName(ArgOS, P, /*WithGlobalNsPrefix*/ true);
+
+    // Print template arguments substituting enumerators
+    ASTContext &Ctx = RD->getASTContext();
+    const TemplateArgumentList &Args = TSD->getTemplateArgs();
+    printTemplateArguments(Ctx, ArgOS, Args.asArray(), P);
+
+    return eraseAnonNamespace(ArgOS.str().str());
+  }
+
+  return getCPPTypeString(T);
+}
+
 void SYCLIntegrationHeader::emit(raw_ostream &O) {
   O << "// This is auto-generated SYCL integration header.\n";
   O << "\n";
@@ -1938,8 +2065,9 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
         O << "', '" << c;
       O << "'> {\n";
     } else {
-      O << "template <> struct KernelInfo<" << getCPPTypeString(K.NameType)
-        << "> {\n";
+
+      O << "template <> struct KernelInfo<"
+        << getKernelNameTypeString(K.NameType) << "> {\n";
     }
     O << "  DLL_LOCAL\n";
     O << "  static constexpr const char* getName() { return \"" << K.Name
diff --git a/clang/test/CodeGenSYCL/int_header1.cpp b/clang/test/CodeGenSYCL/int_header1.cpp
index 646373e7f6adc..a2e6288a7875c 100644
--- a/clang/test/CodeGenSYCL/int_header1.cpp
+++ b/clang/test/CodeGenSYCL/int_header1.cpp
@@ -4,14 +4,14 @@
 // CHECK:template <> struct KernelInfo<class KernelName> {
 // CHECK:template <> struct KernelInfo<::nm1::nm2::KernelName0> {
 // CHECK:template <> struct KernelInfo<::nm1::KernelName1> {
-// CHECK:template <> struct KernelInfo<::nm1::KernelName3< ::nm1::nm2::KernelName0>> {
-// CHECK:template <> struct KernelInfo<::nm1::KernelName3< ::nm1::KernelName1>> {
-// CHECK:template <> struct KernelInfo<::nm1::KernelName4< ::nm1::nm2::KernelName0>> {
-// CHECK:template <> struct KernelInfo<::nm1::KernelName4< ::nm1::KernelName1>> {
+// CHECK:template <> struct KernelInfo<::nm1::KernelName3<::nm1::nm2::KernelName0>> {
+// CHECK:template <> struct KernelInfo<::nm1::KernelName3<::nm1::KernelName1>> {
+// CHECK:template <> struct KernelInfo<::nm1::KernelName4<::nm1::nm2::KernelName0>> {
+// CHECK:template <> struct KernelInfo<::nm1::KernelName4<::nm1::KernelName1>> {
 // CHECK:template <> struct KernelInfo<::nm1::KernelName3<KernelName5>> {
 // CHECK:template <> struct KernelInfo<::nm1::KernelName4<KernelName7>> {
-// CHECK:template <> struct KernelInfo<::nm1::KernelName8< ::nm1::nm2::C>> {
-// CHECK:template <> struct KernelInfo<class TmplClassInAnonNS<class ClassInAnonNS>> {
+// CHECK:template <> struct KernelInfo<::nm1::KernelName8<::nm1::nm2::C>> {
+// CHECK:template <> struct KernelInfo<::TmplClassInAnonNS<ClassInAnonNS>> {
 
 // This test checks if the SYCL device compiler is able to generate correct
 // integration header when the kernel name class is expressed in different
diff --git a/clang/test/CodeGenSYCL/integration_header.cpp b/clang/test/CodeGenSYCL/integration_header.cpp
index 84b35578f48e6..58d0c3addcd8c 100644
--- a/clang/test/CodeGenSYCL/integration_header.cpp
+++ b/clang/test/CodeGenSYCL/integration_header.cpp
@@ -57,8 +57,8 @@
 //
 // CHECK: template <> struct KernelInfo<class first_kernel> {
 // CHECK: template <> struct KernelInfo<::second_namespace::second_kernel<char>> {
-// CHECK: template <> struct KernelInfo<::third_kernel<1, int, ::point<X> >> {
-// CHECK: template <> struct KernelInfo<::fourth_kernel< ::template_arg_ns::namespaced_arg<1> >> {
+// CHECK: template <> struct KernelInfo<::third_kernel<1, int, ::point<X>>> {
+// CHECK: template <> struct KernelInfo<::fourth_kernel<::template_arg_ns::namespaced_arg<1>>> {
 
 #include "sycl.hpp"
 
diff --git a/clang/test/CodeGenSYCL/kernel_name_with_typedefs.cpp b/clang/test/CodeGenSYCL/kernel_name_with_typedefs.cpp
index c124f95770b8f..2ad4c70020281 100644
--- a/clang/test/CodeGenSYCL/kernel_name_with_typedefs.cpp
+++ b/clang/test/CodeGenSYCL/kernel_name_with_typedefs.cpp
@@ -117,16 +117,16 @@ int main() {
   single_task<kernel_name2<const space::long_t, space::a_t>>(f);
   // CHECK: template <> struct KernelInfo<::kernel_name2<const volatile long, const ::space::B>> {
   single_task<kernel_name2<volatile space::clong_t, const space::b_t>>(f);
-  // CHECK: template <> struct KernelInfo<::kernel_name2< ::A, long>> {
+  // CHECK: template <> struct KernelInfo<::kernel_name2<::A, long>> {
   single_task<kernel_name2<space::a_t, space::long_t>>(f);
-  // CHECK: template <> struct KernelInfo<::kernel_name2< ::space::B, int>> {
+  // CHECK: template <> struct KernelInfo<::kernel_name2<::space::B, int>> {
   single_task<kernel_name2<space::b_t, int_t>>(f);
   // full template specialization
   // CHECK: template <> struct KernelInfo<::kernel_name2<int, const unsigned int>> {
   single_task<kernel_name2<int_t, const uint_t>>(f);
-  // CHECK: template <> struct KernelInfo<::kernel_name2<const long, volatile const unsigned long>> {
+  // CHECK: template <> struct KernelInfo<::kernel_name2<const long, const volatile unsigned long>> {
   single_task<kernel_name2<space::clong_t, volatile space::culong_t>>(f);
-  // CHECK: template <> struct KernelInfo<::kernel_name2< ::A, volatile ::space::B>> {
+  // CHECK: template <> struct KernelInfo<::kernel_name2<::A, volatile ::space::B>> {
   single_task<kernel_name2<space::a_t, volatile space::b_t>>(f);
   // CHECK: template <> struct KernelInfo<::kernel_name3<1>> {
   single_task<kernel_name3<1>>(f);
diff --git a/clang/test/CodeGenSYCL/kernelname-enum.cpp b/clang/test/CodeGenSYCL/kernelname-enum.cpp
new file mode 100644
index 0000000000000..8aa170b15a566
--- /dev/null
+++ b/clang/test/CodeGenSYCL/kernelname-enum.cpp
@@ -0,0 +1,157 @@
+// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -fsycl-int-header=%t.h %s -fsyntax-only
+// RUN: FileCheck -input-file=%t.h %s
+
+#include "sycl.hpp"
+
+enum unscoped_enum : int {
+  val_1,
+  val_2
+};
+
+enum class no_namespace_int : int {
+  val_1,
+  val_2
+};
+
+enum class no_namespace_short : short {
+  val_1,
+  val_2
+};
+
+namespace internal {
+enum class namespace_short : short {
+  val_1,
+  val_2
+};
+}
+
+namespace {
+enum class enum_in_anonNS : short {
+  val_1,
+  val_2
+};
+}
+
+enum class no_type_set {
+  val_1,
+  val_2
+};
+
+template <no_namespace_int EnumType>
+class dummy_functor_1 {
+public:
+  void operator()() {}
+};
+
+template <no_namespace_short EnumType>
+class dummy_functor_2 {
+public:
+  void operator()() {}
+};
+
+template <internal::namespace_short EnumType>
+class dummy_functor_3 {
+public:
+  void operator()() {}
+};
+
+template <enum_in_anonNS EnumType>
+class dummy_functor_4 {
+public:
+  void operator()() {}
+};
+
+template <no_type_set EnumType>
+class dummy_functor_5 {
+public:
+  void operator()() {}
+};
+
+template <unscoped_enum EnumType>
+class dummy_functor_6 {
+public:
+  void operator()() {}
+};
+
+template <typename EnumType>
+class dummy_functor_7 {
+public:
+  void operator()() {}
+};
+
+int main() {
+
+  dummy_functor_1<no_namespace_int::val_1> f1;
+  dummy_functor_2<no_namespace_short::val_2> f2;
+  dummy_functor_3<internal::namespace_short::val_2> f3;
+  dummy_functor_4<enum_in_anonNS::val_2> f4;
+  dummy_functor_5<no_type_set::val_1> f5;
+  dummy_functor_6<unscoped_enum::val_1> f6;
+  dummy_functor_7<no_namespace_int> f7;
+  dummy_functor_7<internal::namespace_short> f8;
+
+  cl::sycl::queue q;
+
+  q.submit([&](cl::sycl::handler &cgh) {
+    cgh.single_task(f1);
+  });
+
+  q.submit([&](cl::sycl::handler &cgh) {
+    cgh.single_task(f2);
+  });
+
+  q.submit([&](cl::sycl::handler &cgh) {
+    cgh.single_task(f3);
+  });
+
+  q.submit([&](cl::sycl::handler &cgh) {
+    cgh.single_task(f4);
+  });
+
+  q.submit([&](cl::sycl::handler &cgh) {
+    cgh.single_task(f5);
+  });
+
+  q.submit([&](cl::sycl::handler &cgh) {
+    cgh.single_task(f6);
+  });
+
+  q.submit([&](cl::sycl::handler &cgh) {
+    cgh.single_task(f7);
+  });
+
+  q.submit([&](cl::sycl::handler &cgh) {
+    cgh.single_task(f8);
+  });
+
+  return 0;
+}
+
+// CHECK: Forward declarations of templated kernel function types:
+// CHECK: enum class no_namespace_int : int;
+// CHECK: template <no_namespace_int EnumType> class dummy_functor_1;
+// CHECK: enum class no_namespace_short : short;
+// CHECK: template <no_namespace_short EnumType> class dummy_functor_2;
+// CHECK: namespace internal {
+// CHECK-NEXT: enum class namespace_short : short;
+// CHECK-NEXT: }
+// CHECK: template <internal::namespace_short EnumType> class dummy_functor_3;
+// CHECK: namespace  {
+// CHECK-NEXT: enum class enum_in_anonNS : short;
+// CHECK-NEXT: }
+// CHECK: template <enum_in_anonNS EnumType> class dummy_functor_4;
+// CHECK: enum class no_type_set : int;
+// CHECK: template <no_type_set EnumType> class dummy_functor_5;
+// CHECK: enum unscoped_enum : int;
+// CHECK: template <unscoped_enum EnumType> class dummy_functor_6;
+// CHECK: template <typename EnumType> class dummy_functor_7;
+
+// CHECK: Specializations of KernelInfo for kernel function types:
+// CHECK: template <> struct KernelInfo<::dummy_functor_1<(no_namespace_int)0>>
+// CHECK: template <> struct KernelInfo<::dummy_functor_2<(no_namespace_short)1>>
+// CHECK: template <> struct KernelInfo<::dummy_functor_3<(internal::namespace_short)1>>
+// CHECK: template <> struct KernelInfo<::dummy_functor_4<(enum_in_anonNS)1>>
+// CHECK: template <> struct KernelInfo<::dummy_functor_5<(no_type_set)0>>
+// CHECK: template <> struct KernelInfo<::dummy_functor_6<(unscoped_enum)0>>
+// CHECK: template <> struct KernelInfo<::dummy_functor_7<::no_namespace_int>>
+// CHECK: template <> struct KernelInfo<::dummy_functor_7<::internal::namespace_short>>
diff --git a/clang/test/SemaSYCL/kernelname-enum.cpp b/clang/test/SemaSYCL/kernelname-enum.cpp
new file mode 100644
index 0000000000000..22a9f96acc50a
--- /dev/null
+++ b/clang/test/SemaSYCL/kernelname-enum.cpp
@@ -0,0 +1,77 @@
+// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsyntax-only -verify %s
+
+#include "sycl.hpp"
+
+enum unscoped_enum_int : int {
+  val_1,
+  val_2
+};
+
+// expected-note@+1 {{'unscoped_enum_no_type_set' declared here}}
+enum unscoped_enum_no_type_set {
+  val_3,
+  val_4
+};
+
+enum class scoped_enum_int : int {
+  val_1,
+  val_2
+};
+
+enum class scoped_enum_no_type_set {
+  val_3,
+  val_4
+};
+
+template <unscoped_enum_int EnumType>
+class dummy_functor_1 {
+public:
+  void operator()() {}
+};
+
+// expected-error@+2 {{kernel name is invalid. Unscoped enum requires fixed underlying type}}
+template <unscoped_enum_no_type_set EnumType>
+class dummy_functor_2 {
+public:
+  void operator()() {}
+};
+
+template <scoped_enum_int EnumType>
+class dummy_functor_3 {
+public:
+  void operator()() {}
+};
+
+template <scoped_enum_no_type_set EnumType>
+class dummy_functor_4 {
+public:
+  void operator()() {}
+};
+
+int main() {
+
+  dummy_functor_1<val_1> f1;
+  dummy_functor_2<val_3> f2;
+  dummy_functor_3<scoped_enum_int::val_2> f3;
+  dummy_functor_4<scoped_enum_no_type_set::val_4> f4;
+
+  cl::sycl::queue q;
+
+  q.submit([&](cl::sycl::handler &cgh) {
+    cgh.single_task(f1);
+  });
+
+  q.submit([&](cl::sycl::handler &cgh) {
+    cgh.single_task(f2);
+  });
+
+  q.submit([&](cl::sycl::handler &cgh) {
+    cgh.single_task(f3);
+  });
+
+  q.submit([&](cl::sycl::handler &cgh) {
+    cgh.single_task(f4);
+  });
+
+  return 0;
+}