diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
index c0a5f90e79491..ce58dff5833fd 100644
--- a/clang/include/clang/Basic/Attr.td
+++ b/clang/include/clang/Basic/Attr.td
@@ -1198,6 +1198,14 @@ def SYCLIntelBufferLocation : InheritableAttr {
   let Documentation = [Undocumented];
 }
 
+def SYCLRequiresDecomposition : InheritableAttr {
+  // No spellings, as this is for internal use.
+  let Spellings = [];
+  let Subjects = SubjectList<[Named]>;
+  let LangOpts = [SYCLIsDevice, SYCLIsHost];
+  let Documentation = [Undocumented];
+}
+
 def SYCLIntelKernelArgsRestrict : InheritableAttr {
   let Spellings = [ CXX11<"intel", "kernel_args_restrict"> ];
   let Subjects = SubjectList<[Function], ErrorDiag>;
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index adb982ad24bd0..800e12ba6ef84 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -828,9 +828,9 @@ class KernelObjVisitor {
   // type (which doesn't exist in cases where it is a FieldDecl in the
   // 'root'), and Wrapper is the current struct being unwrapped.
   template <typename ParentTy, typename... HandlerTys>
-  void visitRecord(const CXXRecordDecl *Owner, ParentTy &Parent,
-                   const CXXRecordDecl *Wrapper, QualType RecordTy,
-                   HandlerTys &... Handlers) {
+  void visitComplexRecord(const CXXRecordDecl *Owner, ParentTy &Parent,
+                          const CXXRecordDecl *Wrapper, QualType RecordTy,
+                          HandlerTys &... Handlers) {
     (void)std::initializer_list<int>{
         (Handlers.enterStruct(Owner, Parent, RecordTy), 0)...};
     VisitRecordHelper(Wrapper, Wrapper->bases(), Handlers...);
@@ -839,6 +839,19 @@ class KernelObjVisitor {
         (Handlers.leaveStruct(Owner, Parent, RecordTy), 0)...};
   }
 
+  template <typename ParentTy, typename... HandlerTys>
+  void visitSimpleRecord(const CXXRecordDecl *Owner, ParentTy &Parent,
+                         const CXXRecordDecl *Wrapper, QualType RecordTy,
+                         HandlerTys &... Handlers) {
+    (void)std::initializer_list<int>{
+        (Handlers.handleNonDecompStruct(Owner, Parent, RecordTy), 0)...};
+  }
+
+  template <typename ParentTy, typename... HandlerTys>
+  void visitRecord(const CXXRecordDecl *Owner, ParentTy &Parent,
+                   const CXXRecordDecl *Wrapper, QualType RecordTy,
+                   HandlerTys &... Handlers);
+
   template <typename ParentTy, typename... HandlerTys>
   void VisitUnion(const CXXRecordDecl *Owner, ParentTy &Parent,
                   const CXXRecordDecl *Wrapper, HandlerTys &... Handlers);
@@ -908,8 +921,15 @@ class KernelObjVisitor {
                             HandlerTys &... Handlers);
 
   template <typename... HandlerTys>
-  void visitArray(const CXXRecordDecl *Owner, FieldDecl *Field,
-                  QualType ArrayTy, HandlerTys &... Handlers) {
+  void visitSimpleArray(const CXXRecordDecl *Owner, FieldDecl *Field,
+                        QualType ArrayTy, HandlerTys &... Handlers) {
+    (void)std::initializer_list<int>{
+        (Handlers.handleSimpleArrayType(Field, ArrayTy), 0)...};
+  }
+
+  template <typename... HandlerTys>
+  void visitComplexArray(const CXXRecordDecl *Owner, FieldDecl *Field,
+                         QualType ArrayTy, HandlerTys &... Handlers) {
     // Array workflow is:
     // handleArrayType
     // enterArray
@@ -940,6 +960,10 @@ class KernelObjVisitor {
         (Handlers.leaveArray(Field, ArrayTy, ET), 0)...};
   }
 
+  template <typename... HandlerTys>
+  void visitArray(const CXXRecordDecl *Owner, FieldDecl *Field,
+                  QualType ArrayTy, HandlerTys &... Handlers);
+
   template <typename... HandlerTys>
   void visitField(const CXXRecordDecl *Owner, FieldDecl *Field,
                   QualType FieldTy, HandlerTys &... Handlers) {
@@ -1003,6 +1027,10 @@ class SyclKernelFieldHandlerBase {
 public:
   static constexpr const bool VisitUnionBody = false;
   static constexpr const bool VisitNthArrayElement = true;
+  // Opt-in based on whether we should visit inside simple containers (structs,
+  // arrays). All of the 'check' types should likely be true, the int-header,
+  // and kernel decl creation types should not.
+  static constexpr const bool VisitInsideSimpleContainers = true;
   // Mark these virtual so that we can use override in the implementer classes,
   // despite virtual dispatch never being used.
 
@@ -1039,6 +1067,23 @@ class SyclKernelFieldHandlerBase {
   // Most handlers shouldn't be handling this, just the field checker.
   virtual bool handleOtherType(FieldDecl *, QualType) { return true; }
 
+  // Handle a simple struct that doesn't need to be decomposed, only called on
+  // handlers with VisitInsideSimpleContainers as false.  Replaces
+  // handleStructType, enterStruct, leaveStruct, and visiting of sub-elements.
+  virtual bool handleNonDecompStruct(const CXXRecordDecl *, FieldDecl *,
+                                     QualType) {
+    return true;
+  }
+  virtual bool handleNonDecompStruct(const CXXRecordDecl *,
+                                     const CXXBaseSpecifier &, QualType) {
+    return true;
+  }
+
+  // Instead of handleArrayType, enterArray, leaveArray, and nextElement (plus
+  // descending down the elements), this function gets called in the event of an
+  // array containing simple elements (even in the case of an MD array).
+  virtual bool handleSimpleArrayType(FieldDecl *, QualType) { return true; }
+
   // The following are only used for keeping track of where we are in the base
   // class/field graph. Int Headers use this to calculate offset, most others
   // don't have a need for these.
@@ -1110,6 +1155,14 @@ template <bool B, bool... Rest> struct AnyTrue {
   static constexpr bool Value = B || AnyTrue<Rest...>::Value;
 };
 
+template <bool B, bool... Rest> struct AllTrue;
+
+template <bool B> struct AllTrue<B> { static constexpr bool Value = B; };
+
+template <bool B, bool... Rest> struct AllTrue {
+  static constexpr bool Value = B && AllTrue<Rest...>::Value;
+};
+
 template <typename ParentTy, typename... Handlers>
 void KernelObjVisitor::VisitUnion(const CXXRecordDecl *Owner, ParentTy &Parent,
                                   const CXXRecordDecl *Wrapper,
@@ -1138,6 +1191,64 @@ void KernelObjVisitor::visitNthArrayElement(const CXXRecordDecl *Owner,
             .Handler...);
 }
 
+template <typename ParentTy, typename... HandlerTys>
+void KernelObjVisitor::visitRecord(const CXXRecordDecl *Owner, ParentTy &Parent,
+                                   const CXXRecordDecl *Wrapper,
+                                   QualType RecordTy,
+                                   HandlerTys &... Handlers) {
+  if (RecordTy->getAsRecordDecl()->hasAttr<SYCLRequiresDecompositionAttr>()) {
+    // If this container requires decomposition, we have to visit it as
+    // 'complex', so all handlers are called in this case with the 'complex'
+    // case.
+    visitComplexRecord(Owner, Parent, Wrapper, RecordTy, Handlers...);
+  } else {
+    // "Simple" Containers are those that do NOT need to be decomposed,
+    // "Complex" containers are those that DO. In the case where the container
+    // does NOT need to be decomposed, we can call VisitSimpleRecord on the
+    // handlers that have opted-out of VisitInsideSimpleContainers. The 'if'
+    // makes sure we only do that if at least 1 has opted out.
+    if (!AllTrue<HandlerTys::VisitInsideSimpleContainers...>::Value)
+      visitSimpleRecord(
+          Owner, Parent, Wrapper, RecordTy,
+          HandlerFilter<!HandlerTys::VisitInsideSimpleContainers, HandlerTys>(
+              Handlers)
+              .Handler...);
+
+    // Even though this is a 'simple' container, some handlers (via
+    // VisitInsideSimpleContainers = true) need to treat it as if it needs
+    // decomposing, so we call VisitComplexRecord iif at least one has.
+    if (AnyTrue<HandlerTys::VisitInsideSimpleContainers...>::Value)
+      visitComplexRecord(
+          Owner, Parent, Wrapper, RecordTy,
+          HandlerFilter<HandlerTys::VisitInsideSimpleContainers, HandlerTys>(
+              Handlers)
+              .Handler...);
+  }
+}
+
+template <typename... HandlerTys>
+void KernelObjVisitor::visitArray(const CXXRecordDecl *Owner, FieldDecl *Field,
+                                  QualType ArrayTy, HandlerTys &... Handlers) {
+
+  if (Field->hasAttr<SYCLRequiresDecompositionAttr>()) {
+    visitComplexArray(Owner, Field, ArrayTy, Handlers...);
+  } else {
+    if (!AllTrue<HandlerTys::VisitInsideSimpleContainers...>::Value)
+      visitSimpleArray(
+          Owner, Field, ArrayTy,
+          HandlerFilter<!HandlerTys::VisitInsideSimpleContainers, HandlerTys>(
+              Handlers)
+              .Handler...);
+
+    if (AnyTrue<HandlerTys::VisitInsideSimpleContainers...>::Value)
+      visitComplexArray(
+          Owner, Field, ArrayTy,
+          HandlerFilter<HandlerTys::VisitInsideSimpleContainers, HandlerTys>(
+              Handlers)
+              .Handler...);
+  }
+}
+
 // A type to check the validity of all of the argument types.
 class SyclKernelFieldChecker : public SyclKernelFieldHandler {
   bool IsInvalid = false;
@@ -1373,6 +1484,132 @@ class SyclKernelUnionChecker : public SyclKernelFieldHandler {
   }
 };
 
+// A type to mark whether a collection requires decomposition.
+class SyclKernelDecompMarker : public SyclKernelFieldHandler {
+  llvm::SmallVector<bool, 16> CollectionStack;
+
+public:
+  static constexpr const bool VisitUnionBody = false;
+  static constexpr const bool VisitNthArrayElement = false;
+
+  SyclKernelDecompMarker(Sema &S) : SyclKernelFieldHandler(S) {
+    // In order to prevent checking this over and over, just add a dummy-base
+    // entry.
+    CollectionStack.push_back(true);
+  }
+
+  bool handleSyclAccessorType(const CXXRecordDecl *, const CXXBaseSpecifier &,
+                              QualType) final {
+    CollectionStack.back() = true;
+    return true;
+  }
+  bool handleSyclAccessorType(FieldDecl *, QualType) final {
+    CollectionStack.back() = true;
+    return true;
+  }
+
+  bool handleSyclSamplerType(const CXXRecordDecl *, const CXXBaseSpecifier &,
+                             QualType) final {
+    CollectionStack.back() = true;
+    return true;
+  }
+  bool handleSyclSamplerType(FieldDecl *, QualType) final {
+    CollectionStack.back() = true;
+    return true;
+  }
+  bool handleSyclSpecConstantType(FieldDecl *, QualType) final {
+    CollectionStack.back() = true;
+    return true;
+  }
+  bool handleSyclStreamType(const CXXRecordDecl *, const CXXBaseSpecifier &,
+                            QualType) final {
+    CollectionStack.back() = true;
+    return true;
+  }
+  bool handleSyclStreamType(FieldDecl *, QualType) final {
+    CollectionStack.back() = true;
+    return true;
+  }
+  bool handleSyclHalfType(const CXXRecordDecl *, const CXXBaseSpecifier &,
+                          QualType) final {
+    CollectionStack.back() = true;
+    return true;
+  }
+  bool handleSyclHalfType(FieldDecl *, QualType) final {
+    CollectionStack.back() = true;
+    return true;
+  }
+
+  bool handlePointerType(FieldDecl *, QualType) final {
+    CollectionStack.back() = true;
+    return true;
+  }
+
+  // Stream is always decomposed (and whether it gets decomposed is handled in
+  // handleSyclStreamType), but we need a CollectionStack entry to capture the
+  // accessors that get handled.
+  bool enterStream(const CXXRecordDecl *, FieldDecl *, QualType) final {
+    CollectionStack.push_back(false);
+    return true;
+  }
+  bool leaveStream(const CXXRecordDecl *, FieldDecl *, QualType Ty) final {
+    CollectionStack.pop_back();
+    return true;
+  }
+
+  bool enterStruct(const CXXRecordDecl *, FieldDecl *, QualType) final {
+    CollectionStack.push_back(false);
+    return true;
+  }
+
+  bool leaveStruct(const CXXRecordDecl *, FieldDecl *, QualType Ty) final {
+    if (CollectionStack.pop_back_val()) {
+      RecordDecl *RD = Ty->getAsRecordDecl();
+      if (!RD->hasAttr<SYCLRequiresDecompositionAttr>())
+        RD->addAttr(SYCLRequiresDecompositionAttr::CreateImplicit(
+            SemaRef.getASTContext()));
+      CollectionStack.back() = true;
+    }
+    return true;
+  }
+
+  bool enterStruct(const CXXRecordDecl *, const CXXBaseSpecifier &,
+                   QualType) final {
+    CollectionStack.push_back(false);
+    return true;
+  }
+
+  bool leaveStruct(const CXXRecordDecl *, const CXXBaseSpecifier &,
+                   QualType Ty) final {
+    if (CollectionStack.pop_back_val()) {
+      RecordDecl *RD = Ty->getAsRecordDecl();
+      if (!RD->hasAttr<SYCLRequiresDecompositionAttr>())
+        RD->addAttr(SYCLRequiresDecompositionAttr::CreateImplicit(
+            SemaRef.getASTContext()));
+      CollectionStack.back() = true;
+    }
+
+    return true;
+  }
+
+  bool enterArray(FieldDecl *, QualType ArrayTy, QualType ElementTy) final {
+    CollectionStack.push_back(false);
+    return true;
+  }
+
+  bool leaveArray(FieldDecl *FD, QualType ArrayTy, QualType ElementTy) final {
+    if (CollectionStack.pop_back_val()) {
+      // Cannot assert, since in MD arrays we'll end up marking them multiple
+      // times.
+      if (!FD->hasAttr<SYCLRequiresDecompositionAttr>())
+        FD->addAttr(SYCLRequiresDecompositionAttr::CreateImplicit(
+            SemaRef.getASTContext()));
+      CollectionStack.back() = true;
+    }
+    return true;
+  }
+};
+
 // A type to Create and own the FunctionDecl for the kernel.
 class SyclKernelDeclCreator : public SyclKernelFieldHandler {
   FunctionDecl *KernelDecl;
@@ -1507,6 +1744,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
   }
 
 public:
+  static constexpr const bool VisitInsideSimpleContainers = false;
   SyclKernelDeclCreator(Sema &S, StringRef Name, SourceLocation Loc,
                         bool IsInline, bool IsSIMDKernel)
       : SyclKernelFieldHandler(S),
@@ -1548,6 +1786,18 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
     return true;
   }
 
+  bool enterStruct(const CXXRecordDecl *, const CXXBaseSpecifier &BS,
+                   QualType FieldTy) final {
+    ++StructDepth;
+    return true;
+  }
+
+  bool leaveStruct(const CXXRecordDecl *, const CXXBaseSpecifier &BS,
+                   QualType FieldTy) final {
+    --StructDepth;
+    return true;
+  }
+
   bool handleSyclAccessorType(const CXXRecordDecl *, const CXXBaseSpecifier &BS,
                               QualType FieldTy) final {
     const auto *RecordDecl = FieldTy->getAsCXXRecordDecl();
@@ -1622,11 +1872,32 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
     return true;
   }
 
+  bool handleSimpleArrayType(FieldDecl *FD, QualType FieldTy) final {
+    // Arrays are always wrapped in a struct since they cannot be passed
+    // directly.
+    RecordDecl *WrappedArray = wrapField(FD, FieldTy);
+    QualType ModTy = SemaRef.getASTContext().getRecordType(WrappedArray);
+    addParam(FD, ModTy);
+    return true;
+  }
+
   bool handleScalarType(FieldDecl *FD, QualType FieldTy) final {
     addParam(FD, FieldTy);
     return true;
   }
 
+  bool handleNonDecompStruct(const CXXRecordDecl *, FieldDecl *FD,
+                             QualType Ty) final {
+    addParam(FD, Ty);
+    return true;
+  }
+
+  bool handleNonDecompStruct(const CXXRecordDecl *Base,
+                             const CXXBaseSpecifier &BS, QualType Ty) final {
+    addParam(BS, Ty);
+    return true;
+  }
+
   bool handleUnionType(FieldDecl *FD, QualType FieldTy) final {
     return handleScalarType(FD, FieldTy);
   }
@@ -1658,9 +1929,6 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
   }
   using SyclKernelFieldHandler::handleSyclHalfType;
   using SyclKernelFieldHandler::handleSyclSamplerType;
-  // Required to handle pointers inside structs
-  using SyclKernelFieldHandler::enterStruct;
-  using SyclKernelFieldHandler::leaveStruct;
 };
 
 class SyclKernelArgsSizeChecker : public SyclKernelFieldHandler {
@@ -1683,6 +1951,7 @@ class SyclKernelArgsSizeChecker : public SyclKernelFieldHandler {
   }
 
 public:
+  static constexpr const bool VisitInsideSimpleContainers = false;
   SyclKernelArgsSizeChecker(Sema &S, SourceLocation Loc)
       : SyclKernelFieldHandler(S), KernelLoc(Loc) {}
 
@@ -1720,6 +1989,23 @@ class SyclKernelArgsSizeChecker : public SyclKernelFieldHandler {
     return true;
   }
 
+  bool handleSimpleArrayType(FieldDecl *FD, QualType FieldTy) final {
+    addParam(FieldTy);
+    return true;
+  }
+
+  bool handleNonDecompStruct(const CXXRecordDecl *, FieldDecl *FD,
+                             QualType Ty) final {
+    addParam(Ty);
+    return true;
+  }
+
+  bool handleNonDecompStruct(const CXXRecordDecl *Base,
+                             const CXXBaseSpecifier &BS, QualType Ty) final {
+    addParam(Ty);
+    return true;
+  }
+
   bool handleUnionType(FieldDecl *FD, QualType FieldTy) final {
     return handleScalarType(FD, FieldTy);
   }
@@ -1865,6 +2151,19 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
     return DRE;
   }
 
+  Expr *createSimpleArrayParamReferenceExpr(QualType ArrayTy) {
+    ParmVarDecl *KernelParameter =
+        DeclCreator.getParamVarDeclsForCurrentField()[0];
+    QualType ParamType = KernelParameter->getOriginalType();
+    Expr *DRE = SemaRef.BuildDeclRefExpr(KernelParameter, ParamType, VK_LValue,
+                                         KernelCallerSrcLoc);
+
+    // Unwrap the array.
+    CXXRecordDecl *WrapperStruct = ParamType->getAsCXXRecordDecl();
+    FieldDecl *ArrayField = *(WrapperStruct->field_begin());
+    return buildMemberExpr(DRE, ArrayField);
+  }
+
   // Returns 'true' if the thing we're visiting (Based on the FD/QualType pair)
   // is an element of an array.  This will determine whether we do
   // MemberExprBases in some cases or not, AND determines how we initialize
@@ -1892,8 +2191,11 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
 
   void addFieldInit(FieldDecl *FD, QualType Ty, MultiExprArg ParamRef,
                     InitializationKind InitKind) {
-    InitializedEntity Entity = getFieldEntity(FD, Ty);
+    addFieldInit(FD, Ty, ParamRef, InitKind, getFieldEntity(FD, Ty));
+  }
 
+  void addFieldInit(FieldDecl *FD, QualType Ty, MultiExprArg ParamRef,
+                    InitializationKind InitKind, InitializedEntity Entity) {
     InitializationSequence InitSeq(SemaRef, Entity, InitKind, ParamRef);
     ExprResult Init = InitSeq.Perform(SemaRef, Entity, InitKind, ParamRef);
 
@@ -1914,6 +2216,22 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
                           Init.get());
   }
 
+  void addSimpleBaseInit(const CXXBaseSpecifier &BS, QualType Ty) {
+    InitializationKind InitKind =
+        InitializationKind::CreateCopy(KernelCallerSrcLoc, KernelCallerSrcLoc);
+
+    InitializedEntity Entity = InitializedEntity::InitializeBase(
+        SemaRef.Context, &BS, /*IsInheritedVirtualBase*/ false, &VarEntity);
+
+    Expr *ParamRef = createParamReferenceExpr();
+    InitializationSequence InitSeq(SemaRef, Entity, InitKind, ParamRef);
+    ExprResult Init = InitSeq.Perform(SemaRef, Entity, InitKind, ParamRef);
+
+    InitListExpr *ParentILE = CollectionInitExprs.back();
+    ParentILE->updateInit(SemaRef.getASTContext(), ParentILE->getNumInits(),
+                          Init.get());
+  }
+
   // Adds an initializer that handles a simple initialization of a field.
   void addSimpleFieldInit(FieldDecl *FD, QualType Ty) {
     Expr *ParamRef = createParamReferenceExpr();
@@ -2047,6 +2365,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
   }
 
 public:
+  static constexpr const bool VisitInsideSimpleContainers = false;
   SyclKernelBodyCreator(Sema &S, SyclKernelDeclCreator &DC,
                         const CXXRecordDecl *KernelObj,
                         FunctionDecl *KernelCallerFunc)
@@ -2116,6 +2435,29 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
     return true;
   }
 
+  bool handleSimpleArrayType(FieldDecl *FD, QualType FieldTy) final {
+    Expr *ArrayRef = createSimpleArrayParamReferenceExpr(FieldTy);
+    InitializationKind InitKind = InitializationKind::CreateDirect({}, {}, {});
+
+    InitializedEntity Entity =
+        InitializedEntity::InitializeMember(FD, &VarEntity, /*Implicit*/ true);
+
+    addFieldInit(FD, FieldTy, ArrayRef, InitKind, Entity);
+    return true;
+  }
+
+  bool handleNonDecompStruct(const CXXRecordDecl *, FieldDecl *FD,
+                             QualType Ty) final {
+    addSimpleFieldInit(FD, Ty);
+    return true;
+  }
+
+  bool handleNonDecompStruct(const CXXRecordDecl *Base,
+                             const CXXBaseSpecifier &BS, QualType Ty) final {
+    addSimpleBaseInit(BS, Ty);
+    return true;
+  }
+
   bool handleScalarType(FieldDecl *FD, QualType FieldTy) final {
     addSimpleFieldInit(FD, FieldTy);
     return true;
@@ -2276,10 +2618,9 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
 
   void addParam(const FieldDecl *FD, QualType ArgTy,
                 SYCLIntegrationHeader::kernel_param_kind_t Kind) {
-    addParam(FD, ArgTy, Kind, offsetOf(FD, ArgTy));
+    addParam(ArgTy, Kind, offsetOf(FD, ArgTy));
   }
-  void addParam(const FieldDecl *FD, QualType ArgTy,
-                SYCLIntegrationHeader::kernel_param_kind_t Kind,
+  void addParam(QualType ArgTy, SYCLIntegrationHeader::kernel_param_kind_t Kind,
                 uint64_t OffsetAdj) {
     uint64_t Size;
     Size = SemaRef.getASTContext().getTypeSizeInChars(ArgTy).getQuantity();
@@ -2296,6 +2637,7 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
   }
 
 public:
+  static constexpr const bool VisitInsideSimpleContainers = false;
   SyclKernelIntHeaderCreator(Sema &S, SYCLIntegrationHeader &H,
                              const CXXRecordDecl *KernelObj, QualType NameType,
                              StringRef Name, StringRef StableName)
@@ -2348,7 +2690,7 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
     // offsetOf calculation wouldn't work correctly. Therefore, we need to call
     // a version of addParam where we calculate the offset based on the true
     // FieldDecl/FieldType pair, rather than the SampleArg type.
-    addParam(FD, SamplerArg->getType(), SYCLIntegrationHeader::kind_sampler,
+    addParam(SamplerArg->getType(), SYCLIntegrationHeader::kind_sampler,
              offsetOf(FD, FieldTy));
     return true;
   }
@@ -2381,6 +2723,26 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
     return true;
   }
 
+  bool handleSimpleArrayType(FieldDecl *FD, QualType FieldTy) final {
+    // Arrays are always wrapped inside of structs, so just treat it as a simple
+    // struct.
+    addParam(FD, FieldTy, SYCLIntegrationHeader::kind_std_layout);
+    return true;
+  }
+
+  bool handleNonDecompStruct(const CXXRecordDecl *, FieldDecl *FD,
+                             QualType Ty) final {
+    addParam(FD, Ty, SYCLIntegrationHeader::kind_std_layout);
+    return true;
+  }
+
+  bool handleNonDecompStruct(const CXXRecordDecl *Base,
+                             const CXXBaseSpecifier &, QualType Ty) final {
+    addParam(Ty, SYCLIntegrationHeader::kind_std_layout,
+             offsetOf(Base, Ty->getAsCXXRecordDecl()));
+    return true;
+  }
+
   bool handleUnionType(FieldDecl *FD, QualType FieldTy) final {
     return handleScalarType(FD, FieldTy);
   }
@@ -2583,9 +2945,6 @@ void Sema::CheckSYCLKernelCall(FunctionDecl *KernelFunc, SourceRange CallLoc,
       }
   }
 
-  SyclKernelFieldChecker FieldChecker(*this);
-  SyclKernelUnionChecker UnionChecker(*this);
-  SyclKernelArgsSizeChecker ArgsSizeChecker(*this, Args[0]->getExprLoc());
   // check that calling kernel conforms to spec
   QualType KernelParamTy = KernelFunc->getParamDecl(0)->getType();
   if (KernelParamTy->isReferenceType()) {
@@ -2602,16 +2961,28 @@ void Sema::CheckSYCLKernelCall(FunctionDecl *KernelFunc, SourceRange CallLoc,
   if (KernelObj->isInvalidDecl())
     return;
 
+  SyclKernelDecompMarker DecompMarker(*this);
+  SyclKernelFieldChecker FieldChecker(*this);
+  SyclKernelUnionChecker UnionChecker(*this);
+  SyclKernelArgsSizeChecker ArgsSizeChecker(*this, Args[0]->getExprLoc());
+
   KernelObjVisitor Visitor{*this};
   SYCLKernelNameTypeVisitor KernelTypeVisitor(*this, Args[0]->getExprLoc());
   // Emit diagnostics for SYCL device kernels only
   if (LangOpts.SYCLIsDevice)
     KernelTypeVisitor.Visit(KernelNameType);
   DiagnosingSYCLKernel = true;
-  Visitor.VisitRecordBases(KernelObj, FieldChecker, UnionChecker,
-                           ArgsSizeChecker);
+  Visitor.VisitRecordBases(KernelObj, FieldChecker, UnionChecker, DecompMarker);
   Visitor.VisitRecordFields(KernelObj, FieldChecker, UnionChecker,
-                            ArgsSizeChecker);
+                            DecompMarker);
+  // ArgSizeChecker needs to happen after DecompMarker has completed, since it
+  // cares about the decomp attributes. DecompMarker cannot run before the
+  // others, since it counts on the FieldChecker to make sure it is visiting
+  // valid arrays/etc. Thus, ArgSizeChecker has its own visitation.
+  if (FieldChecker.isValid() && UnionChecker.isValid()) {
+    Visitor.VisitRecordBases(KernelObj, ArgsSizeChecker);
+    Visitor.VisitRecordFields(KernelObj, ArgsSizeChecker);
+  }
   DiagnosingSYCLKernel = false;
   if (!FieldChecker.isValid() || !UnionChecker.isValid())
     KernelFunc->setInvalidDecl();
diff --git a/clang/test/CodeGenSYCL/inheritance.cpp b/clang/test/CodeGenSYCL/inheritance.cpp
index 9cbea0ca8de48..b5007d1885447 100644
--- a/clang/test/CodeGenSYCL/inheritance.cpp
+++ b/clang/test/CodeGenSYCL/inheritance.cpp
@@ -4,7 +4,7 @@
 
 class second_base {
 public:
-  int e;
+  int *e;
 };
 
 class InnerFieldBase {
@@ -40,45 +40,33 @@ int main() {
 }
 
 // Check kernel paramters
-// CHECK: define spir_kernel void @{{.*}}derived(i32 %_arg_b, i32 %_arg_d, i32 %_arg_c, i32 %_arg_e, i32 %_arg_a)
+// CHECK: define spir_kernel void @{{.*}}derived(%struct.{{.*}}.base* byval(%struct.{{.*}}.base) align 4 %_arg__base, %struct.{{.*}}.__wrapper_class* byval(%struct.{{.*}}.__wrapper_class) align 8 %_arg_e, i32 %_arg_a)
 
 // Check alloca for kernel paramters
-// CHECK: %[[ARG_B:[a-zA-Z0-9_.]+]] = alloca i32, align 4
-// CHECK: %[[ARG_D:[a-zA-Z0-9_.]+]] = alloca i32, align 4
-// CHECK: %[[ARG_C:[a-zA-Z0-9_.]+]] = alloca i32, align 4
-// CHECK: %[[ARG_E:[a-zA-Z0-9_.]+]] = alloca i32, align 4
 // CHECK: %[[ARG_A:[a-zA-Z0-9_.]+]] = alloca i32, align 4
-
 // Check alloca for local functor object
-// CHECK: %[[LOCAL_OBJECT:[a-zA-Z0-9_.]+]] = alloca %struct.{{.*}}.derived, align 4
-
-// Initialize field 'b'
-// CHECK: %[[BITCAST1:[0-9]+]] = bitcast %struct.{{.*}}.derived* %[[LOCAL_OBJECT]] to %struct.{{.*}}.base*
-// CHECK: %[[GEP_B:[a-zA-Z0-9]+]] = getelementptr inbounds %struct.{{.*}}.base, %struct.{{.*}}.base* %[[BITCAST1]], i32 0, i32 0
-// CHECK: %[[LOAD_B:[0-9]+]] = load i32, i32* %[[ARG_B]], align 4
-// CHECK: store i32 %[[LOAD_B]], i32* %[[GEP_B]], align 4
-
-// Initialize field 'd'
-// CHECK: %[[GEP_OBJ:[a-zA-Z0-9]+]] = getelementptr inbounds %struct.{{.*}}.base, %struct.{{.*}}.base* %[[BITCAST1]], i32 0, i32 1
-// CHECK: %[[BITCAST2:[0-9]+]] = bitcast %class.{{.*}}.InnerField* %[[GEP_OBJ]] to %class.{{.*}}.InnerFieldBase*
-// CHECK: %[[GEP_D:[a-zA-Z0-9]+]] = getelementptr inbounds %class.{{.*}}.InnerFieldBase, %class.{{.*}}.InnerFieldBase* %[[BITCAST2]], i32 0, i32 0
-// CHECK: %[[LOAD_D:[0-9]+]] = load i32, i32* %[[ARG_D]], align 4
-// CHECK: store i32 %[[LOAD_D]], i32* %[[GEP_D]], align 4
-
-// Initialize field 'c'
-// CHECK: %[[GEP_C:[a-zA-Z0-9]+]] = getelementptr inbounds %class.{{.*}}.InnerField, %class.{{.*}}.InnerField* %[[GEP_OBJ]], i32 0, i32 1
-// CHECK: %[[LOAD_C:[0-9]+]] = load i32, i32* %[[ARG_C]], align 4
-// CHECK: store i32 %[[LOAD_C]], i32* %[[GEP_C]], align 4
-
-// Initialize field 'e'
-// CHECK: %[[BITCAST3:[0-9]+]] = bitcast %struct.{{.*}}.derived* %[[LOCAL_OBJECT]] to i8*
-// CHECK: %[[GEP_DERIVED:[a-zA-Z0-9]+]] = getelementptr inbounds i8, i8* %[[BITCAST3]], i64 12
-// CHECK: %[[BITCAST4:[0-9]+]] = bitcast i8* %[[GEP_DERIVED]] to %class.{{.*}}.second_base*
-// CHECK: %[[GEP_E:[a-zA-Z0-9]+]] = getelementptr inbounds %class.{{.*}}.second_base, %class.{{.*}}.second_base* %[[BITCAST4]], i32 0, i32 0
-// CHECK: %[[LOAD_E:[0-9]+]] = load i32, i32* %[[ARG_E]], align 4
-// CHECK: store i32 %[[LOAD_E]], i32* %[[GEP_E]], align 4
+// CHECK: %[[LOCAL_OBJECT:[a-zA-Z0-9_.]+]] = alloca %struct.{{.*}}.derived, align 8
+// CHECK: store i32 %_arg_a, i32* %[[ARG_A]], align 4
+
+// Initialize 'base' subobject
+// CHECK: %[[DERIVED_TO_BASE:.*]] = bitcast %struct.{{.*}}.derived* %[[LOCAL_OBJECT]] to %struct.{{.*}}.base*
+// CHECK: %[[BASE_TO_PTR:.*]] = bitcast %struct.{{.*}}.base* %[[DERIVED_TO_BASE]] to i8*
+// CHECK: %[[PARAM_TO_PTR:.*]] = bitcast %struct.{{.*}}.base* %_arg__base to i8*
+// CHECK: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 8 %[[BASE_TO_PTR]], i8* align 4 %[[PARAM_TO_PTR]], i64 12, i1 false)
+
+// Initialize 'second_base' subobject
+// First, derived-to-base cast with offset:
+// CHECK: %[[DERIVED_PTR:.*]] = bitcast %struct.{{.*}}.derived* %[[LOCAL_OBJECT]] to i8*
+// CHECK: %[[OFFSET_CALC:.*]] = getelementptr inbounds i8, i8* %[[DERIVED_PTR]], i64 16
+// CHECK: %[[TO_SECOND_BASE:.*]] = bitcast i8* %[[OFFSET_CALC]] to %class.{{.*}}.second_base*
+// Initialize 'second_base::e'
+// CHECK: %[[SECOND_BASE_PTR:.*]] = getelementptr inbounds %class.{{.*}}.second_base, %class.{{.*}}.second_base* %[[TO_SECOND_BASE]], i32 0, i32 0
+// CHECK: %[[PTR_TO_WRAPPER:.*]] = getelementptr inbounds %struct.{{.*}}.__wrapper_class, %struct.{{.*}}.__wrapper_class* %_arg_e, i32 0, i32 0
+// CHECK: %[[LOAD_PTR:.*]] = load i32 addrspace(1)*, i32 addrspace(1)** %[[PTR_TO_WRAPPER]]
+// CHECK: %[[AS_CAST:.*]] = addrspacecast i32 addrspace(1)* %[[LOAD_PTR]] to i32 addrspace(4)*
+// CHECK: store i32 addrspace(4)* %[[AS_CAST]], i32 addrspace(4)** %[[SECOND_BASE_PTR]]
 
 // Initialize field 'a'
-// CHECK: %[[GEP_A:[a-zA-Z0-9]+]] = getelementptr inbounds %struct.{{.*}}.derived, %struct.{{.*}}.derived* %[[LOCAL_OBJECT]], i32 0, i32 2
+// CHECK: %[[GEP_A:[a-zA-Z0-9]+]] = getelementptr inbounds %struct.{{.*}}.derived, %struct.{{.*}}.derived* %[[LOCAL_OBJECT]], i32 0, i32 3
 // CHECK: %[[LOAD_A:[0-9]+]] = load i32, i32* %[[ARG_A]], align 4
-// CHECK: store i32 %[[LOAD_A]], i32* %[[GEP_A]], align 4
+// CHECK: store i32 %[[LOAD_A]], i32* %[[GEP_A]]
diff --git a/clang/test/CodeGenSYCL/integration_header.cpp b/clang/test/CodeGenSYCL/integration_header.cpp
index d5eca9624f3f2..5bc45080d4235 100644
--- a/clang/test/CodeGenSYCL/integration_header.cpp
+++ b/clang/test/CodeGenSYCL/integration_header.cpp
@@ -28,8 +28,7 @@
 // CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = {
 // CHECK-NEXT:   //--- _ZTSZ4mainE12first_kernel
 // CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 0 },
-// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 1, 4 },
-// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 8 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 8, 4 },
 // CHECK-NEXT:   { kernel_param_kind_t::kind_accessor, 4062, 12 },
 // CHECK-NEXT:   { kernel_param_kind_t::kind_accessor, 6112, 24 },
 // CHECK-NEXT:   { kernel_param_kind_t::kind_sampler, 8, 40 },
diff --git a/clang/test/CodeGenSYCL/kernel-param-pod-array-ih.cpp b/clang/test/CodeGenSYCL/kernel-param-pod-array-ih.cpp
index ad2b36524a78e..c5881940a3fa3 100644
--- a/clang/test/CodeGenSYCL/kernel-param-pod-array-ih.cpp
+++ b/clang/test/CodeGenSYCL/kernel-param-pod-array-ih.cpp
@@ -21,41 +21,21 @@
 // CHECK: static constexpr
 // CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = {
 // CHECK-NEXT: //--- _ZTSZ4mainE8kernel_B
-// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 0 },
-// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 4 },
-// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 8 },
-// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 12 },
-// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 16 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 20, 0 },
 // CHECK-EMPTY:
 // CHECK-NEXT: //--- _ZTSZ4mainE8kernel_C
-// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 0 },
-// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 4 },
-// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 8 },
-// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 12 },
-// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 16 },
-// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 20 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 24, 0 },
 // CHECK-EMPTY:
 // CHECK-NEXT: //--- _ZTSZ4mainE8kernel_D
-// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 0 },
-// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 4 },
-// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 8 },
-// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 12 },
-// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 16 },
-// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 20 },
-// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 24 },
-// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 28 },
-// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 32 },
-// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 36 },
-// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 40 },
-// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 44 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 48, 0 },
 // CHECK-EMPTY:
 // CHECK-NEXT: };
 
 // CHECK: static constexpr
 // CHECK-NEXT: const unsigned kernel_signature_start[] = {
 // CHECK-NEXT:  0, // _ZTSZ4mainE8kernel_B
-// CHECK-NEXT:  6, // _ZTSZ4mainE8kernel_C
-// CHECK-NEXT:  13 // _ZTSZ4mainE8kernel_D
+// CHECK-NEXT:  2, // _ZTSZ4mainE8kernel_C
+// CHECK-NEXT:  4 // _ZTSZ4mainE8kernel_D
 // CHECK-NEXT: };
 
 // CHECK: template <> struct KernelInfo<class kernel_B> {
diff --git a/clang/test/CodeGenSYCL/kernel-param-pod-array.cpp b/clang/test/CodeGenSYCL/kernel-param-pod-array.cpp
index 26b28a4a58cf6..c9f602261a8f7 100644
--- a/clang/test/CodeGenSYCL/kernel-param-pod-array.cpp
+++ b/clang/test/CodeGenSYCL/kernel-param-pod-array.cpp
@@ -46,136 +46,87 @@ int main() {
 
 // Check kernel_B parameters
 // CHECK: define spir_kernel void @{{.*}}kernel_B
-// CHECK-SAME: i32 [[ELEM_ARG0:%[a-zA-Z0-9_]+]],
-// CHECK-SAME: i32 [[ELEM_ARG1:%[a-zA-Z_]+_[0-9]+]])
+// CHECK-SAME:(%struct.{{.*}}.__wrapper_class* byval(%struct.{{.*}}.__wrapper_class) align 4 %[[ARR_ARG:.*]])
 
 // Check local lambda object alloca
-// CHECK: [[LOCAL_OBJECT:%[0-9]+]] = alloca %"class.{{.*}}.anon", align 4
-
-// Check local variables created for parameters
-// CHECK: store i32 [[ELEM_ARG0]], i32* [[ELEM_L0:%[a-zA-Z_]+.addr]], align 4
-// CHECK: store i32 [[ELEM_ARG1]], i32* [[ELEM_L1:%[a-zA-Z_]+.addr[0-9]*]], align 4
-
-// Check init of local array
-// CHECK: [[ARRAY:%[0-9]*]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0
-// CHECK: [[ARRAY_BEGIN:%[a-zA-Z_.]+]] = getelementptr inbounds [2 x i32], [2 x i32]* [[ARRAY]], i64 0, i64 0
-// CHECK: [[ARRAY0:%[0-9]*]] = load i32, i32* [[ELEM_L0]], align 4
-// CHECK: store i32 [[ARRAY0]], i32* [[ARRAY_BEGIN]], align 4
-// CHECK: [[ARRAY_ELEMENT:%[a-zA-Z_.]+]] = getelementptr inbounds i32, i32* %arrayinit.begin, i64 1
-// CHECK: [[ARRAY1:%[0-9]*]] = load i32, i32* [[ELEM_L1]], align 4
-// CHECK: store i32 [[ARRAY1]], i32* [[ARRAY_ELEMENT]], align 4
+// CHECK: %[[LOCAL_OBJECT:[0-9]+]] = alloca %"class.{{.*}}.anon", align 4
+
+// Check for Array init loop
+// CHECK: %[[LAMBDA_PTR:.+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* %[[LOCAL_OBJECT]], i32 0, i32 0
+// CHECK: %[[WRAPPER_PTR:.+]] = getelementptr inbounds %struct.{{.*}}.__wrapper_class, %struct.{{.*}}.__wrapper_class* %[[ARR_ARG]], i32 0, i32 0
+// CHECK: %[[ARRAY_BEGIN:.+]] = getelementptr inbounds [2 x i32], [2 x i32]* %[[LAMBDA_PTR]], i64 0, i64 0
+// CHECK: br label %[[ARRAYINITBODY:.+]]
+
+// The loop body itself
+// CHECK: [[ARRAYINITBODY]]:
+// CHECK: %[[ARRAYINDEX:.+]] = phi i64 [ 0, %{{.*}} ], [ %[[NEXTINDEX:.+]], %[[ARRAYINITBODY]] ]
+// CHECK: %[[TARG_ARRAY_ELEM:.+]] = getelementptr inbounds i32, i32* %[[ARRAY_BEGIN]], i64 %[[ARRAYINDEX]]
+// CHECK: %[[SRC_ELEM:.+]] = getelementptr inbounds [2 x i32], [2 x i32]* %[[WRAPPER_PTR]], i64 0, i64 %[[ARRAYINDEX]]
+// CHECK: %[[SRC_VAL:.+]] = load i32, i32* %[[SRC_ELEM]]
+// CHECK: store i32 %[[SRC_VAL]], i32* %[[TARG_ARRAY_ELEM]]
+// CHECK: %[[NEXTINDEX]] = add nuw i64 %[[ARRAYINDEX]], 1
+// CHECK: %[[ISDONE:.+]] = icmp eq i64 %[[NEXTINDEX]], 2
+// CHECK: br i1 %[[ISDONE]], label %{{.*}}, label %[[ARRAYINITBODY]]
 
 // Check kernel_C parameters
 // CHECK: define spir_kernel void @{{.*}}kernel_C
-// CHECK-SAME: i32 [[FOO1_A:%[a-zA-Z0-9_]+]], i32 [[FOO1_B1_X:%[a-zA-Z0-9_]+]], i32 [[FOO1_B1_Y:%[a-zA-Z0-9_]+]], i32 [[FOO1_B2_X:%[a-zA-Z0-9_]+]], i32 [[FOO1_B2_Y:%[a-zA-Z0-9_]+]], i32 [[FOO1_C:%[a-zA-Z0-9_]+]],
-// CHECK-SAME: i32 [[FOO2_A:%[a-zA-Z0-9_]+]], i32 [[FOO2_B1_X:%[a-zA-Z0-9_]+]], i32 [[FOO2_B1_Y:%[a-zA-Z0-9_]+]], i32 [[FOO2_B2_X:%[a-zA-Z0-9_]+]], i32 [[FOO2_B2_Y:%[a-zA-Z0-9_]+]], i32 [[FOO2_C:%[a-zA-Z0-9_]+]]
+// CHECK-SAME:(%struct.{{.*}}.__wrapper_class{{.*}}* byval(%struct.{{.*}}.__wrapper_class{{.*}}) align 4 %[[ARR_ARG:.*]])
 
 // Check local lambda object alloca
-// CHECK: [[KERNEL_OBJ:%[0-9]+]] = alloca %"class.{{.*}}.anon.0", align 4
-
-// Check local stores
-// CHECK: store i32 [[FOO1_A]], i32* [[FOO1_A_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4
-// CHECK: store i32 [[FOO1_B1_X]], i32* [[FOO1_B1_X_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4
-// CHECK: store i32 [[FOO1_B1_Y]], i32* [[FOO1_B1_Y_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4
-// CHECK: store i32 [[FOO1_B2_X]], i32* [[FOO1_B2_X_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4
-// CHECK: store i32 [[FOO1_B2_Y]], i32* [[FOO1_B2_Y_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4
-// CHECK: store i32 [[FOO1_C]], i32* [[FOO1_C_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4
-// CHECK: store i32 [[FOO2_A]], i32* [[FOO2_A_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4
-// CHECK: store i32 [[FOO2_B1_X]], i32* [[FOO2_B1_X_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4
-// CHECK: store i32 [[FOO2_B1_Y]], i32* [[FOO2_B1_Y_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4
-// CHECK: store i32 [[FOO2_B2_X]], i32* [[FOO2_B2_X_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4
-// CHECK: store i32 [[FOO2_B2_Y]], i32* [[FOO2_B2_Y_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4
-// CHECK: store i32 [[FOO2_C]], i32* [[FOO2_C_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4
-
-// Check initialization of local array
-
-// Initialize struct_array[0].foo_a
-// CHECK: [[GEP:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon.0", %"class.{{.*}}.anon.0"* [[KERNEL_OBJ]], i32 0, i32 0
-// CHECK: [[FOO_ARRAY_0:%[a-zA-Z_.]+]] = getelementptr inbounds [2 x %struct.{{.*}}.foo], [2 x %struct.{{.*}}.foo]* [[GEP]], i64 0, i64 0
-// CHECK: [[GEP_FOO1_A:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}.foo, %struct.{{.*}}.foo* [[FOO_ARRAY_0]], i32 0, i32 0
-// CHECK: [[LOAD_FOO1_A:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO1_A_LOCAL]], align 4
-// CHECK: store i32 [[LOAD_FOO1_A]], i32* [[GEP_FOO1_A]], align 4
-
-// Initialize struct_array[0].foo_b[0].x
-// CHECK: [[GEP_FOO1_B:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}.foo, %struct.{{.*}}.foo* [[FOO_ARRAY_0]], i32 0, i32 1
-// CHECK: [[B_ARRAY_0:%[a-zA-Z0-9_.]+]] = getelementptr inbounds [2 x %struct.{{.*}}foo_inner.foo_inner], [2 x %struct.{{.*}}foo_inner.foo_inner]* [[GEP_FOO1_B]], i64 0, i64 0
-// CHECK: [[GEP_FOO1_B1_X:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}foo_inner.foo_inner, %struct.{{.*}}foo_inner.foo_inner* [[B_ARRAY_0]], i32 0, i32 0
-// CHECK: [[LOAD_FOO1_B1_X:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO1_B1_X_LOCAL]], align 4
-// CHECK: store i32 [[LOAD_FOO1_B1_X]], i32* [[GEP_FOO1_B1_X]], align 4
-
-// Initialize struct_array[0].foo_b[0].y
-// CHECK: [[GEP_FOO1_B1_Y:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}foo_inner.foo_inner, %struct.{{.*}}foo_inner.foo_inner* [[B_ARRAY_0]], i32 0, i32 1
-// CHECK: [[LOAD_FOO1_B1_Y:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO1_B1_Y_LOCAL]], align 4
-// CHECK: store i32 [[LOAD_FOO1_B1_Y]], i32* [[GEP_FOO1_B1_Y]], align 4
-
-// Initialize struct_array[0].foo_b[1].x
-// CHECK: [[B_ARRAY_1:%[a-zA-Z0-9_.]+]] = getelementptr inbounds %struct.{{.*}}foo_inner.foo_inner, %struct.{{.*}}foo_inner.foo_inner* [[B_ARRAY_0]], i64 1
-// CHECK: [[GEP_FOO1_B2_X:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}foo_inner.foo_inner, %struct.{{.*}}foo_inner.foo_inner* [[B_ARRAY_1]], i32 0, i32 0
-// CHECK: [[LOAD_FOO1_B2_X:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO1_B2_X_LOCAL]], align 4
-// CHECK: store i32 [[LOAD_FOO1_B2_X]], i32* [[GEP_FOO1_B2_X]], align 4
-
-// Initialize struct_array[0].foo_b[1].y
-// CHECK: [[GEP_FOO1_B2_Y:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}foo_inner.foo_inner, %struct.{{.*}}foo_inner.foo_inner* [[B_ARRAY_1]], i32 0, i32 1
-// CHECK: [[LOAD_FOO1_B2_Y:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO1_B2_Y_LOCAL]], align 4
-// CHECK: store i32 [[LOAD_FOO1_B2_Y]], i32* [[GEP_FOO1_B2_Y]], align 4
-
-// Initialize struct_array[0].foo_c
-// CHECK: [[GEP_FOO1_C:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}foo.foo, %struct.{{.*}}foo.foo* [[FOO_ARRAY_0]], i32 0, i32 2
-// CHECK: [[LOAD_FOO1_C:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO1_C_LOCAL]], align 4
-// CHECK: store i32 [[LOAD_FOO1_C]], i32* [[GEP_FOO1_C]], align 4
-
-// Initialize struct_array[1].foo_a
-// CHECK: [[FOO_ARRAY_1:%[a-zA-Z0-9_.]+]] = getelementptr inbounds %struct._ZTS3foo.foo, %struct._ZTS3foo.foo* [[FOO_ARRAY_0]], i64 1
-// CHECK: [[GEP_FOO2_A:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}foo.foo, %struct.{{.*}}foo.foo* [[FOO_ARRAY_1]], i32 0, i32 0
-// CHECK: [[LOAD_FOO2_A:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO2_A_LOCAL]], align 4
-// CHECK: store i32 [[LOAD_FOO2_A]], i32* [[GEP_FOO2_A]], align 4
-
-// Initialize struct_array[1].foo_b[0].x
-// CHECK: [[GEP_FOO2_B:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}.foo, %struct.{{.*}}.foo* [[FOO_ARRAY_1]], i32 0, i32 1
-// CHECK: [[FOO2_B_ARRAY_0:%[a-zA-Z0-9_.]+]] = getelementptr inbounds [2 x %struct.{{.*}}foo_inner.foo_inner], [2 x %struct.{{.*}}foo_inner.foo_inner]* [[GEP_FOO2_B]], i64 0, i64 0
-// CHECK: [[GEP_FOO2_B1_X:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}foo_inner.foo_inner, %struct.{{.*}}foo_inner.foo_inner* [[FOO2_B_ARRAY_0]], i32 0, i32 0
-// CHECK: [[LOAD_FOO2_B1_X:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO2_B1_X_LOCAL]], align 4
-// CHECK: store i32 [[LOAD_FOO2_B1_X]], i32* [[GEP_FOO2_B1_X]]
-
-// Initialize struct_array[1].foo_b[0].y
-// CHECK: [[GEP_FOO2_B1_Y:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}foo_inner.foo_inner, %struct.{{.*}}foo_inner.foo_inner* [[FOO2_B_ARRAY_0]], i32 0, i32 1
-// CHECK: [[LOAD_FOO2_B1_Y:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO2_B1_Y_LOCAL]], align 4
-// CHECK: store i32 [[LOAD_FOO2_B1_Y]], i32* [[GEP_FOO2_B1_Y]], align 4
-
-// Initialize struct_array[1].foo_b[1].x
-// CHECK: [[FOO2_B_ARRAY_1:%[a-zA-Z0-9_.]+]] = getelementptr inbounds %struct.{{.*}}foo_inner.foo_inner, %struct.{{.*}}foo_inner.foo_inner* [[FOO2_B_ARRAY_0]], i64 1
-// CHECK: [[GEP_FOO2_B2_X:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}foo_inner.foo_inner, %struct.{{.*}}foo_inner.foo_inner* [[FOO2_B_ARRAY_1]], i32 0, i32 0
-// CHECK: [[LOAD_FOO2_B2_X:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO2_B2_X_LOCAL]], align 4
-// CHECK: store i32 [[LOAD_FOO2_B2_X]], i32* [[GEP_FOO2_B2_X]], align 4
-
-// Initialize struct_array[1].foo_b[1].y
-// CHECK: [[GEP_FOO2_B2_Y:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}foo_inner.foo_inner, %struct.{{.*}}foo_inner.foo_inner* [[FOO2_B_ARRAY_1]], i32 0, i32 1
-// CHECK: [[LOAD_FOO2_B2_Y:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO2_B2_Y_LOCAL]], align 4
-// CHECK: store i32 [[LOAD_FOO2_B2_Y]], i32* [[GEP_FOO2_B2_Y]], align 4
-
-// Initialize struct_array[1].foo_c
-// CHECK: [[GEP_FOO2_C:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}foo.foo, %struct.{{.*}}foo.foo* [[FOO_ARRAY_1]], i32 0, i32 2
-// CHECK: [[LOAD_FOO2_C:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO2_C_LOCAL]], align 4
-// CHECK: store i32 [[LOAD_FOO2_C]], i32* [[GEP_FOO2_C]], align 4
+// CHECK: %[[LOCAL_OBJECT:[0-9]+]] = alloca %"class.{{.*}}.anon{{.*}}", align 4
+
+// Check for Array init loop
+// CHECK: %[[LAMBDA_PTR:.+]] = getelementptr inbounds %"class.{{.*}}.anon{{.*}}", %"class.{{.*}}.anon{{.*}}"* %[[LOCAL_OBJECT]], i32 0, i32 0
+// CHECK: %[[WRAPPER_PTR:.+]] = getelementptr inbounds %struct.{{.*}}.__wrapper_class{{.*}}, %struct.{{.*}}.__wrapper_class{{.*}}* %[[ARR_ARG]], i32 0, i32 0
+// CHECK: %[[ARRAY_BEGIN:.+]] = getelementptr inbounds [2 x %struct.{{.*}}.foo], [2 x %struct.{{.*}}.foo]* %[[LAMBDA_PTR]], i64 0, i64 0
+// CHECK: br label %[[ARRAYINITBODY:.+]]
+
+// The loop body itself
+// CHECK: [[ARRAYINITBODY]]:
+// CHECK: %[[ARRAYINDEX:.+]] = phi i64 [ 0, %{{.*}} ], [ %[[NEXTINDEX:.+]], %[[ARRAYINITBODY]] ]
+// CHECK: %[[TARG_ARRAY_ELEM:.+]] = getelementptr inbounds %struct.{{.*}}.foo, %struct.{{.*}}.foo* %[[ARRAY_BEGIN]], i64 %[[ARRAYINDEX]]
+// CHECK: %[[SRC_ELEM:.+]] = getelementptr inbounds [2 x %struct.{{.*}}.foo], [2 x %struct.{{.*}}.foo]* %[[WRAPPER_PTR]], i64 0, i64 %[[ARRAYINDEX]]
+// CHECK: %[[TARG_PTR:.+]] = bitcast %struct.{{.*}}.foo* %[[TARG_ARRAY_ELEM]] to i8*
+// CHECK: %[[SRC_PTR:.+]] = bitcast %struct.{{.*}}.foo* %[[SRC_ELEM]] to i8*
+// call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 4 %[[TARG_PTR]], i8* align %[[SRC_PTR]], i64 24, i1 false)
+// CHECK: %[[NEXTINDEX]] = add nuw i64 %[[ARRAYINDEX]], 1
+// CHECK: %[[ISDONE:.+]] = icmp eq i64 %[[NEXTINDEX]], 2
+// CHECK: br i1 %[[ISDONE]], label %{{.*}}, label %[[ARRAYINITBODY]]
 
 // Check kernel_D parameters
 // CHECK: define spir_kernel void @{{.*}}kernel_D
-// CHECK-SAME: i32 [[ARR_2D_1:%[a-zA-Z0-9_]+]], i32 [[ARR_2D_2:%[a-zA-Z0-9_]+]]
+// CHECK-SAME:(%struct.{{.*}}.__wrapper_class{{.*}}* byval(%struct.{{.*}}.__wrapper_class{{.*}}) align 4 %[[ARR_ARG:.*]])
 
 // Check local lambda object alloca
-// CHECK: [[LAMBDA_OBJ:%[0-9]+]] = alloca %"class.{{.*}}.anon.1", align 4
-
-// Check local stores
-// CHECK: store i32 [[ARR_2D_1]], i32* [[ARR_2D_1_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4
-// CHECK: store i32 [[ARR_2D_2]], i32* [[ARR_2D_2_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4
-
-// Check initialization of local array
-// CHECK: [[GEP_ARR_2D:%[0-9]*]] = getelementptr inbounds %"class._ZTSZ4mainE3$_0.anon.1", %"class._ZTSZ4mainE3$_0.anon.1"* [[LAMBDA_OBJ]], i32 0, i32 0
-// CHECK: [[GEP_ARR_BEGIN1:%[a-zA-Z0-9_.]+]] = getelementptr inbounds [2 x [1 x i32]], [2 x [1 x i32]]* [[GEP_ARR_2D]], i64 0, i64 0
-// CHECK: [[GEP_ARR_ELEM0:%[a-zA-Z0-9_.]+]] = getelementptr inbounds [1 x i32], [1 x i32]* [[GEP_ARR_BEGIN1]], i64 0, i64 0
-// CHECK: [[ARR_2D_ELEM0:%[0-9]*]] = load i32, i32* [[ARR_2D_1_LOCAL]], align 4
-// CHECK: store i32 [[ARR_2D_ELEM0]], i32* [[GEP_ARR_ELEM0]], align 4
-// CHECK: [[GEP_ARR_BEGIN2:%[a-zA-Z_.]+]] = getelementptr inbounds [1 x i32], [1 x i32]* [[GEP_ARR_BEGIN1]], i64 1
-// CHECK: [[GEP_ARR_ELEM1:%[a-zA-Z0-9_.]+]] = getelementptr inbounds [1 x i32], [1 x i32]* [[GEP_ARR_BEGIN2]], i64 0, i64 0
-// CHECK: [[ARR_2D_ELEM1:%[0-9]*]] = load i32, i32* [[ARR_2D_2_LOCAL]], align 4
-// CHECK: store i32 [[ARR_2D_ELEM1]], i32* [[GEP_ARR_ELEM1]], align 4
+// CHECK: %[[LOCAL_OBJECT:[0-9]+]] = alloca %"class.{{.*}}.anon{{.*}}", align 4
+
+// Check for Array init loop
+// CHECK: %[[LAMBDA_PTR:.+]] = getelementptr inbounds %"class.{{.*}}.anon{{.*}}", %"class.{{.*}}.anon{{.*}}"* %[[LOCAL_OBJECT]], i32 0, i32 0
+// CHECK: %[[WRAPPER_PTR:.+]] = getelementptr inbounds %struct.{{.*}}.__wrapper_class{{.*}}, %struct.{{.*}}.__wrapper_class{{.*}}* %[[ARR_ARG]], i32 0, i32 0
+// CHECK: %[[ARRAY_BEGIN:.+]] = getelementptr inbounds [2 x [1 x i32]], [2 x [1 x i32]]* %[[LAMBDA_PTR]], i64 0, i64 0
+// CHECK: br label %[[ARRAYINITBODY:.+]]
+
+// Check Outer loop.
+// CHECK: [[ARRAYINITBODY]]:
+// CHECK: %[[ARRAYINDEX:.+]] = phi i64 [ 0, %{{.*}} ], [ %[[NEXTINDEX:.+]], %[[ARRAYINITEND:.+]] ]
+// CHECK: %[[TARG_OUTER_ELEM:.+]] = getelementptr inbounds [1 x i32], [1 x i32]* %[[ARRAY_BEGIN]], i64 %[[ARRAYINDEX]]
+// CHECK: %[[SRC_OUTER_ELEM:.+]] = getelementptr inbounds [2 x [1 x i32]], [2 x [1 x i32]]* %[[WRAPPER_PTR]], i64 0, i64 %[[ARRAYINDEX]]
+// CHECK: %[[ARRAY_BEGIN_INNER:.+]] = getelementptr inbounds [1 x i32], [1 x i32]* %[[TARG_OUTER_ELEM]], i64 0, i64 0
+// CHECK: br label %[[ARRAYINITBODY_INNER:.+]]
+
+// Check Inner Loop
+// CHECK: [[ARRAYINITBODY_INNER]]:
+// CHECK: %[[ARRAYINDEX_INNER:.+]] = phi i64 [ 0, %{{.*}} ], [ %[[NEXTINDEX_INNER:.+]], %[[ARRAYINITBODY_INNER:.+]] ]
+// CHECK: %[[TARG_INNER_ELEM:.+]] = getelementptr inbounds i32, i32* %[[ARRAY_BEGIN_INNER]], i64 %[[ARRAYINDEX_INNER]]
+// CHECK: %[[SRC_INNER_ELEM:.+]] = getelementptr inbounds [1 x i32], [1 x i32]* %[[SRC_OUTER_ELEM]], i64 0, i64 %[[ARRAYINDEX_INNER]]
+// CHECK: %[[SRC_LOAD:.+]] = load i32, i32* %[[SRC_INNER_ELEM]]
+// CHECK: store i32 %[[SRC_LOAD]], i32* %[[TARG_INNER_ELEM]]
+// CHECK: %[[NEXTINDEX_INNER]] = add nuw i64 %[[ARRAYINDEX_INNER]], 1
+// CHECK: %[[ISDONE_INNER:.+]] = icmp eq i64 %[[NEXTINDEX_INNER]], 1
+// CHECK: br i1 %[[ISDONE_INNER]], label %[[ARRAYINITEND]], label %[[ARRAYINITBODY_INNER]]
+
+// Check Inner loop 'end'
+// CHECK: [[ARRAYINITEND]]:
+// CHECK: %[[NEXTINDEX]] = add nuw i64 %[[ARRAYINDEX]], 1
+// CHECK: %[[ISDONE:.+]] = icmp eq i64 %[[NEXTINDEX]], 2
+// CHECK: br i1 %[[ISDONE]], label %{{.*}}, label %[[ARRAYINITBODY]]
diff --git a/clang/test/CodeGenSYCL/struct_kernel_param.cpp b/clang/test/CodeGenSYCL/struct_kernel_param.cpp
index e67915455539a..11c1526f41040 100644
--- a/clang/test/CodeGenSYCL/struct_kernel_param.cpp
+++ b/clang/test/CodeGenSYCL/struct_kernel_param.cpp
@@ -3,13 +3,16 @@
 
 // CHECK:     const kernel_param_desc_t kernel_signatures[] = {
 // CHECK-NEXT:  //--- _ZTSZZ5test0vENK3$_0clERN2cl4sycl7handlerEE8MyKernel
+// Accessor
 // CHECK-NEXT:  { kernel_param_kind_t::kind_accessor, 4062, 0 },
-// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 12 },
+// FldInt, offset to 16 because the float* causes the alignment of the structs
+// to change.
 // CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 16 },
-// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 20 },
+// FldArr
 // CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 24 },
-// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 28 },
-// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 32 },
+// FldFloat
+// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 8, 32 },
+// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 12, 40 },
 // CHECK-EMPTY:
 // CHECK-NEXT:};
 
@@ -21,7 +24,7 @@ using namespace cl::sycl;
 
 struct MyNestedStruct {
   int FldArr[1];
-  float FldFloat;
+  float *FldFloat;
 };
 
 struct MyStruct {
diff --git a/clang/test/SemaSYCL/array-kernel-param.cpp b/clang/test/SemaSYCL/array-kernel-param.cpp
index 945b7decd9f57..ea1d4ac4a899d 100644
--- a/clang/test/SemaSYCL/array-kernel-param.cpp
+++ b/clang/test/SemaSYCL/array-kernel-param.cpp
@@ -24,6 +24,8 @@ int main() {
 
   Accessor acc[2];
   int a[2];
+  int *a_ptrs[2];
+
   struct struct_acc_t {
     Accessor member_acc[2];
   } struct_acc;
@@ -32,17 +34,25 @@ int main() {
   struct foo_inner {
     int foo_inner_x;
     int foo_inner_y;
-    int foo_inner_z[2];
+    int *foo_inner_z[2];
   };
 
   struct foo {
     int foo_a;
     foo_inner foo_b[2];
+    int *foo_2D[2][1];
+    int foo_c;
+  };
+
+  // Not decomposed.
+  struct foo2 {
+    int foo_a;
     int foo_2D[2][1];
     int foo_c;
   };
 
   foo struct_array[2];
+  foo2 struct_array2[2];
 
   int array_2D[2][3];
 
@@ -56,6 +66,11 @@ int main() {
         int local = a[1];
       });
 
+  a_kernel<class kernel_B_ptrs>(
+      [=]() {
+        int local = *a_ptrs[1];
+      });
+
   a_kernel<class kernel_C>(
       [=]() {
         struct_acc.member_acc[2].use();
@@ -75,6 +90,11 @@ int main() {
       [=]() {
         int local = array_2D[1][1];
       });
+
+  a_kernel<class kernel_G>(
+      [=]() {
+        foo2 local = struct_array2[0];
+      });
 }
 
 // Check kernel_A parameters
@@ -93,19 +113,40 @@ int main() {
 // CHECK-NEXT: MemberExpr {{.*}}__init
 
 // Check kernel_B parameters
-// CHECK: FunctionDecl {{.*}}kernel_B{{.*}} 'void (int, int)'
-// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ 'int'
-// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ 'int'
+// CHECK: FunctionDecl {{.*}}kernel_B{{.*}} 'void (__wrapper_class)'
+// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ '__wrapper_class'
 // Check kernel_B inits
 // CHECK-NEXT: CompoundStmt
 // CHECK-NEXT: DeclStmt
 // CHECK-NEXT: VarDecl {{.*}} cinit
 // CHECK-NEXT: InitListExpr
-// CHECK-NEXT: InitListExpr {{.*}} 'int [2]'
-// CHECK: ImplicitCastExpr
-// CHECK: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_' 'int'
-// CHECK: ImplicitCastExpr
-// CHECK: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_' 'int'
+// CHECK-NEXT: ArrayInitLoopExpr {{.*}} 'int [2]'
+// CHECK-NEXT: OpaqueValueExpr {{.*}} 'int [2]' lvalue
+// CHECK-NEXT: MemberExpr {{.*}} 'int [2]' lvalue .
+// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_' '__wrapper_class'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
+// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'int' lvalue
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <ArrayToPointerDecay>
+// CHECK-NEXT: OpaqueValueExpr {{.*}} 'int [2]' lvalue
+// CHECK-NEXT: MemberExpr {{.*}} 'int [2]' lvalue .
+// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_' '__wrapper_class'
+
+// Check kernel_B_ptrs parameters
+// CHECK: FunctionDecl {{.*}}kernel_B_ptrs{{.*}} 'void (__global int *, __global int *)'
+// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ '__global int *'
+// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ '__global int *'
+// Check kernel_B_ptrs inits
+// CHECK-NEXT: CompoundStmt
+// CHECK-NEXT: DeclStmt
+// CHECK-NEXT: VarDecl {{.*}} cinit
+// CHECK-NEXT: InitListExpr
+// CHECK-NEXT: InitListExpr {{.*}} 'int *[2]'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <AddressSpaceConversion>
+// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr {{.*}} '__global int *' lvalue ParmVar {{.*}} '_arg_' '__global int *'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <AddressSpaceConversion>
+// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr {{.*}} '__global int *' lvalue ParmVar {{.*}} '_arg_' '__global int *'
 
 // Check kernel_C parameters
 // CHECK: FunctionDecl {{.*}}kernel_C{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, __global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)'
@@ -133,30 +174,30 @@ int main() {
 // CHECK-NEXT: MemberExpr {{.*}}__init
 
 // Check kernel_D parameters
-// CHECK: FunctionDecl {{.*}}kernel_D{{.*}} 'void (int, int, int, int, int, int, int, int, int, int, int, int, int, int, int, int, int, int, int, int, int, int, int, int)'
+// CHECK: FunctionDecl {{.*}}kernel_D{{.*}} 'void (int, int, int, __wrapper_class, __wrapper_class, int, int, __wrapper_class, __wrapper_class, __wrapper_class, __wrapper_class, int, int, int, int, __wrapper_class, __wrapper_class, int, int, __wrapper_class, __wrapper_class, __wrapper_class, __wrapper_class, int)'
 // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_a 'int'
 // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_x 'int'
 // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_y 'int'
-// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z 'int'
-// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z 'int'
+// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z '__wrapper_class'
+// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z '__wrapper_class'
 // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_x 'int'
 // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_y 'int'
-// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z 'int'
-// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z 'int'
-// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_2D 'int'
-// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_2D 'int'
+// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z '__wrapper_class'
+// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z '__wrapper_class'
+// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_2D '__wrapper_class'
+// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_2D '__wrapper_class'
 // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_c 'int'
 // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_a 'int'
 // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_x 'int'
 // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_y 'int'
-// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z 'int'
-// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z 'int'
+// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z '__wrapper_class'
+// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z '__wrapper_class'
 // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_x 'int'
 // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_y 'int'
-// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z 'int'
-// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z 'int'
-// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_2D 'int'
-// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_2D 'int'
+// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z '__wrapper_class'
+// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z '__wrapper_class'
+// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_2D '__wrapper_class'
+// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_2D '__wrapper_class'
 // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_c 'int'
 // CHECK-NEXT: CompoundStmt
 // CHECK-NEXT: DeclStmt
@@ -168,120 +209,164 @@ int main() {
 
 // Initializer for first element of struct_array
 // CHECK-NEXT: InitListExpr {{.*}} 'foo'
-// CHECK-NEXT: ImplicitCastExpr
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
 // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_a' 'int'
+
 // Initializer for struct array inside foo i.e. foo_inner foo_b[2]
 // CHECK-NEXT: InitListExpr {{.*}} 'foo_inner [2]'
 // Initializer for first element of inner struct array
 // CHECK-NEXT: InitListExpr {{.*}} 'foo_inner'
-// CHECK-NEXT: ImplicitCastExpr
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
 // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_x' 'int'
-// CHECK-NEXT: ImplicitCastExpr
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
 // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_y' 'int'
-// CHECK-NEXT: InitListExpr {{.*}} 'int [2]'
+// CHECK-NEXT: InitListExpr {{.*}} 'int *[2]'
 // CHECK-NEXT: ImplicitCastExpr
-// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_z' 'int'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' <LValueToRValue>
+// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue .
+// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_foo_inner_z' '__wrapper_class'
 // CHECK-NEXT: ImplicitCastExpr
-// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_z' 'int'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' <LValueToRValue>
+// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue .
+// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_foo_inner_z' '__wrapper_class'
 // Initializer for second element of inner struct array
 // CHECK-NEXT: InitListExpr {{.*}} 'foo_inner'
 // CHECK-NEXT: ImplicitCastExpr
 // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_x' 'int'
 // CHECK-NEXT: ImplicitCastExpr
 // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_y' 'int'
-// CHECK-NEXT: InitListExpr {{.*}} 'int [2]'
+// CHECK-NEXT: InitListExpr {{.*}} 'int *[2]'
 // CHECK-NEXT: ImplicitCastExpr
-// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_z' 'int'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' <LValueToRValue>
+// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue .
+// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_foo_inner_z' '__wrapper_class'
 // CHECK-NEXT: ImplicitCastExpr
-// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_z' 'int'
-// CHECK-NEXT: InitListExpr {{.*}} 'int [2][1]'
-// CHECK-NEXT: InitListExpr {{.*}} 'int [1]'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' <LValueToRValue>
+// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue .
+// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_foo_inner_z' '__wrapper_class'
+// CHECK-NEXT: InitListExpr {{.*}} 'int *[2][1]'
+// CHECK-NEXT: InitListExpr {{.*}} 'int *[1]'
 // CHECK-NEXT: ImplicitCastExpr
-// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_2D' 'int'
-// CHECK-NEXT: InitListExpr {{.*}} 'int [1]'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *'
+// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue .
+// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_foo_2D' '__wrapper_class'
+// CHECK-NEXT: InitListExpr {{.*}} 'int *[1]'
 // CHECK-NEXT: ImplicitCastExpr
-// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_2D' 'int'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *'
+// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue .
+// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_foo_2D' '__wrapper_class'
 // CHECK-NEXT: ImplicitCastExpr
-// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_c' 'int'
+// CHECK-NEXT: DeclRefExpr{{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_c' 'int'
 
 // Initializer for second element of struct_array
 // CHECK-NEXT: InitListExpr {{.*}} 'foo'
-// CHECK-NEXT: ImplicitCastExpr
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
 // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_a' 'int'
+
+// Initializer for struct array inside foo i.e. foo_inner foo_b[2]
 // CHECK-NEXT: InitListExpr {{.*}} 'foo_inner [2]'
+// Initializer for first element of inner struct array
 // CHECK-NEXT: InitListExpr {{.*}} 'foo_inner'
-// CHECK-NEXT: ImplicitCastExpr
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
 // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_x' 'int'
-// CHECK-NEXT: ImplicitCastExpr
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
 // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_y' 'int'
-// CHECK-NEXT: InitListExpr {{.*}} 'int [2]'
+// CHECK-NEXT: InitListExpr {{.*}} 'int *[2]'
 // CHECK-NEXT: ImplicitCastExpr
-// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_z' 'int'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' <LValueToRValue>
+// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue .
+// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_foo_inner_z' '__wrapper_class'
 // CHECK-NEXT: ImplicitCastExpr
-// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_z' 'int'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' <LValueToRValue>
+// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue .
+// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_foo_inner_z' '__wrapper_class'
+// Initializer for second element of inner struct array
 // CHECK-NEXT: InitListExpr {{.*}} 'foo_inner'
 // CHECK-NEXT: ImplicitCastExpr
 // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_x' 'int'
 // CHECK-NEXT: ImplicitCastExpr
 // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_y' 'int'
-// CHECK-NEXT: InitListExpr {{.*}} 'int [2]'
-// CHECK-NEXT: ImplicitCastExpr
-// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_z' 'int'
+// CHECK-NEXT: InitListExpr {{.*}} 'int *[2]'
 // CHECK-NEXT: ImplicitCastExpr
-// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_z' 'int'
-// CHECK-NEXT: InitListExpr {{.*}} 'int [2][1]'
-// CHECK-NEXT: InitListExpr {{.*}} 'int [1]'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' <LValueToRValue>
+// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue .
+// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_foo_inner_z' '__wrapper_class'
 // CHECK-NEXT: ImplicitCastExpr
-// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_2D' 'int'
-// CHECK-NEXT: InitListExpr {{.*}} 'int [1]'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' <LValueToRValue>
+// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue .
+// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_foo_inner_z' '__wrapper_class'
+// CHECK-NEXT: InitListExpr {{.*}} 'int *[2][1]'
+// CHECK-NEXT: InitListExpr {{.*}} 'int *[1]'
 // CHECK-NEXT: ImplicitCastExpr
-// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_2D' 'int'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *'
+// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue .
+// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_foo_2D' '__wrapper_class'
+// CHECK-NEXT: InitListExpr {{.*}} 'int *[1]'
 // CHECK-NEXT: ImplicitCastExpr
-// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_c' 'int'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *'
+// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue .
+// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_foo_2D' '__wrapper_class'
 
 // Check kernel_E parameters
-// CHECK: FunctionDecl {{.*}}kernel_E{{.*}} 'void (int, int, int)'
-// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_a 'int':'int'
-// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_a 'int':'int'
-// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_a 'int':'int'
+// CHECK: FunctionDecl {{.*}}kernel_E{{.*}} 'void (S<int>)'
+// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ 'S<int>':'S<int>'
 // CHECK-NEXT: CompoundStmt
 // CHECK-NEXT: DeclStmt
 // CHECK-NEXT: VarDecl {{.*}} used '(lambda at {{.*}}array-kernel-param.cpp{{.*}})' cinit
 // CHECK-NEXT: InitListExpr {{.*}} '(lambda at {{.*}}array-kernel-param.cpp{{.*}})'
-// CHECK-NEXT: InitListExpr {{.*}} 'S<int>'
-// CHECK-NEXT: InitListExpr {{.*}} 'int [3]'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int':'int'
-// CHECK-NEXT: DeclRefExpr {{.*}} 'int':'int'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int':'int'
-// CHECK-NEXT: DeclRefExpr {{.*}} 'int':'int'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int':'int'
-// CHECK-NEXT: DeclRefExpr {{.*}} 'int':'int'
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'S<int>':'S<int>' 'void (const S<int> &) noexcept'
+// CHECK-NEXT: ImplicitCastExpr
+// CHECK-NEXT: DeclRefExpr {{.*}} 'S<int>':'S<int>' lvalue ParmVar {{.*}} '_arg_' 'S<int>':'S<int>'
 
 // Check kernel_F parameters
-// CHECK: FunctionDecl {{.*}}kernel_F{{.*}} 'void (int, int, int, int, int, int)'
-// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ 'int'
-// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ 'int'
-// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ 'int'
-// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ 'int'
-// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ 'int'
-// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ 'int'
+// CHECK: FunctionDecl {{.*}}kernel_F{{.*}} 'void (__wrapper_class)'
+// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ '__wrapper_class'
 // Check kernel_F inits
 // CHECK-NEXT: CompoundStmt
 // CHECK-NEXT: DeclStmt
 // CHECK-NEXT: VarDecl {{.*}} cinit
 // CHECK-NEXT: InitListExpr
-// CHECK-NEXT: InitListExpr {{.*}} 'int [2][3]'
-// CHECK-NEXT: InitListExpr {{.*}} 'int [3]'
-// CHECK-NEXT: ImplicitCastExpr
-// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_' 'int'
-// CHECK-NEXT: ImplicitCastExpr
-// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_' 'int'
-// CHECK-NEXT: ImplicitCastExpr
-// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_' 'int'
-// CHECK-NEXT: InitListExpr {{.*}} 'int [3]'
-// CHECK-NEXT: ImplicitCastExpr
-// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_' 'int'
-// CHECK-NEXT: ImplicitCastExpr
-// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_' 'int'
-// CHECK-NEXT: ImplicitCastExpr
-// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_' 'int'
+// CHECK-NEXT: ArrayInitLoopExpr {{.*}} 'int [2][3]'
+// CHECK-NEXT: OpaqueValueExpr {{.*}} 'int [2][3]' lvalue
+// CHECK-NEXT: MemberExpr {{.*}} 'int [2][3]' lvalue .
+// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_' '__wrapper_class'
+// CHECK-NEXT: ArrayInitLoopExpr {{.*}} 'int [3]'
+// CHECK-NEXT: OpaqueValueExpr {{.*}} 'int [3]' lvalue
+// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'int [3]' lvalue
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int (*)[3]' <ArrayToPointerDecay>
+// CHECK-NEXT: OpaqueValueExpr {{.*}} 'int [2][3]' lvalue
+// CHECK-NEXT: MemberExpr {{.*}} 'int [2][3]' lvalue .
+// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_' '__wrapper_class'
+// CHECK-NEXT: ArrayInitIndexExpr {{.*}} 'unsigned long
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
+// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'int' lvalue
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <ArrayToPointerDecay>
+// CHECK-NEXT: OpaqueValueExpr {{.*}} 'int [3]' lvalue
+// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'int [3]' lvalue
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int (*)[3]' <ArrayToPointerDecay>
+// CHECK-NEXT: OpaqueValueExpr {{.*}} 'int [2][3]' lvalue
+// CHECK-NEXT: MemberExpr {{.*}} 'int [2][3]' lvalue .
+// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_' '__wrapper_class'
+// CHECK-NEXT: ArrayInitIndexExpr {{.*}} 'unsigned long
+// CHECK-NEXT: ArrayInitIndexExpr {{.*}} 'unsigned long
+
+// Check kernel_G parameters.
+// CHECK: FunctionDecl {{.*}}kernel_G{{.*}} 'void (__wrapper_class)'
+// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ '__wrapper_class'
+// Check kernel_G inits
+// CHECK-NEXT: CompoundStmt
+// CHECK-NEXT: DeclStmt
+// CHECK-NEXT: VarDecl {{.*}} cinit
+// CHECK-NEXT: InitListExpr
+// CHECK-NEXT: ArrayInitLoopExpr {{.*}} 'foo2 [2]'
+// CHECK-NEXT: OpaqueValueExpr {{.*}} 'foo2 [2]' lvalue
+// CHECK-NEXT: MemberExpr {{.*}} 'foo2 [2]' lvalue .
+// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_' '__wrapper_class'
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'foo2' 'void (const foo2 &) noexcept'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const foo2' lvalue <NoOp>
+// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'foo2' lvalue
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'foo2 *' <ArrayToPointerDecay>
+// CHECK-NEXT: OpaqueValueExpr {{.*}} 'foo2 [2]' lvalue
+// CHECK-NEXT: MemberExpr {{.*}} 'foo2 [2]' lvalue .
+// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_' '__wrapper_class'
+// CHECK-NEXT: ArrayInitIndexExpr {{.*}} 'unsigned long
diff --git a/clang/test/SemaSYCL/decomposition.cpp b/clang/test/SemaSYCL/decomposition.cpp
new file mode 100644
index 0000000000000..6830c944f239c
--- /dev/null
+++ b/clang/test/SemaSYCL/decomposition.cpp
@@ -0,0 +1,124 @@
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -ast-dump %s | FileCheck %s
+
+#include "Inputs/sycl.hpp"
+
+using namespace cl::sycl;
+
+struct has_acc {
+  accessor<char, 1, access::mode::read> acc;
+};
+
+struct acc_base : accessor<char, 1, access::mode::read> {
+  int i;
+};
+
+struct has_sampler {
+  sampler sampl;
+};
+
+struct has_spec_const {
+  ONEAPI::experimental::spec_constant<int, class f1> SC;
+};
+
+handler H;
+
+struct has_stream {
+  stream s1{0, 0, H};
+};
+
+struct has_half {
+  half h;
+};
+
+struct non_decomposed {
+  int i;
+  float f;
+  double d;
+};
+
+struct use_non_decomposed : non_decomposed {
+  non_decomposed member;
+  float f;
+  double d;
+};
+
+template <typename T>
+struct Test1 {
+  T a;
+  T b[2];
+  non_decomposed d;
+  int i;
+};
+
+template <typename T>
+struct Test2 : T {
+  non_decomposed d;
+  int i;
+};
+
+template <typename Name, typename Func>
+__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
+  kernelFunc();
+}
+
+int main() {
+
+  non_decomposed d;
+  non_decomposed ds[5];
+  use_non_decomposed d2;
+  use_non_decomposed d2s[5];
+  // Check to ensure that these are not decomposed.
+  kernel<class NonDecomp>([=]() { return d.i + ds[0].i + d2.i + d2s[0].i; });
+  // CHECK: FunctionDecl {{.*}}NonDecomp{{.*}} 'void (non_decomposed, __wrapper_class, use_non_decomposed, __wrapper_class)'
+
+  {
+    Test1<has_acc> t1;
+    kernel<class Acc1>([=]() { return t1.i; });
+    // CHECK: FunctionDecl {{.*}}Acc1{{.*}} 'void (__global char *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, __global char *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, __global char *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, non_decomposed, int)'
+    Test2<has_acc> t2;
+    kernel<class Acc2>([=]() { return t2.i; });
+    // CHECK: FunctionDecl {{.*}}Acc2{{.*}} 'void (__global char *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, non_decomposed, int)'
+    Test1<acc_base> t3;
+    kernel<class Acc3>([=]() { return t3.i; });
+    // CHECK: FunctionDecl {{.*}}Acc3{{.*}} 'void (__global char *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, int, __global char *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, int, __global char *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, int, non_decomposed, int)'
+    Test2<acc_base> t4;
+    kernel<class Acc4>([=]() { return t4.i; });
+    // CHECK: FunctionDecl {{.*}}Acc4{{.*}} 'void (__global char *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, int, non_decomposed, int)'
+  }
+
+  {
+    Test1<has_sampler> t1;
+    kernel<class Sampl1>([=]() { return t1.i; });
+    // CHECK: FunctionDecl {{.*}}Sampl1{{.*}} 'void (sampler_t, sampler_t, sampler_t, non_decomposed, int)'
+    Test2<has_sampler> t2;
+    kernel<class Sampl2>([=]() { return t2.i; });
+    // CHECK: FunctionDecl {{.*}}Sampl2{{.*}} 'void (sampler_t, non_decomposed, int)'
+  }
+
+  {
+    Test1<has_spec_const> t1;
+    kernel<class SpecConst1>([=]() { return t1.i; });
+    // CHECK: FunctionDecl {{.*}}SpecConst{{.*}} 'void (non_decomposed, int)'
+    Test2<has_spec_const> t2;
+    kernel<class SpecConst2>([=]() { return t2.i; });
+    // CHECK: FunctionDecl {{.*}}SpecConst2{{.*}} 'void (non_decomposed, int)'
+  }
+
+  {
+    Test1<has_stream> t1;
+    kernel<class Stream1>([=]() { return t1.i; });
+    // CHECK: FunctionDecl {{.*}}Stream1{{.*}} 'void (cl::sycl::stream, __global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, cl::sycl::stream, __global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, cl::sycl::stream, __global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, non_decomposed, int)'
+    Test2<has_stream> t2;
+    kernel<class Stream2>([=]() { return t2.i; });
+    // CHECK: FunctionDecl {{.*}}Stream2{{.*}} 'void (cl::sycl::stream, __global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, non_decomposed, int)'
+  }
+
+  {
+    Test1<has_half> t1;
+    kernel<class Half1>([=]() { return t1.i; });
+    // CHECK: FunctionDecl {{.*}}Half1{{.*}} 'void (cl::sycl::half, cl::sycl::half, cl::sycl::half, non_decomposed, int)'
+    Test2<has_half> t2;
+    kernel<class Half2>([=]() { return t2.i; });
+    // CHECK: FunctionDecl {{.*}}Half2{{.*}} 'void (cl::sycl::half, non_decomposed, int)'
+  }
+}
diff --git a/clang/test/SemaSYCL/fake-accessors.cpp b/clang/test/SemaSYCL/fake-accessors.cpp
index 3a7a44aa8dc68..d4dd2a0f60209 100644
--- a/clang/test/SemaSYCL/fake-accessors.cpp
+++ b/clang/test/SemaSYCL/fake-accessors.cpp
@@ -51,6 +51,6 @@ int main() {
         });
   return 0;
 }
-// CHECK: fake_accessors{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, int, int)
-// CHECK: accessor_typedef{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, int, int)
-// CHECK: accessor_alias{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, int, int)
+// CHECK: fake_accessors{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor)
+// CHECK: accessor_typedef{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor)
+// CHECK: accessor_alias{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor)
diff --git a/clang/test/SemaSYCL/inheritance.cpp b/clang/test/SemaSYCL/inheritance.cpp
index f8dcfaf84af94..ab1f505619fee 100644
--- a/clang/test/SemaSYCL/inheritance.cpp
+++ b/clang/test/SemaSYCL/inheritance.cpp
@@ -4,7 +4,7 @@
 
 class second_base {
 public:
-  int e;
+  int *e;
 };
 
 class InnerFieldBase {
@@ -40,13 +40,11 @@ int main() {
 }
 
 // Check declaration of the kernel
-// CHECK: derived{{.*}} 'void (int, int, int, int, int)'
+// CHECK: derived{{.*}} 'void (base, __wrapper_class, int)
 
 // Check parameters of the kernel
-// CHECK: ParmVarDecl {{.*}} used _arg_b 'int'
-// CHECK: ParmVarDecl {{.*}} used _arg_d 'int'
-// CHECK: ParmVarDecl {{.*}} used _arg_c 'int'
-// CHECK: ParmVarDecl {{.*}} used _arg_e 'int'
+// CHECK: ParmVarDecl {{.*}} used _arg__base 'base'
+// CHECK: ParmVarDecl {{.*}} used _arg_e '__wrapper_class'
 // CHECK: ParmVarDecl {{.*}} used _arg_a 'int'
 
 // Check initializers for derived and base classes.
@@ -54,17 +52,13 @@ int main() {
 // Base classes should be initialized first.
 // CHECK: VarDecl {{.*}} derived 'derived' cinit
 // CHECK-NEXT: InitListExpr {{.*}} 'derived'
-// CHECK-NEXT: InitListExpr {{.*}} 'base'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
-// CHECK-NEXT: DeclRefExpr {{.*}} lvalue ParmVar {{.*}} '_arg_b' 'int'
-// CHECK-NEXT: InitListExpr {{.*}} 'InnerField'
-// CHECK-NEXT: InitListExpr {{.*}} 'InnerFieldBase'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
-// CHECK-NEXT: DeclRefExpr {{.*}} lvalue ParmVar {{.*}} '_arg_d' 'int'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
-// CHECK-NEXT: DeclRefExpr {{.*}} lvalue ParmVar {{.*}} '_arg_c' 'int'
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'base' 'void (const base &) noexcept'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const base' lvalue <NoOp>
+// CHECK-NEXT: DeclRefExpr {{.*}} lvalue ParmVar {{.*}} '_arg__base' 'base'
 // CHECK-NEXT: InitListExpr {{.*}} 'second_base'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
-// CHECK-NEXT: DeclRefExpr {{.*}} lvalue ParmVar {{.*}} '_arg_e' 'int'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <AddressSpaceConversion>
+// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' <LValueToRValue>
+// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue .
+// CHECK-NEXT: DeclRefExpr {{.*}} lvalue ParmVar {{.*}} '_arg_e' '__wrapper_class'
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
 // CHECK-NEXT: DeclRefExpr {{.*}} lvalue ParmVar {{.*}} '_arg_a' 'int'
diff --git a/clang/test/SemaSYCL/union-kernel-param2.cpp b/clang/test/SemaSYCL/union-kernel-param2.cpp
index 6fdd90f672993..0fd340b343b5e 100644
--- a/clang/test/SemaSYCL/union-kernel-param2.cpp
+++ b/clang/test/SemaSYCL/union-kernel-param2.cpp
@@ -25,7 +25,7 @@ int main() {
       float b;
       char c;
     } union_mem;
-    int d;
+    int *d;
   } struct_mem;
 
   a_kernel<class kernel_A>(
@@ -53,9 +53,9 @@ int main() {
 // CHECK-NEXT: DeclRefExpr {{.*}} 'union MyUnion':'MyUnion' lvalue ParmVar {{.*}} '_arg_' 'union MyUnion':'MyUnion'
 
 // Check kernel_B parameters
-// CHECK: FunctionDecl {{.*}}kernel_B{{.*}} 'void (union MyUnion, int)'
+// CHECK: FunctionDecl {{.*}}kernel_B{{.*}} 'void (union MyUnion, __wrapper_class)'
 // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_union_mem 'union MyUnion':'MyStruct::MyUnion'
-// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_d 'int'
+// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_d '__wrapper_class'
 
 // Check kernel_B inits
 // CHECK-NEXT: CompoundStmt
@@ -67,4 +67,6 @@ int main() {
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'const MyStruct::MyUnion'
 // CHECK-NEXT: DeclRefExpr {{.*}} 'union MyUnion':'MyStruct::MyUnion' lvalue ParmVar {{.*}} '_arg_union_mem' 'union MyUnion':'MyStruct::MyUnion'
 // CHECK-NEXT: ImplicitCastExpr
-// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_d' 'int'
+// CHECK-NEXT: ImplicitCastExpr
+// CHECK-NEXT: MemberExpr
+// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_d' '__wrapper_class'