diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 8b01c3de185a9..02f688aa9774f 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -12433,6 +12433,12 @@ def err_sycl_kernel_incorrectly_named : Error<
   "'-fsycl-unnamed-lambda' to enable unnamed kernel lambdas"
   "}0">;
 
+// SYCL free function kernels extension.
+def note_free_function_kernel_param_type_not_fwd_declarable : Note<
+  "%0 is not forward declarable">;
+def note_free_function_kernel_param_type_not_supported : Note<
+  "%0 is not yet supported as a free function kernel parameter">;
+
 def err_sycl_kernel_not_function_object
     : Error<"kernel parameter must be a lambda or function object">;
 def err_sycl_restrict : Error<
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index ebc58fa1fb21f..99d9cbe1f4f60 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -490,6 +490,84 @@ void SemaSYCL::checkSYCLDeviceVarDecl(VarDecl *Var) {
   checkSYCLType(*this, Ty, Loc, Visited);
 }
 
+enum NotForwardDeclarableReason {
+  UnscopedEnum,
+  StdNamespace,
+  UnnamedTag,
+  NotAtNamespaceScope,
+  None
+};
+
+// This is a helper function which is used to check if a class declaration is:
+//   * declared within namespace 'std' (at any level)
+//     e.g., namespace std { namespace literals { class Whatever; } }
+//     h.single_task<std::literals::Whatever>([]() {});
+//   * declared within a function
+//     e.g., void foo() { struct S { int i; };
+//     h.single_task<S>([]() {}); }
+//   * declared within another tag
+//     e.g., struct S { struct T { int i } t; };
+//     h.single_task<S::T>([]() {});
+//  User for kernel name types and class/struct types used in free function
+//  kernel arguments.
+static NotForwardDeclarableReason
+isForwardDeclarable(const NamedDecl *DeclToCheck, SemaSYCL &S,
+                    bool DiagForFreeFunction = false) {
+  if (const auto *ED = dyn_cast<EnumDecl>(DeclToCheck);
+      ED && !ED->isScoped() && !ED->isFixed())
+    return NotForwardDeclarableReason::UnscopedEnum;
+
+  const DeclContext *DeclCtx = DeclToCheck->getDeclContext();
+  if (DeclCtx) {
+    while (!DeclCtx->isTranslationUnit() &&
+           (isa<NamespaceDecl>(DeclCtx) || isa<LinkageSpecDecl>(DeclCtx))) {
+      const auto *NSDecl = dyn_cast<NamespaceDecl>(DeclCtx);
+      // We don't report free function kernel parameter case because the
+      // restriction for the type used there to be forward declarable comes from
+      // the need to forward declare it in the integration header. We're safe
+      // to do so because the integration header is an implemention detail and
+      // is generated by the compiler.
+      // We do diagnose case with kernel name type since the spec requires us to
+      // do so.
+      if (!DiagForFreeFunction && NSDecl && NSDecl->isStdNamespace())
+        return NotForwardDeclarableReason::StdNamespace;
+      DeclCtx = DeclCtx->getParent();
+    }
+  }
+
+  // Check if the we've met a Tag declaration local to a non-namespace scope
+  // (i.e. Inside a function or within another Tag etc).
+  if (const auto *Tag = dyn_cast<TagDecl>(DeclToCheck)) {
+    if (Tag->getIdentifier() == nullptr)
+      return NotForwardDeclarableReason::UnnamedTag;
+    if (!DeclCtx->isTranslationUnit()) {
+      // Diagnose used types without complete definition i.e.
+      //   int main() {
+      //     class KernelName1;
+      //     parallel_for<class KernelName1>(..);
+      //   }
+      // For kernel name type This case can only be diagnosed during host
+      // compilation because the integration header is required to distinguish
+      // between the invalid code (above) and the following valid code:
+      //   int main() {
+      //     parallel_for<class KernelName2>(..);
+      //   }
+      // The device compiler forward declares both KernelName1 and
+      // KernelName2 in the integration header as ::KernelName1 and
+      // ::KernelName2. The problem with the former case is the additional
+      // declaration 'class KernelName1' in non-global scope. Lookup in this
+      // case will resolve to ::main::KernelName1 (instead of
+      // ::KernelName1). Since this is not visible to runtime code that
+      // submits kernels, this is invalid.
+      if (Tag->isCompleteDefinition() ||
+          S.getLangOpts().SYCLEnableIntHeaderDiags || DiagForFreeFunction)
+        return NotForwardDeclarableReason::NotAtNamespaceScope;
+    }
+  }
+
+  return NotForwardDeclarableReason::None;
+}
+
 // Tests whether given function is a lambda function or '()' operator used as
 // SYCL kernel body function (e.g. in parallel_for).
 // NOTE: This is incomplete implemenation. See TODO in the FE TODO list for the
@@ -1446,9 +1524,12 @@ class KernelObjVisitor {
                   HandlerTys &...Handlers) {
     if (isSyclSpecialType(ParamTy, SemaSYCLRef))
       KP_FOR_EACH(handleOtherType, Param, ParamTy);
-    else if (ParamTy->isStructureOrClassType())
-      KP_FOR_EACH(handleOtherType, Param, ParamTy);
-    else if (ParamTy->isUnionType())
+    else if (ParamTy->isStructureOrClassType()) {
+      if (KP_FOR_EACH(handleStructType, Param, ParamTy)) {
+        CXXRecordDecl *RD = ParamTy->getAsCXXRecordDecl();
+        visitRecord(nullptr, Param, RD, ParamTy, Handlers...);
+      }
+    } else if (ParamTy->isUnionType())
       KP_FOR_EACH(handleOtherType, Param, ParamTy);
     else if (ParamTy->isReferenceType())
       KP_FOR_EACH(handleOtherType, Param, ParamTy);
@@ -1960,8 +2041,21 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler {
   }
 
   bool handleStructType(ParmVarDecl *PD, QualType ParamTy) final {
-    Diag.Report(PD->getLocation(), diag::err_bad_kernel_param_type) << ParamTy;
-    IsInvalid = true;
+    CXXRecordDecl *RD = ParamTy->getAsCXXRecordDecl();
+    // For free functions all struct/class kernel arguments are forward declared
+    // in integration header, that adds additional restrictions for kernel
+    // arguments.
+    NotForwardDeclarableReason NFDR =
+        isForwardDeclarable(RD, SemaSYCLRef, /*DiagForFreeFunction=*/true);
+    if (NFDR != NotForwardDeclarableReason::None) {
+      Diag.Report(PD->getLocation(),
+                  diag::err_bad_kernel_param_type)
+          << ParamTy;
+      Diag.Report(PD->getLocation(),
+                  diag::note_free_function_kernel_param_type_not_fwd_declarable)
+          << ParamTy;
+      IsInvalid = true;
+    }
     return isValid();
   }
 
@@ -2040,15 +2134,31 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler {
   }
 
   bool enterStruct(const CXXRecordDecl *, ParmVarDecl *, QualType) final {
-    // TODO
-    unsupportedFreeFunctionParamType();
+    // TODO manipulate struct depth once special types are supported for free
+    // function kernels.
+    // ++StructFieldDepth;
     return true;
   }
 
-  bool leaveStruct(const CXXRecordDecl *, ParmVarDecl *, QualType) final {
-    // TODO
-    unsupportedFreeFunctionParamType();
-    return true;
+  bool leaveStruct(const CXXRecordDecl *, ParmVarDecl *PD,
+                   QualType ParamTy) final {
+    // TODO manipulate struct depth once special types are supported for free
+    // function kernels.
+    // --StructFieldDepth;
+    // TODO We don't yet support special types and therefore structs that
+    // require decomposition and leaving/entering. Diagnose for better user
+    // experience.
+    CXXRecordDecl *RD = ParamTy->getAsCXXRecordDecl();
+    if (RD->hasAttr<SYCLRequiresDecompositionAttr>()) {
+      Diag.Report(PD->getLocation(),
+                  diag::err_bad_kernel_param_type)
+          << ParamTy;
+      Diag.Report(PD->getLocation(),
+                  diag::note_free_function_kernel_param_type_not_supported)
+          << ParamTy;
+      IsInvalid = true;
+    }
+    return isValid();
   }
 
   bool enterStruct(const CXXRecordDecl *, const CXXBaseSpecifier &BS,
@@ -2154,8 +2264,9 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler {
   }
 
   bool handleSyclSpecialType(ParmVarDecl *, QualType) final {
-    // TODO
-    unsupportedFreeFunctionParamType();
+    // TODO We don't support special types in free function kernel parameters,
+    // but track them to diagnose the case properly.
+    CollectionStack.back() = true;
     return true;
   }
 
@@ -2165,8 +2276,7 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler {
   }
 
   bool handlePointerType(ParmVarDecl *, QualType) final {
-    // TODO
-    unsupportedFreeFunctionParamType();
+    PointerStack.back() = targetRequiresNewType(SemaSYCLRef.getASTContext());
     return true;
   }
 
@@ -2197,8 +2307,8 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler {
   }
 
   bool enterStruct(const CXXRecordDecl *, ParmVarDecl *, QualType) final {
-    // TODO
-    unsupportedFreeFunctionParamType();
+    CollectionStack.push_back(false);
+    PointerStack.push_back(false);
     return true;
   }
 
@@ -2224,10 +2334,22 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler {
     return true;
   }
 
-  bool leaveStruct(const CXXRecordDecl *RD, ParmVarDecl *PD,
+  bool leaveStruct(const CXXRecordDecl *, ParmVarDecl *,
                    QualType ParamTy) final {
-    // TODO
-    unsupportedFreeFunctionParamType();
+    CXXRecordDecl *RD = ParamTy->getAsCXXRecordDecl();
+    assert(RD && "should not be null.");
+    if (CollectionStack.pop_back_val()) {
+      if (!RD->hasAttr<SYCLRequiresDecompositionAttr>())
+        RD->addAttr(SYCLRequiresDecompositionAttr::CreateImplicit(
+            SemaSYCLRef.getASTContext()));
+      CollectionStack.back() = true;
+      PointerStack.pop_back();
+    } else if (PointerStack.pop_back_val()) {
+      PointerStack.back() = true;
+      if (!RD->hasAttr<SYCLGenerateNewTypeAttr>())
+        RD->addAttr(SYCLGenerateNewTypeAttr::CreateImplicit(
+            SemaSYCLRef.getASTContext()));
+    }
     return true;
   }
 
@@ -2826,7 +2948,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
 
   bool enterStruct(const CXXRecordDecl *, ParmVarDecl *, QualType) final {
     // TODO
-    unsupportedFreeFunctionParamType();
+    // ++StructDepth;
     return true;
   }
 
@@ -2837,7 +2959,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
 
   bool leaveStruct(const CXXRecordDecl *, ParmVarDecl *, QualType) final {
     // TODO
-    unsupportedFreeFunctionParamType();
+    // --StructDepth;
     return true;
   }
 
@@ -2977,8 +3099,15 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
 
   bool handleNonDecompStruct(const CXXRecordDecl *RD, ParmVarDecl *PD,
                              QualType ParamTy) final {
-    // TODO
-    unsupportedFreeFunctionParamType();
+    // This is a struct parameter which should not be decomposed.
+    CXXRecordDecl *ParamRecordDecl = ParamTy->getAsCXXRecordDecl();
+    assert(ParamRecordDecl && "Type must be a C++ record type");
+    // Check if we need to generate a new type for this record,
+    // i.e. this record contains pointers.
+    if (ParamRecordDecl->hasAttr<SYCLGenerateNewTypeAttr>())
+      addParam(PD, GenerateNewRecordType(ParamRecordDecl));
+    else
+      addParam(PD, ParamTy);
     return true;
   }
 
@@ -3206,8 +3335,7 @@ class SyclKernelArgsSizeChecker : public SyclKernelFieldHandler {
 
   bool handleNonDecompStruct(const CXXRecordDecl *, ParmVarDecl *,
                              QualType ParamTy) final {
-    // TODO
-    unsupportedFreeFunctionParamType();
+    addParam(ParamTy);
     return true;
   }
 
@@ -4197,7 +4325,7 @@ class FreeFunctionKernelBodyCreator : public SyclKernelFieldHandler {
 
   // Creates a DeclRefExpr to the ParmVar that represents the current pointer
   // parameter.
-  Expr *createPointerParamReferenceExpr(QualType PointerTy, bool Wrapped) {
+  Expr *createPointerParamReferenceExpr(QualType PointerTy) {
     ParmVarDecl *FreeFunctionParameter =
         DeclCreator.getParamVarDeclsForCurrentField()[0];
 
@@ -4215,6 +4343,50 @@ class FreeFunctionKernelBodyCreator : public SyclKernelFieldHandler {
     return DRE;
   }
 
+  Expr *createGetAddressOf(Expr *E) {
+    return UnaryOperator::Create(
+        SemaSYCLRef.getASTContext(), E, UO_AddrOf,
+        SemaSYCLRef.getASTContext().getPointerType(E->getType()), VK_PRValue,
+        OK_Ordinary, SourceLocation(), false,
+        SemaSYCLRef.SemaRef.CurFPFeatureOverrides());
+  }
+
+  Expr *createDerefOp(Expr *E) {
+    return UnaryOperator::Create(SemaSYCLRef.getASTContext(), E, UO_Deref,
+                                 E->getType()->getPointeeType(), VK_LValue,
+                                 OK_Ordinary, SourceLocation(), false,
+                                 SemaSYCLRef.SemaRef.CurFPFeatureOverrides());
+  }
+
+  Expr *createReinterpretCastExpr(Expr *E, QualType To) {
+    return CXXReinterpretCastExpr::Create(
+        SemaSYCLRef.getASTContext(), To, VK_PRValue, CK_BitCast, E,
+        /*Path=*/nullptr,
+        SemaSYCLRef.getASTContext().getTrivialTypeSourceInfo(To),
+        SourceLocation(), SourceLocation(), SourceRange());
+  }
+
+  Expr *createCopyInitExpr(ParmVarDecl *OrigFunctionParameter) {
+    Expr *DRE = createParamReferenceExpr();
+
+    assert(OrigFunctionParameter && "no parameter?");
+
+    CXXRecordDecl *RD = OrigFunctionParameter->getType()->getAsCXXRecordDecl();
+    InitializedEntity Entity = InitializedEntity::InitializeParameter(
+        SemaSYCLRef.getASTContext(), OrigFunctionParameter);
+
+    if (RD->hasAttr<SYCLGenerateNewTypeAttr>()) {
+      DRE = createReinterpretCastExpr(
+          createGetAddressOf(DRE), SemaSYCLRef.getASTContext().getPointerType(
+                                       OrigFunctionParameter->getType()));
+      DRE = createDerefOp(DRE);
+    }
+
+    ExprResult ArgE = SemaSYCLRef.SemaRef.PerformCopyInitialization(
+        Entity, SourceLocation(), DRE, false, false);
+    return ArgE.getAs<Expr>();
+  }
+
   // For a free function such as:
   // void f(int i, int* p, struct Simple S) { ... }
   //
@@ -4284,7 +4456,7 @@ class FreeFunctionKernelBodyCreator : public SyclKernelFieldHandler {
   }
 
   bool handlePointerType(ParmVarDecl *PD, QualType ParamTy) final {
-    Expr *PointerRef = createPointerParamReferenceExpr(ParamTy, false);
+    Expr *PointerRef = createPointerParamReferenceExpr(ParamTy);
     ArgExprs.push_back(PointerRef);
     return true;
   }
@@ -4302,10 +4474,10 @@ class FreeFunctionKernelBodyCreator : public SyclKernelFieldHandler {
     return true;
   }
 
-  bool handleNonDecompStruct(const CXXRecordDecl *, ParmVarDecl *,
+  bool handleNonDecompStruct(const CXXRecordDecl *, ParmVarDecl *PD,
                              QualType) final {
-    // TODO
-    unsupportedFreeFunctionParamType();
+    Expr *TempCopy = createCopyInitExpr(PD);
+    ArgExprs.push_back(TempCopy);
     return true;
   }
 
@@ -4591,8 +4763,7 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
 
   bool handleNonDecompStruct(const CXXRecordDecl *, ParmVarDecl *PD,
                              QualType ParamTy) final {
-    // TODO
-    unsupportedFreeFunctionParamType();
+    addParam(PD, ParamTy, SYCLIntegrationHeader::kind_std_layout);
     return true;
   }
 
@@ -4769,89 +4940,37 @@ class SYCLKernelNameTypeVisitor
   }
 
   void DiagnoseKernelNameType(const NamedDecl *DeclNamed) {
-    /*
-    This is a helper function which throws an error if the kernel name
-    declaration is:
-      * declared within namespace 'std' (at any level)
-        e.g., namespace std { namespace literals { class Whatever; } }
-        h.single_task<std::literals::Whatever>([]() {});
-      * declared within a function
-        e.g., void foo() { struct S { int i; };
-        h.single_task<S>([]() {}); }
-      * declared within another tag
-        e.g., struct S { struct T { int i } t; };
-        h.single_task<S::T>([]() {});
-    */
-
-    if (const auto *ED = dyn_cast<EnumDecl>(DeclNamed)) {
-      if (!ED->isScoped() && !ED->isFixed()) {
+    if (!IsUnnamedKernel) {
+      NotForwardDeclarableReason NFDR = isForwardDeclarable(DeclNamed, S);
+      switch (NFDR) {
+      case NotForwardDeclarableReason::UnscopedEnum:
         S.Diag(KernelInvocationFuncLoc, diag::err_sycl_kernel_incorrectly_named)
             << /* unscoped enum requires fixed underlying type */ 1
             << DeclNamed;
         IsInvalid = true;
-      }
-    }
-
-    const DeclContext *DeclCtx = DeclNamed->getDeclContext();
-    if (DeclCtx && !IsUnnamedKernel) {
-
-      // Check if the kernel name declaration is declared within namespace
-      // "std" (at any level).
-      while (!DeclCtx->isTranslationUnit() && isa<NamespaceDecl>(DeclCtx)) {
-        const auto *NSDecl = cast<NamespaceDecl>(DeclCtx);
-        if (NSDecl->isStdNamespace()) {
-          S.Diag(KernelInvocationFuncLoc,
-                 diag::err_invalid_std_type_in_sycl_kernel)
-              << KernelNameType << DeclNamed;
-          IsInvalid = true;
-          return;
-        }
-        DeclCtx = DeclCtx->getParent();
-      }
-
-      // Check if the kernel name is a Tag declaration
-      // local to a non-namespace scope (i.e. Inside a function or within
-      // another Tag etc).
-      if (!DeclCtx->isTranslationUnit() && !isa<NamespaceDecl>(DeclCtx)) {
-        if (const auto *Tag = dyn_cast<TagDecl>(DeclNamed)) {
-          bool UnnamedLambdaUsed = Tag->getIdentifier() == nullptr;
-
-          if (UnnamedLambdaUsed) {
-            S.Diag(KernelInvocationFuncLoc,
-                   diag::err_sycl_kernel_incorrectly_named)
-                << /* unnamed type is invalid */ 2 << KernelNameType;
-            IsInvalid = true;
-            return;
-          }
-
-          // Diagnose used types without complete definition i.e.
-          //   int main() {
-          //     class KernelName1;
-          //     parallel_for<class KernelName1>(..);
-          //   }
-          // This case can only be diagnosed during host compilation because the
-          // integration header is required to distinguish between the invalid
-          // code (above) and the following valid code:
-          //   int main() {
-          //     parallel_for<class KernelName2>(..);
-          //   }
-          // The device compiler forward declares both KernelName1 and
-          // KernelName2 in the integration header as ::KernelName1 and
-          // ::KernelName2. The problem with the former case is the additional
-          // declaration 'class KernelName1' in non-global scope. Lookup in this
-          // case will resolve to ::main::KernelName1 (instead of
-          // ::KernelName1). Since this is not visible to runtime code that
-          // submits kernels, this is invalid.
-          if (Tag->isCompleteDefinition() ||
-              S.getLangOpts().SYCLEnableIntHeaderDiags) {
-            S.Diag(KernelInvocationFuncLoc,
-                   diag::err_sycl_kernel_incorrectly_named)
-                << /* kernel name should be forward declarable at namespace
-                      scope */
-                0 << KernelNameType;
-            IsInvalid = true;
-          }
-        }
+        return;
+      case NotForwardDeclarableReason::StdNamespace:
+        S.Diag(KernelInvocationFuncLoc,
+               diag::err_invalid_std_type_in_sycl_kernel)
+            << KernelNameType << DeclNamed;
+        IsInvalid = true;
+        return;
+      case NotForwardDeclarableReason::UnnamedTag:
+        S.Diag(KernelInvocationFuncLoc, diag::err_sycl_kernel_incorrectly_named)
+            << /* unnamed type is invalid */ 2 << KernelNameType;
+        IsInvalid = true;
+        return;
+      case NotForwardDeclarableReason::NotAtNamespaceScope:
+        S.Diag(KernelInvocationFuncLoc, diag::err_sycl_kernel_incorrectly_named)
+            << /* kernel name should be forward declarable at namespace
+                  scope */
+            0 << KernelNameType;
+        IsInvalid = true;
+        return;
+      case NotForwardDeclarableReason::None:
+      default:
+        // Do nothing, we're fine.
+        break;
       }
     }
   }
@@ -5438,6 +5557,7 @@ void SemaSYCL::MarkDevices() {
 
 void SemaSYCL::ProcessFreeFunction(FunctionDecl *FD) {
   if (isFreeFunction(*this, FD)) {
+    SyclKernelDecompMarker DecompMarker(*this);
     SyclKernelFieldChecker FieldChecker(*this);
     SyclKernelUnionChecker UnionChecker(*this);
 
@@ -5446,7 +5566,8 @@ void SemaSYCL::ProcessFreeFunction(FunctionDecl *FD) {
     DiagnosingSYCLKernel = true;
 
     // Check parameters of free function.
-    Visitor.VisitFunctionParameters(FD, FieldChecker, UnionChecker);
+    Visitor.VisitFunctionParameters(FD, DecompMarker, FieldChecker,
+                                    UnionChecker);
 
     DiagnosingSYCLKernel = false;
 
@@ -5892,6 +6013,14 @@ class SYCLFwdDeclEmitter
   void VisitPackTemplateArgument(const TemplateArgument &TA) {
     VisitTemplateArgs(TA.getPackAsArray());
   }
+
+  void VisitFunctionProtoType(const FunctionProtoType *T) {
+    for (const auto Ty : T->getParamTypes())
+      Visit(Ty.getCanonicalType());
+    // So far this visitor method is only used for free function kernels whose
+    // return type is void anyway, so it is not visited. Otherwise, add if
+    // required.
+  }
 };
 
 class SYCLKernelNameTypePrinter
@@ -6328,10 +6457,15 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
   for (const KernelDesc &K : KernelDescs) {
     if (!isFreeFunction(S, K.SyclKernel))
       continue;
-
     ++FreeFunctionCount;
     // Generate forward declaration for free function.
     O << "\n// Definition of " << K.Name << " as a free function kernel\n";
+
+    O << "\n";
+    O << "// Forward declarations of kernel and its argument types:\n";
+    FwdDeclEmitter.Visit(K.SyclKernel->getType());
+    O << "\n";
+
     if (K.SyclKernel->getLanguageLinkage() == CLanguageLinkage)
       O << "extern \"C\" ";
     std::string ParmList;
diff --git a/clang/test/CodeGenSYCL/free_function_int_header.cpp b/clang/test/CodeGenSYCL/free_function_int_header.cpp
index ab61b06d66728..147b4ff903418 100755
--- a/clang/test/CodeGenSYCL/free_function_int_header.cpp
+++ b/clang/test/CodeGenSYCL/free_function_int_header.cpp
@@ -1,8 +1,8 @@
 // RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -sycl-std=2020 -fsycl-int-header=%t.h %s
 // RUN: FileCheck -input-file=%t.h %s
 // 
-// This test checks integration header contents for free functions with scalar
-// and pointer parameters.
+// This test checks integration header contents for free functions with scalar,
+// pointer and non-decomposed struct parameters.
 
 #include "mock_properties.hpp"
 #include "sycl.hpp"
@@ -43,12 +43,47 @@ template <> void ff_3<double>(double *ptr, double start, double end) {
     ptr[i] = end;
 }
 
+struct NoPointers {
+  int f;
+};
+
+struct Pointers {
+  int * a;
+  float * b;
+};
+
+struct Agg {
+  NoPointers F1;
+  int F2;
+  int *F3;
+  Pointers F4;
+};
+
+struct Derived : Agg {
+  int a;
+};
+
+__attribute__((sycl_device))
+[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 0)]]
+void ff_4(NoPointers S1, Pointers S2, Agg S3) {
+}
+
+template <typename T1, typename T2>
+__attribute__((sycl_device))
+[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 0)]]
+  void ff_6(T1 S1, T2 S2, int end) {
+}
+
+template void ff_6(Agg S1, Derived S2, int);
+
 // CHECK:      const char* const kernel_names[] = {
 // CHECK-NEXT:   {{.*}}__sycl_kernel_ff_2Piii
 // CHECK-NEXT:   {{.*}}__sycl_kernel_ff_2Piiii
 // CHECK-NEXT:   {{.*}}__sycl_kernel_ff_3IiEvPT_S0_S0_
 // CHECK-NEXT:   {{.*}}__sycl_kernel_ff_3IfEvPT_S0_S0_
 // CHECK-NEXT:   {{.*}}__sycl_kernel_ff_3IdEvPT_S0_S0_
+// CHECK-NEXT:   {{.*}}__sycl_kernel_ff_410NoPointers8Pointers3Agg
+// CHECK-NEXT:   {{.*}}__sycl_kernel_ff_6I3Agg7DerivedEvT_T0_i
 // CHECK-NEXT:   ""
 // CHECK-NEXT: };
 
@@ -79,11 +114,22 @@ template <> void ff_3<double>(double *ptr, double start, double end) {
 // CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 8, 8 },
 // CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 8, 16 },
 
+// CHECK:  //--- _Z18__sycl_kernel_ff_410NoPointers8Pointers3Agg
+// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 0 },
+// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 16, 4 },
+// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 32, 20 },
+
+// CHECK:  //--- _Z18__sycl_kernel_ff_6I3Agg7DerivedEvT_T0_i
+// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 32, 0 },
+// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 40, 32 },
+// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 72 },
+
 // CHECK:        { kernel_param_kind_t::kind_invalid, -987654321, -987654321 },
 // CHECK-NEXT: };
 
 // CHECK: Definition of _Z18__sycl_kernel_ff_2Piii as a free function kernel
-// CHECK-NEXT: void ff_2(int *ptr, int start, int end);
+// CHECK: Forward declarations of kernel and its argument types:
+// CHECK: void ff_2(int *ptr, int start, int end);
 // CHECK-NEXT: static constexpr auto __sycl_shim1() {
 // CHECK-NEXT:   return (void (*)(int *, int, int))ff_2;
 // CHECK-NEXT: }
@@ -99,7 +145,8 @@ template <> void ff_3<double>(double *ptr, double start, double end) {
 // CHECK-NEXT: }
 
 // CHECK: Definition of _Z18__sycl_kernel_ff_2Piiii as a free function kernel
-// CHECK-NEXT: void ff_2(int *ptr, int start, int end, int value);
+// CHECK: Forward declarations of kernel and its argument types:
+// CHECK: void ff_2(int *ptr, int start, int end, int value);
 // CHECK-NEXT: static constexpr auto __sycl_shim2() {
 // CHECK-NEXT:   return (void (*)(int *, int, int, int))ff_2;
 // CHECK-NEXT: }
@@ -115,7 +162,8 @@ template <> void ff_3<double>(double *ptr, double start, double end) {
 // CHECK-NEXT: }
 
 // CHECK: Definition of _Z18__sycl_kernel_ff_3IiEvPT_S0_S0_ as a free function kernel
-// CHECK-NEXT: template <typename T> void ff_3(T *ptr, T start, T end);
+// CHECK: Forward declarations of kernel and its argument types:
+// CHECK: template <typename T> void ff_3(T *ptr, T start, T end);
 // CHECK-NEXT: static constexpr auto __sycl_shim3() {
 // CHECK-NEXT:   return (void (*)(int *, int, int))ff_3<int>;
 // CHECK-NEXT: }
@@ -129,9 +177,10 @@ template <> void ff_3<double>(double *ptr, double start, double end) {
 // CHECK-NEXT:     static constexpr bool value = true;
 // CHECK-NEXT:   };
 // CHECK-NEXT: }
- 
+
 // CHECK: Definition of _Z18__sycl_kernel_ff_3IfEvPT_S0_S0_ as a free function kernel
-// CHECK-NEXT: template <typename T> void ff_3(T *ptr, T start, T end);
+// CHECK: Forward declarations of kernel and its argument types:
+// CHECK: template <typename T> void ff_3(T *ptr, T start, T end);
 // CHECK-NEXT: static constexpr auto __sycl_shim4() {
 // CHECK-NEXT:   return (void (*)(float *, float, float))ff_3<float>;
 // CHECK-NEXT: }
@@ -147,7 +196,8 @@ template <> void ff_3<double>(double *ptr, double start, double end) {
 // CHECK-NEXT: }
 
 // CHECK: Definition of _Z18__sycl_kernel_ff_3IdEvPT_S0_S0_ as a free function kernel
-// CHECK-NEXT: template <typename T> void ff_3(T *ptr, T start, T end);
+// CHECK: Forward declarations of kernel and its argument types:
+// CHECK: template <typename T> void ff_3(T *ptr, T start, T end);
 // CHECK-NEXT: static constexpr auto __sycl_shim5() {
 // CHECK-NEXT:   return (void (*)(double *, double, double))ff_3<double>;
 // CHECK-NEXT: }
@@ -162,6 +212,44 @@ template <> void ff_3<double>(double *ptr, double start, double end) {
 // CHECK-NEXT:   };
 // CHECK-NEXT: }
 
+// CHECK: Definition of _Z18__sycl_kernel_ff_410NoPointers8Pointers3Agg as a free function kernel
+// CHECK: Forward declarations of kernel and its argument types:
+// CHECK-NEXT: struct NoPointers;
+// CHECK-NEXT: struct Pointers;
+// CHECK-NEXT: struct Agg;
+// CHECK: void ff_4(NoPointers S1, Pointers S2, Agg S3);
+// CHECK-NEXT: static constexpr auto __sycl_shim6() {
+// CHECK-NEXT:   return (void (*)(struct NoPointers, struct Pointers, struct Agg))ff_4;
+// CHECK-NEXT: }
+// CHECK-NEXT: namespace sycl {
+// CHECK-NEXT: template <>
+// CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim6()> {
+// CHECK-NEXT:   static constexpr bool value = true;
+// CHECK-NEXT: };
+// CHECK-NEXT: template <>
+// CHECK-NEXT: struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim6()> {
+// CHECK-NEXT:   static constexpr bool value = true;
+// CHECK-NEXT: };
+// CHECK-NEXT: }
+
+// CHECK: Definition of _Z18__sycl_kernel_ff_6I3Agg7DerivedEvT_T0_i as a free function kernel
+// CHECK: Forward declarations of kernel and its argument types:
+// CHECK: struct Derived;
+// CHECK: template <typename T1, typename T2> void ff_6(T1 S1, T2 S2, int end);
+// CHECK-NEXT: static constexpr auto __sycl_shim7() {
+// CHECK-NEXT:   return (void (*)(struct Agg, struct Derived, int))ff_6<struct Agg, struct Derived>;
+// CHECK-NEXT: }
+// CHECK-NEXT: namespace sycl {
+// CHECK-NEXT: template <>
+// CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim7()> {
+// CHECK-NEXT:   static constexpr bool value = true;
+// CHECK-NEXT: };
+// CHECK-NEXT: template <>
+// CHECK-NEXT: struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim7()> {
+// CHECK-NEXT:   static constexpr bool value = true;
+// CHECK-NEXT: };
+// CHECK-NEXT: }
+
 // CHECK: #include <sycl/kernel_bundle.hpp>
 
 // CHECK: Definition of kernel_id of _Z18__sycl_kernel_ff_2Piii
@@ -201,5 +289,21 @@ template <> void ff_3<double>(double *ptr, double start, double end) {
 // CHECK-NEXT:   template <>
 // CHECK-NEXT:   kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim5()>() {
 // CHECK-NEXT:     return sycl::detail::get_kernel_id_impl(std::string_view{"_Z18__sycl_kernel_ff_3IdEvPT_S0_S0_"});
-// CHECK-NEXT:   }
+// CHECK-NEXT: }
+// CHECK-NEXT: }
+
+// CHECK: Definition of kernel_id of _Z18__sycl_kernel_ff_410NoPointers8Pointers3Agg
+// CHECK-NEXT: namespace sycl {
+// CHECK-NEXT: template <>
+// CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim6()>() {
+// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{"_Z18__sycl_kernel_ff_410NoPointers8Pointers3Agg"});
+// CHECK-NEXT: }
+// CHECK-NEXT: }
+
+// CHECK: Definition of kernel_id of _Z18__sycl_kernel_ff_6I3Agg7DerivedEvT_T0_i
+// CHECK-NEXT: namespace sycl {
+// CHECK-NEXT: template <>
+// CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim7()>() {
+// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{"_Z18__sycl_kernel_ff_6I3Agg7DerivedEvT_T0_i"});
+// CHECK-NEXT: }
 // CHECK-NEXT: }
diff --git a/clang/test/CodeGenSYCL/free_function_kernel_params.cpp b/clang/test/CodeGenSYCL/free_function_kernel_params.cpp
new file mode 100755
index 0000000000000..97a8d29b95bc2
--- /dev/null
+++ b/clang/test/CodeGenSYCL/free_function_kernel_params.cpp
@@ -0,0 +1,35 @@
+// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -triple spir64 \
+// RUN: -emit-llvm %s -o - | FileCheck %s
+// This test checks parameter IR generation for free functions with parameters
+// of non-decomposed struct type.
+
+#include "sycl.hpp"
+
+struct NoPointers {
+  int f;
+};
+
+struct Pointers {
+  int * a;
+  float * b;
+};
+
+struct Agg {
+  NoPointers F1;
+  int F2;
+  int *F3;
+  Pointers F4;
+};
+
+__attribute__((sycl_device))
+[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 0)]]
+void ff_4(NoPointers S1, Pointers S2, Agg S3) {
+}
+
+// CHECK: %struct.NoPointers = type { i32 }
+// CHECK: %struct.Pointers = type { ptr addrspace(4), ptr addrspace(4) }
+// CHECK: %struct.Agg = type { %struct.NoPointers, i32, ptr addrspace(4), %struct.Pointers }
+// CHECK: %struct.__generated_Pointers = type { ptr addrspace(1), ptr addrspace(1) }
+// CHECK: %struct.__generated_Agg = type { %struct.NoPointers, i32, ptr addrspace(1), %struct.__generated_Pointers.0 }
+// CHECK: %struct.__generated_Pointers.0 = type { ptr addrspace(1), ptr addrspace(1) }
+// CHECK: define dso_local spir_kernel void @{{.*}}__sycl_kernel{{.*}}(ptr noundef byval(%struct.NoPointers) align 4 %__arg_S1, ptr noundef byval(%struct.__generated_Pointers) align 8 %__arg_S2, ptr noundef byval(%struct.__generated_Agg) align 8 %__arg_S3)
diff --git a/clang/test/SemaSYCL/free_function_kernel_params.cpp b/clang/test/SemaSYCL/free_function_kernel_params.cpp
index 6f73d6f172aa6..2de4f896a1513 100755
--- a/clang/test/SemaSYCL/free_function_kernel_params.cpp
+++ b/clang/test/SemaSYCL/free_function_kernel_params.cpp
@@ -1,7 +1,7 @@
 // RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -ast-dump \
 // RUN: %s -o - | FileCheck %s
 // This test checks parameter rewriting for free functions with parameters
-// of type scalar and pointer.
+// of type scalar, pointer and non-decomposed struct.
 
 #include "sycl.hpp"
 
@@ -55,3 +55,119 @@ template void ff_3(int* ptr, int start, int end);
 // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '__arg_start' 'int'
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
 // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '__arg_end' 'int'
+
+struct NoPointers {
+  int f;
+};
+
+struct Pointers {
+  int * a;
+  float * b;
+};
+
+struct Agg {
+  NoPointers F1;
+  int F2;
+  int *F3;
+  Pointers F4;
+};
+
+struct Agg1 {
+  NoPointers F1;
+  int F2;
+};
+
+struct Derived : Agg {
+  int a;
+};
+
+class Derived1 : Pointers {
+  int a;
+};
+
+__attribute__((sycl_device))
+[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 0)]]
+void ff_4(NoPointers S1, Pointers S2, Agg S3) {
+}
+// CHECK: FunctionDecl {{.*}}__sycl_kernel{{.*}}'void (NoPointers, __generated_Pointers, __generated_Agg)'
+// CHECK-NEXT: ParmVarDecl {{.*}} used __arg_S1 'NoPointers'
+// CHECK-NEXT: ParmVarDecl {{.*}} used __arg_S2 '__generated_Pointers'
+// CHECK-NEXT: ParmVarDecl {{.*}} used __arg_S3 '__generated_Agg'
+// CHECK-NEXT: CompoundStmt
+// CHECK-NEXT: CallExpr
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'void (*)(NoPointers, Pointers, Agg)' <FunctionToPointerDecay>
+// CHECK-NEXT: DeclRefExpr {{.*}} 'void (NoPointers, Pointers, Agg)' lvalue Function {{.*}} 'ff_4' 'void (NoPointers, Pointers, Agg)'
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'NoPointers' 'void (const NoPointers &) noexcept'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const NoPointers' lvalue <NoOp>
+// CHECK-NEXT: DeclRefExpr {{.*}} 'NoPointers' lvalue ParmVar {{.*}} '__arg_S1' 'NoPointers'
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'Pointers' 'void (const Pointers &) noexcept'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const Pointers' lvalue <NoOp>
+// CHECK-NEXT: UnaryOperator {{.*}} 'Pointers' lvalue prefix '*' cannot overflow
+// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'Pointers *' reinterpret_cast<Pointers *> <BitCast>
+// CHECK-NEXT: UnaryOperator {{.*}} '__generated_Pointers *' prefix '&' cannot overflow
+// CHECK-NEXT: DeclRefExpr {{.*}} '__generated_Pointers' lvalue ParmVar {{.*}} '__arg_S2' '__generated_Pointers'
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'Agg' 'void (const Agg &) noexcept'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const Agg' lvalue <NoOp>
+// CHECK-NEXT: UnaryOperator {{.*}} 'Agg' lvalue prefix '*' cannot overflow
+// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'Agg *' reinterpret_cast<Agg *> <BitCast>
+// CHECK-NEXT: UnaryOperator {{.*}} '__generated_Agg *' prefix '&' cannot overflow
+// CHECK-NEXT: DeclRefExpr {{.*}} '__generated_Agg' lvalue ParmVar {{.*}} '__arg_S3' '__generated_Agg'
+
+__attribute__((sycl_device))
+[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 0)]]
+void ff_5(Agg1 S1, Derived S2, Derived1 S3) {
+}
+// CHECK: FunctionDecl {{.*}}__sycl_kernel{{.*}}'void (Agg1, __generated_Derived, __generated_Derived1)'
+// CHECK-NEXT: ParmVarDecl {{.*}} used __arg_S1 'Agg1'
+// CHECK-NEXT: ParmVarDecl {{.*}} used __arg_S2 '__generated_Derived'
+// CHECK-NEXT: ParmVarDecl {{.*}} used __arg_S3 '__generated_Derived1'
+// CHECK-NEXT: CompoundStmt
+// CHECK-NEXT: CallExpr
+// CHECK-NEXT: ImplicitCastExpr{{.*}}'void (*)(Agg1, Derived, Derived1)' <FunctionToPointerDecay>
+// CHECK-NEXT: DeclRefExpr{{.*}}'void (Agg1, Derived, Derived1)' lvalue Function {{.*}} 'ff_5' 'void (Agg1, Derived, Derived1)'
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'Agg1' 'void (const Agg1 &) noexcept'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const Agg1' lvalue <NoOp>
+// CHECK-NEXT: DeclRefExpr {{.*}} 'Agg1' lvalue ParmVar {{.*}} '__arg_S1' 'Agg1'
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'Derived' 'void (const Derived &) noexcept'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const Derived' lvalue <NoOp>
+// CHECK-NEXT: UnaryOperator {{.*}} 'Derived' lvalue prefix '*' cannot overflow
+// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'Derived *' reinterpret_cast<Derived *> <BitCast>
+// CHECK-NEXT: UnaryOperator {{.*}} '__generated_Derived *' prefix '&' cannot overflow
+// CHECK-NEXT: DeclRefExpr {{.*}} '__generated_Derived' lvalue ParmVar {{.*}} '__arg_S2' '__generated_Derived'
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'Derived1' 'void (const Derived1 &) noexcept'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const Derived1' lvalue <NoOp>
+// CHECK-NEXT: UnaryOperator {{.*}} 'Derived1' lvalue prefix '*' cannot overflow
+// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'Derived1 *' reinterpret_cast<Derived1 *> <BitCast>
+// CHECK-NEXT: UnaryOperator {{.*}} '__generated_Derived1 *' prefix '&' cannot overflow
+// CHECK-NEXT: DeclRefExpr {{.*}} '__generated_Derived1' lvalue ParmVar {{.*}} '__arg_S3' '__generated_Derived1'
+
+template <typename T1, typename T2>
+__attribute__((sycl_device))
+[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 0)]]
+  void ff_6(T1 S1, T2 S2, int end) {
+}
+
+// Explicit instantiation.
+template void ff_6(Agg S1, Derived1 S2, int);
+// CHECK: FunctionDecl {{.*}}__sycl_kernel{{.*}}'void (__generated_Agg, __generated_Derived1, int)'
+// CHECK-NEXT: ParmVarDecl {{.*}} used __arg_S1 '__generated_Agg'
+// CHECK-NEXT: ParmVarDecl {{.*}} used __arg_S2 '__generated_Derived1'
+// CHECK-NEXT: ParmVarDecl {{.*}} used __arg_end 'int'
+// CHECK-NEXT: CompoundStmt
+// CHECK-NEXT: CallExpr {{.*}} 'void'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'void (*)(Agg, Derived1, int)' <FunctionToPointerDecay>
+// CHECK-NEXT: DeclRefExpr {{.*}} 'void (Agg, Derived1, int)' lvalue Function {{.*}} 'ff_6' 'void (Agg, Derived1, int)'
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'Agg' 'void (const Agg &) noexcept'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const Agg' lvalue <NoOp>
+// CHECK-NEXT: UnaryOperator {{.*}} 'Agg' lvalue prefix '*' cannot overflow
+// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'Agg *' reinterpret_cast<struct Agg *> <BitCast>
+// CHECK-NEXT: UnaryOperator {{.*}} '__generated_Agg *' prefix '&' cannot overflow
+// CHECK-NEXT: DeclRefExpr {{.*}} '__generated_Agg' lvalue ParmVar {{.*}} '__arg_S1' '__generated_Agg'
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'Derived1' 'void (const Derived1 &) noexcept'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const Derived1' lvalue <NoOp>
+// CHECK-NEXT: UnaryOperator {{.*}} 'Derived1' lvalue prefix '*' cannot overflow
+// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'Derived1 *' reinterpret_cast<class Derived1 *> <BitCast>
+// CHECK-NEXT: UnaryOperator {{.*}} '__generated_Derived1 *' prefix '&' cannot overflow
+// CHECK-NEXT: DeclRefExpr {{.*}} '__generated_Derived1' lvalue ParmVar {{.*}} '__arg_S2' '__generated_Derived1'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '__arg_end' 'int'
diff --git a/clang/test/SemaSYCL/free_function_kernel_params_restrictions.cpp b/clang/test/SemaSYCL/free_function_kernel_params_restrictions.cpp
new file mode 100755
index 0000000000000..d1bdc0e3da475
--- /dev/null
+++ b/clang/test/SemaSYCL/free_function_kernel_params_restrictions.cpp
@@ -0,0 +1,61 @@
+// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -verify %s
+// This test checks that compiler correctly diagnoses violations of restrictions
+// applied to free function kernel parameters defined by the spec.
+
+#include "sycl.hpp"
+
+class Outer {
+public:
+  class DefinedWithinAClass {
+    int f;
+  };
+};
+
+__attribute__((sycl_device))
+[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 0)]]
+void ff_4(Outer::DefinedWithinAClass S1) { // expected-error {{'Outer::DefinedWithinAClass' cannot be used as the type of a kernel parameter}}
+                                           // expected-note@-1 {{'Outer::DefinedWithinAClass' is not forward declarable}}
+}
+
+template <typename T1>
+__attribute__((sycl_device))
+[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 0)]]
+  void ff_6(T1 S1) { // expected-error 2{{cannot be used as the type of a kernel parameter}}
+                     // expected-note@-1 2{{is not forward declarable}}
+}
+
+void bar() {
+  ff_6([=](){});
+}
+
+auto Glob = [](int P){ return P + 1;};
+
+template void ff_6(typeof(Glob) S1);
+
+extern "C" {
+  struct A {
+    int a;
+  };
+}
+
+__attribute__((sycl_device))
+[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 0)]]
+void ff_5(A S1) {
+}
+
+
+
+struct StructWithAccessor {
+  sycl::accessor<char, 1, sycl::access::mode::read> acc;
+  int *ptr;
+};
+
+struct Wrapper {
+  StructWithAccessor SWA;
+
+};
+
+[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 0)]]
+void ff_6(Wrapper S1) { // expected-error {{cannot be used as the type of a kernel parameter}}
+                        // expected-note@-1 {{'Wrapper' is not yet supported as a free function kernel parameter}}
+}
diff --git a/clang/test/SemaSYCL/kernelname-enum.cpp b/clang/test/SemaSYCL/kernelname-enum.cpp
index 34c201cd7f1c2..726d339c84e16 100644
--- a/clang/test/SemaSYCL/kernelname-enum.cpp
+++ b/clang/test/SemaSYCL/kernelname-enum.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -sycl-std=2020 -verify %s
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -sycl-std=2020 -fno-sycl-unnamed-lambda -verify %s
 
 // This test verifies that kernel names containing unscoped enums are diagnosed correctly.
 
diff --git a/sycl/test-e2e/KernelAndProgram/free_function_kernels.cpp b/sycl/test-e2e/KernelAndProgram/free_function_kernels.cpp
index 36974cd372fd3..c260a281f596c 100644
--- a/sycl/test-e2e/KernelAndProgram/free_function_kernels.cpp
+++ b/sycl/test-e2e/KernelAndProgram/free_function_kernels.cpp
@@ -257,6 +257,120 @@ bool test_3(queue Queue) {
   return PassA && PassB;
 }
 
+struct KArg {
+  KArg(int *_p, int _s) : ptr(_p), start(_s) {}
+  int *ptr;
+  int start;
+};
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(
+    (ext::oneapi::experimental::nd_range_kernel<2>))
+void ff_4(KArg Arg) {
+  int(&ptr2D)[4][4] = *reinterpret_cast<int(*)[4][4]>(Arg.ptr);
+  nd_item<2> Item = ext::oneapi::this_work_item::get_nd_item<2>();
+  id<2> GId = Item.get_global_id();
+  id<2> LId = Item.get_local_id();
+  ptr2D[GId.get(0)][GId.get(1)] = LId.get(0) + LId.get(1) + Arg.start;
+}
+
+bool test_4(queue Queue) {
+  constexpr int Range = 16;
+  int *usmPtr = malloc_shared<int>(Range, Queue);
+  int value = 55;
+  int Result[Range] = {55, 56, 55, 56, 56, 57, 56, 57,
+                       55, 56, 55, 56, 56, 57, 56, 57};
+  nd_range<2> R2{range<2>{4, 4}, range<2>{2, 2}};
+
+  memset(usmPtr, 0, Range * sizeof(int));
+  Queue.submit([&](handler &Handler) {
+    Handler.parallel_for(R2, [=](nd_item<2> Item) {
+      int(&ptr2D)[4][4] = *reinterpret_cast<int(*)[4][4]>(usmPtr);
+      id<2> GId = Item.get_global_id();
+      id<2> LId = Item.get_local_id();
+      ptr2D[GId.get(0)][GId.get(1)] = LId.get(0) + LId.get(1) + value;
+    });
+  });
+  Queue.wait();
+  bool PassA = checkUSM(usmPtr, Range, Result);
+  std::cout << "Test 4a: " << (PassA ? "PASS" : "FAIL") << std::endl;
+
+  bool PassB = false;
+#ifndef __SYCL_DEVICE_ONLY__
+  kernel_bundle Bundle =
+      get_kernel_bundle<bundle_state::executable>(Queue.get_context());
+  kernel_id Kernel_id =
+      ext::oneapi::experimental::get_kernel_id<(void (*)(KArg))ff_4>();
+  kernel Kernel = Bundle.get_kernel(Kernel_id);
+  memset(usmPtr, 0, Range * sizeof(int));
+  Queue.submit([&](handler &Handler) {
+    Handler.set_arg(0, KArg(usmPtr, value));
+    Handler.parallel_for(R2, Kernel);
+  });
+  Queue.wait();
+  PassB = checkUSM(usmPtr, Range, Result);
+  std::cout << "Test 4b: " << (PassB ? "PASS" : "FAIL") << std::endl;
+
+  free(usmPtr, Queue);
+#endif
+  return PassA && PassB;
+}
+
+template <typename T>
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(
+    (ext::oneapi::experimental::nd_range_kernel<2>))
+void ff_5(T Arg) {
+  int(&ptr2D)[4][4] = *reinterpret_cast<int(*)[4][4]>(Arg.ptr);
+  nd_item<2> Item = ext::oneapi::this_work_item::get_nd_item<2>();
+  id<2> GId = Item.get_global_id();
+  id<2> LId = Item.get_local_id();
+  ptr2D[GId.get(0)][GId.get(1)] = LId.get(0) + LId.get(1) + Arg.start;
+}
+
+// Explicit instantiation with "KArg".
+template void ff_5(KArg Arg);
+
+bool test_5(queue Queue) {
+  constexpr int Range = 16;
+  int *usmPtr = malloc_shared<int>(Range, Queue);
+  int value = 55;
+  int Result[Range] = {55, 56, 55, 56, 56, 57, 56, 57,
+                       55, 56, 55, 56, 56, 57, 56, 57};
+  nd_range<2> R2{range<2>{4, 4}, range<2>{2, 2}};
+
+  memset(usmPtr, 0, Range * sizeof(int));
+  Queue.submit([&](handler &Handler) {
+    Handler.parallel_for(R2, [=](nd_item<2> Item) {
+      int(&ptr2D)[4][4] = *reinterpret_cast<int(*)[4][4]>(usmPtr);
+      id<2> GId = Item.get_global_id();
+      id<2> LId = Item.get_local_id();
+      ptr2D[GId.get(0)][GId.get(1)] = LId.get(0) + LId.get(1) + value;
+    });
+  });
+  Queue.wait();
+  bool PassA = checkUSM(usmPtr, Range, Result);
+  std::cout << "Test 5a: " << (PassA ? "PASS" : "FAIL") << std::endl;
+
+  bool PassB = false;
+#ifndef __SYCL_DEVICE_ONLY__
+  kernel_bundle Bundle =
+      get_kernel_bundle<bundle_state::executable>(Queue.get_context());
+  kernel_id Kernel_id =
+      ext::oneapi::experimental::get_kernel_id<(void (*)(KArg))ff_5<KArg>>();
+  kernel Kernel = Bundle.get_kernel(Kernel_id);
+  memset(usmPtr, 0, Range * sizeof(int));
+  Queue.submit([&](handler &Handler) {
+    Handler.set_arg(0, KArg(usmPtr, value));
+    Handler.parallel_for(R2, Kernel);
+  });
+  Queue.wait();
+  PassB = checkUSM(usmPtr, Range, Result);
+  std::cout << "Test 5b: " << (PassB ? "PASS" : "FAIL") << std::endl;
+
+  free(usmPtr, Queue);
+#endif
+  return PassA && PassB;
+}
+
 int main() {
   queue Queue;
 
@@ -265,6 +379,8 @@ int main() {
   Pass &= test_1(Queue);
   Pass &= test_2(Queue);
   Pass &= test_3(Queue);
+  Pass &= test_4(Queue);
+  Pass &= test_5(Queue);
 
   return Pass ? 0 : 1;
 }
diff --git a/sycl/test/extensions/free_function_errors.cpp b/sycl/test/extensions/free_function_errors.cpp
index 344bf17bdcc32..1838ec9ec072e 100755
--- a/sycl/test/extensions/free_function_errors.cpp
+++ b/sycl/test/extensions/free_function_errors.cpp
@@ -25,7 +25,6 @@ union U {
 
 using accType = accessor<int, 1, access::mode::read_write>;
 
-// expected-error@+3 {{'struct S' cannot be used as the type of a kernel parameter}}
 SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(
     (ext::oneapi::experimental::single_task_kernel))
 void ff(struct S s) {}
@@ -40,7 +39,6 @@ SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(
     (ext::oneapi::experimental::single_task_kernel))
 void ff(accType acc) {}
 
-// expected-error@+3 {{'std::array<int, 10>' cannot be used as the type of a kernel parameter}}
 SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(
     (ext::oneapi::experimental::single_task_kernel))
 void ff(std::array<int, 10> a) {}