diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index d1e2a22702bb5..bfafb3da5c30d 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -3195,6 +3195,9 @@ def warn_dllimport_dropped_from_inline_function : Warning<
   InGroup<IgnoredAttributes>;
 def warn_attribute_ignored : Warning<"%0 attribute ignored">,
   InGroup<IgnoredAttributes>;
+def warn_attribute_on_direct_kernel_callee_only : Warning<"%0 attribute allowed"
+  " only on a function directly called from a SYCL kernel function; attribute ignored">,
+  InGroup<IgnoredAttributes>;
 def warn_nothrow_attribute_ignored : Warning<"'nothrow' attribute conflicts with"
   " exception specification; attribute ignored">,
   InGroup<IgnoredAttributes>;
diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index c4c4a8dae34a6..990066008bc3b 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -13017,6 +13017,7 @@ class Sema final {
 
   bool isKnownGoodSYCLDecl(const Decl *D);
   void checkSYCLDeviceVarDecl(VarDecl *Var);
+  void copySYCLKernelAttrs(const CXXRecordDecl *KernelObj);
   void ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc, MangleContext &MC);
   void MarkDevice();
 
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index d6a43df48e54b..16b393f2be1fc 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -306,6 +306,37 @@ static int64_t getIntExprValue(const Expr *E, ASTContext &Ctx) {
   return E->getIntegerConstantExpr(Ctx)->getSExtValue();
 }
 
+// Collect function attributes related to SYCL.
+static void collectSYCLAttributes(Sema &S, FunctionDecl *FD,
+                                  llvm::SmallVector<Attr *, 4> &Attrs,
+                                  bool DirectlyCalled = true) {
+  if (!FD->hasAttrs())
+    return;
+
+  llvm::copy_if(FD->getAttrs(), std::back_inserter(Attrs), [](Attr *A) {
+    // FIXME: Make this list self-adapt as new SYCL attributes are added.
+    return isa<IntelReqdSubGroupSizeAttr, ReqdWorkGroupSizeAttr,
+               SYCLIntelKernelArgsRestrictAttr, SYCLIntelNumSimdWorkItemsAttr,
+               SYCLIntelSchedulerTargetFmaxMhzAttr,
+               SYCLIntelMaxWorkGroupSizeAttr, SYCLIntelMaxGlobalWorkDimAttr,
+               SYCLIntelNoGlobalWorkOffsetAttr, SYCLSimdAttr>(A);
+  });
+
+  // Allow the kernel attribute "use_stall_enable_clusters" only on lambda
+  // functions and function objects called directly from a kernel.
+  // For all other cases, emit a warning and ignore.
+  if (auto *A = FD->getAttr<SYCLIntelUseStallEnableClustersAttr>()) {
+    if (DirectlyCalled) {
+      Attrs.push_back(A);
+    } else {
+      S.Diag(A->getLocation(),
+             diag::warn_attribute_on_direct_kernel_callee_only)
+          << A;
+      FD->dropAttr<SYCLIntelUseStallEnableClustersAttr>();
+    }
+  }
+}
+
 class MarkDeviceFunction : public RecursiveASTVisitor<MarkDeviceFunction> {
   // Used to keep track of the constexpr depth, so we know whether to skip
   // diagnostics.
@@ -477,7 +508,7 @@ class MarkDeviceFunction : public RecursiveASTVisitor<MarkDeviceFunction> {
   // Returns the kernel body function found during traversal.
   FunctionDecl *
   CollectPossibleKernelAttributes(FunctionDecl *SYCLKernel,
-                                  llvm::SmallPtrSet<Attr *, 4> &Attrs) {
+                                  llvm::SmallVector<Attr *, 4> &Attrs) {
     typedef std::pair<FunctionDecl *, FunctionDecl *> ChildParentPair;
     llvm::SmallPtrSet<FunctionDecl *, 16> Visited;
     llvm::SmallVector<ChildParentPair, 16> WorkList;
@@ -508,55 +539,23 @@ class MarkDeviceFunction : public RecursiveASTVisitor<MarkDeviceFunction> {
                               "function can be called");
         KernelBody = FD;
       }
+
       WorkList.pop_back();
       if (!Visited.insert(FD).second)
         continue; // We've already seen this Decl
 
-      if (auto *A = FD->getAttr<IntelReqdSubGroupSizeAttr>())
-        Attrs.insert(A);
-
-      if (auto *A = FD->getAttr<ReqdWorkGroupSizeAttr>())
-        Attrs.insert(A);
-
-      if (auto *A = FD->getAttr<SYCLIntelKernelArgsRestrictAttr>())
-        Attrs.insert(A);
-
-      if (auto *A = FD->getAttr<SYCLIntelNumSimdWorkItemsAttr>())
-        Attrs.insert(A);
-
-      if (auto *A = FD->getAttr<SYCLIntelSchedulerTargetFmaxMhzAttr>())
-        Attrs.insert(A);
-
-      if (auto *A = FD->getAttr<SYCLIntelMaxWorkGroupSizeAttr>())
-        Attrs.insert(A);
-
-      if (auto *A = FD->getAttr<SYCLIntelMaxGlobalWorkDimAttr>())
-        Attrs.insert(A);
-
-      if (auto *A = FD->getAttr<SYCLIntelNoGlobalWorkOffsetAttr>())
-        Attrs.insert(A);
-
-      if (auto *A = FD->getAttr<SYCLSimdAttr>())
-        Attrs.insert(A);
-
-      // Allow the kernel attribute "use_stall_enable_clusters" only on lambda
-      // functions and function objects that are called directly from a kernel
-      // (i.e. the one passed to the single_task or parallel_for functions).
-      // For all other cases, emit a warning and ignore.
-      if (auto *A = FD->getAttr<SYCLIntelUseStallEnableClustersAttr>()) {
-        if (ParentFD == SYCLKernel) {
-          Attrs.insert(A);
-        } else {
-          SemaRef.Diag(A->getLocation(), diag::warn_attribute_ignored) << A;
-          FD->dropAttr<SYCLIntelUseStallEnableClustersAttr>();
-        }
-      }
+      // Gather all attributes of FD that are SYCL related.
+      // Some attributes are allowed only on lambda functions and function
+      // objects called directly from a kernel (i.e. the one passed to the
+      // single_task or parallel_for functions).
+      bool DirectlyCalled = (ParentFD == SYCLKernel);
+      collectSYCLAttributes(SemaRef, FD, Attrs, DirectlyCalled);
 
       // Attribute "loop_fuse" can be applied explicitly on kernel function.
       // Attribute should not be propagated from device functions to kernel.
       if (auto *A = FD->getAttr<SYCLIntelLoopFuseAttr>()) {
         if (ParentFD == SYCLKernel) {
-          Attrs.insert(A);
+          Attrs.push_back(A);
         }
       }
 
@@ -2058,8 +2057,8 @@ class SyclKernelArgsSizeChecker : public SyclKernelFieldHandler {
   using SyclKernelFieldHandler::handleSyclHalfType;
 };
 
-static const CXXMethodDecl *getOperatorParens(const CXXRecordDecl *Rec) {
-  for (const auto *MD : Rec->methods()) {
+static CXXMethodDecl *getOperatorParens(const CXXRecordDecl *Rec) {
+  for (auto *MD : Rec->methods()) {
     if (MD->getOverloadedOperator() == OO_Call)
       return MD;
   }
@@ -3149,6 +3148,56 @@ void Sema::CheckSYCLKernelCall(FunctionDecl *KernelFunc, SourceRange CallLoc,
     KernelFunc->setInvalidDecl();
 }
 
+// For a wrapped parallel_for, copy attributes from original
+// kernel to wrapped kernel.
+void Sema::copySYCLKernelAttrs(const CXXRecordDecl *KernelObj) {
+  // Get the operator() function of the wrapper.
+  CXXMethodDecl *OpParens = getOperatorParens(KernelObj);
+  assert(OpParens && "invalid kernel object");
+
+  typedef std::pair<FunctionDecl *, FunctionDecl *> ChildParentPair;
+  llvm::SmallPtrSet<FunctionDecl *, 16> Visited;
+  llvm::SmallVector<ChildParentPair, 16> WorkList;
+  WorkList.push_back({OpParens, nullptr});
+  FunctionDecl *KernelBody = nullptr;
+
+  CallGraph SYCLCG;
+  SYCLCG.addToCallGraph(getASTContext().getTranslationUnitDecl());
+  while (!WorkList.empty()) {
+    FunctionDecl *FD = WorkList.back().first;
+    FunctionDecl *ParentFD = WorkList.back().second;
+
+    if ((ParentFD == OpParens) && isSYCLKernelBodyFunction(FD)) {
+      KernelBody = FD;
+      break;
+    }
+
+    WorkList.pop_back();
+    if (!Visited.insert(FD).second)
+      continue; // We've already seen this Decl
+
+    CallGraphNode *N = SYCLCG.getNode(FD);
+    if (!N)
+      continue;
+
+    for (const CallGraphNode *CI : *N) {
+      if (auto *Callee = dyn_cast<FunctionDecl>(CI->getDecl())) {
+        Callee = Callee->getMostRecentDecl();
+        if (!Visited.count(Callee))
+          WorkList.push_back({Callee, FD});
+      }
+    }
+  }
+
+  assert(KernelBody && "improper parallel_for wrap");
+  if (KernelBody) {
+    llvm::SmallVector<Attr *, 4> Attrs;
+    collectSYCLAttributes(*this, KernelBody, Attrs);
+    if (!Attrs.empty())
+      llvm::for_each(Attrs, [OpParens](Attr *A) { OpParens->addAttr(A); });
+  }
+}
+
 // Generates the OpenCL kernel using KernelCallerFunc (kernel caller
 // function) defined is SYCL headers.
 // Generated OpenCL kernel contains the body of the kernel caller function,
@@ -3181,14 +3230,20 @@ void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc,
   if (KernelObj->isInvalidDecl())
     return;
 
-  bool IsSIMDKernel = isESIMDKernelType(KernelObj);
-
   // Calculate both names, since Integration headers need both.
   std::string CalculatedName, StableName;
   std::tie(CalculatedName, StableName) =
       constructKernelName(*this, KernelCallerFunc, MC);
   StringRef KernelName(getLangOpts().SYCLUnnamedLambda ? StableName
                                                        : CalculatedName);
+
+  // Attributes of a user-written SYCL kernel must be copied to the internally
+  // generated alternative kernel, identified by a known string in its name.
+  if (StableName.find("__pf_kernel_wrapper") != std::string::npos)
+    copySYCLKernelAttrs(KernelObj);
+
+  bool IsSIMDKernel = isESIMDKernelType(KernelObj);
+
   SyclKernelDeclCreator kernel_decl(*this, KernelName, KernelObj->getLocation(),
                                     KernelCallerFunc->isInlined(),
                                     IsSIMDKernel);
@@ -3226,7 +3281,7 @@ void Sema::MarkDevice(void) {
       Marker.CollectKernelSet(SYCLKernel, SYCLKernel, VisitedSet);
 
       // Let's propagate attributes from device functions to a SYCL kernels
-      llvm::SmallPtrSet<Attr *, 4> Attrs;
+      llvm::SmallVector<Attr *, 4> Attrs;
       // This function collects all kernel attributes which might be applied to
       // a device functions, but need to be propagated down to callers, i.e.
       // SYCL kernels
diff --git a/clang/test/SemaSYCL/Inputs/sycl.hpp b/clang/test/SemaSYCL/Inputs/sycl.hpp
index 6162171e84001..a1a8d626b7c40 100644
--- a/clang/test/SemaSYCL/Inputs/sycl.hpp
+++ b/clang/test/SemaSYCL/Inputs/sycl.hpp
@@ -206,11 +206,24 @@ template <typename Type>
 struct get_kernel_name_t<auto_name, Type> {
   using name = Type;
 };
+
+// Used when parallel_for range is rounded-up.
+template <typename Type> class __pf_kernel_wrapper;
+
+template <typename Type> struct get_kernel_wrapper_name_t {
+  using name =
+      __pf_kernel_wrapper<typename get_kernel_name_t<auto_name, Type>::name>;
+};
+
 #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();
 }
+template <typename KernelName = auto_name, typename KernelType>
+ATTR_SYCL_KERNEL void kernel_parallel_for(const KernelType &kernelFunc) {
+  kernelFunc();
+}
 class handler {
 public:
   template <typename KernelName = auto_name, typename KernelType>
@@ -220,6 +233,16 @@ class handler {
     kernel_single_task<NameT>(kernelFunc);
 #else
     kernelFunc();
+#endif
+  }
+  template <typename KernelName = auto_name, typename KernelType>
+  void parallel_for(const KernelType &kernelObj) {
+    using NameT = typename get_kernel_name_t<KernelName, KernelType>::name;
+    using NameWT = typename get_kernel_wrapper_name_t<NameT>::name;
+#ifdef __SYCL_DEVICE_ONLY__
+    kernel_parallel_for<NameT>(kernelObj);
+#else
+    kernelObj();
 #endif
   }
 };
diff --git a/clang/test/SemaSYCL/args-size-overflow.cpp b/clang/test/SemaSYCL/args-size-overflow.cpp
index 3b442f583b415..f94822f2be49e 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:220 {{size of kernel arguments (7994 bytes) may exceed the supported maximum of 2048 bytes on some devices}}
+// expected-warning@Inputs/sycl.hpp:233 {{size of kernel arguments (7994 bytes) may exceed the supported maximum of 2048 bytes on some devices}}
 #elif SPIR32
-// expected-warning@Inputs/sycl.hpp:220 {{size of kernel arguments (7986 bytes) may exceed the supported maximum of 2048 bytes on some devices}}
+// expected-warning@Inputs/sycl.hpp:233 {{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 904b28bda7c87..23bbbf5c0d29e 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:212 {{called by 'kernel_single_task<AName, (lambda}}
+    // expected-note@Inputs/sycl.hpp:221 {{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}}
diff --git a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp
index f634e7bc6933c..5a7c03f5e9546 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:212 {{called by 'kernel_single_task<AName, (lambda}}
+    // expected-note@Inputs/sycl.hpp:221 {{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:212 2 {{called by 'kernel_single_task<AName, (lambda}}
+    // expected-note@Inputs/sycl.hpp:221 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 0f9d6180d5652..b3b5d06ed4cad 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:212 {{called by 'kernel_single_task<variables, (lambda}}
+    // expected-note@Inputs/sycl.hpp:221 {{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:212 4{{called by 'kernel_single_task<functions, (lambda}}
+    // expected-note@Inputs/sycl.hpp:221 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:212 {{called by 'kernel_single_task<ok, (lambda}}
+    // expected-note@Inputs/sycl.hpp:221 {{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 3f4239855055d..c458859ea1be5 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:220 {{'InvalidKernelName1' is an invalid kernel name type}}
-  // expected-note@Inputs/sycl.hpp:220 {{'InvalidKernelName1' should be globally-visible}}
+  // 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-note@+8 {{in instantiation of function template specialization}}
 #elif defined(ERROR)
-  // expected-error@Inputs/sycl.hpp:220 {{'InvalidKernelName1' is an invalid kernel name type}}
-  // expected-note@Inputs/sycl.hpp:220 {{'InvalidKernelName1' should be globally-visible}}
+  // 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-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:220 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}}
+  // 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}}
 #elif defined(ERROR)
-  // expected-error@Inputs/sycl.hpp:220 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}}
+  // 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}}
 #endif
 
   q.submit([&](handler &h) {
@@ -53,9 +53,9 @@ int main() {
   });
 
 #if defined(WARN)
-  // expected-warning@Inputs/sycl.hpp:220 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}}
+  // 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}}
 #elif defined(ERROR)
-  // expected-error@Inputs/sycl.hpp:220 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}}
+  // 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}}
 #endif
 
   q.submit([&](handler &h) {
diff --git a/clang/test/SemaSYCL/int128.cpp b/clang/test/SemaSYCL/int128.cpp
index 2eea7b8734715..ea9d708083ea4 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:212 {{called by 'kernel_single_task<variables, (lambda}}
+    // expected-note@Inputs/sycl.hpp:221 {{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:212 5{{called by 'kernel_single_task<functions, (lambda}}
+    // expected-note@Inputs/sycl.hpp:221 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:212 {{called by 'kernel_single_task<ok, (lambda}}
+    // expected-note@Inputs/sycl.hpp:221 {{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/intel-max-global-work-dim-device.cpp b/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp
index 48e9305d6ecbf..3d8a4be3e64a3 100644
--- a/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp
+++ b/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp
@@ -86,7 +86,11 @@ int main() {
 
     h.single_task<class test_kernel4>(TRIFuncObjGood1());
     // CHECK-LABEL: FunctionDecl {{.*}}test_kernel4
-    // CHECK:       ReqdWorkGroupSizeAttr {{.*}}
+    // CHECK:       SYCLIntelMaxGlobalWorkDimAttr {{.*}}
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 0
+    // CHECK-NEXT:  IntegerLiteral{{.*}}0{{$}}
+    // CHECK:       SYCLIntelMaxWorkGroupSizeAttr {{.*}}
     // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
     // CHECK-NEXT:  value: Int 1
     // CHECK-NEXT:  IntegerLiteral{{.*}}1{{$}}
@@ -96,7 +100,7 @@ int main() {
     // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
     // CHECK-NEXT:  value: Int 1
     // CHECK-NEXT:  IntegerLiteral{{.*}}1{{$}}
-    // CHECK:       SYCLIntelMaxWorkGroupSizeAttr {{.*}}
+    // CHECK:       ReqdWorkGroupSizeAttr {{.*}}
     // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
     // CHECK-NEXT:  value: Int 1
     // CHECK-NEXT:  IntegerLiteral{{.*}}1{{$}}
@@ -106,37 +110,33 @@ int main() {
     // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
     // CHECK-NEXT:  value: Int 1
     // CHECK-NEXT:  IntegerLiteral{{.*}}1{{$}}
-    // CHECK:       SYCLIntelMaxGlobalWorkDimAttr {{.*}}
-    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
-    // CHECK-NEXT:  value: Int 0
-    // CHECK-NEXT:  IntegerLiteral{{.*}}0{{$}}
 
     h.single_task<class test_kernel5>(TRIFuncObjGood2());
     // CHECK-LABEL: FunctionDecl {{.*}}test_kernel5
-    // CHECK:       ReqdWorkGroupSizeAttr {{.*}}
+    // CHECK:       SYCLIntelMaxGlobalWorkDimAttr {{.*}}
     // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
-    // CHECK-NEXT:  value: Int 4
-    // CHECK-NEXT:  IntegerLiteral{{.*}}4{{$}}
+    // CHECK-NEXT:  value: Int 3
+    // CHECK-NEXT:  IntegerLiteral{{.*}}3{{$}}
+    // CHECK:       SYCLIntelMaxWorkGroupSizeAttr {{.*}}
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 8
+    // CHECK-NEXT:  IntegerLiteral{{.*}}8{{$}}
     // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
     // CHECK-NEXT:  value: Int 1
     // CHECK-NEXT:  IntegerLiteral{{.*}}1{{$}}
     // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
     // CHECK-NEXT:  value: Int 1
     // CHECK-NEXT:  IntegerLiteral{{.*}}1{{$}}
-    // CHECK:       SYCLIntelMaxWorkGroupSizeAttr {{.*}}
+    // CHECK:       ReqdWorkGroupSizeAttr {{.*}}
     // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
-    // CHECK-NEXT:  value: Int 8
-    // CHECK-NEXT:  IntegerLiteral{{.*}}8{{$}}
+    // CHECK-NEXT:  value: Int 4
+    // CHECK-NEXT:  IntegerLiteral{{.*}}4{{$}}
     // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
     // CHECK-NEXT:  value: Int 1
     // CHECK-NEXT:  IntegerLiteral{{.*}}1{{$}}
     // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
     // CHECK-NEXT:  value: Int 1
     // CHECK-NEXT:  IntegerLiteral{{.*}}1{{$}}
-    // CHECK:       SYCLIntelMaxGlobalWorkDimAttr {{.*}}
-    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
-    // CHECK-NEXT:  value: Int 3
-    // CHECK-NEXT:  IntegerLiteral{{.*}}3{{$}}
 #ifdef TRIGGER_ERROR
     [[intel::max_global_work_dim(1)]] int Var = 0; // expected-error{{'max_global_work_dim' attribute only applies to functions}}
 
diff --git a/clang/test/SemaSYCL/kernelname-enum.cpp b/clang/test/SemaSYCL/kernelname-enum.cpp
index 6a7156c646bac..bb00330af2a05 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:220 {{'dummy_functor_2<val_3>' is an invalid kernel name type}}
-    // expected-note@Inputs/sycl.hpp:220 {{unscoped enum 'unscoped_enum_no_type_set' requires fixed underlying type}}
+    // 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-note@+1{{in instantiation of function template specialization}}
     cgh.single_task(f2);
   });
 
   q.submit([&](cl::sycl::handler &cgh) {
-    // expected-error@Inputs/sycl.hpp:220 {{'templated_functor<dummy_functor_2>' is an invalid kernel name type}}
-    // expected-note@Inputs/sycl.hpp:220 {{unscoped enum 'unscoped_enum_no_type_set' requires fixed underlying type}}
+    // 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-note@+1{{in instantiation of function template specialization}}
     cgh.single_task(f5);
   });
diff --git a/clang/test/SemaSYCL/num_simd_work_items_device.cpp b/clang/test/SemaSYCL/num_simd_work_items_device.cpp
index b7fc2a8c9475f..c099407d7780b 100644
--- a/clang/test/SemaSYCL/num_simd_work_items_device.cpp
+++ b/clang/test/SemaSYCL/num_simd_work_items_device.cpp
@@ -243,6 +243,10 @@ int main() {
 
     h.single_task<class test_kernel4>(TRIFuncObjGood1());
     // CHECK-LABEL: FunctionDecl {{.*}}test_kernel4
+    // CHECK:       SYCLIntelNumSimdWorkItemsAttr {{.*}}
+    // CHECK-NEXT:  ConstantExpr{{.*}}'int'
+    // CHECK-NEXT:  value: Int 4
+    // CHECK-NEXT:  IntegerLiteral{{.*}}4{{$}}
     // CHECK:       ReqdWorkGroupSizeAttr {{.*}}
     // CHECK-NEXT:  ConstantExpr{{.*}}'int'
     // CHECK-NEXT:  value: Int 64
@@ -253,10 +257,6 @@ int main() {
     // CHECK-NEXT:  ConstantExpr{{.*}}'int'
     // CHECK-NEXT:  value: Int 5
     // CHECK-NEXT:  IntegerLiteral{{.*}}5{{$}}
-    // CHECK:       SYCLIntelNumSimdWorkItemsAttr {{.*}}
-    // CHECK-NEXT:  ConstantExpr{{.*}}'int'
-    // CHECK-NEXT:  value: Int 4
-    // CHECK-NEXT:  IntegerLiteral{{.*}}4{{$}}
 
     h.single_task<class test_kernel5>(TRIFuncObjGood2());
     // CHECK-LABEL: FunctionDecl {{.*}}test_kernel5
@@ -277,6 +277,10 @@ int main() {
 
     h.single_task<class test_kernel6>(TRIFuncObjGood3());
     // CHECK-LABEL: FunctionDecl {{.*}}test_kernel6
+    // CHECK:       SYCLIntelNumSimdWorkItemsAttr {{.*}}
+    // CHECK-NEXT:  ConstantExpr{{.*}}'int'
+    // CHECK-NEXT:  value: Int 4
+    // CHECK-NEXT:  IntegerLiteral{{.*}}4{{$}}
     // CHECK:       ReqdWorkGroupSizeAttr {{.*}}
     // CHECK-NEXT:  ConstantExpr{{.*}}'int'
     // CHECK-NEXT:  value: Int 64
@@ -287,10 +291,6 @@ int main() {
     // CHECK-NEXT:  ConstantExpr{{.*}}'int'
     // CHECK-NEXT:  value: Int 5
     // CHECK-NEXT:  IntegerLiteral{{.*}}5{{$}}
-    // CHECK:       SYCLIntelNumSimdWorkItemsAttr {{.*}}
-    // CHECK-NEXT:  ConstantExpr{{.*}}'int'
-    // CHECK-NEXT:  value: Int 4
-    // CHECK-NEXT:  IntegerLiteral{{.*}}4{{$}}
 
     h.single_task<class test_kernel7>(TRIFuncObjGood4());
     // CHECK-LABEL: FunctionDecl {{.*}}test_kernel7
@@ -311,6 +311,10 @@ int main() {
 
     h.single_task<class test_kernel8>(TRIFuncObjGood5());
     // CHECK-LABEL: FunctionDecl {{.*}}test_kernel8
+    // CHECK:       SYCLIntelNumSimdWorkItemsAttr {{.*}}
+    // CHECK-NEXT:  ConstantExpr{{.*}}'int'
+    // CHECK-NEXT:  value: Int 4
+    // CHECK-NEXT:  IntegerLiteral{{.*}}4{{$}}
     // CHECK:       ReqdWorkGroupSizeAttr {{.*}}
     // CHECK-NEXT:  ConstantExpr{{.*}}'int'
     // CHECK-NEXT:  value: Int 64
@@ -321,10 +325,6 @@ int main() {
     // CHECK-NEXT:  ConstantExpr{{.*}}'int'
     // CHECK-NEXT:  value: Int 1
     // CHECK-NEXT:  IntegerLiteral{{.*}}1{{$}}
-    // CHECK:       SYCLIntelNumSimdWorkItemsAttr {{.*}}
-    // CHECK-NEXT:  ConstantExpr{{.*}}'int'
-    // CHECK-NEXT:  value: Int 4
-    // CHECK-NEXT:  IntegerLiteral{{.*}}4{{$}}
 
     h.single_task<class test_kernel9>(TRIFuncObjGood6());
     // CHECK-LABEL: FunctionDecl {{.*}}test_kernel9
@@ -345,6 +345,10 @@ int main() {
 
     h.single_task<class test_kernel10>(TRIFuncObjGood7());
     // CHECK-LABEL: FunctionDecl {{.*}}test_kernel10
+    // CHECK:       SYCLIntelNumSimdWorkItemsAttr {{.*}}
+    // CHECK-NEXT:  ConstantExpr{{.*}}'int'
+    // CHECK-NEXT:  value: Int 4
+    // CHECK-NEXT:  IntegerLiteral{{.*}}4{{$}}
     // CHECK:       ReqdWorkGroupSizeAttr {{.*}}
     // CHECK-NEXT:  ConstantExpr{{.*}}'int'
     // CHECK-NEXT:  value: Int 64
@@ -355,10 +359,6 @@ int main() {
     // CHECK-NEXT:  ConstantExpr{{.*}}'int'
     // CHECK-NEXT:  value: Int 1
     // CHECK-NEXT:  IntegerLiteral{{.*}}1{{$}}
-    // CHECK:       SYCLIntelNumSimdWorkItemsAttr {{.*}}
-    // CHECK-NEXT:  ConstantExpr{{.*}}'int'
-    // CHECK-NEXT:  value: Int 4
-    // CHECK-NEXT:  IntegerLiteral{{.*}}4{{$}}
 
     h.single_task<class test_kernel11>(TRIFuncObjGood8());
     // CHECK-LABEL: FunctionDecl {{.*}}test_kernel11
diff --git a/clang/test/SemaSYCL/parallel_for_wrapper_attr.cpp b/clang/test/SemaSYCL/parallel_for_wrapper_attr.cpp
new file mode 100755
index 0000000000000..79c1c4f6d92a1
--- /dev/null
+++ b/clang/test/SemaSYCL/parallel_for_wrapper_attr.cpp
@@ -0,0 +1,54 @@
+// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl -fsycl-is-device -triple spir64 | FileCheck %s
+
+#include "Inputs/sycl.hpp"
+
+template <typename T> class Fobj {
+public:
+  Fobj() {}
+  void operator()() const {
+    auto L0 = []() [[intel::reqd_sub_group_size(4)]]{};
+    L0();
+  }
+};
+
+void invoke() {
+  sycl::queue q;
+  q.submit([&](sycl::handler &h) {
+    Fobj<int> fobj1;
+    h.parallel_for<class __pf_kernel_wrapper>(fobj1);
+  });
+  q.submit([&](sycl::handler &h) {
+    Fobj<short> fobj2;
+    h.parallel_for<class PPP>(fobj2);
+  });
+}
+
+// CHECK-LABEL: ClassTemplateSpecializationDecl {{.*}} class Fobj definition
+// CHECK:       TemplateArgument type 'int'
+// CHECK:       CXXMethodDecl {{.*}} used operator() 'void () const'
+// CHECK:       CXXMethodDecl {{.*}} used operator() 'void () const' inline
+// CHECK-NEXT:  CompoundStmt
+// CHECK-NEXT:  IntelReqdSubGroupSizeAttr {{.*}}
+// CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+// CHECK-NEXT:  value: Int 4
+// CHECK-NEXT:  IntegerLiteral{{.*}}4{{$}}
+// CHECK:       CXXOperatorCallExpr {{.*}} 'void':'void' '()'
+// CHECK:       IntelReqdSubGroupSizeAttr {{.*}}
+// CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+// CHECK-NEXT:  value: Int 4
+// CHECK-NEXT:  IntegerLiteral{{.*}}4{{$}}
+// CHECK:       CXXConstructorDecl
+// CHECK:       CXXConstructorDecl
+
+// CHECK-LABEL: ClassTemplateSpecializationDecl {{.*}} class Fobj definition
+// CHECK:       TemplateArgument type 'short'
+// CHECK:       CXXMethodDecl {{.*}} used operator() 'void () const'
+// CHECK:       CXXMethodDecl {{.*}} used operator() 'void () const' inline
+// CHECK-NEXT:  CompoundStmt
+// CHECK-NEXT:  IntelReqdSubGroupSizeAttr {{.*}}
+// CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+// CHECK-NEXT:  value: Int 4
+// CHECK-NEXT:  IntegerLiteral{{.*}}4{{$}}
+// CHECK:       CXXOperatorCallExpr {{.*}} 'void':'void' '()'
+// CHECK-NOT:   IntelReqdSubGroupSizeAttr {{.*}}
+// CHECK:       CXXConstructorDecl
diff --git a/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp b/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp
index c1fa227b9f73a..082fe8bf64594 100644
--- a/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp
+++ b/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp
@@ -53,6 +53,10 @@ int main() {
 #ifndef TRIGGER_ERROR
     // CHECK-LABEL:  FunctionDecl {{.*}} main 'int ()'
     // CHECK:  `-FunctionDecl {{.*}}test_kernel1 'void ()'
+    // CHECK:  -SYCLIntelNoGlobalWorkOffsetAttr {{.*}}
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 1
+    // CHECK-NEXT:  IntegerLiteral{{.*}}1{{$}}
     // CHECK:  -SYCLIntelMaxWorkGroupSizeAttr {{.*}} Inherited
     // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
     // CHECK-NEXT:  value: Int 4
@@ -63,10 +67,6 @@ int main() {
     // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
     // CHECK-NEXT:  value: Int 4
     // CHECK-NEXT:  IntegerLiteral{{.*}}4{{$}}
-    // CHECK:  -SYCLIntelNoGlobalWorkOffsetAttr {{.*}}
-    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
-    // CHECK-NEXT:  value: Int 1
-    // CHECK-NEXT:  IntegerLiteral{{.*}}1{{$}}
     // CHECK:  `-ReqdWorkGroupSizeAttr {{.*}}
     // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
     // CHECK-NEXT:  value: Int 2
diff --git a/clang/test/SemaSYCL/stall_enable.cpp b/clang/test/SemaSYCL/stall_enable.cpp
index 45ba33a58a906..ec496f441469f 100644
--- a/clang/test/SemaSYCL/stall_enable.cpp
+++ b/clang/test/SemaSYCL/stall_enable.cpp
@@ -6,7 +6,7 @@
 using namespace cl::sycl;
 queue q;
 
-[[intel::use_stall_enable_clusters]] void test() {} //expected-warning{{'use_stall_enable_clusters' attribute ignored}}
+[[intel::use_stall_enable_clusters]] void test() {} // expected-warning{{'use_stall_enable_clusters' attribute allowed only on a function directly called from a SYCL kernel}}
 
 #ifdef TRIGGER_ERROR
 [[intel::use_stall_enable_clusters(1)]] void bar1() {} // expected-error{{'use_stall_enable_clusters' attribute takes no arguments}}
diff --git a/clang/test/SemaSYCL/stdtypes_kernel_type.cpp b/clang/test/SemaSYCL/stdtypes_kernel_type.cpp
index 59d212f64ded7..0d44415bf2394 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:220 {{'nullptr_t' is an invalid kernel name type}}
-    // expected-note@Inputs/sycl.hpp:220 {{type 'nullptr_t' cannot be in the "std" namespace}}
+    // 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-note@+1{{in instantiation of function template specialization}}
     h.single_task<std::nullptr_t>([=] {});
   });
   q.submit([&](handler &h) {
-    // expected-error@Inputs/sycl.hpp:220 {{'std::T' is an invalid kernel name type}}
-    // expected-note@Inputs/sycl.hpp:220 {{type 'std::T' cannot be in the "std" namespace}}
+    // 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-note@+1{{in instantiation of function template specialization}}
     h.single_task<std::T>([=] {});
   });
   q.submit([&](handler &h) {
-    // expected-error@Inputs/sycl.hpp:220 {{'Templated_kernel_name<nullptr_t>' is an invalid kernel name type}}
-    // expected-note@Inputs/sycl.hpp:220 {{type 'nullptr_t' cannot be in the "std" namespace}}
+    // 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-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:220 {{'Templated_kernel_name<std::U>' is an invalid kernel name type}}
-    // expected-note@Inputs/sycl.hpp:220 {{type 'std::U' cannot be in the "std" namespace}}
+    // 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-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:220 {{'Templated_kernel_name2<Templated_kernel_name<std::Foo>>' is an invalid kernel name type}}
-    // expected-note@Inputs/sycl.hpp:220{{type 'std::Foo' cannot be in the "std" namespace}}
+    // 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-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:220 {{'TemplParamPack<int, float, nullptr_t, double>' is an invalid kernel name type}}
-    // expected-note@Inputs/sycl.hpp:220 {{type 'nullptr_t' cannot be in the "std" namespace}}
+    // 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-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 c4e3746f78679..97edfca682745 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:220 {{'InvalidKernelName1' is an invalid kernel name type}}
-    // expected-note@Inputs/sycl.hpp:220 {{'InvalidKernelName1' should be globally-visible}}
+    // 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-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:220 {{'namespace1::KernelName<InvalidKernelName2>' is an invalid kernel name type}}
-    // expected-note@Inputs/sycl.hpp:220 {{'InvalidKernelName2' should be globally-visible}}
+    // 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-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:220 {{'MyWrapper::InvalidKernelName0' is an invalid kernel name type}}
-    // expected-note@Inputs/sycl.hpp:220 {{'MyWrapper::InvalidKernelName0' should be globally-visible}}
+    // 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-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:220 {{'namespace1::KernelName<MyWrapper::InvalidKernelName3>' is an invalid kernel name type}}
-    // expected-note@Inputs/sycl.hpp:220 {{'MyWrapper::InvalidKernelName3' should be globally-visible}}
+    // 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-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:220 {{'std::max_align_t' is an invalid kernel name type}}
-    // expected-note@Inputs/sycl.hpp:220 {{type 'std::max_align_t' cannot be in the "std" namespace}}
+    // 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-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:220 {{'MyWrapper::InvalidKernelName4' is an invalid kernel name type}}
-    // expected-note@Inputs/sycl.hpp:220 {{'MyWrapper::InvalidKernelName4' should be globally-visible}}
+    // 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-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:220 {{'namespace1::KernelName<MyWrapper::InvalidKernelName5>' is an invalid kernel name type}}
-    // expected-note@Inputs/sycl.hpp:220 {{'MyWrapper::InvalidKernelName5' should be globally-visible}}
+    // 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-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:220 {{'Templated_kernel_name2<Templated_kernel_name<InvalidKernelName1>>' is an invalid kernel name type}}
-    // expected-note@Inputs/sycl.hpp:220 {{'InvalidKernelName1' should be globally-visible}}
+    // 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-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:220 {{'(lambda at {{.*}}unnamed-kernel.cpp{{.*}}' is an invalid kernel name type}}
-  // expected-note@Inputs/sycl.hpp:220 {{unnamed type used in a SYCL kernel name}}
+  // 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-note@+2{{in instantiation of function template specialization}}
 #endif
   q.submit([&](cl::sycl::handler &h) { h.single_task([] {}); });
diff --git a/sycl/test/kernel_param/attr.cpp b/sycl/test/kernel_param/attr.cpp
new file mode 100755
index 0000000000000..1037d23f54686
--- /dev/null
+++ b/sycl/test/kernel_param/attr.cpp
@@ -0,0 +1,22 @@
+// RUN: %clangxx -fsycl-device-only -Xclang -fsycl-is-device -emit-llvm %s -S -o %t.ll -I %sycl_include
+// RUN: FileCheck %s --input-file %t.ll
+
+// Check copying of parallel_for kernel attributes to wrapper kernel.
+
+#include <CL/sycl.hpp>
+using namespace cl::sycl;
+
+int main() {
+  range<1> Size{10};
+  {
+    queue myQueue;
+    myQueue.submit([&](handler &cgh) {
+      cgh.parallel_for<class C>(Size, [=](item<1> ITEM)
+                                          [[intel::reqd_work_group_size(4)]]{});
+    });
+  }
+
+  return 0;
+}
+
+// CHECK: define {{.*}}spir_kernel void @{{.*}}__pf_kernel_wrapper{{.*}}reqd_work_group_size